[SYCL][SYCLLowerWGLocalMemoryPass] Remove implicit dependency on AlwaysInlinerPass and move to PipelineStart #16356
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -9,15 +9,17 @@ | |
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h" | ||
| #include "llvm/Demangle/Demangle.h" | ||
| #include "llvm/IR/Function.h" | ||
| #include "llvm/IR/IRBuilder.h" | ||
| #include "llvm/IR/InstIterator.h" | ||
| #include "llvm/Pass.h" | ||
| #include "llvm/TargetParser/Triple.h" | ||
| #include "llvm/Transforms/Utils/Cloning.h" | ||
|
|
||
| using namespace llvm; | ||
|
|
||
| #define DEBUG_TYPE "sycllowerwglocalmemory" | ||
|
|
||
| static constexpr char SYCL_ALLOCLOCALMEM_CALL[] = "__sycl_allocateLocalMemory"; | ||
| static constexpr char SYCL_DYNAMIC_LOCALMEM_CALL[] = | ||
|
|
@@ -84,6 +86,36 @@ ModulePass *llvm::createSYCLLowerWGLocalMemoryLegacyPass() { | |
| return new SYCLLowerWGLocalMemoryLegacy(); | ||
| } | ||
|
|
||
| // In sycl header __sycl_allocateLocalMemory builtin call is wrapped in | ||
| // group_local_memory/group_local_memory_for_overwrite functions, which must be | ||
| // inlined first before each __sycl_allocateLocalMemory call can be lowered to a | ||
| // unique global variable. Inlining them here so that this pass doesn't have | ||
| // implicit dependency on AlwaysInlinerPass. | ||
| static bool inlineGroupLocalMemoryFunc(Module &M) { | ||
| Function *ALMFunc = M.getFunction(SYCL_ALLOCLOCALMEM_CALL); | ||
| if (!ALMFunc || ALMFunc->use_empty()) | ||
| return false; | ||
|
|
||
| bool Changed = false; | ||
| for (auto *U : ALMFunc->users()) { | ||
|
There was a problem hiding this comment. we need to use a work list here rather than the simple loop. This function https://github.com/intel/llvm/blob/sycl/sycl/include/syclcompat/memory.hpp#L71 needs to be updated as well, and this function won't be able to handle the nesting. The CI is currently green because there is no test requesting 2 distinct local memory objects using this function in the same kernel. There was a problem hiding this comment. done, thank you for the suggestion. Now I understand what you mean by |
||
| auto *Caller = cast<CallInst>(U)->getFunction(); | ||
| if (!Caller->hasFnAttribute("sycl_forceinline")) { | ||
| // Already inlined. | ||
| continue; | ||
| } | ||
| for (auto *U2 : make_early_inc_range(Caller->users())) { | ||
| auto *CI = cast<CallInst>(U2); | ||
| InlineFunctionInfo IFI; | ||
| [[maybe_unused]] auto Result = InlineFunction(*CI, IFI); | ||
| assert(Result.isSuccess() && "inlining failed"); | ||
| } | ||
| Caller->eraseFromParent(); | ||
| Changed = true; | ||
| } | ||
|
|
||
| return Changed; | ||
| } | ||
|
|
||
| // TODO: It should be checked that __sycl_allocateLocalMemory (or its source | ||
| // form - group_local_memory) does not occur: | ||
| // - in a function (other than user lambda/functor) | ||
|
|
@@ -317,9 +349,8 @@ static bool dynamicWGLocalMemory(Module &M) { | |
|
|
||
| PreservedAnalyses SYCLLowerWGLocalMemoryPass::run(Module &M, | ||
| ModuleAnalysisManager &) { | ||
| bool Changed = inlineGroupLocalMemoryFunc(M); | ||
| Changed |= allocaWGLocalMemory(M); | ||
| Changed |= dynamicWGLocalMemory(M); | ||
| return Changed ? PreservedAnalyses::none() : PreservedAnalyses::all(); | ||
jsji marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| } | ||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,66 @@ | ||
| ; RUN: opt < %s -passes=sycllowerwglocalmemory -S | FileCheck %s | ||
|
|
||
| ; Check group_local_memory_for_overwrite and group_local_memory functions are inlined. | ||
| ; Check __sycl_allocateLocalMemory calls are lowered to four separate allocations. | ||
|
|
||
| target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" | ||
| target triple = "spir64-unknown-unknown" | ||
|
|
||
| %"class.sycl::_V1::multi_ptr" = type { ptr addrspace(3) } | ||
| %"class.sycl::_V1::group" = type { %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::id" } | ||
| %"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" } | ||
| %"class.sycl::_V1::detail::array" = type { [1 x i64] } | ||
| %"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" } | ||
|
|
||
| ; CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4 | ||
| ; CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4 | ||
| ; CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4 | ||
| ; CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4 | ||
|
|
||
| ; Function Attrs: alwaysinline | ||
| define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlNS0_7nd_itemILi1EEEE_clES5_() #0 { | ||
| entry: | ||
| ; CHECK: define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlNS0_7nd_itemILi1EEEE_clES5_( | ||
| ; CHECK: store ptr addrspace(3) @WGLocalMem{{.*}}, ptr addrspace(4) %AllocatedMem{{.*}}, align 8 | ||
| ; CHECK: store ptr addrspace(3) @WGLocalMem{{.*}}, ptr addrspace(4) %AllocatedMem{{.*}}, align 8 | ||
| ; CHECK: store ptr addrspace(3) @WGLocalMem{{.*}}, ptr addrspace(4) %AllocatedMem{{.*}}, align 8 | ||
| ; CHECK: store ptr addrspace(3) @WGLocalMem{{.*}}, ptr addrspace(4) %AllocatedMem{{.*}}, align 8 | ||
|
|
||
| %Ptr = alloca %"class.sycl::_V1::multi_ptr", align 8 | ||
| %agg = alloca %"class.sycl::_V1::group", align 8 | ||
| %Ptr.ascast = addrspacecast ptr %Ptr to ptr addrspace(4) | ||
| call 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 %Ptr.ascast, ptr byval(%"class.sycl::_V1::group") align 8 %agg) | ||
| call 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 %Ptr.ascast, ptr byval(%"class.sycl::_V1::group") align 8 %agg) | ||
| call 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 %Ptr.ascast, ptr byval(%"class.sycl::_V1::group") align 8 %agg) | ||
| call 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 %Ptr.ascast, ptr byval(%"class.sycl::_V1::group") align 8 %agg) | ||
| ret void | ||
| } | ||
|
|
||
| ; 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) #1 { | ||
| entry: | ||
| %AllocatedMem = alloca ptr addrspace(3), align 8 | ||
| %AllocatedMem.ascast = addrspacecast ptr %AllocatedMem to ptr addrspace(4) | ||
| %call = call spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 4, i64 4) | ||
| store ptr addrspace(3) %call, ptr addrspace(4) %AllocatedMem.ascast, align 8 | ||
| ret void | ||
| } | ||
|
|
||
| ; 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) #1 { | ||
| entry: | ||
| %AllocatedMem = alloca ptr addrspace(3), align 8 | ||
| %AllocatedMem.ascast = addrspacecast ptr %AllocatedMem to ptr addrspace(4) | ||
| %call = call spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 4, i64 4) | ||
| store ptr addrspace(3) %call, ptr addrspace(4) %AllocatedMem.ascast, align 8 | ||
| ret void | ||
| } | ||
|
|
||
| declare spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 noundef, i64 noundef) | ||
|
|
||
| attributes #0 = { alwaysinline } | ||
| attributes #1 = { "sycl_forceinline"="true" } |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,38 @@ | ||
| // RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s | ||
| // RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm %s -fno-sycl-early-optimizations -o - | FileCheck %s | ||
| // RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm %s -O0 -o - | FileCheck %s | ||
|
|
||
| // The test checks that multiple calls to the same template instantiation of a | ||
| // group local memory function result in separate allocations. | ||
|
|
||
| #include <sycl/detail/core.hpp> | ||
| #include <sycl/ext/oneapi/group_local_memory.hpp> | ||
| #include <sycl/usm.hpp> | ||
|
|
||
| using namespace sycl; | ||
|
|
||
| int main() { | ||
| queue Q; | ||
|
|
||
| int **Out = malloc_shared<int *>(4, Q); | ||
|
|
||
| Q.submit([&](handler &Cgh) { | ||
| Cgh.parallel_for(nd_range<1>({1}, {1}), [=](nd_item<1> Item) { | ||
| auto Ptr0 = | ||
| ext::oneapi::group_local_memory_for_overwrite<int>(Item.get_group()); | ||
| auto Ptr1 = | ||
| ext::oneapi::group_local_memory_for_overwrite<int>(Item.get_group()); | ||
| auto Ptr2 = ext::oneapi::group_local_memory<int>(Item.get_group()); | ||
| auto Ptr3 = ext::oneapi::group_local_memory<int>(Item.get_group()); | ||
| Out[0] = Ptr0; | ||
| Out[1] = Ptr1; | ||
| Out[2] = Ptr2; | ||
| Out[3] = Ptr3; | ||
| }); | ||
| }); | ||
| } | ||
|
|
||
| // CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4 | ||
| // CHECK-NEXT: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4 | ||
| // CHECK-NEXT: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4 | ||
| // CHECK-NEXT: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4 |
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Why can we not rewrite the SYCL headers to 'inline' these calls? Is there a specific reason? Thanks
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
We can't ask users to call
__sycl_allocateLocalMemoryinternal intrsinsic when documented interface issycl::ext::something::something::group_local_memory<T>