From fb018e596d039d2cef8a802520d8363451ed9f32 Mon Sep 17 00:00:00 2001 From: Gregory Fine Date: Wed, 21 Feb 2024 21:14:37 -0800 Subject: [PATCH 01/17] Initial Implementation --- llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp | 42 ++++++++++++++++++++ 1 file changed, 42 insertions(+) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index 88c95c8005fa2..4361b66e3b465 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -12,6 +12,7 @@ //===----------------------------------------------------------------------===// #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h" +#include "llvm/ADT/SmallSet.h" #include "llvm/Demangle/Demangle.h" #include "llvm/Demangle/ItaniumDemangle.h" #include "llvm/IR/InstIterator.h" @@ -20,6 +21,7 @@ #include "llvm/InitializePasses.h" #include "llvm/Pass.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" +#include "llvm/SYCLLowerIR/SYCLUtils.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" #include "llvm/Support/Regex.h" @@ -97,6 +99,7 @@ class ESIMDVerifierImpl { void verify() { SmallPtrSet Visited; SmallVector Worklist; + SmallSet SLMKernelsSeen; auto Add2Worklist = [&Worklist, &Visited](const Function *F) { if (Visited.insert(F).second) @@ -112,6 +115,45 @@ class ESIMDVerifierImpl { // for invalid calls. while (!Worklist.empty()) { const Function *F = Worklist.pop_back_val(); + if (isSlmInit(*F)) { + // Filter function for graph traversal when propagating ESIMD attribute. + // While traversing the call graph, non-call use of the traversed + // function is not added to the graph. The reason is that it is + // impossible to gurantee correct inference of use of that function, in + // particular to determine if that function is used as an argument for + // invoke_simd. As a result, any use of function pointers requires + // explicit marking of the functions as ESIMD_FUNCTION if needed. + auto filterInvokeSimdUse = [](const Instruction *, const Function *) { + return false; + }; + + sycl::utils::traverseCallgraphUp( + const_cast(F), + [&](Function *GraphNode) { + if (llvm::esimd::isESIMDKernel(*GraphNode)) { + StringRef KernelName = GraphNode->getName(); + if (SLMKernelsSeen.count(KernelName) != 0) { + std::string ErrorMsg = + std::string("slm_init is called more than once from kernel '") + + demangle(KernelName.str()) + "'."; + GraphNode->getContext().emitError(ErrorMsg); + } else { + SLMKernelsSeen.insert(KernelName); + } + } else { + if (!GraphNode->hasFnAttribute(llvm::Attribute::AlwaysInline)) { + StringRef MangledName = GraphNode->getName(); + std::string ErrorMsg = std::string("slm_init is called from function '") + + demangle(MangledName.str()) + + "' which is not guaranteed to be inlined."; + GraphNode->getContext().emitError(ErrorMsg); + } + } + }, + false, filterInvokeSimdUse); + } + + for (const Instruction &I : instructions(F)) { if (auto *CB = dyn_cast(&I)) { Function *Callee = CB->getCalledFunction(); From 10fceb87393868d23bbd09bb2c398fd83c742d8a Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Thu, 22 Feb 2024 12:12:03 -0800 Subject: [PATCH 02/17] rework the solution and add test --- llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp | 40 ------------------- .../ESIMD/LowerESIMDSlmReservation.cpp | 13 ++++-- sycl/test/esimd/slm_init_check.cpp | 22 ++++++++++ 3 files changed, 32 insertions(+), 43 deletions(-) create mode 100644 sycl/test/esimd/slm_init_check.cpp diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index 4361b66e3b465..46be89de9d1a0 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -12,7 +12,6 @@ //===----------------------------------------------------------------------===// #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h" -#include "llvm/ADT/SmallSet.h" #include "llvm/Demangle/Demangle.h" #include "llvm/Demangle/ItaniumDemangle.h" #include "llvm/IR/InstIterator.h" @@ -99,7 +98,6 @@ class ESIMDVerifierImpl { void verify() { SmallPtrSet Visited; SmallVector Worklist; - SmallSet SLMKernelsSeen; auto Add2Worklist = [&Worklist, &Visited](const Function *F) { if (Visited.insert(F).second) @@ -115,44 +113,6 @@ class ESIMDVerifierImpl { // for invalid calls. while (!Worklist.empty()) { const Function *F = Worklist.pop_back_val(); - if (isSlmInit(*F)) { - // Filter function for graph traversal when propagating ESIMD attribute. - // While traversing the call graph, non-call use of the traversed - // function is not added to the graph. The reason is that it is - // impossible to gurantee correct inference of use of that function, in - // particular to determine if that function is used as an argument for - // invoke_simd. As a result, any use of function pointers requires - // explicit marking of the functions as ESIMD_FUNCTION if needed. - auto filterInvokeSimdUse = [](const Instruction *, const Function *) { - return false; - }; - - sycl::utils::traverseCallgraphUp( - const_cast(F), - [&](Function *GraphNode) { - if (llvm::esimd::isESIMDKernel(*GraphNode)) { - StringRef KernelName = GraphNode->getName(); - if (SLMKernelsSeen.count(KernelName) != 0) { - std::string ErrorMsg = - std::string("slm_init is called more than once from kernel '") + - demangle(KernelName.str()) + "'."; - GraphNode->getContext().emitError(ErrorMsg); - } else { - SLMKernelsSeen.insert(KernelName); - } - } else { - if (!GraphNode->hasFnAttribute(llvm::Attribute::AlwaysInline)) { - StringRef MangledName = GraphNode->getName(); - std::string ErrorMsg = std::string("slm_init is called from function '") + - demangle(MangledName.str()) + - "' which is not guaranteed to be inlined."; - GraphNode->getContext().emitError(ErrorMsg); - } - } - }, - false, filterInvokeSimdUse); - } - for (const Instruction &I : instructions(F)) { if (auto *CB = dyn_cast(&I)) { diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp index 323e16faa2171..29aa63c231ace 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp @@ -64,6 +64,8 @@ #include "llvm/ADT/SmallPtrSet.h" #include "llvm/ADT/SmallVector.h" +#include "llvm/Demangle/Demangle.h" +#include "llvm/Demangle/ItaniumDemangle.h" #include "llvm/IR/Function.h" #include "llvm/IR/Instruction.h" #include "llvm/IR/Module.h" @@ -362,9 +364,14 @@ class ScopedCallGraph { if (auto *CB = dyn_cast(&I)) { if (isSlmInitCall(dyn_cast(CB))) { auto *CI = dyn_cast(CB); - esimd::assert_and_diag(!SlmInitCall, - "multiple slm_init calls in function ", - F.getName()); + if (SlmInitCall) { + std::string ErrorMsg = + std::string( + "slm_init is called more than once from function '") + + demangle(F.getName().str()) + "'."; + F.getContext().emitError(ErrorMsg); + } + // TODO: this diagnostics incorrectly fires on functor's // operator() marked as SYCL_ESIMD_KERNEL, because becomes neither // spir_kernel nor SYCL_EXERNAL function in IR. It rather becomes diff --git a/sycl/test/esimd/slm_init_check.cpp b/sycl/test/esimd/slm_init_check.cpp new file mode 100644 index 0000000000000..43a442b19195d --- /dev/null +++ b/sycl/test/esimd/slm_init_check.cpp @@ -0,0 +1,22 @@ +// RUN: not %clangxx -fsycl %s 2>&1 | FileCheck %s + +// This test verifies more than 1 call to slm_init triggers an error. + +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; + +int main() { + queue Q; + nd_range<1> NDR{range<1>{2}, range<1>{2}}; + Q.parallel_for(NDR, [=](nd_item<1> NDI) SYCL_ESIMD_KERNEL { + slm_init(1024); + slm_init(1024); + }).wait(); + // CHECK: error: slm_init is called more than once from function 'typeinfo name for main::'lambda'(sycl::_V1::nd_item<1>)'. + + return 0; +} From 95febd75042ee84c8fb578e686e69fd7e355decd Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Thu, 22 Feb 2024 12:15:25 -0800 Subject: [PATCH 03/17] Remove unnecessary changes --- llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index 46be89de9d1a0..88c95c8005fa2 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -20,7 +20,6 @@ #include "llvm/InitializePasses.h" #include "llvm/Pass.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" -#include "llvm/SYCLLowerIR/SYCLUtils.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" #include "llvm/Support/Regex.h" @@ -113,7 +112,6 @@ class ESIMDVerifierImpl { // for invalid calls. while (!Worklist.empty()) { const Function *F = Worklist.pop_back_val(); - for (const Instruction &I : instructions(F)) { if (auto *CB = dyn_cast(&I)) { Function *Callee = CB->getCalledFunction(); From 26c5eaf6d283a099fa2aa050f76694ace86cdbc3 Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Sun, 25 Feb 2024 10:43:48 -0800 Subject: [PATCH 04/17] Refactor the checks to cover more cases --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 38 ++++++++++++++++++- .../ESIMD/LowerESIMDSlmReservation.cpp | 12 ++---- sycl/test/esimd/slm_init_check.cpp | 2 +- sycl/test/esimd/slm_init_noinline_check.cpp | 31 +++++++++++++++ 4 files changed, 73 insertions(+), 10 deletions(-) create mode 100644 sycl/test/esimd/slm_init_noinline_check.cpp diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index e370afb494fdf..760f2e42cb334 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1750,6 +1750,41 @@ void lowerGlobalsToVector(Module &M) { } // namespace +static void checkSLMInit(Module &M) { + SmallPtrSet Callers; + for (auto &F : M) { + if (isSlmInit(F)) { + auto filterInvokeSimdUse = [](const Instruction *, const Function *) { + return false; + }; + for (User *U : F.users()) { + auto *FCall = dyn_cast(U); + if (FCall && FCall->getCalledFunction() == &F) { + Function *GenF = FCall->getFunction(); + + sycl::utils::traverseCallgraphUp( + GenF, + [&](Function *GraphNode) { + if (llvm::esimd::isESIMDKernel(*GraphNode)) { + if (Callers.contains(GraphNode)) { + StringRef KernelName = GraphNode->getName(); + std::string ErrorMsg = + std::string( + "slm_init is called more than once from kernel '") + + demangle(KernelName.str()) + "'."; + GraphNode->getContext().emitError(ErrorMsg); + } else { + Callers.insert(GraphNode); + } + } + }, + false, filterInvokeSimdUse); + } + } + } + } +} + bool SYCLLowerESIMDPass::prepareForAlwaysInliner(Module &M) { auto markAlwaysInlined = [](Function &F) -> bool { @@ -1851,7 +1886,7 @@ bool SYCLLowerESIMDPass::prepareForAlwaysInliner(Module &M) { if (FCall && FCall->getCalledFunction() == &F) { Function *GenF = FCall->getFunction(); // The original kernel (UserK) if often automatically separated into - // a spir_func (GenF) that is then called from spir_kernel (GenK). + // a spir_func (GenF) that is then cal led from spir_kernel (GenK). // When that happens, the calls of slm_init() originally placed // in 'UserK' get moved to spir_func 'GenF', which creates wrong IR // because slm_init() must be called only from a kernel. @@ -1908,6 +1943,7 @@ static void fixFunctionReadWriteAttributes(Module &M) { PreservedAnalyses SYCLLowerESIMDPass::run(Module &M, ModuleAnalysisManager &MAM) { + checkSLMInit(M); // AlwaysInlinerPass is required for correctness. bool ForceInline = prepareForAlwaysInliner(M); if (ForceInline) { diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp index 29aa63c231ace..db35f852f0878 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp @@ -364,14 +364,10 @@ class ScopedCallGraph { if (auto *CB = dyn_cast(&I)) { if (isSlmInitCall(dyn_cast(CB))) { auto *CI = dyn_cast(CB); - if (SlmInitCall) { - std::string ErrorMsg = - std::string( - "slm_init is called more than once from function '") + - demangle(F.getName().str()) + "'."; - F.getContext().emitError(ErrorMsg); - } - + esimd::assert_and_diag(!SlmInitCall, + "multiple slm_init calls in function ", + F.getName()); + // TODO: this diagnostics incorrectly fires on functor's // operator() marked as SYCL_ESIMD_KERNEL, because becomes neither // spir_kernel nor SYCL_EXERNAL function in IR. It rather becomes diff --git a/sycl/test/esimd/slm_init_check.cpp b/sycl/test/esimd/slm_init_check.cpp index 43a442b19195d..67ad468550b6b 100644 --- a/sycl/test/esimd/slm_init_check.cpp +++ b/sycl/test/esimd/slm_init_check.cpp @@ -16,7 +16,7 @@ int main() { slm_init(1024); slm_init(1024); }).wait(); - // CHECK: error: slm_init is called more than once from function 'typeinfo name for main::'lambda'(sycl::_V1::nd_item<1>)'. + // CHECK: error: slm_init is called more than once from kernel 'typeinfo name for main::'lambda'(sycl::_V1::nd_item<1>)'. return 0; } diff --git a/sycl/test/esimd/slm_init_noinline_check.cpp b/sycl/test/esimd/slm_init_noinline_check.cpp new file mode 100644 index 0000000000000..0edda013fc1f5 --- /dev/null +++ b/sycl/test/esimd/slm_init_noinline_check.cpp @@ -0,0 +1,31 @@ +// RUN: not %clangxx -fsycl %s 2>&1 | FileCheck %s + +// This test verifies call to slm_init from a function not marked as +// always_inline triggers an error. + +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; + +#ifdef _MSC_VER +#define __SYCL_NOINLINE __declspec(noinline) +#else +#define __SYCL_NOINLINE __attribute__((noinline)) +#endif + +__SYCL_NOINLINE void bar() { slm_init(1024); } +__SYCL_NOINLINE void foo() { + slm_init(1024); + bar(); +} + +int main() { + queue Q; + nd_range<1> NDR{range<1>{2}, range<1>{2}}; + Q.parallel_for(NDR, [=](nd_item<1> NDI) SYCL_ESIMD_KERNEL { foo(); }).wait(); + return 0; +} +// CHECK: error: slm_init is called more than once from kernel 'typeinfo name for main::'lambda'(sycl::_V1::nd_item<1>)'. \ No newline at end of file From be39333eee8378cefcc931ec05a44867c131bb5b Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Sun, 25 Feb 2024 10:45:56 -0800 Subject: [PATCH 05/17] Remove unnecessary changes --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp index db35f852f0878..ffe282adb9b3a 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp @@ -364,10 +364,9 @@ class ScopedCallGraph { if (auto *CB = dyn_cast(&I)) { if (isSlmInitCall(dyn_cast(CB))) { auto *CI = dyn_cast(CB); - esimd::assert_and_diag(!SlmInitCall, - "multiple slm_init calls in function ", + esimd::assert_and_diag(!SlmInitCall, + "multiple slm_init calls in function ", F.getName()); - // TODO: this diagnostics incorrectly fires on functor's // operator() marked as SYCL_ESIMD_KERNEL, because becomes neither // spir_kernel nor SYCL_EXERNAL function in IR. It rather becomes From a0f13d1554f9a9da26d339a2868a3f67120e4679 Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Sun, 25 Feb 2024 10:47:14 -0800 Subject: [PATCH 06/17] Remove unneeded includes --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp index ffe282adb9b3a..323e16faa2171 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDSlmReservation.cpp @@ -64,8 +64,6 @@ #include "llvm/ADT/SmallPtrSet.h" #include "llvm/ADT/SmallVector.h" -#include "llvm/Demangle/Demangle.h" -#include "llvm/Demangle/ItaniumDemangle.h" #include "llvm/IR/Function.h" #include "llvm/IR/Instruction.h" #include "llvm/IR/Module.h" From 5edbbd521b2c7581d72ad3d9b0438538821da323 Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Sun, 25 Feb 2024 10:59:08 -0800 Subject: [PATCH 07/17] Add comments --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 760f2e42cb334..d23489ca2e6b8 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1943,6 +1943,7 @@ static void fixFunctionReadWriteAttributes(Module &M) { PreservedAnalyses SYCLLowerESIMDPass::run(Module &M, ModuleAnalysisManager &MAM) { + // Check validity of slm_init calls. checkSLMInit(M); // AlwaysInlinerPass is required for correctness. bool ForceInline = prepareForAlwaysInliner(M); From 1fa24a17e428d1f1a818c735ea4151f25f4f7144 Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Mon, 26 Feb 2024 22:27:19 -0800 Subject: [PATCH 08/17] Add checks to prevent use of slm_init in functions called using invoke_simd --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 25 +++++++++++++---- sycl/test/esimd/slm_init_invoke_simd.cpp | 34 +++++++++++++++++++++++ 2 files changed, 53 insertions(+), 6 deletions(-) create mode 100644 sycl/test/esimd/slm_init_invoke_simd.cpp diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index d23489ca2e6b8..6a278fcdcf39e 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1754,14 +1754,11 @@ static void checkSLMInit(Module &M) { SmallPtrSet Callers; for (auto &F : M) { if (isSlmInit(F)) { - auto filterInvokeSimdUse = [](const Instruction *, const Function *) { - return false; - }; for (User *U : F.users()) { auto *FCall = dyn_cast(U); if (FCall && FCall->getCalledFunction() == &F) { Function *GenF = FCall->getFunction(); - + SmallPtrSet Visited; sycl::utils::traverseCallgraphUp( GenF, [&](Function *GraphNode) { @@ -1778,7 +1775,21 @@ static void checkSLMInit(Module &M) { } } }, - false, filterInvokeSimdUse); + Visited, false); + bool VisitedKernel = false; + for (const Function *Caller : Visited) { + if (llvm::esimd::isESIMDKernel(*Caller)) { + VisitedKernel = true; + break; + } + } + if (!VisitedKernel) { + F.getContext().emitError( + "slm_init must be called directly from ESIMD kernel."); + } + } else { + F.getContext().emitError( + "slm_init can only be used as a direct call."); } } } @@ -1886,7 +1897,7 @@ bool SYCLLowerESIMDPass::prepareForAlwaysInliner(Module &M) { if (FCall && FCall->getCalledFunction() == &F) { Function *GenF = FCall->getFunction(); // The original kernel (UserK) if often automatically separated into - // a spir_func (GenF) that is then cal led from spir_kernel (GenK). + // a spir_func (GenF) that is then called from spir_kernel (GenK). // When that happens, the calls of slm_init() originally placed // in 'UserK' get moved to spir_func 'GenF', which creates wrong IR // because slm_init() must be called only from a kernel. @@ -1943,8 +1954,10 @@ static void fixFunctionReadWriteAttributes(Module &M) { PreservedAnalyses SYCLLowerESIMDPass::run(Module &M, ModuleAnalysisManager &MAM) { + // Check validity of slm_init calls. checkSLMInit(M); + // AlwaysInlinerPass is required for correctness. bool ForceInline = prepareForAlwaysInliner(M); if (ForceInline) { diff --git a/sycl/test/esimd/slm_init_invoke_simd.cpp b/sycl/test/esimd/slm_init_invoke_simd.cpp new file mode 100644 index 0000000000000..bc57240644002 --- /dev/null +++ b/sycl/test/esimd/slm_init_invoke_simd.cpp @@ -0,0 +1,34 @@ +// This test verifies call to slm_init from a function called through +// invoke_simd triggers an error. + +// RUN: not %clangxx -fsycl -fno-sycl-device-code-split-esimd -Xclang -fsycl-allow-func-ptr %s 2>&1 | FileCheck %s + +#include +#include +#include +#include + +#include +#include +#include + +using namespace sycl::ext::oneapi::experimental; +using namespace sycl; +namespace esimd = sycl::ext::intel::esimd; + +SYCL_EXTERNAL +[[intel::device_indirectly_callable]] void __regcall SIMD_CALLEE_VOID() + SYCL_ESIMD_FUNCTION { + esimd::slm_init<1024>(); +} + +int main() { + queue Q; + nd_range<1> NDR{range<1>{2}, range<1>{2}}; + Q.parallel_for(NDR, [=](nd_item<1> NDI) [[intel::reqd_sub_group_size(16)]] { + sub_group sg = NDI.get_sub_group(); + invoke_simd(sg, SIMD_CALLEE_VOID); + }).wait(); + return 0; +} +// CHECK: slm_init must be called directly from ESIMD kernel. \ No newline at end of file From 4685babd4170ce4f0ba0be7d842155b378a9f50b Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Wed, 28 Feb 2024 12:59:37 -0800 Subject: [PATCH 09/17] Address PR comments --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 67 +++++++++++------------ 1 file changed, 33 insertions(+), 34 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 6a278fcdcf39e..a0cfa2167a161 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1753,44 +1753,43 @@ void lowerGlobalsToVector(Module &M) { static void checkSLMInit(Module &M) { SmallPtrSet Callers; for (auto &F : M) { - if (isSlmInit(F)) { - for (User *U : F.users()) { - auto *FCall = dyn_cast(U); - if (FCall && FCall->getCalledFunction() == &F) { - Function *GenF = FCall->getFunction(); - SmallPtrSet Visited; - sycl::utils::traverseCallgraphUp( - GenF, - [&](Function *GraphNode) { - if (llvm::esimd::isESIMDKernel(*GraphNode)) { - if (Callers.contains(GraphNode)) { - StringRef KernelName = GraphNode->getName(); - std::string ErrorMsg = - std::string( - "slm_init is called more than once from kernel '") + - demangle(KernelName.str()) + "'."; - GraphNode->getContext().emitError(ErrorMsg); - } else { - Callers.insert(GraphNode); - } + if (!isSlmInit(F)) + continue; + for (User *U : F.users()) { + auto *FCall = dyn_cast(U); + if (FCall && FCall->getCalledFunction() == &F) { + Function *GenF = FCall->getFunction(); + SmallPtrSet Visited; + sycl::utils::traverseCallgraphUp( + GenF, + [&](Function *GraphNode) { + if (llvm::esimd::isESIMDKernel(*GraphNode)) { + if (Callers.contains(GraphNode)) { + StringRef KernelName = GraphNode->getName(); + std::string ErrorMsg = + std::string( + "slm_init is called more than once from kernel '") + + demangle(KernelName.str()) + "'."; + GraphNode->getContext().emitError(ErrorMsg); + } else { + Callers.insert(GraphNode); } - }, - Visited, false); - bool VisitedKernel = false; - for (const Function *Caller : Visited) { - if (llvm::esimd::isESIMDKernel(*Caller)) { - VisitedKernel = true; - break; - } - } - if (!VisitedKernel) { - F.getContext().emitError( - "slm_init must be called directly from ESIMD kernel."); + } + }, + Visited, false); + bool VisitedKernel = false; + for (const Function *Caller : Visited) { + if (llvm::esimd::isESIMDKernel(*Caller)) { + VisitedKernel = true; + break; } - } else { + } + if (!VisitedKernel) { F.getContext().emitError( - "slm_init can only be used as a direct call."); + "slm_init must be called directly from ESIMD kernel."); } + } else { + F.getContext().emitError("slm_init can only be used as a direct call."); } } } From 3bbc5b2c099dce7d7831fd504e78db52d3b88340 Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Mon, 4 Mar 2024 22:07:11 -0800 Subject: [PATCH 10/17] Add checks to detect use of local_accessor and slm_init --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 105 ++++++++++++------ .../esimd/slm_init_local_accessor_check.cpp | 24 ++++ .../slm_init_local_accessor_parameter.cpp | 25 +++++ .../slm_init_local_accessor_subscript.cpp | 25 +++++ 4 files changed, 146 insertions(+), 33 deletions(-) create mode 100644 sycl/test/esimd/slm_init_local_accessor_check.cpp create mode 100644 sycl/test/esimd/slm_init_local_accessor_parameter.cpp create mode 100644 sycl/test/esimd/slm_init_local_accessor_subscript.cpp diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index a0cfa2167a161..a3a4e486560fd 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -136,6 +136,8 @@ static constexpr char ESIMD_INTRIN_PREF0[] = "_Z"; static constexpr char ESIMD_INTRIN_PREF1[] = "__esimd_"; static constexpr char ESIMD_INSERTED_VSTORE_FUNC_NAME[] = "_Z14__esimd_vstorev"; static constexpr char SPIRV_INTRIN_PREF[] = "__spirv_BuiltIn"; +static constexpr char SPIRV_LOCAL_ACCESSOR_PREF[] = + "_ZN4sycl3_V114local_accessor"; struct ESIMDIntrinDesc { // Denotes argument translation rule kind. enum GenXArgRuleKind { @@ -1752,46 +1754,83 @@ void lowerGlobalsToVector(Module &M) { static void checkSLMInit(Module &M) { SmallPtrSet Callers; + bool Kernel_Has_slm_init = false; + bool Kernel_Has_local_accessor = false; + for (auto &F : M) { - if (!isSlmInit(F)) - continue; - for (User *U : F.users()) { - auto *FCall = dyn_cast(U); - if (FCall && FCall->getCalledFunction() == &F) { - Function *GenF = FCall->getFunction(); - SmallPtrSet Visited; - sycl::utils::traverseCallgraphUp( - GenF, - [&](Function *GraphNode) { - if (llvm::esimd::isESIMDKernel(*GraphNode)) { - if (Callers.contains(GraphNode)) { - StringRef KernelName = GraphNode->getName(); - std::string ErrorMsg = - std::string( - "slm_init is called more than once from kernel '") + - demangle(KernelName.str()) + "'."; - GraphNode->getContext().emitError(ErrorMsg); - } else { - Callers.insert(GraphNode); - } - } - }, - Visited, false); - bool VisitedKernel = false; - for (const Function *Caller : Visited) { - if (llvm::esimd::isESIMDKernel(*Caller)) { - VisitedKernel = true; - break; + if (!isSlmInit(F)) { + if (Kernel_Has_local_accessor) { + continue; + } + if (F.getName().starts_with(SPIRV_LOCAL_ACCESSOR_PREF)) { + Kernel_Has_local_accessor = true; + continue; + } + unsigned Idx = 0; + for (const Argument &Arg : F.args()) { + if (Arg.getType()->isPointerTy()) { + auto *KernelArgAccPtrs = F.getMetadata("kernel_arg_accessor_ptr"); + + if (KernelArgAccPtrs) { + auto *AccMD = + cast(KernelArgAccPtrs->getOperand(Idx)); + auto AccMDVal = cast(AccMD->getValue())->getValue(); + bool IsAcc = static_cast(AccMDVal.getZExtValue()); + + constexpr unsigned LocalAS{3}; + if (IsAcc && cast(Arg.getType())->getAddressSpace() == + LocalAS) { + Kernel_Has_local_accessor = true; + break; + } } } - if (!VisitedKernel) { + Idx++; + } + } else { + Kernel_Has_slm_init = true; + for (User *U : F.users()) { + auto *FCall = dyn_cast(U); + if (FCall && FCall->getCalledFunction() == &F) { + Function *GenF = FCall->getFunction(); + SmallPtrSet Visited; + sycl::utils::traverseCallgraphUp( + GenF, + [&](Function *GraphNode) { + if (llvm::esimd::isESIMDKernel(*GraphNode)) { + if (Callers.contains(GraphNode)) { + StringRef KernelName = GraphNode->getName(); + std::string ErrorMsg = + std::string("slm_init is called more than once " + "from kernel '") + + demangle(KernelName.str()) + "'."; + GraphNode->getContext().emitError(ErrorMsg); + } else { + Callers.insert(GraphNode); + } + } + }, + Visited, false); + bool VisitedKernel = false; + for (const Function *Caller : Visited) { + if (llvm::esimd::isESIMDKernel(*Caller)) { + VisitedKernel = true; + break; + } + } + if (!VisitedKernel) { + F.getContext().emitError( + "slm_init must be called directly from ESIMD kernel."); + } + } else { F.getContext().emitError( - "slm_init must be called directly from ESIMD kernel."); + "slm_init can only be used as a direct call."); } - } else { - F.getContext().emitError("slm_init can only be used as a direct call."); } } + if (Kernel_Has_slm_init && Kernel_Has_local_accessor) { + F.getContext().emitError("slm_init can not be used with local_accessor."); + } } } diff --git a/sycl/test/esimd/slm_init_local_accessor_check.cpp b/sycl/test/esimd/slm_init_local_accessor_check.cpp new file mode 100644 index 0000000000000..2f55002666253 --- /dev/null +++ b/sycl/test/esimd/slm_init_local_accessor_check.cpp @@ -0,0 +1,24 @@ +// RUN: not %clangxx -O0 -fsycl %s 2>&1 | FileCheck %s + +// This test verifies usage of slm_init and local_accessor triggers an error. + +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; + +int main() { + queue Q; + nd_range<1> NDR{range<1>{2}, range<1>{2}}; + Q.submit([&](handler &CGH) { + CGH.parallel_for(NDR, [=](nd_item<1> NDI) SYCL_ESIMD_KERNEL { + auto InAcc = local_accessor(); + slm_init(1024); + }); + }).wait(); + // CHECK: error: slm_init can not be used with local_accessor. + + return 0; +} diff --git a/sycl/test/esimd/slm_init_local_accessor_parameter.cpp b/sycl/test/esimd/slm_init_local_accessor_parameter.cpp new file mode 100644 index 0000000000000..7da7e6fb7936d --- /dev/null +++ b/sycl/test/esimd/slm_init_local_accessor_parameter.cpp @@ -0,0 +1,25 @@ +// RUN: not %clangxx -fsycl %s 2>&1 | FileCheck %s + +// This test verifies usage of slm_init and local_accessor triggers an error. + +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; + +int main() { + queue Q; + nd_range<1> NDR{range<1>{2}, range<1>{2}}; + Q.submit([&](handler &CGH) { + auto InAcc = local_accessor(5, CGH); + CGH.parallel_for(NDR, [=](nd_item<1> NDI) SYCL_ESIMD_KERNEL { + slm_init(1024); + scalar_load(InAcc, 0); + }); + }).wait(); + // CHECK: error: slm_init can not be used with local_accessor. + + return 0; +} diff --git a/sycl/test/esimd/slm_init_local_accessor_subscript.cpp b/sycl/test/esimd/slm_init_local_accessor_subscript.cpp new file mode 100644 index 0000000000000..e99d72ce6a9af --- /dev/null +++ b/sycl/test/esimd/slm_init_local_accessor_subscript.cpp @@ -0,0 +1,25 @@ +// RUN: not %clangxx -fsycl %s 2>&1 | FileCheck %s + +// This test verifies usage of slm_init and local_accessor triggers an error. + +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; + +int main() { + queue Q; + nd_range<1> NDR{range<1>{2}, range<1>{2}}; + Q.submit([&](handler &CGH) { + auto InAcc = local_accessor(5, CGH); + CGH.parallel_for(NDR, [=](nd_item<1> NDI) SYCL_ESIMD_KERNEL { + slm_init(1024); + InAcc[0] = 5; + }); + }).wait(); + // CHECK: error: slm_init can not be used with local_accessor. + + return 0; +} From 836c1ed84dbda2f3230442bbb0e4af33833dbf1c Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Tue, 5 Mar 2024 11:23:24 -0800 Subject: [PATCH 11/17] Fix test failure --- sycl/test-e2e/ESIMD/unified_memory_api/Inputs/block_store.hpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/block_store.hpp b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/block_store.hpp index ebd2b6c957842..607840cdb7db0 100644 --- a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/block_store.hpp +++ b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/block_store.hpp @@ -344,9 +344,6 @@ bool testLocalAccSLM(queue Q, uint32_t Groups, auto OutPtr = Out.data(); CGH.parallel_for(Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL { - constexpr uint32_t SLMSize = (GroupSize * N) * sizeof(T); - slm_init(); - uint16_t GlobalID = ndi.get_global_id(0); uint16_t LocalID = ndi.get_local_id(0); uint32_t LocalElemOffset = LocalID * N * sizeof(T); From 8853912a2378171e81ada81e83060604cd4cd487 Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Wed, 6 Mar 2024 15:12:01 -0800 Subject: [PATCH 12/17] Address PR comments --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 20 ++++++++++--------- .../esimd/slm_init_local_accessor_check.cpp | 2 +- .../slm_init_local_accessor_parameter.cpp | 2 +- .../slm_init_local_accessor_subscript.cpp | 2 +- 4 files changed, 14 insertions(+), 12 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index a3a4e486560fd..e7eb591846ba9 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1754,16 +1754,17 @@ void lowerGlobalsToVector(Module &M) { static void checkSLMInit(Module &M) { SmallPtrSet Callers; - bool Kernel_Has_slm_init = false; - bool Kernel_Has_local_accessor = false; + bool KernelHasSLMInit = false; + bool KernelHasLocalAccessor = false; for (auto &F : M) { + F.dump(); if (!isSlmInit(F)) { - if (Kernel_Has_local_accessor) { + if (KernelHasLocalAccessor) continue; - } + if (F.getName().starts_with(SPIRV_LOCAL_ACCESSOR_PREF)) { - Kernel_Has_local_accessor = true; + KernelHasLocalAccessor = true; continue; } unsigned Idx = 0; @@ -1780,7 +1781,7 @@ static void checkSLMInit(Module &M) { constexpr unsigned LocalAS{3}; if (IsAcc && cast(Arg.getType())->getAddressSpace() == LocalAS) { - Kernel_Has_local_accessor = true; + KernelHasLocalAccessor = true; break; } } @@ -1788,7 +1789,7 @@ static void checkSLMInit(Module &M) { Idx++; } } else { - Kernel_Has_slm_init = true; + KernelHasSLMInit = true; for (User *U : F.users()) { auto *FCall = dyn_cast(U); if (FCall && FCall->getCalledFunction() == &F) { @@ -1828,8 +1829,9 @@ static void checkSLMInit(Module &M) { } } } - if (Kernel_Has_slm_init && Kernel_Has_local_accessor) { - F.getContext().emitError("slm_init can not be used with local_accessor."); + if (KernelHasSLMInit && KernelHasLocalAccessor) { + F.getContext().emitError( + "slm_init can not be used with local accessors."); } } } diff --git a/sycl/test/esimd/slm_init_local_accessor_check.cpp b/sycl/test/esimd/slm_init_local_accessor_check.cpp index 2f55002666253..b1fd27049d003 100644 --- a/sycl/test/esimd/slm_init_local_accessor_check.cpp +++ b/sycl/test/esimd/slm_init_local_accessor_check.cpp @@ -18,7 +18,7 @@ int main() { slm_init(1024); }); }).wait(); - // CHECK: error: slm_init can not be used with local_accessor. + // CHECK: error: slm_init can not be used with local accessors. return 0; } diff --git a/sycl/test/esimd/slm_init_local_accessor_parameter.cpp b/sycl/test/esimd/slm_init_local_accessor_parameter.cpp index 7da7e6fb7936d..c32c7fa880bd9 100644 --- a/sycl/test/esimd/slm_init_local_accessor_parameter.cpp +++ b/sycl/test/esimd/slm_init_local_accessor_parameter.cpp @@ -19,7 +19,7 @@ int main() { scalar_load(InAcc, 0); }); }).wait(); - // CHECK: error: slm_init can not be used with local_accessor. + // CHECK: error: slm_init can not be used with local accessors. return 0; } diff --git a/sycl/test/esimd/slm_init_local_accessor_subscript.cpp b/sycl/test/esimd/slm_init_local_accessor_subscript.cpp index e99d72ce6a9af..53a2c10f25b6e 100644 --- a/sycl/test/esimd/slm_init_local_accessor_subscript.cpp +++ b/sycl/test/esimd/slm_init_local_accessor_subscript.cpp @@ -19,7 +19,7 @@ int main() { InAcc[0] = 5; }); }).wait(); - // CHECK: error: slm_init can not be used with local_accessor. + // CHECK: error: slm_init can not be used with local accessors. return 0; } From 955f44af33b5f9de60fed68a8e45ffac6daf693f Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Wed, 6 Mar 2024 15:14:55 -0800 Subject: [PATCH 13/17] Remove debug dump --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index e7eb591846ba9..0a14285a228be 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1758,7 +1758,6 @@ static void checkSLMInit(Module &M) { bool KernelHasLocalAccessor = false; for (auto &F : M) { - F.dump(); if (!isSlmInit(F)) { if (KernelHasLocalAccessor) continue; From 909bf52ec71b98daa0f0159ea8494ae32f198a1f Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" Date: Thu, 7 Mar 2024 16:40:59 -0800 Subject: [PATCH 14/17] Address PR comments, Improve handling of multiple kernels --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 72 ++++++++++++--------- sycl/test/esimd/slm_init_local_accessor.cpp | 29 +++++++++ sycl/test/esimd/slm_init_noinline_check.cpp | 4 +- 3 files changed, 71 insertions(+), 34 deletions(-) create mode 100644 sycl/test/esimd/slm_init_local_accessor.cpp diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 0a14285a228be..44a00c0dee3a5 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1753,42 +1753,49 @@ void lowerGlobalsToVector(Module &M) { } // namespace static void checkSLMInit(Module &M) { - SmallPtrSet Callers; - bool KernelHasSLMInit = false; - bool KernelHasLocalAccessor = false; + SmallPtrSet SLMInitKernels; + SmallPtrSet LocalAccessorKernels; for (auto &F : M) { if (!isSlmInit(F)) { - if (KernelHasLocalAccessor) - continue; - + bool LocalAccessorUsed = false; if (F.getName().starts_with(SPIRV_LOCAL_ACCESSOR_PREF)) { - KernelHasLocalAccessor = true; - continue; - } - unsigned Idx = 0; - for (const Argument &Arg : F.args()) { - if (Arg.getType()->isPointerTy()) { - auto *KernelArgAccPtrs = F.getMetadata("kernel_arg_accessor_ptr"); - - if (KernelArgAccPtrs) { - auto *AccMD = - cast(KernelArgAccPtrs->getOperand(Idx)); - auto AccMDVal = cast(AccMD->getValue())->getValue(); - bool IsAcc = static_cast(AccMDVal.getZExtValue()); - - constexpr unsigned LocalAS{3}; - if (IsAcc && cast(Arg.getType())->getAddressSpace() == - LocalAS) { - KernelHasLocalAccessor = true; - break; + LocalAccessorUsed = true; + } else { + unsigned Idx = 0; + for (const Argument &Arg : F.args()) { + if (Arg.getType()->isPointerTy()) { + auto *KernelArgAccPtrs = F.getMetadata("kernel_arg_accessor_ptr"); + + if (KernelArgAccPtrs) { + auto *AccMD = + cast(KernelArgAccPtrs->getOperand(Idx)); + auto AccMDVal = cast(AccMD->getValue())->getValue(); + bool IsAcc = static_cast(AccMDVal.getZExtValue()); + + constexpr unsigned LocalAS{3}; + if (IsAcc && + cast(Arg.getType())->getAddressSpace() == + LocalAS) { + LocalAccessorUsed = true; + break; + } } } + Idx++; } - Idx++; + } + if (LocalAccessorUsed) { + sycl::utils::traverseCallgraphUp( + &F, + [&](Function *GraphNode) { + if (llvm::esimd::isESIMDKernel(*GraphNode)) { + LocalAccessorKernels.insert(GraphNode); + } + }, + false); } } else { - KernelHasSLMInit = true; for (User *U : F.users()) { auto *FCall = dyn_cast(U); if (FCall && FCall->getCalledFunction() == &F) { @@ -1798,7 +1805,7 @@ static void checkSLMInit(Module &M) { GenF, [&](Function *GraphNode) { if (llvm::esimd::isESIMDKernel(*GraphNode)) { - if (Callers.contains(GraphNode)) { + if (SLMInitKernels.contains(GraphNode)) { StringRef KernelName = GraphNode->getName(); std::string ErrorMsg = std::string("slm_init is called more than once " @@ -1806,7 +1813,7 @@ static void checkSLMInit(Module &M) { demangle(KernelName.str()) + "'."; GraphNode->getContext().emitError(ErrorMsg); } else { - Callers.insert(GraphNode); + SLMInitKernels.insert(GraphNode); } } }, @@ -1828,9 +1835,10 @@ static void checkSLMInit(Module &M) { } } } - if (KernelHasSLMInit && KernelHasLocalAccessor) { - F.getContext().emitError( - "slm_init can not be used with local accessors."); + for (const Function *Kernel : LocalAccessorKernels) { + if (SLMInitKernels.contains(Kernel)) + F.getContext().emitError( + "slm_init can not be used with local accessors."); } } } diff --git a/sycl/test/esimd/slm_init_local_accessor.cpp b/sycl/test/esimd/slm_init_local_accessor.cpp new file mode 100644 index 0000000000000..3ca7408011781 --- /dev/null +++ b/sycl/test/esimd/slm_init_local_accessor.cpp @@ -0,0 +1,29 @@ +// RUN: %clangxx -fsycl %s + +// This test verifies usage of slm_init and local_accessor in different kernels +// passes. + +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; + +int main() { + queue Q; + nd_range<1> NDR{range<1>{2}, range<1>{2}}; + Q.submit([&](handler &CGH) { + auto InAcc = local_accessor(5, CGH); + CGH.parallel_for(NDR, [=](nd_item<1> NDI) SYCL_ESIMD_KERNEL { + scalar_load(InAcc, 0); + }); + }).wait(); + + Q.submit([&](handler &CGH) { + CGH.parallel_for(NDR, [=](nd_item<1> NDI) + SYCL_ESIMD_KERNEL { slm_init(1024); }); + }).wait(); + + return 0; +} diff --git a/sycl/test/esimd/slm_init_noinline_check.cpp b/sycl/test/esimd/slm_init_noinline_check.cpp index 0edda013fc1f5..117090a424524 100644 --- a/sycl/test/esimd/slm_init_noinline_check.cpp +++ b/sycl/test/esimd/slm_init_noinline_check.cpp @@ -1,7 +1,7 @@ // RUN: not %clangxx -fsycl %s 2>&1 | FileCheck %s -// This test verifies call to slm_init from a function not marked as -// always_inline triggers an error. +// This test verifies call to slm_init from a function marked as +// noinline triggers an error. #include #include From a909ff27f5d67fe31a1dc8694fe56bcb941296ca Mon Sep 17 00:00:00 2001 From: gregory Date: Mon, 18 Mar 2024 12:30:33 -0700 Subject: [PATCH 15/17] Remove incorrect test --- .../esimd/slm_init_local_accessor_check.cpp | 24 ------------------- 1 file changed, 24 deletions(-) delete mode 100644 sycl/test/esimd/slm_init_local_accessor_check.cpp diff --git a/sycl/test/esimd/slm_init_local_accessor_check.cpp b/sycl/test/esimd/slm_init_local_accessor_check.cpp deleted file mode 100644 index b1fd27049d003..0000000000000 --- a/sycl/test/esimd/slm_init_local_accessor_check.cpp +++ /dev/null @@ -1,24 +0,0 @@ -// RUN: not %clangxx -O0 -fsycl %s 2>&1 | FileCheck %s - -// This test verifies usage of slm_init and local_accessor triggers an error. - -#include -#include -#include - -using namespace sycl; -using namespace sycl::ext::intel::esimd; - -int main() { - queue Q; - nd_range<1> NDR{range<1>{2}, range<1>{2}}; - Q.submit([&](handler &CGH) { - CGH.parallel_for(NDR, [=](nd_item<1> NDI) SYCL_ESIMD_KERNEL { - auto InAcc = local_accessor(); - slm_init(1024); - }); - }).wait(); - // CHECK: error: slm_init can not be used with local accessors. - - return 0; -} From 36e8273e78a962ffcd7fc03b0253b438e686fdac Mon Sep 17 00:00:00 2001 From: gregory Date: Mon, 18 Mar 2024 17:18:11 -0700 Subject: [PATCH 16/17] Simplify the vlidation logic --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 52 +++++++++-------------- 1 file changed, 19 insertions(+), 33 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 44a00c0dee3a5..8fbfd17373793 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1758,42 +1758,28 @@ static void checkSLMInit(Module &M) { for (auto &F : M) { if (!isSlmInit(F)) { - bool LocalAccessorUsed = false; - if (F.getName().starts_with(SPIRV_LOCAL_ACCESSOR_PREF)) { - LocalAccessorUsed = true; - } else { - unsigned Idx = 0; - for (const Argument &Arg : F.args()) { - if (Arg.getType()->isPointerTy()) { - auto *KernelArgAccPtrs = F.getMetadata("kernel_arg_accessor_ptr"); - - if (KernelArgAccPtrs) { - auto *AccMD = - cast(KernelArgAccPtrs->getOperand(Idx)); - auto AccMDVal = cast(AccMD->getValue())->getValue(); - bool IsAcc = static_cast(AccMDVal.getZExtValue()); - - constexpr unsigned LocalAS{3}; - if (IsAcc && - cast(Arg.getType())->getAddressSpace() == - LocalAS) { - LocalAccessorUsed = true; - break; - } + if (!llvm::esimd::isESIMDKernel(F)) + continue; + unsigned Idx = 0; + for (const Argument &Arg : F.args()) { + if (Arg.getType()->isPointerTy()) { + auto *KernelArgAccPtrs = F.getMetadata("kernel_arg_accessor_ptr"); + + if (KernelArgAccPtrs) { + auto *AccMD = + cast(KernelArgAccPtrs->getOperand(Idx)); + auto AccMDVal = cast(AccMD->getValue())->getValue(); + bool IsAcc = static_cast(AccMDVal.getZExtValue()); + + constexpr unsigned LocalAS{3}; + if (IsAcc && cast(Arg.getType())->getAddressSpace() == + LocalAS) { + LocalAccessorKernels.insert(&F); + break; } } - Idx++; } - } - if (LocalAccessorUsed) { - sycl::utils::traverseCallgraphUp( - &F, - [&](Function *GraphNode) { - if (llvm::esimd::isESIMDKernel(*GraphNode)) { - LocalAccessorKernels.insert(GraphNode); - } - }, - false); + Idx++; } } else { for (User *U : F.users()) { From 795a0857a169db06b29eb060b5160567d416f477 Mon Sep 17 00:00:00 2001 From: fineg74 <61437305+fineg74@users.noreply.github.com> Date: Mon, 18 Mar 2024 17:49:53 -0700 Subject: [PATCH 17/17] Update llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp Co-authored-by: Vyacheslav Klochkov --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 8fbfd17373793..430d571c4abd1 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -136,8 +136,6 @@ static constexpr char ESIMD_INTRIN_PREF0[] = "_Z"; static constexpr char ESIMD_INTRIN_PREF1[] = "__esimd_"; static constexpr char ESIMD_INSERTED_VSTORE_FUNC_NAME[] = "_Z14__esimd_vstorev"; static constexpr char SPIRV_INTRIN_PREF[] = "__spirv_BuiltIn"; -static constexpr char SPIRV_LOCAL_ACCESSOR_PREF[] = - "_ZN4sycl3_V114local_accessor"; struct ESIMDIntrinDesc { // Denotes argument translation rule kind. enum GenXArgRuleKind {