Skip to content

Commit ec8346e

Browse files
AllanZynekbenzie
AllanZyne
and
kbenzie
authored
[DeviceMSAN] Fix gpu crashed on device global variable (#16566)
UR: oneapi-src/unified-runtime#2534 --------- Co-authored-by: Kenneth Benzie (Benie) <[email protected]>
1 parent df91a6b commit ec8346e

File tree

4 files changed

+174
-46
lines changed

4 files changed

+174
-46
lines changed

‎llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp

Lines changed: 99 additions & 39 deletions
Original file line numberDiff line numberDiff line change
@@ -767,50 +767,110 @@ Constant *getOrCreateGlobalString(Module &M, StringRef Name, StringRef Value,
767767
});
768768
}
769769

770-
static void extendSpirKernelArgs(Module &M) {
771-
SmallVector<Constant *, 8> SpirKernelsMetadata;
770+
static bool isUnsupportedDeviceGlobal(const GlobalVariable *G) {
771+
// Skip instrumenting on "__MsanKernelMetadata" etc.
772+
if (G->getName().starts_with("__Msan"))
773+
return true;
774+
if (G->getName().starts_with("__spirv_BuiltIn"))
775+
return true;
776+
if (G->getName().starts_with("__usid_str"))
777+
return true;
778+
if (G->getAddressSpace() == kSpirOffloadLocalAS ||
779+
G->getAddressSpace() == kSpirOffloadConstantAS)
780+
return true;
781+
return false;
782+
}
783+
784+
static void instrumentSPIRModule(Module &M) {
772785

773786
const auto &DL = M.getDataLayout();
774787
Type *IntptrTy = DL.getIntPtrType(M.getContext());
775788

776-
// SpirKernelsMetadata only saves fixed kernels, and is described by
777-
// following structure:
778-
// uptr unmangled_kernel_name
779-
// uptr unmangled_kernel_name_size
780-
StructType *StructTy = StructType::get(IntptrTy, IntptrTy);
781-
for (Function &F : M) {
782-
if (F.getCallingConv() != CallingConv::SPIR_KERNEL)
783-
continue;
789+
// Instrument __MsanKernelMetadata, which records information of sanitized
790+
// kernel
791+
{
792+
SmallVector<Constant *, 8> SpirKernelsMetadata;
793+
794+
// SpirKernelsMetadata only saves fixed kernels, and is described by
795+
// following structure:
796+
// uptr unmangled_kernel_name
797+
// uptr unmangled_kernel_name_size
798+
StructType *StructTy = StructType::get(IntptrTy, IntptrTy);
799+
for (Function &F : M) {
800+
if (F.getCallingConv() != CallingConv::SPIR_KERNEL)
801+
continue;
784802

785-
if (!F.hasFnAttribute(Attribute::SanitizeMemory) ||
786-
F.hasFnAttribute(Attribute::DisableSanitizerInstrumentation))
787-
continue;
803+
if (!F.hasFnAttribute(Attribute::SanitizeMemory) ||
804+
F.hasFnAttribute(Attribute::DisableSanitizerInstrumentation))
805+
continue;
788806

789-
auto KernelName = F.getName();
790-
auto *KernelNameGV = getOrCreateGlobalString(M, "__msan_kernel", KernelName,
791-
kSpirOffloadConstantAS);
792-
SpirKernelsMetadata.emplace_back(ConstantStruct::get(
793-
StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy),
794-
ConstantInt::get(IntptrTy, KernelName.size())));
795-
}
796-
797-
// Create global variable to record spirv kernels' information
798-
ArrayType *ArrayTy = ArrayType::get(StructTy, SpirKernelsMetadata.size());
799-
Constant *MetadataInitializer =
800-
ConstantArray::get(ArrayTy, SpirKernelsMetadata);
801-
GlobalVariable *MsanSpirKernelMetadata = new GlobalVariable(
802-
M, MetadataInitializer->getType(), false, GlobalValue::AppendingLinkage,
803-
MetadataInitializer, "__MsanKernelMetadata", nullptr,
804-
GlobalValue::NotThreadLocal, 1);
805-
MsanSpirKernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
806-
// Add device global attributes
807-
MsanSpirKernelMetadata->addAttribute(
808-
"sycl-device-global-size", std::to_string(DL.getTypeAllocSize(ArrayTy)));
809-
MsanSpirKernelMetadata->addAttribute("sycl-device-image-scope");
810-
MsanSpirKernelMetadata->addAttribute("sycl-host-access", "0"); // read only
811-
MsanSpirKernelMetadata->addAttribute("sycl-unique-id",
812-
"_Z20__MsanKernelMetadata");
813-
MsanSpirKernelMetadata->setDSOLocal(true);
807+
auto KernelName = F.getName();
808+
auto *KernelNameGV = getOrCreateGlobalString(
809+
M, "__msan_kernel", KernelName, kSpirOffloadConstantAS);
810+
SpirKernelsMetadata.emplace_back(ConstantStruct::get(
811+
StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy),
812+
ConstantInt::get(IntptrTy, KernelName.size())));
813+
}
814+
815+
// Create global variable to record spirv kernels' information
816+
ArrayType *ArrayTy = ArrayType::get(StructTy, SpirKernelsMetadata.size());
817+
Constant *MetadataInitializer =
818+
ConstantArray::get(ArrayTy, SpirKernelsMetadata);
819+
GlobalVariable *MsanSpirKernelMetadata = new GlobalVariable(
820+
M, MetadataInitializer->getType(), false, GlobalValue::AppendingLinkage,
821+
MetadataInitializer, "__MsanKernelMetadata", nullptr,
822+
GlobalValue::NotThreadLocal, 1);
823+
MsanSpirKernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
824+
// Add device global attributes
825+
MsanSpirKernelMetadata->addAttribute(
826+
"sycl-device-global-size",
827+
std::to_string(DL.getTypeAllocSize(ArrayTy)));
828+
MsanSpirKernelMetadata->addAttribute("sycl-device-image-scope");
829+
MsanSpirKernelMetadata->addAttribute("sycl-host-access",
830+
"0"); // read only
831+
MsanSpirKernelMetadata->addAttribute("sycl-unique-id",
832+
"_Z20__MsanKernelMetadata");
833+
MsanSpirKernelMetadata->setDSOLocal(true);
834+
}
835+
836+
// Handle global variables:
837+
// - Skip sanitizing unsupported variables
838+
// - Instrument __MsanDeviceGlobalMetadata for device globals
839+
do {
840+
SmallVector<Constant *, 8> DeviceGlobalMetadata;
841+
842+
// Device global meta data is described by a structure
843+
// size_t device_global_size
844+
// size_t beginning address of the device global
845+
StructType *StructTy = StructType::get(IntptrTy, IntptrTy);
846+
847+
for (auto &G : M.globals()) {
848+
if (isUnsupportedDeviceGlobal(&G)) {
849+
for (auto *User : G.users())
850+
if (auto *Inst = dyn_cast<Instruction>(User))
851+
Inst->setNoSanitizeMetadata();
852+
continue;
853+
}
854+
855+
DeviceGlobalMetadata.push_back(ConstantStruct::get(
856+
StructTy,
857+
ConstantInt::get(IntptrTy, DL.getTypeAllocSize(G.getValueType())),
858+
ConstantExpr::getPointerCast(&G, IntptrTy)));
859+
}
860+
861+
if (DeviceGlobalMetadata.empty())
862+
break;
863+
864+
// Create meta data global to record device globals' information
865+
ArrayType *ArrayTy = ArrayType::get(StructTy, DeviceGlobalMetadata.size());
866+
Constant *MetadataInitializer =
867+
ConstantArray::get(ArrayTy, DeviceGlobalMetadata);
868+
GlobalVariable *MsanDeviceGlobalMetadata = new GlobalVariable(
869+
M, MetadataInitializer->getType(), false, GlobalValue::AppendingLinkage,
870+
MetadataInitializer, "__MsanDeviceGlobalMetadata", nullptr,
871+
GlobalValue::NotThreadLocal, 1);
872+
MsanDeviceGlobalMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
873+
} while (false);
814874
}
815875

