[SYCL][TSAN] AOT support for device thread sanitizer

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -603,7 +603,15 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
       // For DG2, we just use libsycl-msan as placeholder.
       {"libsycl-msan", "internal"},
       {"libsycl-msan-pvc", "internal"}};
-  const SYCLDeviceLibsList SYCLDeviceTsanLibs = {{"libsycl-tsan", "internal"}};
+  const SYCLDeviceLibsList SYCLDeviceTsanLibs = {
+      {"libsycl-tsan", "internal"},
+      {"libsycl-tsan-cpu", "internal"},
+      // Currently, we only provide aot tsan libdevice for PVC and CPU.
+      // For DG2, we just use libsycl-tsan as placeholder.
+      // TODO: replace "libsycl-tsan" with "libsycl-tsan-dg2" when DG2
+      // AOT support is added.
+      {"libsycl-tsan", "internal"},
+      {"libsycl-tsan-pvc", "internal"}};
 #endif
 
   const SYCLDeviceLibsList SYCLNativeCpuDeviceLibs = {
@@ -759,7 +767,7 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
   else if (SanitizeVal == "memory")
     addSingleLibrary(SYCLDeviceMsanLibs[sanitizer_lib_idx]);
   else if (SanitizeVal == "thread")
-    addLibraries(SYCLDeviceTsanLibs);
+    addSingleLibrary(SYCLDeviceTsanLibs[sanitizer_lib_idx]);
 #endif
 
   if (isSYCLNativeCPU(TargetTriple))
@@ -883,6 +891,8 @@ static llvm::SmallVector<StringRef, 16> SYCLDeviceLibList{
     "msan-pvc",
     "msan-cpu",
     "tsan",
+    "tsan-pvc",
+    "tsan-cpu",
 #endif
     "imf",
     "imf-fp64",

clang/test/Driver/sycl-device-lib-old-model.cpp

@@ -415,3 +415,26 @@
 // SYCL_DEVICE_TSAN_MACRO-SAME: "USE_SYCL_DEVICE_TSAN"
 // SYCL_DEVICE_TSAN_MACRO: llvm-link{{.*}} "-only-needed"
 // SYCL_DEVICE_TSAN_MACRO-SAME: "{{.*}}libsycl-tsan.bc"
+
+/// ###########################################################################
+/// test behavior of linking libsycl-tsan-pvc for PVC target AOT compilation when tsan flag is applied.
+// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_pvc --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL \
+// RUN: -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device pvc" --no-offload-new-driver %s \
+// RUN: --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device pvc" --no-offload-new-driver %s \
+// RUN: --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device 12.60.7" --no-offload-new-driver %s \
+// RUN: --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xs "-device 12.60.7" --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL \
+// RUN: -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
+// SYCL_DEVICE_LIB_TSAN_PVC: llvm-link{{.*}} "-only-needed" "{{.*}}libsycl-crt.bc"
+// SYCL_DEVICE_LIB_TSAN_PVC-SAME: "{{.*}}libsycl-tsan-pvc.bc"
+
+
+/// ###########################################################################
+/// test behavior of linking libsycl-tsan-cpu for CPU target AOT compilation when tsan flag is applied.
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL \
+// RUN: -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_CPU
+// SYCL_DEVICE_LIB_TSAN_CPU: llvm-link{{.*}} "-only-needed" "{{.*}}libsycl-crt.bc"
+// SYCL_DEVICE_LIB_TSAN_CPU-SAME: "{{.*}}libsycl-tsan-cpu.bc"

clang/test/Driver/sycl-device-lib.cpp

@@ -401,3 +401,15 @@
 // SYCL_DEVICE_TSAN_MACRO: "-cc1"
 // SYCL_DEVICE_TSAN_MACRO-SAME: "USE_SYCL_DEVICE_TSAN"
 // SYCL_DEVICE_TSAN_MACRO: libsycl-tsan.new.o
+
+/// test behavior of tsan libdevice linking when -fsanitize=thread is available for AOT targets
+// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_pvc --offload-new-driver %s --sysroot=%S/Inputs/SYCL \
+// RUN: -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
+// SYCL_DEVICE_LIB_TSAN_PVC: clang-linker-wrapper{{.*}} "-sycl-device-libraries
+// SYCL_DEVICE_LIB_TSAN_PVC-SAME: {{.*}}libsycl-tsan-pvc.new.o
+
+/// test behavior of tsan libdevice linking when -fsanitize=thread is available for AOT targets
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 --offload-new-driver %s --sysroot=%S/Inputs/SYCL \
+// RUN: -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_CPU
+// SYCL_DEVICE_LIB_TSAN_CPU: clang-linker-wrapper{{.*}} "-sycl-device-libraries
+// SYCL_DEVICE_LIB_TSAN_CPU-SAME: {{.*}}libsycl-tsan-cpu.new.o

libdevice/cmake/modules/SYCLLibdevice.cmake

@@ -433,6 +433,19 @@ else()
       EXTRA_OPTS -fno-sycl-instrument-device-code
                  -I${UR_SANITIZER_INCLUDE_DIR}
                  -I${CMAKE_CURRENT_SOURCE_DIR})
+
+    set(tsan_devicetypes pvc cpu)
+
+    foreach(tsan_ft IN LISTS sanitizer_filetypes)
+      foreach(tsan_device IN LISTS tsan_devicetypes)
+        compile_lib_ext(libsycl-tsan-${tsan_device}
+                        SRC sanitizer/tsan_rtl.cpp
+                        FILETYPE ${tsan_ft}
+                        DEPENDENCIES ${tsan_obj_deps}
+                        OPTS ${sanitizer_${tsan_device}_compile_opts_${tsan_ft}})
+      endforeach()
+    endforeach()
+
   endif()
 endif()
 

libdevice/sanitizer/tsan_rtl.cpp

@@ -115,6 +115,11 @@ inline __SYCL_GLOBAL__ RawShadow *MemToShadow_PVC(uptr addr, uint32_t as) {
 inline __SYCL_GLOBAL__ RawShadow *MemToShadow(uptr addr, uint32_t as) {
   __SYCL_GLOBAL__ RawShadow *shadow_ptr = nullptr;
 
+#if defined(__LIBDEVICE_CPU__)
+  shadow_ptr = MemToShadow_CPU(addr, as);
+#elif defined(__LIBDEVICE_PVC__)
+  shadow_ptr = MemToShadow_PVC(addr, as);
+#else
   if (TsanLaunchInfo->DeviceTy == DeviceType::CPU) {
     shadow_ptr = MemToShadow_CPU(addr, as);
   } else if (TsanLaunchInfo->DeviceTy == DeviceType::GPU_PVC) {
@@ -124,6 +129,7 @@ inline __SYCL_GLOBAL__ RawShadow *MemToShadow(uptr addr, uint32_t as) {
                                   (int)TsanLaunchInfo->DeviceTy));
     return nullptr;
   }
+#endif
 
   return shadow_ptr;
 }
@@ -148,6 +154,11 @@ inline Sid GetCurrentSid_GPU() {
 }
 
 inline Sid GetCurrentSid() {
+#if defined(__LIBDEVICE_CPU__)
+  return GetCurrentSid_CPU();
+#elif defined(__LIBDEVICE_PVC__)
+  return GetCurrentSid_GPU();
+#else
   if (TsanLaunchInfo->DeviceTy == DeviceType::CPU) {
     return GetCurrentSid_CPU();
   } else if (TsanLaunchInfo->DeviceTy != DeviceType::UNKNOWN) {
@@ -157,6 +168,7 @@ inline Sid GetCurrentSid() {
                                   (int)TsanLaunchInfo->DeviceTy));
     return 0;
   }
+#endif
 }
 
 inline RawShadow LoadShadow(const __SYCL_GLOBAL__ RawShadow *p) {
@@ -426,10 +438,7 @@ __tsan_unaligned_read16(uptr addr, uint32_t as,
   __tsan_unaligned_read8(addr + 8, as, file, line, func);
 }
 
-DEVICE_EXTERN_C_NOINLINE void __tsan_cleanup_private(uptr addr, uint32_t size) {
-  if (TsanLaunchInfo->DeviceTy != DeviceType::CPU)
-    return;
-
+static inline void __tsan_cleanup_private_cpu_impl(uptr addr, uint32_t size) {
   if (size) {
     addr = RoundDownTo(addr, kShadowCell);
     size = RoundUpTo(size, kShadowCell);
@@ -443,6 +452,19 @@ DEVICE_EXTERN_C_NOINLINE void __tsan_cleanup_private(uptr addr, uint32_t size) {
   }
 }
 
+DEVICE_EXTERN_C_NOINLINE void __tsan_cleanup_private(uptr addr, uint32_t size) {
+#if defined(__LIBDEVICE_CPU__)
+  __tsan_cleanup_private_cpu_impl(addr, size);
+#elif defined(__LIBDEVICE_PVC__)
+  return;
+#else
+  if (TsanLaunchInfo->DeviceTy != DeviceType::CPU)
+    return;
+
+  __tsan_cleanup_private_cpu_impl(addr, size);
+#endif
+}
+
 DEVICE_EXTERN_C_INLINE void __tsan_device_barrier() {
   Sid sid = GetCurrentSid();
   __spirv_ControlBarrier(__spv::Scope::Device, __spv::Scope::Device,
@@ -470,10 +492,7 @@ DEVICE_EXTERN_C_INLINE void __tsan_device_barrier() {
                              __spv::MemorySemanticsMask::WorkgroupMemory);
 }
 
-DEVICE_EXTERN_C_INLINE void __tsan_group_barrier() {
-  if (TsanLaunchInfo->DeviceTy == DeviceType::CPU)
-    return;
-
+static inline void __tsan_group_barrier_impl() {
   Sid sid = GetCurrentSid();
   __spirv_ControlBarrier(__spv::Scope::Workgroup, __spv::Scope::Workgroup,
                          __spv::MemorySemanticsMask::SequentiallyConsistent |
@@ -500,4 +519,16 @@ DEVICE_EXTERN_C_INLINE void __tsan_group_barrier() {
                              __spv::MemorySemanticsMask::WorkgroupMemory);
 }
 
+DEVICE_EXTERN_C_INLINE void __tsan_group_barrier() {
+#if defined(__LIBDEVICE_CPU__)
+  return;
+#elif defined(__LIBDEVICE_PVC__)
+  __tsan_group_barrier_impl();
+#else
+  if (TsanLaunchInfo->DeviceTy == DeviceType::CPU)
+    return;
+  __tsan_group_barrier_impl();
+#endif
+}
+
 #endif // __SPIR__ || __SPIRV__

sycl/test-e2e/ThreadSanitizer/aot/Inputs/usm_data_race.cpp

@@ -0,0 +1,17 @@
+#include "sycl/detail/core.hpp"
+#include "sycl/usm.hpp"
+
+int main() {
+  sycl::queue Q;
+  auto *array = sycl::malloc_device<char>(1, Q);
+  Q.submit([&](sycl::handler &h) {
+     h.parallel_for<class Test>(sycl::nd_range<1>(32, 8),
+                                [=](sycl::nd_item<1>) { array[0]++; });
+   }).wait();
+  // CHECK: DeviceSanitizer: data race
+  // CHECK-NEXT: When write of size 1 at 0x{{.*}} in kernel <{{.*}}Test>
+  // CHECK-NEXT: #0 {{.*}}usm_data_race.cpp:[[@LINE-4]]
+
+  sycl::free(array, Q);
+  return 0;
+}

sycl/test-e2e/ThreadSanitizer/aot/cpu.cpp

@@ -0,0 +1,14 @@
+// REQUIRES: linux, opencl-aot, cpu
+// ALLOW_RETRIES: 10
+
+// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O0 -g %S/Inputs/usm_data_race.cpp -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
+
+// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O1 -g %S/Inputs/usm_data_race.cpp -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
+
+// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O2 -g %S/Inputs/usm_data_race.cpp -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
+
+// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O3 -g %S/Inputs/usm_data_race.cpp -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp

sycl/test-e2e/ThreadSanitizer/aot/gpu.cpp

@@ -0,0 +1,15 @@
+// REQUIRES: linux, gpu && level_zero
+// REQUIRES: arch-intel_gpu_pvc
+// ALLOW_RETRIES: 10
+
+// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O0 -g %S/Inputs/usm_data_race.cpp -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
+
+// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O1 -g %S/Inputs/usm_data_race.cpp -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
+
+// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O2 -g %S/Inputs/usm_data_race.cpp -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
+
+// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O3 -g %S/Inputs/usm_data_race.cpp -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp

sycl/test-e2e/ThreadSanitizer/lit.local.cfg

@@ -21,6 +21,11 @@ config.substitutions.append(
     ("%device_tsan_flags", "-Xarch_device -fsanitize=thread")
 )
 
+# For GPU, DTSAN only support pvc currently
+config.substitutions.append(
+    ("%device_tsan_aot_flags", "-Xarch_device -fsanitize=thread %if cpu %{ -fsycl-targets=spir64_x86_64 %} %if gpu %{ -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen '-devices pvc' %}")
+)
+
 unsupported_san_flags = [
     "-fsanitize=address",
     "-fsanitize=memory",