Skip to content

Commit 2ae0331

Browse files
AllanZynekbenzie
authored and
KornevNikita
committed
[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 7057fde commit 2ae0331

File tree

3 files changed

+167
-39
lines changed

3 files changed

+167
-39
lines changed

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

Lines changed: 99 additions & 39 deletions
Original file line numberDiff line numberDiff line change
@@ -750,50 +750,110 @@ Constant *getOrCreateGlobalString(Module &M, StringRef Name, StringRef Value,
750750
});
751751
}
752752

753-
static void extendSpirKernelArgs(Module &M) {
754-
SmallVector<Constant *, 8> SpirKernelsMetadata;
753+
static bool isUnsupportedDeviceGlobal(const GlobalVariable *G) {
754+
// Skip instrumenting on "__MsanKernelMetadata" etc.
755+
if (G->getName().starts_with("__Msan"))
756+
return true;
757+
if (G->getName().starts_with("__spirv_BuiltIn"))
758+
return true;
759+
if (G->getName().starts_with("__usid_str"))
760+
return true;
761+
if (G->getAddressSpace() == kSpirOffloadLocalAS ||
762+
G->getAddressSpace() == kSpirOffloadConstantAS)
763+
return true;
764+
return false;
765+
}
766+
767+
static void instrumentSPIRModule(Module &M) {
755768

756769
const auto &DL = M.getDataLayout();
757770
Type *IntptrTy = DL.getIntPtrType(M.getContext());
758771

759-
// SpirKernelsMetadata only saves fixed kernels, and is described by
760-
// following structure:
761-
// uptr unmangled_kernel_name
762-
// uptr unmangled_kernel_name_size
763-
StructType *StructTy = StructType::get(IntptrTy, IntptrTy);
764-
for (Function &F : M) {
765-
if (F.getCallingConv() != CallingConv::SPIR_KERNEL)
766-
continue;
772+
// Instrument __MsanKernelMetadata, which records information of sanitized
773+
// kernel
774+
{
775+
SmallVector<Constant *, 8> SpirKernelsMetadata;
776+
777+
// SpirKernelsMetadata only saves fixed kernels, and is described by
778+
// following structure:
779+
// uptr unmangled_kernel_name
780+
// uptr unmangled_kernel_name_size
781+
StructType *StructTy = StructType::get(IntptrTy, IntptrTy);
782+
for (Function &F : M) {
783+
if (F.getCallingConv() != CallingConv::SPIR_KERNEL)
784+
continue;
767785

768-
if (!F.hasFnAttribute(Attribute::SanitizeMemory) ||
769-
F.hasFnAttribute(Attribute::DisableSanitizerInstrumentation))
770-
continue;
786+
if (!F.hasFnAttribute(Attribute::SanitizeMemory) ||
787+
F.hasFnAttribute(Attribute::DisableSanitizerInstrumentation))
788+
continue;
771789

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

799859
PreservedAnalyses MemorySanitizerPass::run(Module &M,
@@ -810,7 +870,7 @@ PreservedAnalyses MemorySanitizerPass::run(Module &M,
810870
}
811871

812872
if (TargetTriple.isSPIROrSPIRV()) {
813-
extendSpirKernelArgs(M);
873+
instrumentSPIRModule(M);
814874
Modified = true;
815875
}
816876

‎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/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)