Skip to content

[DeviceSanitizer] Fix device global type of KernelMetadata #16357

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

Merged
merged 3 commits into from
Dec 23, 2024
Merged

[DeviceSanitizer] Fix device global type of KernelMetadata #16357

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
662920e
Fix type of KernelMetada, by wrapping a structure type
AllanZyne Dec 13, 2024
99cbc03
fix window crash
AllanZyne Dec 18, 2024
4e31af0
Address comments
AllanZyne Dec 20, 2024
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
23 changes: 23 additions & 0 deletions llvm/lib/SYCLLowerIR/SanitizerKernelMetadata.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,6 +38,29 @@ PreservedAnalyses SanitizerKernelMetadataPass::run(Module &M,
auto &DL = M.getDataLayout();
auto &Ctx = M.getContext();

// Fix device global type, by wrapping a structure type
{
assert(KernelMetadata->getValueType()->isArrayTy());

auto *KernelMetadataOld = KernelMetadata;

StructType *StructTypeWithArray = StructType::create(Ctx);
StructTypeWithArray->setBody(KernelMetadataOld->getValueType());

KernelMetadata = new GlobalVariable(
M, StructTypeWithArray, false, GlobalValue::ExternalLinkage,
ConstantStruct::get(StructTypeWithArray,
KernelMetadataOld->getInitializer()),
"", nullptr, GlobalValue::NotThreadLocal, 1); // Global AddressSpace
KernelMetadata->takeName(KernelMetadataOld);
KernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
KernelMetadata->setDSOLocal(true);
KernelMetadata->copyAttributesFrom(KernelMetadataOld);
KernelMetadata->copyMetadata(KernelMetadataOld, 0);

KernelMetadataOld->eraseFromParent();
Copy link
Contributor

@wenju-he wenju-he Dec 16, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

why not generate KernelMetadataOld in struct type in the first place?

Copy link
Contributor Author

@AllanZyne AllanZyne Dec 16, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is a good question.

"KernelMetadata" is an array with appending linkage, so it can't be struct type at beginning.

I also considered to generate "KernelMetadata" at sycl-post-link step completely, but I use "KernelMetadata" to distinguish if the device image is a sanitized module at the same time. If we don't add "KernelMetadata" at sanitizer pass, it's possible that a sanitized module (but without any sanitized kernels) won't enable sanitizer layer at UR.

This is a design decision that if user uses "-fsanitize=" to compile program, we will always enable sanitizer layer in UR.

}

// Fix attributes
KernelMetadata->addAttribute(
"sycl-device-global-size",
Expand Down
2 changes: 1 addition & 1 deletion llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any

@__asan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00"
@__AsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__asan_kernel to i64), i64 54 }] #2
; CHECK-IR: @__AsanKernelMetadata {{.*}} !spirv.Decorations
; CHECK-IR: @__AsanKernelMetadata = 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"

Expand Down
2 changes: 1 addition & 1 deletion 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 @@ -18,7 +18,7 @@ $_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 {{.*}} !spirv.Decorations
; 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"

Expand Down
Loading