From 976b1ea0e9be72a4c46f1209b2d32195fdbd0be7 Mon Sep 17 00:00:00 2001 From: "Zhao, Maosu" Date: Fri, 5 Dec 2025 06:09:00 +0100 Subject: [PATCH 1/2] [DevSAN] Don't keep original meta data for new SanitizerKernelMetadata Otherwise, IGC compiler may not export the SanitizerKernelMetadata as device globals. --- llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp b/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp index ce75bc888a5e8..5500f83dea8a7 100644 --- a/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp +++ b/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp @@ -84,7 +84,6 @@ static bool FixSanitizerKernelMetadata(Module &M) { KernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local); KernelMetadata->setDSOLocal(true); KernelMetadata->copyAttributesFrom(KernelMetadataOld); - KernelMetadata->copyMetadata(KernelMetadataOld, 0); KernelMetadataOld->eraseFromParent(); } From ec84ef70fc6faff8ad91341a6c8ed22907b6a187 Mon Sep 17 00:00:00 2001 From: "Zhao, Maosu" Date: Thu, 11 Dec 2025 07:52:58 +0100 Subject: [PATCH 2/2] update related test --- llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll b/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll index fc4d6ff8e78af..763e49219a5ba 100644 --- a/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll +++ b/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll @@ -17,8 +17,10 @@ target triple = "spir64-unknown-unknown" $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any @__msan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00" -@__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 +@__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 }], !spirv.Decorations !9 #0 +; CHECK-IR: @__MsanKernelMetadata = dso_local local_unnamed_addr addrspace(1) global %0 { {{.*}} }, !spirv.Decorations [[MD1:![0-9]+]] #{{.*}} +; CHECK-IR: [[MD1]] = !{[[MD2:![0-9]+]]} +; CHECK-IR: [[MD2]] = !{i32 6147, i32 0, !"_Z20__SanitizerKernelMetadata"} @__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 @__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" @@ -85,3 +87,5 @@ attributes #2 = { mustprogress noinline norecurse nounwind sanitize_memory uwtab !6 = !{i32 563} !7 = !{i32 -1} !8 = !{} +!9 = !{!10} +!10 = !{i32 41, !"__MsanKernelMetadata", i32 0}