Skip to content

Commit 6249f30

Browse files
jinge90jinge90
authored
[SYCL][TSAN] AOT support for device thread sanitizer (#18130)
Signed-off-by: jinge90 <[email protected]>
1 parent ccb2df5 commit 6249f30

File tree

11 files changed

+150
-10
lines changed

11 files changed

+150
-10
lines changed

‎clang/lib/Driver/ToolChains/SYCL.cpp‎

Lines changed: 12 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -603,7 +603,15 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
603603
// For DG2, we just use libsycl-msan as placeholder.
604604
{"libsycl-msan", "internal"},
605605
{"libsycl-msan-pvc", "internal"}};
606-
const SYCLDeviceLibsList SYCLDeviceTsanLibs = {{"libsycl-tsan", "internal"}};
606+
const SYCLDeviceLibsList SYCLDeviceTsanLibs = {
607+
{"libsycl-tsan", "internal"},
608+
{"libsycl-tsan-cpu", "internal"},
609+
// Currently, we only provide aot tsan libdevice for PVC and CPU.
610+
// For DG2, we just use libsycl-tsan as placeholder.
611+
// TODO: replace "libsycl-tsan" with "libsycl-tsan-dg2" when DG2
612+
// AOT support is added.
613+
{"libsycl-tsan", "internal"},
614+
{"libsycl-tsan-pvc", "internal"}};
607615
#endif
608616

