From c4bbac5c63e5f99c5729c7bcdd5ff172d07253e6 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Thu, 9 Jan 2025 04:14:14 +0100 Subject: [PATCH 1/9] add __MsanDeviceGlobalMetadata --- .../Instrumentation/MemorySanitizer.cpp | 130 +++++++++++++----- 1 file changed, 93 insertions(+), 37 deletions(-) diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 3d5fede606f9f..24f82b3f3e1cb 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -767,50 +767,106 @@ Constant *getOrCreateGlobalString(Module &M, StringRef Name, StringRef Value, }); } +static bool isUnsupportedDeviceGlobal(GlobalVariable *G) { + // Non image scope device globals are implemented by device USM, and the + // out-of-bounds check for them will be done by sanitizer USM part. So we + // exclude them here. + if (!G->hasAttribute("sycl-device-image-scope")) + return true; + + // Skip instrumenting on "__MsanKernelMetadata" etc. + if (G->getName().starts_with("__Msan")) + return true; + + Attribute Attr = G->getAttribute("sycl-device-image-scope"); + return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false"); +} + static void extendSpirKernelArgs(Module &M) { - SmallVector SpirKernelsMetadata; const auto &DL = M.getDataLayout(); Type *IntptrTy = DL.getIntPtrType(M.getContext()); - // SpirKernelsMetadata only saves fixed kernels, and is described by - // following structure: - // uptr unmangled_kernel_name - // uptr unmangled_kernel_name_size - StructType *StructTy = StructType::get(IntptrTy, IntptrTy); - for (Function &F : M) { - if (F.getCallingConv() != CallingConv::SPIR_KERNEL) - continue; + // Instrument __MsanKernelMetadata, which records information of sanitized + // kernel + { + SmallVector SpirKernelsMetadata; + + // SpirKernelsMetadata only saves fixed kernels, and is described by + // following structure: + // uptr unmangled_kernel_name + // uptr unmangled_kernel_name_size + StructType *StructTy = StructType::get(IntptrTy, IntptrTy); + for (Function &F : M) { + if (F.getCallingConv() != CallingConv::SPIR_KERNEL) + continue; - if (!F.hasFnAttribute(Attribute::SanitizeMemory) || - F.hasFnAttribute(Attribute::DisableSanitizerInstrumentation)) - continue; + if (!F.hasFnAttribute(Attribute::SanitizeMemory) || + F.hasFnAttribute(Attribute::DisableSanitizerInstrumentation)) + continue; - auto KernelName = F.getName(); - auto *KernelNameGV = getOrCreateGlobalString(M, "__msan_kernel", KernelName, - kSpirOffloadConstantAS); - SpirKernelsMetadata.emplace_back(ConstantStruct::get( - StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy), - ConstantInt::get(IntptrTy, KernelName.size()))); - } - - // Create global variable to record spirv kernels' information - ArrayType *ArrayTy = ArrayType::get(StructTy, SpirKernelsMetadata.size()); - Constant *MetadataInitializer = - ConstantArray::get(ArrayTy, SpirKernelsMetadata); - GlobalVariable *MsanSpirKernelMetadata = new GlobalVariable( - M, MetadataInitializer->getType(), false, GlobalValue::AppendingLinkage, - MetadataInitializer, "__MsanKernelMetadata", nullptr, - GlobalValue::NotThreadLocal, 1); - MsanSpirKernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local); - // Add device global attributes - MsanSpirKernelMetadata->addAttribute( - "sycl-device-global-size", std::to_string(DL.getTypeAllocSize(ArrayTy))); - MsanSpirKernelMetadata->addAttribute("sycl-device-image-scope"); - MsanSpirKernelMetadata->addAttribute("sycl-host-access", "0"); // read only - MsanSpirKernelMetadata->addAttribute("sycl-unique-id", - "_Z20__MsanKernelMetadata"); - MsanSpirKernelMetadata->setDSOLocal(true); + auto KernelName = F.getName(); + auto *KernelNameGV = getOrCreateGlobalString( + M, "__msan_kernel", KernelName, kSpirOffloadConstantAS); + SpirKernelsMetadata.emplace_back(ConstantStruct::get( + StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy), + ConstantInt::get(IntptrTy, KernelName.size()))); + } + + // Create global variable to record spirv kernels' information + ArrayType *ArrayTy = ArrayType::get(StructTy, SpirKernelsMetadata.size()); + Constant *MetadataInitializer = + ConstantArray::get(ArrayTy, SpirKernelsMetadata); + GlobalVariable *MsanSpirKernelMetadata = new GlobalVariable( + M, MetadataInitializer->getType(), false, GlobalValue::AppendingLinkage, + MetadataInitializer, "__MsanKernelMetadata", nullptr, + GlobalValue::NotThreadLocal, 1); + MsanSpirKernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local); + // Add device global attributes + MsanSpirKernelMetadata->addAttribute( + "sycl-device-global-size", + std::to_string(DL.getTypeAllocSize(ArrayTy))); + MsanSpirKernelMetadata->addAttribute("sycl-device-image-scope"); + MsanSpirKernelMetadata->addAttribute("sycl-host-access", + "0"); // read only + MsanSpirKernelMetadata->addAttribute("sycl-unique-id", + "_Z20__MsanKernelMetadata"); + MsanSpirKernelMetadata->setDSOLocal(true); + } + + // Instrument __MsanDeviceGlobalMetadata, which records information of device + // global + do { + SmallVector DeviceGlobalMetadata; + + // Device global meta data is described by a structure + // size_t device_global_size + // size_t beginning address of the device global + StructType *StructTy = StructType::get(IntptrTy, IntptrTy); + + for (auto &G : M.globals()) { + if (isUnsupportedDeviceGlobal(&G)) + continue; + + DeviceGlobalMetadata.push_back(ConstantStruct::get( + StructTy, + ConstantInt::get(IntptrTy, DL.getTypeAllocSize(G.getValueType())), + ConstantExpr::getPointerCast(&G, IntptrTy))); + } + + if (DeviceGlobalMetadata.empty()) + break; + + // Create meta data global to record device globals' information + ArrayType *ArrayTy = ArrayType::get(StructTy, DeviceGlobalMetadata.size()); + Constant *MetadataInitializer = + ConstantArray::get(ArrayTy, DeviceGlobalMetadata); + GlobalVariable *MsanDeviceGlobalMetadata = new GlobalVariable( + M, MetadataInitializer->getType(), false, GlobalValue::AppendingLinkage, + MetadataInitializer, "__MsanDeviceGlobalMetadata", nullptr, + GlobalValue::NotThreadLocal, 1); + MsanDeviceGlobalMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local); + } while (false); } PreservedAnalyses MemorySanitizerPass::run(Module &M, From 3a051844f35ab18b7b3dececa1455a8491f85eea Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Thu, 9 Jan 2025 07:09:02 +0100 Subject: [PATCH 2/9] add test --- .../MemorySanitizer/check_device_global.cpp | 60 +++++++++++++++++++ 1 file changed, 60 insertions(+) create mode 100644 sycl/test-e2e/MemorySanitizer/check_device_global.cpp diff --git a/sycl/test-e2e/MemorySanitizer/check_device_global.cpp b/sycl/test-e2e/MemorySanitizer/check_device_global.cpp new file mode 100644 index 0000000000000..d54c15767e677 --- /dev/null +++ b/sycl/test-e2e/MemorySanitizer/check_device_global.cpp @@ -0,0 +1,60 @@ +// REQUIRES: linux, cpu || (gpu && level_zero) +// RUN: %{build} %device_msan_flags -O0 -g -o %t1.out +// RUN: %{run} not %t1.out 2>&1 | FileCheck %s +// RUN: %{build} %device_msan_flags -O1 -g -o %t2.out +// RUN: %{run} not %t2.out 2>&1 | FileCheck %s +// RUN: %{build} %device_msan_flags -O2 -g -o %t3.out +// RUN: %{run} not %t3.out 2>&1 | FileCheck %s + +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::oneapi; +using namespace sycl::ext::oneapi::experimental; + +sycl::ext::oneapi::experimental::device_global< + int[4], decltype(properties(device_image_scope, host_access_read_write))> + dev_global; + +__attribute__((noinline)) int check(int data) { + return data + 1; +} + +int main() { + sycl::queue Q; + int *array = sycl::malloc_device(4, Q); + + Q.submit([&](sycl::handler &h) { + h.single_task([=]() { + dev_global[0] = 42; + array[0] = check(dev_global[1]); + array[1] = dev_global[1]; + }); + }).wait(); + + int val[4]; + Q.copy(dev_global, val).wait(); + assert(val[0] == 42); + + Q.submit([&](sycl::handler &h) { + h.single_task([=]() { + array[0] = check(array[1]); + dev_global[1] = array[2]; // uninitialzed value + }); + }).wait(); + + Q.submit([&](sycl::handler &h) { + h.single_task([=]() { + array[0] = dev_global[1]; + check(array[0]); + }); + }).wait(); + // CHECK: use-of-uninitialized-value + // CHECK-NEXT: kernel <{{.*Test3}}> + + sycl::free(array, Q); + + return 0; +} From 704c7a1c7b347f8e2121a5dc0f17df88fb568e54 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Thu, 9 Jan 2025 07:11:22 +0100 Subject: [PATCH 3/9] update ur tag --- sycl/cmake/modules/FetchUnifiedRuntime.cmake | 2 +- sycl/cmake/modules/UnifiedRuntimeTag.cmake | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 72841724fa01d..1ed066b389aa3 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -116,7 +116,7 @@ if(SYCL_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") + set(UNIFIED_RUNTIME_REPO "https://github.com/AllanZyne/unified-runtime.git") include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/UnifiedRuntimeTag.cmake) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake index a5c42a284ac40..cb7b13e8d6914 100644 --- a/sycl/cmake/modules/UnifiedRuntimeTag.cmake +++ b/sycl/cmake/modules/UnifiedRuntimeTag.cmake @@ -4,4 +4,4 @@ # Date: Tue Jan 7 10:40:35 2025 +0000 # Merge pull request #2524 from nrspruit/fix_enqueue_wait_out_event # [L0]: Fix Out Event in Enqueue Wait Events to prevent reuse -set(UNIFIED_RUNTIME_TAG da04d13807044aaf17615b66577fb0e832011ab1) +set(UNIFIED_RUNTIME_TAG review/yang/msan_device_global) From 7a426cc2f9efbb18d8a2073ab3824c8521262aa9 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Fri, 10 Jan 2025 03:30:05 +0100 Subject: [PATCH 4/9] fix format --- sycl/test-e2e/MemorySanitizer/check_device_global.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/sycl/test-e2e/MemorySanitizer/check_device_global.cpp b/sycl/test-e2e/MemorySanitizer/check_device_global.cpp index d54c15767e677..f8b47569deb9b 100644 --- a/sycl/test-e2e/MemorySanitizer/check_device_global.cpp +++ b/sycl/test-e2e/MemorySanitizer/check_device_global.cpp @@ -18,9 +18,7 @@ sycl::ext::oneapi::experimental::device_global< int[4], decltype(properties(device_image_scope, host_access_read_write))> dev_global; -__attribute__((noinline)) int check(int data) { - return data + 1; -} +__attribute__((noinline)) int check(int data) { return data + 1; } int main() { sycl::queue Q; From 139d1abce5118e1937d16289bec9ed4c08be2213 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Fri, 10 Jan 2025 03:40:04 +0100 Subject: [PATCH 5/9] skip unsupported variables --- .../Instrumentation/MemorySanitizer.cpp | 16 ++++++++++++---- 1 file changed, 12 insertions(+), 4 deletions(-) diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 24f82b3f3e1cb..ecd379f4b8454 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -782,7 +782,7 @@ static bool isUnsupportedDeviceGlobal(GlobalVariable *G) { return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false"); } -static void extendSpirKernelArgs(Module &M) { +static void instrumentSPIRModule(Module &M) { const auto &DL = M.getDataLayout(); Type *IntptrTy = DL.getIntPtrType(M.getContext()); @@ -834,8 +834,9 @@ static void extendSpirKernelArgs(Module &M) { MsanSpirKernelMetadata->setDSOLocal(true); } - // Instrument __MsanDeviceGlobalMetadata, which records information of device - // global + // Handle global variables: + // - Skip sanitizing unsupported variables + // - Instrument __MsanDeviceGlobalMetadata for device globals do { SmallVector DeviceGlobalMetadata; @@ -845,6 +846,13 @@ static void extendSpirKernelArgs(Module &M) { StructType *StructTy = StructType::get(IntptrTy, IntptrTy); for (auto &G : M.globals()) { + // FIXME: temporarily disable local variables + if (G.isConstant() || G.getAddressSpace() == kSpirOffloadLocalAS) { + for (auto *User : G.users()) + if (auto *Inst = dyn_cast(User)) + Inst->setNoSanitizeMetadata(); + } + if (isUnsupportedDeviceGlobal(&G)) continue; @@ -883,7 +891,7 @@ PreservedAnalyses MemorySanitizerPass::run(Module &M, } if (TargetTriple.isSPIROrSPIRV()) { - extendSpirKernelArgs(M); + instrumentSPIRModule(M); Modified = true; } From 0e94f813156af31557d26782d2659267265edffd Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Mon, 13 Jan 2025 07:28:52 +0100 Subject: [PATCH 6/9] fix --- .../Instrumentation/MemorySanitizer.cpp | 24 +++++++------------ 1 file changed, 9 insertions(+), 15 deletions(-) diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index ecd379f4b8454..5341910973271 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -767,19 +767,16 @@ Constant *getOrCreateGlobalString(Module &M, StringRef Name, StringRef Value, }); } -static bool isUnsupportedDeviceGlobal(GlobalVariable *G) { - // Non image scope device globals are implemented by device USM, and the - // out-of-bounds check for them will be done by sanitizer USM part. So we - // exclude them here. - if (!G->hasAttribute("sycl-device-image-scope")) - return true; - +static bool isUnsupportedDeviceGlobal(const GlobalVariable *G) { // Skip instrumenting on "__MsanKernelMetadata" etc. if (G->getName().starts_with("__Msan")) return true; - - Attribute Attr = G->getAttribute("sycl-device-image-scope"); - return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false"); + if (G->getName().starts_with("__spirv_BuiltIn")) + return true; + if (G->getAddressSpace() == kSpirOffloadLocalAS || + G->getAddressSpace() == kSpirOffloadConstantAS) + return true; + return false; } static void instrumentSPIRModule(Module &M) { @@ -846,15 +843,12 @@ static void instrumentSPIRModule(Module &M) { StructType *StructTy = StructType::get(IntptrTy, IntptrTy); for (auto &G : M.globals()) { - // FIXME: temporarily disable local variables - if (G.isConstant() || G.getAddressSpace() == kSpirOffloadLocalAS) { + if (isUnsupportedDeviceGlobal(&G)) { for (auto *User : G.users()) if (auto *Inst = dyn_cast(User)) Inst->setNoSanitizeMetadata(); - } - - if (isUnsupportedDeviceGlobal(&G)) continue; + } DeviceGlobalMetadata.push_back(ConstantStruct::get( StructTy, From a69a4439c1c9b65ff6cc889817610cafa3372032 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Mon, 13 Jan 2025 08:49:04 +0100 Subject: [PATCH 7/9] add lit test --- .../MemorySanitizer/SPIRV/instrument_device_global.ll | 10 ++++++++++ 1 file changed, 10 insertions(+) create mode 100644 llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_device_global.ll diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_device_global.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_device_global.ll new file mode 100644 index 0000000000000..39c2775a923c2 --- /dev/null +++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_device_global.ll @@ -0,0 +1,10 @@ +; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -S | FileCheck %s +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" +target triple = "spir64-unknown-unknown" + +@.str = external addrspace(1) constant [59 x i8] +@__spirv_BuiltInGlobalInvocationId = external addrspace(1) constant <3 x i64> + +; CHECK: @__MsanDeviceGlobalMetadata +; CHECK-NOT: @__spirv_BuiltInGlobalInvocationId +; CHECK-SAME: @.str From 66e18aaef27a43f213ad7132ef07cb93f1f7c360 Mon Sep 17 00:00:00 2001 From: "Zhao, Yang2" Date: Tue, 14 Jan 2025 03:19:21 +0100 Subject: [PATCH 8/9] fix private alloca --- llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index 5341910973271..2386fc83fa3c9 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp @@ -773,6 +773,8 @@ static bool isUnsupportedDeviceGlobal(const GlobalVariable *G) { return true; if (G->getName().starts_with("__spirv_BuiltIn")) return true; + if (G->getName().starts_with("__usid_str")) + return true; if (G->getAddressSpace() == kSpirOffloadLocalAS || G->getAddressSpace() == kSpirOffloadConstantAS) return true; From 356d1e90e873e070215a45a3c8dbc480dd7c4087 Mon Sep 17 00:00:00 2001 From: "Kenneth Benzie (Benie)" Date: Wed, 15 Jan 2025 15:01:06 +0000 Subject: [PATCH 9/9] [UR] Bump main tag to 9e48f543 --- sycl/cmake/modules/FetchUnifiedRuntime.cmake | 2 +- sycl/cmake/modules/UnifiedRuntimeTag.cmake | 14 +++++++------- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 1ed066b389aa3..72841724fa01d 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -116,7 +116,7 @@ if(SYCL_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/AllanZyne/unified-runtime.git") + set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/UnifiedRuntimeTag.cmake) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake index 99805d1a0a556..526683c9cdf97 100644 --- a/sycl/cmake/modules/UnifiedRuntimeTag.cmake +++ b/sycl/cmake/modules/UnifiedRuntimeTag.cmake @@ -1,7 +1,7 @@ -# commit b2ac58f27c63b8ff714e8b0c39b79aaab05a3faf -# Merge: 3472b5bd ead3d07d -# Author: Kenneth Benzie (Benie) -# Date: Tue Jan 7 10:40:35 2025 +0000 -# Merge pull request #2524 from nrspruit/fix_enqueue_wait_out_event -# [L0]: Fix Out Event in Enqueue Wait Events to prevent reuse -set(UNIFIED_RUNTIME_TAG review/yang/msan_device_global) +# commit 9e48f543b8dd39d45563169433bb529583625dfe +# Merge: 6a3fece6 1a1108b3 +# Author: Martin Grant +# Date: Wed Jan 15 14:33:29 2025 +0000 +# Merge pull request #2540 from martygrant/martin/program-info-unswitch +# Move urProgramGetInfo success test from a switch to individual tests. +set(UNIFIED_RUNTIME_TAG 9e48f543b8dd39d45563169433bb529583625dfe)