[DeviceSanitizer] Move sanitizer predicates to new SanitizerUtils.h

clang/lib/Driver/OffloadBundler.cpp

@@ -23,17 +23,18 @@
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/StringMap.h"
 #include "llvm/ADT/StringRef.h"
+#include "llvm/BinaryFormat/Magic.h"
 #include "llvm/Bitcode/BitcodeWriter.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/LLVMContext.h"
 #include "llvm/IRReader/IRReader.h"
-#include "llvm/BinaryFormat/Magic.h"
 #include "llvm/Object/Archive.h"
 #include "llvm/Object/ArchiveWriter.h"
 #include "llvm/Object/Binary.h"
 #include "llvm/Object/Error.h"
 #include "llvm/Object/ObjectFile.h"
 #include "llvm/Object/OffloadBundle.h"
+#include "llvm/SYCLLowerIR/SanitizerUtils.h"
 #include "llvm/Support/Casting.h"
 #include "llvm/Support/Compiler.h"
 #include "llvm/Support/Compression.h"
@@ -49,12 +50,12 @@
 #include "llvm/Support/Path.h"
 #include "llvm/Support/Program.h"
 #include "llvm/Support/Signals.h"
+#include "llvm/Support/SourceMgr.h"
 #include "llvm/Support/StringSaver.h"
 #include "llvm/Support/TargetSelect.h"
 #include "llvm/Support/Timer.h"
 #include "llvm/Support/WithColor.h"
 #include "llvm/Support/raw_ostream.h"
-#include "llvm/Support/SourceMgr.h"
 #include "llvm/TargetParser/Host.h"
 #include "llvm/TargetParser/Triple.h"
 #include <algorithm>
@@ -756,8 +757,9 @@ class ObjectFileHandler final : public FileHandler {
              Name == "__AsanDeviceGlobalMetadata" ||
              Name == "__MsanDeviceGlobalMetadata" ||
              Name == "__TsanDeviceGlobalMetadata" ||
-             Name == "__AsanKernelMetadata" || Name == "__MsanKernelMetadata" ||
-             Name == "__TsanKernelMetadata"))
+             Name == llvm::sycl::utils::ASAN_KERNEL_METADATA_PREFIX ||
+             Name == llvm::sycl::utils::MSAN_KERNEL_METADATA_PREFIX ||
+             Name == llvm::sycl::utils::TSAN_KERNEL_METADATA_PREFIX))
           continue;
 
         // Add symbol name with the target prefix to the buffer.

llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

