Skip to content

[DevSAN] Don't keep original meta data for new SanitizerKernelMetadata #20838

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
zhaomaosu wants to merge 4 commits into
base: sycl
Choose a base branch
from
Open

[DevSAN] Don't keep original meta data for new SanitizerKernelMetadata #20838

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
976b1ea
[DevSAN] Don't keep original meta data for new SanitizerKernelMetadata
zhaomaosu Dec 5, 2025
031015c
Merge branch 'sycl' into devsan-dont-copy-metadata
zhaomaosu Dec 11, 2025
ec84ef7
update related test
zhaomaosu Dec 11, 2025
b1152fc
Merge branch 'sycl' into devsan-dont-copy-metadata
zhaomaosu Dec 12, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 0 additions & 1 deletion llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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();
}
Expand Down
8 changes: 6 additions & 2 deletions llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"

Expand Down Expand Up @@ -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}
Loading