609617
const SYCLDeviceLibsList SYCLNativeCpuDeviceLibs = {
@@ -759,7 +767,7 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
759767
else if (SanitizeVal == "memory")
760768
addSingleLibrary(SYCLDeviceMsanLibs[sanitizer_lib_idx]);
761769
else if (SanitizeVal == "thread")
762-
addLibraries(SYCLDeviceTsanLibs);
770+
addSingleLibrary(SYCLDeviceTsanLibs[sanitizer_lib_idx]);
763771
#endif
764772

765773
if (isSYCLNativeCPU(TargetTriple))
@@ -883,6 +891,8 @@ static llvm::SmallVector<StringRef, 16> SYCLDeviceLibList{
883891
"msan-pvc",
884892
"msan-cpu",
885893
"tsan",
894+
"tsan-pvc",
895+
"tsan-cpu",
886896
#endif
887897
"imf",
888898
"imf-fp64",

‎clang/test/Driver/Inputs/SYCL/lib/libsycl-tsan-cpu.bc‎

Whitespace-only changes.

‎clang/test/Driver/Inputs/SYCL/lib/libsycl-tsan-pvc.bc‎

Whitespace-only changes.

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

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -415,3 +415,26 @@
415415
// SYCL_DEVICE_TSAN_MACRO-SAME: "USE_SYCL_DEVICE_TSAN"
416416
// SYCL_DEVICE_TSAN_MACRO: llvm-link{{.*}} "-only-needed"
417417
// SYCL_DEVICE_TSAN_MACRO-SAME: "{{.*}}libsycl-tsan.bc"
418+
419+
/// ###########################################################################
420+
/// test behavior of linking libsycl-tsan-pvc for PVC target AOT compilation when tsan flag is applied.
421+
// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_pvc --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL \
422+
// RUN: -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
423+
// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device pvc" --no-offload-new-driver %s \
424+
// RUN: --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
425+
// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device pvc" --no-offload-new-driver %s \
426+
// RUN: --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
427+
// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device 12.60.7" --no-offload-new-driver %s \
428+
// RUN: --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
429+
// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xs "-device 12.60.7" --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL \
430+
// RUN: -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
431+
// SYCL_DEVICE_LIB_TSAN_PVC: llvm-link{{.*}} "-only-needed" "{{.*}}libsycl-crt.bc"
432+
// SYCL_DEVICE_LIB_TSAN_PVC-SAME: "{{.*}}libsycl-tsan-pvc.bc"
433+
434+
435+
/// ###########################################################################
436+
/// test behavior of linking libsycl-tsan-cpu for CPU target AOT compilation when tsan flag is applied.
437+
// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL \
438+
// RUN: -Xarch_device -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_CPU
439+
// SYCL_DEVICE_LIB_TSAN_CPU: llvm-link{{.*}} "-only-needed" "{{.*}}libsycl-crt.bc"
440+
// SYCL_DEVICE_LIB_TSAN_CPU-SAME: "{{.*}}libsycl-tsan-cpu.bc"

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

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -401,3 +401,15 @@
401401
// SYCL_DEVICE_TSAN_MACRO: "-cc1"
402402
// SYCL_DEVICE_TSAN_MACRO-SAME: "USE_SYCL_DEVICE_TSAN"
403403
// SYCL_DEVICE_TSAN_MACRO: libsycl-tsan.new.o
404+
405+
/// test behavior of tsan libdevice linking when -fsanitize=thread is available for AOT targets
406+
// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_pvc --offload-new-driver %s --sysroot=%S/Inputs/SYCL \
407+
// RUN: -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_PVC
408+
// SYCL_DEVICE_LIB_TSAN_PVC: clang-linker-wrapper{{.*}} "-sycl-device-libraries
409+
// SYCL_DEVICE_LIB_TSAN_PVC-SAME: {{.*}}libsycl-tsan-pvc.new.o
410+
411+
/// test behavior of tsan libdevice linking when -fsanitize=thread is available for AOT targets
412+
// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 --offload-new-driver %s --sysroot=%S/Inputs/SYCL \
413+
// RUN: -fsanitize=thread -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_TSAN_CPU
414+
// SYCL_DEVICE_LIB_TSAN_CPU: clang-linker-wrapper{{.*}} "-sycl-device-libraries
415+
// SYCL_DEVICE_LIB_TSAN_CPU-SAME: {{.*}}libsycl-tsan-cpu.new.o

‎libdevice/cmake/modules/SYCLLibdevice.cmake‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -433,6 +433,19 @@ else()
433433
EXTRA_OPTS -fno-sycl-instrument-device-code
434434
-I${UR_SANITIZER_INCLUDE_DIR}
435435
-I${CMAKE_CURRENT_SOURCE_DIR})
436+
437+
set(tsan_devicetypes pvc cpu)
438+
439+
foreach(tsan_ft IN LISTS sanitizer_filetypes)
440+
foreach(tsan_device IN LISTS tsan_devicetypes)
441+
compile_lib_ext(libsycl-tsan-${tsan_device}
442+
SRC sanitizer/tsan_rtl.cpp
443+
FILETYPE ${tsan_ft}
444+
DEPENDENCIES ${tsan_obj_deps}
445+
OPTS ${sanitizer_${tsan_device}_compile_opts_${tsan_ft}})
446+
endforeach()
447+
endforeach()
448+
436449
endif()
437450
endif()
438451

‎libdevice/sanitizer/tsan_rtl.cpp‎

Lines changed: 39 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -115,6 +115,11 @@ inline __SYCL_GLOBAL__ RawShadow *MemToShadow_PVC(uptr addr, uint32_t as) {
115115
inline __SYCL_GLOBAL__ RawShadow *MemToShadow(uptr addr, uint32_t as) {
116116
__SYCL_GLOBAL__ RawShadow *shadow_ptr = nullptr;
117117

118+
#if defined(__LIBDEVICE_CPU__)
119+
shadow_ptr = MemToShadow_CPU(addr, as);
120+
#elif defined(__LIBDEVICE_PVC__)
121+
shadow_ptr = MemToShadow_PVC(addr, as);
122+
#else
118123
if (TsanLaunchInfo->DeviceTy == DeviceType::CPU) {
119124
shadow_ptr = MemToShadow_CPU(addr, as);
120125
} else if (TsanLaunchInfo->DeviceTy == DeviceType::GPU_PVC) {
@@ -124,6 +129,7 @@ inline __SYCL_GLOBAL__ RawShadow *MemToShadow(uptr addr, uint32_t as) {
124129
(int)TsanLaunchInfo->DeviceTy));
125130
return nullptr;
126131
}
132+
#endif
127133

128134
return shadow_ptr;
129135
}
@@ -148,6 +154,11 @@ inline Sid GetCurrentSid_GPU() {
148154
}
149155

150156
inline Sid GetCurrentSid() {
157+
#if defined(__LIBDEVICE_CPU__)
158+
return GetCurrentSid_CPU();
159+
#elif defined(__LIBDEVICE_PVC__)
160+
return GetCurrentSid_GPU();
161+
#else
151162
if (TsanLaunchInfo->DeviceTy == DeviceType::CPU) {
152163
return GetCurrentSid_CPU();
153164
} else if (TsanLaunchInfo->DeviceTy != DeviceType::UNKNOWN) {
@@ -157,6 +168,7 @@ inline Sid GetCurrentSid() {
157168
(int)TsanLaunchInfo->DeviceTy));
158169
return 0;
159170
}
171+
#endif
160172
}
161173

