add ir attribute sycl_forceinline to group_local_memory

Author: wenju-he

llvm/lib/SYCLLowerIR/LowerWGLocalMemory.cpp

@@ -99,13 +99,7 @@ static bool inlineGroupLocalMemoryFunc(Module &M) {
   bool Changed = false;
   for (auto *U : ALMFunc->users()) {
     auto *Caller = cast<CallInst>(U)->getFunction();
-    if (!Caller->hasFnAttribute(Attribute::AlwaysInline)) {
-      // Already inlined.
-      continue;
-    }
-    std::string FName = llvm::demangle(Caller->getName());
-    if (FName.find("sycl::_V1::ext::oneapi::group_local_memory") ==
-        std::string::npos) {
+    if (!Caller->hasFnAttribute("sycl_forceinline")) {
       // Already inlined.
       continue;
     }

llvm/test/SYCLLowerIR/group_local_memory_inline.ll

@@ -39,7 +39,7 @@ entry:
 ; CHECK-NOT: define {{.*}} @_ZN4sycl3_V13ext6oneapi32group_local_memory_for_overwriteIiNS0_5groupILi1EEEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_(
 
 ; Function Attrs: alwaysinline
-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 {
+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 {
 entry:
   %AllocatedMem = alloca ptr addrspace(3), align 8
   %AllocatedMem.ascast = addrspacecast ptr %AllocatedMem to ptr addrspace(4)
@@ -51,7 +51,7 @@ entry:
 ; CHECK-NOT: define {{.*}} @_ZN4sycl3_V13ext6oneapi18group_local_memoryIiNS0_5groupILi1EEEJEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_DpOT1_(
 
 ; Function Attrs: alwaysinline
-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 {
+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 {
 entry:
   %AllocatedMem = alloca ptr addrspace(3), align 8
   %AllocatedMem.ascast = addrspacecast ptr %AllocatedMem to ptr addrspace(4)
@@ -63,3 +63,4 @@ entry:
 declare spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 noundef, i64 noundef)
 
 attributes #0 = { alwaysinline }
+attributes #1 = { "sycl_forceinline"="true" }

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

@@ -21,10 +21,13 @@ namespace sycl {
 inline namespace _V1 {
 namespace ext::oneapi {
 template <typename T, typename Group>
+#ifdef __SYCL_DEVICE_ONLY__
+[[__sycl_detail__::add_ir_attributes_function("sycl_forceinline", true)]]
+#endif
 std::enable_if_t<
     std::is_trivially_destructible_v<T> && sycl::detail::is_group<Group>::value,
     multi_ptr<T, access::address_space::local_space, access::decorated::legacy>>
-    __SYCL_ALWAYS_INLINE group_local_memory_for_overwrite(Group g) {
+group_local_memory_for_overwrite(Group g) {
   (void)g;
 #ifdef __SYCL_DEVICE_ONLY__
   __attribute__((opencl_local)) std::uint8_t *AllocatedMem =
@@ -44,10 +47,13 @@ std::enable_if_t<
 }
 
 template <typename T, typename Group, typename... Args>
+#ifdef __SYCL_DEVICE_ONLY__
+[[__sycl_detail__::add_ir_attributes_function("sycl_forceinline", true)]]
+#endif
 std::enable_if_t<
     std::is_trivially_destructible_v<T> && sycl::detail::is_group<Group>::value,
     multi_ptr<T, access::address_space::local_space, access::decorated::legacy>>
-    __SYCL_ALWAYS_INLINE group_local_memory(Group g, Args &&...args) {
+group_local_memory(Group g, Args &&...args) {
 #ifdef __SYCL_DEVICE_ONLY__
   __attribute__((opencl_local)) std::uint8_t *AllocatedMem =
       __sycl_allocateLocalMemory(sizeof(T), alignof(T));