Add checks to detect use of local_accessor and slm_init

intel · Mar 5, 2024 · 3bbc5b2 · 3bbc5b2
1 parent 4685bab
commit 3bbc5b2
Show file tree

Hide file tree

Showing 4 changed files with 146 additions and 33 deletions.
diff --git 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<const Function *, 8u> 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<CallInst>(U);
-      if (FCall && FCall->getCalledFunction() == &F) {
-        Function *GenF = FCall->getFunction();
-        SmallPtrSet<Function *, 32> 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<ConstantAsMetadata>(KernelArgAccPtrs->getOperand(Idx));
+            auto AccMDVal = cast<ConstantInt>(AccMD->getValue())->getValue();
+            bool IsAcc = static_cast<unsigned>(AccMDVal.getZExtValue());
+
+            constexpr unsigned LocalAS{3};
+            if (IsAcc && cast<PointerType>(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<CallInst>(U);
+        if (FCall && FCall->getCalledFunction() == &F) {
+          Function *GenF = FCall->getFunction();
+          SmallPtrSet<Function *, 32> 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
@@ -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 <iostream>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+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<int, 1>();
+       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
@@ -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 <iostream>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+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<int, 1>(5, CGH);
+     CGH.parallel_for(NDR, [=](nd_item<1> NDI) SYCL_ESIMD_KERNEL {
+       slm_init(1024);
+       scalar_load<int>(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
@@ -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 <iostream>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+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<int, 1>(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;
+}