Skip to content

Commit a4fe915

Browse files
committed
add ir attribute sycl_forceinline to group_local_memory
1 parent 97add86 commit a4fe915

File tree

3 files changed

+12
-11
lines changed

3 files changed

+12
-11
lines changed

llvm/lib/SYCLLowerIR/LowerWGLocalMemory.cpp

Lines changed: 1 addition & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -99,13 +99,7 @@ static bool inlineGroupLocalMemoryFunc(Module &M) {
9999
bool Changed = false;
100100
for (auto *U : ALMFunc->users()) {
101101
auto *Caller = cast<CallInst>(U)->getFunction();
102-
if (!Caller->hasFnAttribute(Attribute::AlwaysInline)) {
103-
// Already inlined.
104-
continue;
105-
}
106-
std::string FName = llvm::demangle(Caller->getName());
107-
if (FName.find("sycl::_V1::ext::oneapi::group_local_memory") ==
108-
std::string::npos) {
102+
if (!Caller->hasFnAttribute("sycl_forceinline")) {
109103
// Already inlined.
110104
continue;
111105
}

llvm/test/SYCLLowerIR/group_local_memory_inline.ll

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -39,7 +39,7 @@ entry:
3939
; CHECK-NOT: define {{.*}} @_ZN4sycl3_V13ext6oneapi32group_local_memory_for_overwriteIiNS0_5groupILi1EEEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_(
4040

4141
; Function Attrs: alwaysinline
42-
define spir_func void @_ZN4sycl3_V13ext6oneapi32group_local_memory_for_overwriteIiNS0_5groupILi1EEEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %agg.result, ptr byval(%"class.sycl::_V1::group") align 8 %g) #0 {
42+
define spir_func void @_ZN4sycl3_V13ext6oneapi32group_local_memory_for_overwriteIiNS0_5groupILi1EEEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %agg.result, ptr byval(%"class.sycl::_V1::group") align 8 %g) #1 {
4343
entry:
4444
%AllocatedMem = alloca ptr addrspace(3), align 8
4545
%AllocatedMem.ascast = addrspacecast ptr %AllocatedMem to ptr addrspace(4)
@@ -51,7 +51,7 @@ entry:
5151
; CHECK-NOT: define {{.*}} @_ZN4sycl3_V13ext6oneapi18group_local_memoryIiNS0_5groupILi1EEEJEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_DpOT1_(
5252

5353
; Function Attrs: alwaysinline
54-
define spir_func void @_ZN4sycl3_V13ext6oneapi18group_local_memoryIiNS0_5groupILi1EEEJEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_DpOT1_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %agg.result, ptr byval(%"class.sycl::_V1::group") align 8 %g) #0 {
54+
define spir_func void @_ZN4sycl3_V13ext6oneapi18group_local_memoryIiNS0_5groupILi1EEEJEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_DpOT1_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %agg.result, ptr byval(%"class.sycl::_V1::group") align 8 %g) #1 {
5555
entry:
5656
%AllocatedMem = alloca ptr addrspace(3), align 8
5757
%AllocatedMem.ascast = addrspacecast ptr %AllocatedMem to ptr addrspace(4)
@@ -63,3 +63,4 @@ entry:
6363
declare spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 noundef, i64 noundef)
6464

6565
attributes #0 = { alwaysinline }
66+
attributes #1 = { "sycl_forceinline"="true" }

sycl/include/sycl/ext/oneapi/group_local_memory.hpp

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,10 +21,13 @@ namespace sycl {
2121
inline namespace _V1 {
2222
namespace ext::oneapi {
2323
template <typename T, typename Group>
24+
#ifdef __SYCL_DEVICE_ONLY__
25+
[[__sycl_detail__::add_ir_attributes_function("sycl_forceinline", true)]]
26+
#endif
2427
std::enable_if_t<
2528
std::is_trivially_destructible_v<T> && sycl::detail::is_group<Group>::value,
2629
multi_ptr<T, access::address_space::local_space, access::decorated::legacy>>
27-
__SYCL_ALWAYS_INLINE group_local_memory_for_overwrite(Group g) {
30+
group_local_memory_for_overwrite(Group g) {
2831
(void)g;
2932
#ifdef __SYCL_DEVICE_ONLY__
3033
__attribute__((opencl_local)) std::uint8_t *AllocatedMem =
@@ -44,10 +47,13 @@ std::enable_if_t<
4447
}
4548

4649
template <typename T, typename Group, typename... Args>
50+
#ifdef __SYCL_DEVICE_ONLY__
51+
[[__sycl_detail__::add_ir_attributes_function("sycl_forceinline", true)]]
52+
#endif
4753
std::enable_if_t<
4854
std::is_trivially_destructible_v<T> && sycl::detail::is_group<Group>::value,
4955
multi_ptr<T, access::address_space::local_space, access::decorated::legacy>>
50-
__SYCL_ALWAYS_INLINE group_local_memory(Group g, Args &&...args) {
56+
group_local_memory(Group g, Args &&...args) {
5157
#ifdef __SYCL_DEVICE_ONLY__
5258
__attribute__((opencl_local)) std::uint8_t *AllocatedMem =
5359
__sycl_allocateLocalMemory(sizeof(T), alignof(T));

0 commit comments

Comments
 (0)