[SYCL][ESIMD] Report an error when slm_init is called more than once in the kernel

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<CallBase>(&I)) {
  if (isSlmInitCall(dyn_cast<CallInst>(CB))) {
  auto *CI = dyn_cast<CallInst>(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

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 <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.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;
+}