Add checks to prevent use of slm_init in functions called using invoke_simd

Author: fineg74

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

@@ -1754,14 +1754,11 @@ static void checkSLMInit(Module &M) {
   SmallPtrSet<const Function *, 8u> 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<CallInst>(U);
         if (FCall && FCall->getCalledFunction() == &F) {
           Function *GenF = FCall->getFunction();
-
+          SmallPtrSet<Function *, 32> 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<N>() 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) {

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 <sycl/ext/intel/esimd.hpp>
+#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>
+#include <sycl/ext/oneapi/experimental/uniform.hpp>
+#include <sycl/sycl.hpp>
+
+#include <functional>
+#include <iostream>
+#include <type_traits>
+
+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.