[DevTSAN] Support device thread sanitizer for device globals

llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp

@@ -118,13 +118,19 @@ struct ThreadSanitizerOnSpirv {
 
   void initialize();
 
-  void instrumentKernelsMetadata();
+  void instrumentModule();
 
   void appendDebugInfoToArgs(Instruction *I, SmallVectorImpl<Value *> &Args);
 
 private:
+  void instrumentGlobalVariables();
+
+  void instrumentKernelsMetadata();
+
   bool isSupportedSPIRKernel(Function &F);
 
+  bool isUnsupportedDeviceGlobal(const GlobalVariable &G);
+
   GlobalVariable *GetOrCreateGlobalString(StringRef Name, StringRef Value,
                                           unsigned AddressSpace);
 
@@ -243,7 +249,7 @@ PreservedAnalyses ModuleThreadSanitizerPass::run(Module &M,
     return PreservedAnalyses::all();
   if (Triple(M.getTargetTriple()).isSPIROrSPIRV()) {
     ThreadSanitizerOnSpirv Spirv(M);
-    Spirv.instrumentKernelsMetadata();
+    Spirv.instrumentModule();
   } else
     insertModuleCtor(M);
   return PreservedAnalyses::none();
@@ -327,6 +333,70 @@ bool ThreadSanitizerOnSpirv::isSupportedSPIRKernel(Function &F) {
   return true;
 }
 
+bool ThreadSanitizerOnSpirv::isUnsupportedDeviceGlobal(
+    const GlobalVariable &G) {
+  if (G.user_empty())
+    return true;
+  // Skip instrumenting on "__TsanKernelMetadata" etc.
+  if (G.getName().starts_with("__Tsan"))
+    return true;
+  if (G.getName().starts_with("__tsan_"))
+    return true;
+  if (G.getName().starts_with("__spirv_BuiltIn"))
+    return true;
+  if (G.getName().starts_with("__usid_str"))
+    return true;
+  // TODO: Will support global variable with local address space later.
+  if (G.getAddressSpace() == kSpirOffloadLocalAS)
+    return true;
+  // Global variables have constant value or constant address space will not
+  // trigger race condition.
+  if (G.isConstant() || G.getAddressSpace() == kSpirOffloadConstantAS)
+    return true;
+  return false;
+}
+
+void ThreadSanitizerOnSpirv::instrumentModule() {
+  instrumentGlobalVariables();
+  instrumentKernelsMetadata();
+}
+
+void ThreadSanitizerOnSpirv::instrumentGlobalVariables() {
+  SmallVector<Constant *, 8> DeviceGlobalMetadata;
+
+  // Device global metadata 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)) {
+      for (auto *User : G.users())
+        if (auto *Inst = dyn_cast<Instruction>(User))
+          Inst->setNoSanitizeMetadata();
+      continue;
+    }
+
+    DeviceGlobalMetadata.push_back(ConstantStruct::get(
+        StructTy,
+        ConstantInt::get(IntptrTy, DL.getTypeAllocSize(G.getValueType())),
+        ConstantExpr::getPointerCast(&G, IntptrTy)));
+  }
+
+  if (DeviceGlobalMetadata.empty())
+    return;
+
+  // 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, "__TsanDeviceGlobalMetadata", nullptr,
+      GlobalValue::NotThreadLocal, 1);
+  MsanDeviceGlobalMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
+}
+
 void ThreadSanitizerOnSpirv::instrumentKernelsMetadata() {
   SmallVector<Constant *, 8> SpirKernelsMetadata;
 

llvm/test/Instrumentation/ThreadSanitizer/SPIRV/instrument_device_global.ll

@@ -0,0 +1,20 @@
+; RUN: opt < %s -passes='function(tsan),module(tsan-module)' -tsan-instrument-func-entry-exit=0 -tsan-instrument-memintrinsics=0 -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"
+
+@dev_global = external addrspace(1) global { [4 x i32] }
+@dev_global_no_users = dso_local addrspace(1) global { [4 x i32] } zeroinitializer
+@.str = external addrspace(1) constant [59 x i8]
+@__spirv_BuiltInGlobalInvocationId = external addrspace(1) constant <3 x i64>
+
+; CHECK: @__TsanDeviceGlobalMetadata
+; CHECK-NOT: @dev_global_no_users
+; CHECK-NOT: @.str
+; CHECK-NOT: @__spirv_BuiltInGlobalInvocationId
+; CHECK-SAME: @dev_global
+
+define spir_func void @test() {
+entry:
+  %call = call spir_func ptr addrspace(4) null(ptr addrspace(4) addrspacecast (ptr addrspace(1) @dev_global to ptr addrspace(4)), i64 0)
+  ret void
+}

sycl/test-e2e/ThreadSanitizer/check_device_global.cpp

@@ -0,0 +1,30 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// ALLOW_RETRIES: 10
+// RUN: %{build} %device_tsan_flags -O0 -g -o %t1.out
+// RUN: %{run} %t1.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_tsan_flags -O2 -g -o %t2.out
+// RUN: %{run} %t2.out 2>&1 | FileCheck %s
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/device_global/device_global.hpp>
+
+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;
+
+int main() {
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+     h.parallel_for<class Test>(sycl::nd_range<1>(32, 8),
+                                [=](sycl::nd_item<1>) { dev_global[0]++; });
+   }).wait();
+  // CHECK: WARNING: DeviceSanitizer: data race
+  // CHECK-NEXT: When write of size 4 at 0x{{.*}} in kernel <{{.*}}Test>
+  // CHECK-NEXT: #0 {{.*}}check_device_global.cpp:[[@LINE-4]]
+
+  return 0;
+}

