[SYCL][ESIMD] Add ESIMD-specific IR verification pass (#4965)

Signed-off-by: Sergey Dmitriev <[email protected]>

clang/lib/CodeGen/BackendUtil.cpp

@@ -42,6 +42,7 @@
 #include "llvm/Passes/PassBuilder.h"
 #include "llvm/Passes/PassPlugin.h"
 #include "llvm/Passes/StandardInstrumentations.h"
+#include "llvm/SYCLLowerIR/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/Support/BuryPointer.h"
 #include "llvm/Support/CommandLine.h"
@@ -849,6 +850,9 @@ void EmitAssemblyHelper::CreatePasses(legacy::PassManager &MPM,
     FPM.add(createVerifierPass());
 
   // Set up the per-module pass manager.
+  if (LangOpts.SYCLIsDevice)
+    MPM.add(createESIMDVerifierPass());
+
   if (!CodeGenOpts.RewriteMapFiles.empty())
     addSymbolRewriterPass(CodeGenOpts, &MPM);
 

llvm/include/llvm/InitializePasses.h

@@ -438,6 +438,7 @@ void initializeSYCLLowerESIMDLegacyPassPass(PassRegistry &);
 void initializeSPIRITTAnnotationsLegacyPassPass(PassRegistry &);
 void initializeESIMDLowerLoadStorePass(PassRegistry &);
 void initializeESIMDLowerVecArgLegacyPassPass(PassRegistry &);
+void initializeESIMDVerifierPass(PassRegistry &);
 void initializeSYCLLowerWGLocalMemoryLegacyPass(PassRegistry &);
 void initializeTailCallElimPass(PassRegistry&);
 void initializeTailDuplicatePass(PassRegistry&);

llvm/include/llvm/LinkAllPasses.h

@@ -37,6 +37,7 @@
 #include "llvm/CodeGen/Passes.h"
 #include "llvm/IR/Function.h"
 #include "llvm/IR/IRPrintingPasses.h"
+#include "llvm/SYCLLowerIR/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/SYCLLowerIR/LowerWGScope.h"
@@ -208,6 +209,7 @@ namespace {
       (void)llvm::createSYCLLowerESIMDPass();
       (void)llvm::createESIMDLowerLoadStorePass();
       (void)llvm::createESIMDLowerVecArgPass();
+      (void)llvm::createESIMDVerifierPass();
       (void)llvm::createSPIRITTAnnotationsLegacyPass();
       (void)llvm::createSYCLLowerWGLocalMemoryLegacyPass();
       std::string buf;

llvm/include/llvm/SYCLLowerIR/ESIMDVerifier.h

@@ -0,0 +1,30 @@
+//===--------- ESIMDVerifier.h - ESIMD-specific IR verification -----------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// ESIMD verification pass.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_SYCLLOWERIR_ESIMDVERIFIER_H
+#define LLVM_SYCLLOWERIR_ESIMDVERIFIER_H
+
+#include "llvm/IR/PassManager.h"
+
+namespace llvm {
+
+struct ESIMDVerifierPass : public PassInfoMixin<ESIMDVerifierPass> {
+  ESIMDVerifierPass() {}
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+  static bool isRequired() { return true; }
+};
+
+ModulePass *createESIMDVerifierPass();
+
+} // namespace llvm
+
+#endif // LLVM_SYCLLOWERIR_ESIMDVERIFIER_H

llvm/lib/Passes/PassBuilder.cpp

@@ -75,6 +75,7 @@
 #include "llvm/IR/PrintPasses.h"
 #include "llvm/IR/SafepointIRVerifier.h"
 #include "llvm/IR/Verifier.h"
+#include "llvm/SYCLLowerIR/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerWGScope.h"
 #include "llvm/Support/CommandLine.h"

llvm/lib/Passes/PassRegistry.def

@@ -117,6 +117,7 @@ MODULE_PASS("poison-checking", PoisonCheckingPass())
 MODULE_PASS("pseudo-probe-update", PseudoProbeUpdatePass())
 MODULE_PASS("LowerESIMD", SYCLLowerESIMDPass())
 MODULE_PASS("ESIMDLowerVecArg", ESIMDLowerVecArgPass())
+MODULE_PASS("esimd-verifier", ESIMDVerifierPass())
 MODULE_PASS("SPIRITTAnnotations", SPIRITTAnnotationsPass())
 MODULE_PASS("deadargelim-sycl", DeadArgumentEliminationSYCLPass())
 #undef MODULE_PASS

llvm/lib/SYCLLowerIR/CMakeLists.txt

@@ -53,6 +53,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
   LowerESIMDVLoadVStore.cpp
   LowerESIMDVecArg.cpp
   LowerWGLocalMemory.cpp
+  ESIMDVerifier.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLLowerIR

llvm/lib/SYCLLowerIR/ESIMDVerifier.cpp

@@ -0,0 +1,122 @@
+//===---------- ESIMDVerifier.cpp - ESIMD-specific IR verification --------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements ESIMD specific IR verification pass. So far it only
+// detects invalid API calls in ESIMD context.
+//
+//===----------------------------------------------------------------------===//
+
+#include "llvm/SYCLLowerIR/ESIMDVerifier.h"
+#include "llvm/Demangle/Demangle.h"
+#include "llvm/IR/InstIterator.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/Module.h"
+#include "llvm/InitializePasses.h"
+#include "llvm/Pass.h"
+#include "llvm/Support/Debug.h"
+#include "llvm/Support/Regex.h"
+
+using namespace llvm;
+
+#define DEBUG_TYPE "esimd-verifier"
+
+// A list of unsupported functions in ESIMD context.
+static const char *IllegalFunctions[] = {
+    "^cl::sycl::multi_ptr<.+> cl::sycl::accessor<.+>::get_pointer<.+>\\(\\) "
+    "const"};
+
+namespace {
+
+class ESIMDVerifierImpl {
+  const Module &M;
+
+public:
+  ESIMDVerifierImpl(const Module &M) : M(M) {}
+
+  void verify() {
+    SmallPtrSet<const Function *, 8u> Visited;
+    SmallVector<const Function *, 8u> Worklist;
+
+    auto Add2Worklist = [&Worklist, &Visited](const Function *F) {
+      if (Visited.insert(F).second)
+        Worklist.push_back(F);
+    };
+
+    // Start with adding all ESIMD functions to the work list.
+    for (const Function &F : M)
+      if (F.hasMetadata("sycl_explicit_simd"))
+        Add2Worklist(&F);
+
+    // Then check ESIMD functions and all functions called from ESIMD context
+    // for invalid calls.
+    while (!Worklist.empty()) {
+      const Function *F = Worklist.pop_back_val();
+      for (const Instruction &I : instructions(F)) {
+        if (auto *CB = dyn_cast<CallBase>(&I)) {
+          Function *Callee = CB->getCalledFunction();
+          if (!Callee)
+            continue;
+
+          // Demangle called function name and check if it matches any illegal
+          // function name. Report an error if there is a match.
+          std::string DemangledName = demangle(Callee->getName().str());
+          for (const char *Name : IllegalFunctions) {
+            Regex NameRE(Name);
+            assert(NameRE.isValid() && "invalid function name regex");
+            if (NameRE.match(DemangledName)) {
+              std::string ErrorMsg = std::string("function '") + DemangledName +
+                                     "' is not supported in ESIMD context";
+              F->getContext().emitError(&I, ErrorMsg);
+            }
+          }
+
+          // Add callee to the list to be analyzed if it is not a declaration.
+          if (!Callee->isDeclaration())
+            Add2Worklist(Callee);
+        }
+      }
+    }
+  }
+};
+
+} // end anonymous namespace
+
+PreservedAnalyses ESIMDVerifierPass::run(Module &M, ModuleAnalysisManager &AM) {
+  ESIMDVerifierImpl(M).verify();
+  return PreservedAnalyses::all();
+}
+
+namespace {
+
+struct ESIMDVerifier : public ModulePass {
+  static char ID;
+
+  ESIMDVerifier() : ModulePass(ID) {
+    initializeESIMDVerifierPass(*PassRegistry::getPassRegistry());
+  }
+
+  void getAnalysisUsage(AnalysisUsage &AU) const override {
+    AU.setPreservesAll();
+  }
+
+  bool runOnModule(Module &M) override {
+    ESIMDVerifierImpl(M).verify();
+    return false;
+  }
+};
+
+} // end anonymous namespace
+
+char ESIMDVerifier::ID = 0;
+
+INITIALIZE_PASS_BEGIN(ESIMDVerifier, DEBUG_TYPE, "ESIMD-specific IR verifier",
+                      false, false)
+INITIALIZE_PASS_END(ESIMDVerifier, DEBUG_TYPE, "ESIMD-specific IR verifier",
+                    false, false)
+
+ModulePass *llvm::createESIMDVerifierPass() { return new ESIMDVerifier(); }

llvm/tools/opt/opt.cpp

@@ -579,6 +579,7 @@ int main(int argc, char **argv) {
   initializeSPIRITTAnnotationsLegacyPassPass(Registry);
   initializeESIMDLowerLoadStorePass(Registry);
   initializeESIMDLowerVecArgLegacyPassPass(Registry);
+  initializeESIMDVerifierPass(Registry);
   initializeSYCLLowerWGLocalMemoryLegacyPass(Registry);
 
 #ifdef BUILD_EXAMPLES

sycl/test/esimd/esimd_verify.cpp

@@ -0,0 +1,14 @@
+// RUN: not %clangxx -fsycl -fsycl-device-only -S %s -o %t 2>&1 | FileCheck %s
+
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+using namespace cl::sycl;
+using namespace sycl::ext::intel::experimental::esimd;
+
+// CHECK: error: function 'cl::sycl::multi_ptr<{{.+}}> cl::sycl::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
+
+SYCL_EXTERNAL auto
+test(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
+    SYCL_ESIMD_FUNCTION {
+  return acc.get_pointer();
+}