diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 56c73c7d93c34..d308587d2523b 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/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 SYCLDeviceLibList{ "msan-pvc", "msan-cpu", "tsan", + "tsan-pvc", + "tsan-cpu", #endif "imf", "imf-fp64", diff --git a/clang/test/Driver/Inputs/SYCL/lib/libsycl-tsan-cpu.bc b/clang/test/Driver/Inputs/SYCL/lib/libsycl-tsan-cpu.bc new file mode 100644 index 0000000000000..e69de29bb2d1d diff --git a/clang/test/Driver/Inputs/SYCL/lib/libsycl-tsan-pvc.bc b/clang/test/Driver/Inputs/SYCL/lib/libsycl-tsan-pvc.bc new file mode 100644 index 0000000000000..e69de29bb2d1d diff --git a/clang/test/Driver/sycl-device-lib-old-model.cpp b/clang/test/Driver/sycl-device-lib-old-model.cpp index fdabe3b6c5838..ffc752d73abd1 100644 --- a/clang/test/Driver/sycl-device-lib-old-model.cpp +++ b/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" diff --git a/clang/test/Driver/sycl-device-lib.cpp b/clang/test/Driver/sycl-device-lib.cpp index 2fbf7411bae3d..fd59997afda4b 100644 --- a/clang/test/Driver/sycl-device-lib.cpp +++ b/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 diff --git a/libdevice/cmake/modules/SYCLLibdevice.cmake b/libdevice/cmake/modules/SYCLLibdevice.cmake index c3d9be1a338b7..ecbb2b3ca697c 100644 --- a/libdevice/cmake/modules/SYCLLibdevice.cmake +++ b/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() diff --git a/libdevice/sanitizer/tsan_rtl.cpp b/libdevice/sanitizer/tsan_rtl.cpp index 50f8fec1720ac..6c7b6f26bb9bd 100644 --- a/libdevice/sanitizer/tsan_rtl.cpp +++ b/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__ diff --git a/sycl/test-e2e/ThreadSanitizer/aot/Inputs/usm_data_race.cpp b/sycl/test-e2e/ThreadSanitizer/aot/Inputs/usm_data_race.cpp new file mode 100644 index 0000000000000..8a7935d4510e1 --- /dev/null +++ b/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(1, Q); + Q.submit([&](sycl::handler &h) { + h.parallel_for(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; +} diff --git a/sycl/test-e2e/ThreadSanitizer/aot/cpu.cpp b/sycl/test-e2e/ThreadSanitizer/aot/cpu.cpp new file mode 100644 index 0000000000000..a63ca0e348844 --- /dev/null +++ b/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 diff --git a/sycl/test-e2e/ThreadSanitizer/aot/gpu.cpp b/sycl/test-e2e/ThreadSanitizer/aot/gpu.cpp new file mode 100644 index 0000000000000..3c0935cad742b --- /dev/null +++ b/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 diff --git a/sycl/test-e2e/ThreadSanitizer/lit.local.cfg b/sycl/test-e2e/ThreadSanitizer/lit.local.cfg index daa768272e79c..fe03e06b8d89a 100644 --- a/sycl/test-e2e/ThreadSanitizer/lit.local.cfg +++ b/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",