diff --git a/clang/lib/Driver/OffloadBundler.cpp b/clang/lib/Driver/OffloadBundler.cpp index f51707bd88b0a..7a32a0c97c88a 100644 --- a/clang/lib/Driver/OffloadBundler.cpp +++ b/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 @@ -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. diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h index 6238b494e9124..c9ebcdae53f4b 100644 --- a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h +++ b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h @@ -19,15 +19,7 @@ #include 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"; diff --git a/llvm/include/llvm/SYCLLowerIR/SanitizerUtils.h b/llvm/include/llvm/SYCLLowerIR/SanitizerUtils.h new file mode 100644 index 0000000000000..d74c5b470eeff --- /dev/null +++ b/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 diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index badb3fd0972fd..72a49d6d5d494 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/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 diff --git a/llvm/lib/SYCLLowerIR/SYCLUtils.cpp b/llvm/lib/SYCLLowerIR/SYCLUtils.cpp index 86ecf2cb61ee0..84354acf7d416 100644 --- a/llvm/lib/SYCLLowerIR/SYCLUtils.cpp +++ b/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; diff --git a/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp b/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp index c0318044d89b6..32fe93ea6095f 100644 --- a/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp +++ b/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 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(); diff --git a/llvm/lib/SYCLLowerIR/SanitizerUtils.cpp b/llvm/lib/SYCLLowerIR/SanitizerUtils.cpp new file mode 100644 index 0000000000000..0c1214db88897 --- /dev/null +++ b/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 diff --git a/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp b/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp index 0edee0156eda5..d0f25b03b4a8c 100644 --- a/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp +++ b/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 @@ -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"); } diff --git a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp b/llvm/lib/SYCLPostLink/ModuleSplitter.cpp index 26c3c0e9eb6be..9d174068b28c5 100644 --- a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp +++ b/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 diff --git a/llvm/lib/SYCLPostLink/SanitizerPostSplitProcessing.cpp b/llvm/lib/SYCLPostLink/SanitizerPostSplitProcessing.cpp index 2e0cd13994abd..3de2b95b84c5f 100644 --- a/llvm/lib/SYCLPostLink/SanitizerPostSplitProcessing.cpp +++ b/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 ModuleIDs; for (GlobalVariable &GV : M.globals()) { auto GVName = GV.getName(); diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index 6cfb99e9f3319..7290917868a68 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -45,6 +45,7 @@ #include #include #include +#include #include #include #include @@ -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