[DevMSAN] Cleanup dead warning check after optimizations

llvm/include/llvm/SYCLLowerIR/SanitizerKernelMetadata.h → llvm/include/llvm/SYCLLowerIR/SanitizerPostOptimizer.h

@@ -15,8 +15,8 @@
 
 namespace llvm {
 
-class SanitizerKernelMetadataPass
-    : public PassInfoMixin<SanitizerKernelMetadataPass> {
+class SanitizerPostOptimizerPass
+    : public PassInfoMixin<SanitizerPostOptimizerPass> {
 public:
   PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM);
 };

llvm/lib/SYCLLowerIR/CMakeLists.txt

@@ -62,7 +62,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
   GlobalOffset.cpp
   TargetHelpers.cpp
 
-  SanitizerKernelMetadata.cpp
+  SanitizerPostOptimizer.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLLowerIR

llvm/lib/SYCLLowerIR/SanitizerKernelMetadata.cpp → llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp

@@ -12,11 +12,12 @@
 // "spirv.Decorations" is removed by llvm-link, so we add it here again.
 //===----------------------------------------------------------------------===//
 
-#include "llvm/SYCLLowerIR/SanitizerKernelMetadata.h"
+#include "llvm/SYCLLowerIR/SanitizerPostOptimizer.h"
 
 #include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/InstVisitor.h"
 
-#define DEBUG_TYPE "SanitizerKernelMetadata"
+#define DEBUG_TYPE "SanitizerPostOptimizer"
 
 using namespace llvm;
 
@@ -25,8 +26,32 @@ namespace llvm {
 constexpr StringRef SPIRV_DECOR_MD_KIND = "spirv.Decorations";
 constexpr uint32_t SPIRV_HOST_ACCESS_DECOR = 6147;
 
-PreservedAnalyses SanitizerKernelMetadataPass::run(Module &M,
-                                                   ModuleAnalysisManager &MAM) {
+struct EliminateDeadCheck : public InstVisitor<EliminateDeadCheck> {
+  void visitCallInst(CallInst &CI) {
+    // If the shadow value is constant zero, the check instruction can be safely
+    // erased.
+    auto *Func = CI.getCalledFunction();
+    if (!Func)
+      return;
+    auto FuncName = Func->getName();
+    if (!FuncName.contains("__msan_maybe_warning_"))
+      return;
+    auto *Shadow = CI.getArgOperand(0);
+    if (isa<ConstantInt>(Shadow) && cast<ConstantInt>(Shadow)->isZeroValue())
+      InstToErase.push_back(&CI);
+  }
+
+  void eraseDeadCheck() {
+    for (auto *CI : InstToErase)
+      CI->eraseFromParent();
+    InstToErase.clear();
+  }
+
+private:
+  SmallVector<CallInst *, 8> InstToErase;
+};
+
+static bool FixSanitizerKernelMetadata(Module &M) {
   auto *KernelMetadata = M.getNamedGlobal("__AsanKernelMetadata");
 
   if (!KernelMetadata)
@@ -36,7 +61,7 @@ PreservedAnalyses SanitizerKernelMetadataPass::run(Module &M,
     KernelMetadata = M.getNamedGlobal("__TsanKernelMetadata");
 
   if (!KernelMetadata)
-    return PreservedAnalyses::all();
+    return false;
 
   auto &DL = M.getDataLayout();
   auto &Ctx = M.getContext();
@@ -86,6 +111,20 @@ PreservedAnalyses SanitizerKernelMetadataPass::run(Module &M,
 
   KernelMetadata->addMetadata(MDKindID, *MDNode::get(Ctx, MDOps));
 
+  return true;
+}
+
+PreservedAnalyses SanitizerPostOptimizerPass::run(Module &M,
+                                                  ModuleAnalysisManager &MAM) {
+  if (!FixSanitizerKernelMetadata(M))
+    return PreservedAnalyses::all();
+
+  if (M.getNamedGlobal("__MsanKernelMetadata")) {
+    EliminateDeadCheck V;
+    V.visit(M);
+    V.eraseDeadCheck();
+  }
+
   return PreservedAnalyses::none();
 }
 

llvm/lib/SYCLPostLink/ModuleSplitter.cpp

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

llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll

@@ -20,7 +20,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
 @__MsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__msan_kernel to i64), i64 54 }] #0
 ; CHECK-IR: @__MsanKernelMetadata = dso_local local_unnamed_addr addrspace(1) global %0 { {{.*}} }, !spirv.Decorations
 @__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
-@__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"
+@__msan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"
 
 ; Function Attrs: mustprogress norecurse nounwind sanitize_memory uwtable
 define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel(ptr addrspace(1) noundef align 4 %_arg_array) local_unnamed_addr #1 comdat !srcloc !6 !kernel_arg_buffer_location !7 !sycl_fixed_targets !8 {
@@ -42,14 +42,16 @@ entry:
   %_mscmp3 = icmp ne i64 %_msprop, 0
   %_msor = or i1 %_mscmp, %_mscmp3
   %8 = zext i1 %_msor to i8
-  call void @__msan_maybe_warning_1(i8 zeroext %8, i32 zeroext 0, ptr addrspace(2) null, i32 0, ptr addrspace(2) @__asan_func)
+  call void @__msan_maybe_warning_1(i8 zeroext %8, i32 zeroext 0, ptr addrspace(2) null, i32 0, ptr addrspace(2) @__msan_func)
   %call.i = tail call spir_func noundef i64 @_Z3fooix(i32 noundef %0, i64 noundef %conv.i) #4
   %conv4.i = trunc i64 %call.i to i32
   %9 = ptrtoint ptr addrspace(1) %_arg_array to i64
   %10 = call i64 @__msan_get_shadow(i64 %9, i32 1)
   %11 = inttoptr i64 %10 to ptr addrspace(1)
   store i32 0, ptr addrspace(1) %11, align 4
   store i32 %conv4.i, ptr addrspace(1) %_arg_array, align 4
+  call void @__msan_maybe_warning_8(i8 zeroext 0, i32 zeroext 0, ptr addrspace(2) null, i32 0, ptr addrspace(2) @__msan_func)
+  ; CHECK-IR-NOT: call void @__msan_maybe_warning_8
   ret void
 }
 
@@ -63,6 +65,7 @@ entry:
 
 declare i64 @__msan_get_shadow(i64, i32)
 declare void @__msan_maybe_warning_1(i8, i32, ptr addrspace(2), i32, ptr addrspace(2))
+declare void @__msan_maybe_warning_8(i8, i32, ptr addrspace(2), i32, ptr addrspace(2))
 
 attributes #0 = { "sycl-device-global-size"="16" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__MsanKernelMetadata" }
 attributes #1 = { mustprogress norecurse nounwind sanitize_memory uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="check_call.cpp" "sycl-single-task" "uniform-work-group-size"="true" }