From 3f43d4702543d974e80db704aa8c85951dd31e97 Mon Sep 17 00:00:00 2001 From: fineg74 <61437305+fineg74@users.noreply.github.com> Date: Tue, 19 Mar 2024 09:33:22 -0700 Subject: [PATCH] [SYCL][ESIMD] Report an error when slm_init is called more than once in the kernel (#12804) The patch reports an error if: - sim_init() is used together with local_accessor - slm_init() is called not from ESIMD kernel - slm_init() is called more than once in ESIMD kernel Co-authored-by: Vyacheslav Klochkov --- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 81 +++++++++++++++++++ .../unified_memory_api/Inputs/block_store.hpp | 3 - sycl/test/esimd/slm_init_check.cpp | 22 +++++ sycl/test/esimd/slm_init_invoke_simd.cpp | 34 ++++++++ sycl/test/esimd/slm_init_local_accessor.cpp | 29 +++++++ .../slm_init_local_accessor_parameter.cpp | 25 ++++++ .../slm_init_local_accessor_subscript.cpp | 25 ++++++ sycl/test/esimd/slm_init_noinline_check.cpp | 31 +++++++ 8 files changed, 247 insertions(+), 3 deletions(-) create mode 100644 sycl/test/esimd/slm_init_check.cpp create mode 100644 sycl/test/esimd/slm_init_invoke_simd.cpp create mode 100644 sycl/test/esimd/slm_init_local_accessor.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 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 27575bb643a25..0f6bfc932c1e5 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1745,6 +1745,83 @@ void lowerGlobalsToVector(Module &M) { } // namespace +static void checkSLMInit(Module &M) { + SmallPtrSet SLMInitKernels; + SmallPtrSet LocalAccessorKernels; + + for (auto &F : M) { + if (!isSlmInit(F)) { + 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++; + } + } else { + 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 (SLMInitKernels.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 { + SLMInitKernels.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 can only be used as a direct call."); + } + } + } + for (const Function *Kernel : LocalAccessorKernels) { + if (SLMInitKernels.contains(Kernel)) + F.getContext().emitError( + "slm_init can not be used with local accessors."); + } + } +} + bool SYCLLowerESIMDPass::prepareForAlwaysInliner(Module &M) { auto markAlwaysInlined = [](Function &F) -> bool { @@ -1912,6 +1989,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-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); diff --git a/sycl/test/esimd/slm_init_check.cpp b/sycl/test/esimd/slm_init_check.cpp new file mode 100644 index 0000000000000..67ad468550b6b --- /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 kernel 'typeinfo name for main::'lambda'(sycl::_V1::nd_item<1>)'. + + return 0; +} 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 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_local_accessor_parameter.cpp b/sycl/test/esimd/slm_init_local_accessor_parameter.cpp new file mode 100644 index 0000000000000..c32c7fa880bd9 --- /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 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 new file mode 100644 index 0000000000000..53a2c10f25b6e --- /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 accessors. + + 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..117090a424524 --- /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 marked as +// noinline 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