@@ -19,15 +19,7 @@
 #include <functional>
 
 namespace llvm {
-
-class Module;
-
 namespace sycl {
-
-bool isModuleUsingAsan(const Module &M);
-bool isModuleUsingMsan(const Module &M);
-bool isModuleUsingTsan(const Module &M);
-
 namespace utils {
 constexpr char ATTR_SYCL_MODULE_ID[] = "sycl-module-id";
 constexpr char ATTR_SYCL_OPTLEVEL[] = "sycl-optlevel";

llvm/include/llvm/SYCLLowerIR/SanitizerUtils.h

@@ -0,0 +1,31 @@
+//===------------ SanitizerUtils.h - sanitizer utility functions --------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// Utility functions for device sanitizers.
+//===----------------------------------------------------------------------===//
+#pragma once
+
+#include "llvm/ADT/StringRef.h"
+
+namespace llvm {
+
+class Module;
+
+namespace sycl {
+namespace utils {
+
+constexpr StringRef ASAN_KERNEL_METADATA_PREFIX = "__AsanKernelMetadata";
+constexpr StringRef MSAN_KERNEL_METADATA_PREFIX = "__MsanKernelMetadata";
+constexpr StringRef TSAN_KERNEL_METADATA_PREFIX = "__TsanKernelMetadata";
+
+bool isModuleUsingAsan(const Module &M);
+bool isModuleUsingMsan(const Module &M);
+bool isModuleUsingTsan(const Module &M);
+
+} // namespace utils
+} // namespace sycl
+} // namespace llvm

llvm/lib/SYCLLowerIR/CMakeLists.txt

@@ -64,6 +64,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
   TargetHelpers.cpp
 
   SanitizerPostOptimizer.cpp
+  SanitizerUtils.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLLowerIR

llvm/lib/SYCLLowerIR/SYCLUtils.cpp

@@ -14,25 +14,6 @@
 
 namespace llvm {
 namespace sycl {
-
-bool isModuleUsingAsan(const Module &M) {
-  return any_of(M.globals(), [](const GlobalVariable &GV) {
-    return GV.getName().starts_with("__AsanKernelMetadata");
-  });
-}
-
-bool isModuleUsingMsan(const Module &M) {
-  return any_of(M.globals(), [](const GlobalVariable &GV) {
-    return GV.getName().starts_with("__MsanKernelMetadata");
-  });
-}
-
-bool isModuleUsingTsan(const Module &M) {
-  return any_of(M.globals(), [](const GlobalVariable &GV) {
-    return GV.getName().starts_with("__TsanKernelMetadata");
-  });
-}
-
 namespace utils {
 
 using namespace llvm::esimd;

llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp

@@ -13,7 +13,7 @@
 //===----------------------------------------------------------------------===//
 
 #include "llvm/SYCLLowerIR/SanitizerPostOptimizer.h"
-#include "llvm/SYCLLowerIR/SYCLUtils.h"
+#include "llvm/SYCLLowerIR/SanitizerUtils.h"
 
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/InstVisitor.h"
@@ -58,9 +58,9 @@ static bool FixSanitizerKernelMetadata(Module &M) {
   SmallVector<GlobalVariable *, 4> KernelMetadatas;
   for (GlobalVariable &GV : M.globals()) {
     auto GVName = GV.getName();
-    if (GVName.starts_with("__AsanKernelMetadata") ||
-        GVName.starts_with("__MsanKernelMetadata") ||
-        GVName.starts_with("__TsanKernelMetadata")) {
+    if (GVName.starts_with(sycl::utils::ASAN_KERNEL_METADATA_PREFIX) ||
+        GVName.starts_with(sycl::utils::MSAN_KERNEL_METADATA_PREFIX) ||
+        GVName.starts_with(sycl::utils::TSAN_KERNEL_METADATA_PREFIX)) {
       KernelMetadatas.push_back(&GV);
     }
   }
@@ -131,7 +131,7 @@ PreservedAnalyses SanitizerPostOptimizerPass::run(Module &M,
   if (!FixSanitizerKernelMetadata(M))
     return PreservedAnalyses::all();
 
-  if (sycl::isModuleUsingMsan(M)) {
+  if (sycl::utils::isModuleUsingMsan(M)) {
     EliminateDeadCheck V;
     V.visit(M);
     V.eraseDeadCheck();

llvm/lib/SYCLLowerIR/SanitizerUtils.cpp

@@ -0,0 +1,38 @@
+//===------------ SanitizerUtils.cpp - sanitizer utility functions ------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// Utility functions for device sanitizers.
+//===----------------------------------------------------------------------===//
+#include "llvm/SYCLLowerIR/SanitizerUtils.h"
+#include "llvm/IR/GlobalVariable.h"
+#include "llvm/IR/Module.h"
+
+namespace llvm {
+namespace sycl {
+namespace utils {
+
+bool isModuleUsingAsan(const Module &M) {
+  return any_of(M.globals(), [](const GlobalVariable &GV) {
+    return GV.getName().starts_with(ASAN_KERNEL_METADATA_PREFIX);
+  });
+}
+
+bool isModuleUsingMsan(const Module &M) {
+  return any_of(M.globals(), [](const GlobalVariable &GV) {
+    return GV.getName().starts_with(MSAN_KERNEL_METADATA_PREFIX);
+  });
+}
+
+bool isModuleUsingTsan(const Module &M) {
+  return any_of(M.globals(), [](const GlobalVariable &GV) {
+    return GV.getName().starts_with(TSAN_KERNEL_METADATA_PREFIX);
+  });
+}
+
+} // namespace utils
+} // namespace sycl
+} // namespace llvm

llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp

@@ -17,6 +17,7 @@
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/SYCLLowerIR/SYCLKernelParamOptInfo.h"
 #include "llvm/SYCLLowerIR/SYCLUtils.h"
+#include "llvm/SYCLLowerIR/SanitizerUtils.h"
 #include "llvm/SYCLLowerIR/SpecConstants.h"
 #include "llvm/SYCLPostLink/ModuleSplitter.h"
 #include <queue>
@@ -391,11 +392,11 @@ PropSetRegTy computeModuleProperties(const Module &M,
   }
 
   {
-    if (isModuleUsingAsan(M))
+    if (utils::isModuleUsingAsan(M))
       PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "sanUsed", "asan");
-    else if (isModuleUsingMsan(M))
+    else if (utils::isModuleUsingMsan(M))
       PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "sanUsed", "msan");
-    else if (isModuleUsingTsan(M))
+    else if (utils::isModuleUsingTsan(M))
       PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "sanUsed", "tsan");
   }
 

llvm/lib/SYCLPostLink/ModuleSplitter.cpp

@@ -30,6 +30,7 @@
 #include "llvm/SYCLLowerIR/SYCLJointMatrixTransform.h"
 #include "llvm/SYCLLowerIR/SYCLUtils.h"
 #include "llvm/SYCLLowerIR/SanitizerPostOptimizer.h"
+#include "llvm/SYCLLowerIR/SanitizerUtils.h"
 #include "llvm/SYCLLowerIR/SpecConstants.h"
 #include "llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h"
 #include "llvm/Support/CommandLine.h"
@@ -1215,8 +1216,8 @@ bool runPreSplitProcessingPipeline(Module &M) {
     MPM.addPass(RemoveDeviceGlobalFromLLVMCompilerUsed());
 
   // Sanitizer specific passes.
-  if (sycl::isModuleUsingAsan(M) || sycl::isModuleUsingMsan(M) ||
-      sycl::isModuleUsingTsan(M))
+  if (sycl::utils::isModuleUsingAsan(M) || sycl::utils::isModuleUsingMsan(M) ||
+      sycl::utils::isModuleUsingTsan(M))
     MPM.addPass(SanitizerPostOptimizerPass());
 
   // Transform Joint Matrix builtin calls to align them with SPIR-V friendly

llvm/lib/SYCLPostLink/SanitizerPostSplitProcessing.cpp

@@ -14,14 +14,16 @@
 #include "llvm/IR/GlobalVariable.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/Module.h"
+#include "llvm/SYCLLowerIR/SanitizerUtils.h"
 
 using namespace llvm;
 
 namespace {
 
 bool createSanitizerModuleID(Module &M) {
-  constexpr StringRef Prefixes[] = {
-      "__AsanKernelMetadata", "__MsanKernelMetadata", "__TsanKernelMetadata"};
+  constexpr StringRef Prefixes[] = {sycl::utils::ASAN_KERNEL_METADATA_PREFIX,
+                                    sycl::utils::MSAN_KERNEL_METADATA_PREFIX,
+                                    sycl::utils::TSAN_KERNEL_METADATA_PREFIX};
   SmallVector<StringRef, 3> ModuleIDs;
   for (GlobalVariable &GV : M.globals()) {
     auto GVName = GV.getName();

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

@@ -45,6 +45,7 @@
 #include <llvm/SYCLLowerIR/SYCLDeviceLibBF16.h>
 #include <llvm/SYCLLowerIR/SYCLJointMatrixTransform.h>
 #include <llvm/SYCLLowerIR/SYCLUtils.h>
+#include <llvm/SYCLLowerIR/SanitizerUtils.h>
 #include <llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h>
 #include <llvm/SYCLPostLink/ModuleSplitter.h>
 #include <llvm/Support/BLAKE3.h>
@@ -1040,8 +1041,9 @@ jit_compiler::performPostLink(ModuleUPtr Module,
   // Otherwise: Port over the `removeSYCLKernelsConstRefArray` and
   // `removeDeviceGlobalFromCompilerUsed` methods.
 
-  assert(!(isModuleUsingAsan(*Module) || isModuleUsingMsan(*Module) ||
-           isModuleUsingTsan(*Module)));
+  assert(!(utils::isModuleUsingAsan(*Module) ||
+           utils::isModuleUsingMsan(*Module) ||
+           utils::isModuleUsingTsan(*Module)));
   // Otherwise: Run `SanitizerKernelMetadataPass`.
 
   // Transform Joint Matrix builtin calls to align them with SPIR-V friendly