162174
inline RawShadow LoadShadow(const __SYCL_GLOBAL__ RawShadow *p) {
@@ -426,10 +438,7 @@ __tsan_unaligned_read16(uptr addr, uint32_t as,
426438
__tsan_unaligned_read8(addr + 8, as, file, line, func);
427439
}
428440

429-
DEVICE_EXTERN_C_NOINLINE void __tsan_cleanup_private(uptr addr, uint32_t size) {
430-
if (TsanLaunchInfo->DeviceTy != DeviceType::CPU)
431-
return;
432-
441+
static inline void __tsan_cleanup_private_cpu_impl(uptr addr, uint32_t size) {
433442
if (size) {
434443
addr = RoundDownTo(addr, kShadowCell);
435444
size = RoundUpTo(size, kShadowCell);
@@ -443,6 +452,19 @@ DEVICE_EXTERN_C_NOINLINE void __tsan_cleanup_private(uptr addr, uint32_t size) {
443452
}
444453
}
445454

455+
DEVICE_EXTERN_C_NOINLINE void __tsan_cleanup_private(uptr addr, uint32_t size) {
456+
#if defined(__LIBDEVICE_CPU__)
457+
__tsan_cleanup_private_cpu_impl(addr, size);
458+
#elif defined(__LIBDEVICE_PVC__)
459+
return;
460+
#else
461+
if (TsanLaunchInfo->DeviceTy != DeviceType::CPU)
462+
return;
463+
464+
__tsan_cleanup_private_cpu_impl(addr, size);
465+
#endif
466+
}
467+
446468
DEVICE_EXTERN_C_INLINE void __tsan_device_barrier() {
447469
Sid sid = GetCurrentSid();
448470
__spirv_ControlBarrier(__spv::Scope::Device, __spv::Scope::Device,
@@ -470,10 +492,7 @@ DEVICE_EXTERN_C_INLINE void __tsan_device_barrier() {
470492
__spv::MemorySemanticsMask::WorkgroupMemory);
471493
}
472494

473-
DEVICE_EXTERN_C_INLINE void __tsan_group_barrier() {
474-
if (TsanLaunchInfo->DeviceTy == DeviceType::CPU)
475-
return;
476-
495+
static inline void __tsan_group_barrier_impl() {
477496
Sid sid = GetCurrentSid();
478497
__spirv_ControlBarrier(__spv::Scope::Workgroup, __spv::Scope::Workgroup,
479498
__spv::MemorySemanticsMask::SequentiallyConsistent |
@@ -500,4 +519,16 @@ DEVICE_EXTERN_C_INLINE void __tsan_group_barrier() {
500519
__spv::MemorySemanticsMask::WorkgroupMemory);
501520
}
502521

522+
DEVICE_EXTERN_C_INLINE void __tsan_group_barrier() {
523+
#if defined(__LIBDEVICE_CPU__)
524+
return;
525+
#elif defined(__LIBDEVICE_PVC__)
526+
__tsan_group_barrier_impl();
527+
#else
528+
if (TsanLaunchInfo->DeviceTy == DeviceType::CPU)
529+
return;
530+
__tsan_group_barrier_impl();
531+
#endif
532+
}
533+
503534
#endif // __SPIR__ || __SPIRV__

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

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
#include "sycl/detail/core.hpp"
2+
#include "sycl/usm.hpp"
3+
4+
int main() {
5+
sycl::queue Q;
6+
auto *array = sycl::malloc_device<char>(1, Q);
7+
Q.submit([&](sycl::handler &h) {
8+
h.parallel_for<class Test>(sycl::nd_range<1>(32, 8),
9+
[=](sycl::nd_item<1>) { array[0]++; });
10+
}).wait();
11+
// CHECK: DeviceSanitizer: data race
12+
// CHECK-NEXT: When write of size 1 at 0x{{.*}} in kernel <{{.*}}Test>
13+
// CHECK-NEXT: #0 {{.*}}usm_data_race.cpp:[[@LINE-4]]
14+
15+
sycl::free(array, Q);
16+
return 0;
17+
}

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

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
// REQUIRES: linux, opencl-aot, cpu
2+
// ALLOW_RETRIES: 10
3+
4+
// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O0 -g %S/Inputs/usm_data_race.cpp -o %t.out
5+
// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
6+
7+
// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O1 -g %S/Inputs/usm_data_race.cpp -o %t.out
8+
// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
9+
10+
// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O2 -g %S/Inputs/usm_data_race.cpp -o %t.out
11+
// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
12+
13+
// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O3 -g %S/Inputs/usm_data_race.cpp -o %t.out
14+
// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp

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

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,15 @@
1+
// REQUIRES: linux, gpu && level_zero
2+
// REQUIRES: arch-intel_gpu_pvc
3+
// ALLOW_RETRIES: 10
4+
5+
// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O0 -g %S/Inputs/usm_data_race.cpp -o %t.out
6+
// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
7+
8+
// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O1 -g %S/Inputs/usm_data_race.cpp -o %t.out
9+
// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
10+
11+
// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O2 -g %S/Inputs/usm_data_race.cpp -o %t.out
12+
// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp
13+
14+
// RUN: %{run-aux} %{build} %device_tsan_aot_flags -O3 -g %S/Inputs/usm_data_race.cpp -o %t.out
15+
// RUN: %{run} %t.out 2>&1 | FileCheck %S/Inputs/usm_data_race.cpp

0 commit comments

Comments
 (0)