816876
PreservedAnalyses MemorySanitizerPass::run(Module &M,
@@ -827,7 +887,7 @@ PreservedAnalyses MemorySanitizerPass::run(Module &M,
827887
}
828888

829889
if (TargetTriple.isSPIROrSPIRV()) {
830-
extendSpirKernelArgs(M);
890+
instrumentSPIRModule(M);
831891
Modified = true;
832892
}
833893

‎llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_device_global.ll

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -S | FileCheck %s
2+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
3+
target triple = "spir64-unknown-unknown"
4+
5+
@.str = external addrspace(1) constant [59 x i8]
6+
@__spirv_BuiltInGlobalInvocationId = external addrspace(1) constant <3 x i64>
7+
8+
; CHECK: @__MsanDeviceGlobalMetadata
9+
; CHECK-NOT: @__spirv_BuiltInGlobalInvocationId
10+
; CHECK-SAME: @.str

‎sycl/cmake/modules/UnifiedRuntimeTag.cmake

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
# commit afbb289aa8d4f3b27b1536ba33ca618b0aba65c7
2-
# Merge: ef70004f d7c33f88
3-
# Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
4-
# Date: Wed Jan 15 11:54:25 2025 +0000
5-
# Merge pull request #2520 from zhaomaosu/fix-buffer-shadow
6-
# [DevMSAN] Propagate shadow memory in buffer related APIs
7-
set(UNIFIED_RUNTIME_TAG afbb289aa8d4f3b27b1536ba33ca618b0aba65c7)
1+
# commit 9e48f543b8dd39d45563169433bb529583625dfe
2+
# Merge: 6a3fece6 1a1108b3
3+
# Author: Martin Grant <martin.morrisongrant@codeplay.com>
4+
# Date: Wed Jan 15 14:33:29 2025 +0000
5+
# Merge pull request #2540 from martygrant/martin/program-info-unswitch
6+
# Move urProgramGetInfo success test from a switch to individual tests.
7+
set(UNIFIED_RUNTIME_TAG 9e48f543b8dd39d45563169433bb529583625dfe)

‎sycl/test-e2e/MemorySanitizer/check_device_global.cpp

Lines changed: 58 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,58 @@
1+
// REQUIRES: linux, cpu || (gpu && level_zero)
2+
// RUN: %{build} %device_msan_flags -O0 -g -o %t1.out
3+
// RUN: %{run} not %t1.out 2>&1 | FileCheck %s
4+
// RUN: %{build} %device_msan_flags -O1 -g -o %t2.out
5+
// RUN: %{run} not %t2.out 2>&1 | FileCheck %s
6+
// RUN: %{build} %device_msan_flags -O2 -g -o %t3.out
7+
// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
8+
9+
#include <sycl/detail/core.hpp>
10+
#include <sycl/ext/oneapi/device_global/device_global.hpp>
11+
#include <sycl/usm.hpp>
12+
13+
using namespace sycl;
14+
using namespace sycl::ext::oneapi;
15+
using namespace sycl::ext::oneapi::experimental;
16+
17+
sycl::ext::oneapi::experimental::device_global<
18+
int[4], decltype(properties(device_image_scope, host_access_read_write))>
19+
dev_global;
20+
21+
__attribute__((noinline)) int check(int data) { return data + 1; }
22+
23+
int main() {
24+
sycl::queue Q;
25+
int *array = sycl::malloc_device<int>(4, Q);
26+
27+
Q.submit([&](sycl::handler &h) {
28+
h.single_task<class Test1>([=]() {
29+
dev_global[0] = 42;
30+
array[0] = check(dev_global[1]);
31+
array[1] = dev_global[1];
32+
});
33+
}).wait();
34+
35+
int val[4];
36+
Q.copy(dev_global, val).wait();
37+
assert(val[0] == 42);
38+
39+
Q.submit([&](sycl::handler &h) {
40+
h.single_task<class Test2>([=]() {
41+
array[0] = check(array[1]);
42+
dev_global[1] = array[2]; // uninitialzed value
43+
});
44+
}).wait();
45+
46+
Q.submit([&](sycl::handler &h) {
47+
h.single_task<class Test3>([=]() {
48+
array[0] = dev_global[1];
49+
check(array[0]);
50+
});
51+
}).wait();
52+
// CHECK: use-of-uninitialized-value
53+
// CHECK-NEXT: kernel <{{.*Test3}}>
54+
55+
sycl::free(array, Q);
56+
57+
return 0;
58+
}

0 commit comments

Comments
 (0)