unified-runtime/source/loader/layers/sanitizer/tsan/tsan_ddi.cpp

@@ -130,6 +130,95 @@ ur_result_t urContextRelease(
   return UR_RESULT_SUCCESS;
 }
 
+///////////////////////////////////////////////////////////////////////////////
+/// @brief Intercept function for urProgramBuild
+ur_result_t urProgramBuild(
+    /// [in] handle of the context object
+    ur_context_handle_t hContext,
+    /// [in] handle of the program object
+    ur_program_handle_t hProgram,
+    /// [in] string of build options
+    const char *pOptions) {
+  getContext()->logger.debug("==== urProgramBuild");
+
+  UR_CALL(
+      getContext()->urDdiTable.Program.pfnBuild(hContext, hProgram, pOptions));
+
+  UR_CALL(getTsanInterceptor()->registerProgram(hProgram));
+
+  return UR_RESULT_SUCCESS;
+}
+
+///////////////////////////////////////////////////////////////////////////////
+/// @brief Intercept function for urProgramBuildExp
+ur_result_t urProgramBuildExp(
+    /// [in] Handle of the program to build.
+    ur_program_handle_t hProgram,
+    /// [in] number of devices
+    uint32_t numDevices,
+    /// [in][range(0, numDevices)] pointer to array of device handles
+    ur_device_handle_t *phDevices,
+    /// [in][optional] pointer to build options null-terminated string.
+    const char *pOptions) {
+  getContext()->logger.debug("==== urProgramBuildExp");
+
+  UR_CALL(getContext()->urDdiTable.ProgramExp.pfnBuildExp(hProgram, numDevices,
+                                                          phDevices, pOptions));
+  UR_CALL(getTsanInterceptor()->registerProgram(hProgram));
+
+  return UR_RESULT_SUCCESS;
+}
+
+///////////////////////////////////////////////////////////////////////////////
+/// @brief Intercept function for urProgramLink
+ur_result_t urProgramLink(
+    /// [in] handle of the context instance.
+    ur_context_handle_t hContext,
+    /// [in] number of program handles in `phPrograms`.
+    uint32_t count,
+    /// [in][range(0, count)] pointer to array of program handles.
+    const ur_program_handle_t *phPrograms,
+    /// [in][optional] pointer to linker options null-terminated string.
+    const char *pOptions,
+    /// [out] pointer to handle of program object created.
+    ur_program_handle_t *phProgram) {
+  getContext()->logger.debug("==== urProgramLink");
+
+  UR_CALL(getContext()->urDdiTable.Program.pfnLink(hContext, count, phPrograms,
+                                                   pOptions, phProgram));
+
+  UR_CALL(getTsanInterceptor()->registerProgram(*phProgram));
+
+  return UR_RESULT_SUCCESS;
+}
+
+///////////////////////////////////////////////////////////////////////////////
+/// @brief Intercept function for urProgramLinkExp
+ur_result_t urProgramLinkExp(
+    /// [in] handle of the context instance.
+    ur_context_handle_t hContext,
+    /// [in] number of devices
+    uint32_t numDevices,
+    /// [in][range(0, numDevices)] pointer to array of device handles
+    ur_device_handle_t *phDevices,
+    /// [in] number of program handles in `phPrograms`.
+    uint32_t count,
+    /// [in][range(0, count)] pointer to array of program handles.
+    const ur_program_handle_t *phPrograms,
+    /// [in][optional] pointer to linker options null-terminated string.
+    const char *pOptions,
+    /// [out] pointer to handle of program object created.
+    ur_program_handle_t *phProgram) {
+  getContext()->logger.debug("==== urProgramLinkExp");
+
+  UR_CALL(getContext()->urDdiTable.ProgramExp.pfnLinkExp(
+      hContext, numDevices, phDevices, count, phPrograms, pOptions, phProgram));
+
+  UR_CALL(getTsanInterceptor()->registerProgram(*phProgram));
+
+  return UR_RESULT_SUCCESS;
+}
+
 ///////////////////////////////////////////////////////////////////////////////
 /// @brief Intercept function for urUSMDeviceAlloc
 __urdlllocal ur_result_t UR_APICALL urUSMDeviceAlloc(
@@ -283,6 +372,39 @@ __urdlllocal ur_result_t UR_APICALL urGetContextProcAddrTable(
   return result;
 }
 
+///////////////////////////////////////////////////////////////////////////////
+/// @brief Exported function for filling application's Program table
+///        with current process' addresses
+///
+/// @returns
+///     - ::UR_RESULT_SUCCESS
+///     - ::UR_RESULT_ERROR_INVALID_NULL_POINTER
+ur_result_t urGetProgramProcAddrTable(
+    /// [in,out] pointer to table of DDI function pointers
+    ur_program_dditable_t *pDdiTable) {
+  pDdiTable->pfnBuild = ur_sanitizer_layer::tsan::urProgramBuild;
+  pDdiTable->pfnLink = ur_sanitizer_layer::tsan::urProgramLink;
+
+  return UR_RESULT_SUCCESS;
+}
+
+/// @brief Exported function for filling application's ProgramExp table
+///        with current process' addresses
+///
+/// @returns
+///     - ::UR_RESULT_SUCCESS
+///     - ::UR_RESULT_ERROR_INVALID_NULL_POINTER
+ur_result_t urGetProgramExpProcAddrTable(
+    /// [in,out] pointer to table of DDI function pointers
+    ur_program_exp_dditable_t *pDdiTable) {
+  ur_result_t result = UR_RESULT_SUCCESS;
+
+  pDdiTable->pfnBuildExp = ur_sanitizer_layer::tsan::urProgramBuildExp;
+  pDdiTable->pfnLinkExp = ur_sanitizer_layer::tsan::urProgramLinkExp;
+
+  return result;
+}
+
 ///////////////////////////////////////////////////////////////////////////////
 /// @brief Exported function for filling application's USM table
 ///        with current process' addresses
@@ -363,6 +485,16 @@ ur_result_t initTsanDDITable(ur_dditable_t *dditable) {
         UR_API_VERSION_CURRENT, &dditable->Context);
   }
 
+  if (UR_RESULT_SUCCESS == result) {
+    result =
+        ur_sanitizer_layer::tsan::urGetProgramProcAddrTable(&dditable->Program);
+  }
+
+  if (UR_RESULT_SUCCESS == result) {
+    result = ur_sanitizer_layer::tsan::urGetProgramExpProcAddrTable(
+        &dditable->ProgramExp);
+  }
+
   if (UR_RESULT_SUCCESS == result) {
     result = ur_sanitizer_layer::tsan::urGetUSMProcAddrTable(
         UR_API_VERSION_CURRENT, &dditable->USM);