Skip to content

Commit 78a557d

Browse files
zhaomaosuzhaomaosu
authored
[DevTSAN] Fix missing symbols __tsan[_unaligned]_[read|write]16 (#17924)
*Implement tsan builts __tsan_[read|write]16 and __tsan_unaligned_[read|write] in libdevice *Fix one minor issue
1 parent db0da3b commit 78a557d

File tree

6 files changed

+260
-6
lines changed

6 files changed

+260
-6
lines changed

‎libdevice/sanitizer/tsan_rtl.cpp

Lines changed: 85 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -58,6 +58,8 @@ inline constexpr uptr RoundDownTo(uptr x, uptr boundary) {
5858
return x & ~(boundary - 1);
5959
}
6060

61+
inline constexpr uptr Min(uptr a, uptr b) { return a < b ? a : b; }
62+
6163
inline void ConvertGenericPointer(uptr &addr, uint32_t &as) {
6264
auto old = addr;
6365
if ((addr = (uptr)ToPrivate((void *)old))) {
@@ -341,6 +343,89 @@ TSAN_CHECK(write, true, 2)
341343
TSAN_CHECK(write, true, 4)
342344
TSAN_CHECK(write, true, 8)
343345

346+
DEVICE_EXTERN_C_NOINLINE void
347+
__tsan_write16(uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,
348+
uint32_t line, const char __SYCL_CONSTANT__ *func) {
349+
__tsan_write8(addr, as, file, line, func);
350+
__tsan_write8(addr + 8, as, file, line, func);
351+
}
352+
353+
DEVICE_EXTERN_C_NOINLINE void
354+
__tsan_read16(uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,
355+
uint32_t line, const char __SYCL_CONSTANT__ *func) {
356+
__tsan_read8(addr, as, file, line, func);
357+
__tsan_read8(addr + 8, as, file, line, func);
358+
}
359+
360+
#define TSAN_UNALIGNED_CHECK(type, is_write, size) \
361+
DEVICE_EXTERN_C_NOINLINE void __tsan_unaligned_##type##size( \
362+
uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file, \
363+
uint32_t line, const char __SYCL_CONSTANT__ *func) { \
364+
__SYCL_GLOBAL__ RawShadow *shadow_mem = MemToShadow(addr, as); \
365+
if (!shadow_mem) \
366+
return; \
367+
Sid sid = GetCurrentSid(); \
368+
uint16_t current_clock = IncrementEpoch(sid) + 1; \
369+
AccessType type = is_write ? kAccessWrite : kAccessRead; \
370+
uptr size1 = Min(size, RoundUpTo(addr + 1, kShadowCell) - addr); \
371+
{ \
372+
TSAN_DEBUG(__spirv_ocl_printf( \
373+
__tsan_print_raw_shadow, (void *)addr, as, (void *)shadow_mem, \
374+
shadow_mem[0], shadow_mem[1], shadow_mem[2], shadow_mem[3])); \
375+
Shadow cur(addr, size1, current_clock, sid, type); \
376+
TSAN_DEBUG(__spirv_ocl_printf(__tsan_print_shadow_value, (void *)addr, \
377+
as, size1, cur.access(), cur.sid(), \
378+
cur.clock(), is_write)); \
379+
if (ContainsSameAccess(shadow_mem, cur, type)) \
380+
goto SECOND; \
381+
if (CheckRace(shadow_mem, cur, type, addr, size1, file, line, func)) \
382+
return; \
383+
} \
384+
SECOND: \
385+
uptr size2 = size - size1; \
386+
if (size2 == 0) \
387+
return; \
388+
shadow_mem += kShadowCnt; \
389+
{ \
390+
TSAN_DEBUG( \
391+
__spirv_ocl_printf(__tsan_print_raw_shadow, (void *)(addr + size1), \
392+
as, (void *)shadow_mem, shadow_mem[0], \
393+
shadow_mem[1], shadow_mem[2], shadow_mem[3])); \
394+
Shadow cur(0, size2, current_clock, sid, type); \
395+
TSAN_DEBUG(__spirv_ocl_printf( \
396+
__tsan_print_shadow_value, (void *)(addr + size1), as, size2, \
397+
cur.access(), cur.sid(), cur.clock(), is_write)); \
398+
if (ContainsSameAccess(shadow_mem, cur, type)) \
399+
return; \
400+
CheckRace(shadow_mem, cur, type, addr + size1, size2, file, line, func); \
401+
} \
402+
}
403+
404+
TSAN_UNALIGNED_CHECK(read, false, 1)
405+
TSAN_UNALIGNED_CHECK(read, false, 2)
406+
TSAN_UNALIGNED_CHECK(read, false, 4)
407+
TSAN_UNALIGNED_CHECK(read, false, 8)
408+
TSAN_UNALIGNED_CHECK(write, true, 1)
409+
TSAN_UNALIGNED_CHECK(write, true, 2)
410+
TSAN_UNALIGNED_CHECK(write, true, 4)
411+
TSAN_UNALIGNED_CHECK(write, true, 8)
412+
413+
DEVICE_EXTERN_C_NOINLINE void
414+
__tsan_unaligned_write16(uptr addr, uint32_t as,
415+
const char __SYCL_CONSTANT__ *file, uint32_t line,
416+
const char __SYCL_CONSTANT__ *func) {
417+
__tsan_unaligned_write8(addr, as, file, line, func);
418+
__tsan_unaligned_write8(addr + 8, as, file, line, func);
419+
}
420+
421+
DEVICE_EXTERN_C_NOINLINE void
422+
__tsan_unaligned_read16(uptr addr, uint32_t as,
423+
const char __SYCL_CONSTANT__ *file, uint32_t line,
424+
const char __SYCL_CONSTANT__ *func) {
425+
__tsan_unaligned_read8(addr, as, file, line, func);
426+
__tsan_unaligned_read8(addr + 8, as, file, line, func);
427+
}
428+
344429
DEVICE_EXTERN_C_NOINLINE void __tsan_cleanup_private(uptr addr, uint32_t size) {
345430
if (TsanLaunchInfo->DeviceTy != DeviceType::CPU)
346431
return;

‎llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp

Lines changed: 19 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -154,6 +154,8 @@ struct ThreadSanitizerOnSpirv {
154154
FunctionCallee TsanGroupBarrier;
155155
FunctionCallee TsanRead[kNumberOfAccessSizes];
156156
FunctionCallee TsanWrite[kNumberOfAccessSizes];
157+
FunctionCallee TsanUnalignedRead[kNumberOfAccessSizes];
158+
FunctionCallee TsanUnalignedWrite[kNumberOfAccessSizes];
157159

158160
friend struct ThreadSanitizer;
159161
};
@@ -299,6 +301,16 @@ void ThreadSanitizerOnSpirv::initialize() {
299301
TsanWrite[i] = M.getOrInsertFunction(WriteName, Attr, IRB.getVoidTy(),
300302
IntptrTy, IRB.getInt32Ty(), Int8PtrTy,
301303
IRB.getInt32Ty(), Int8PtrTy);
304+
305+
SmallString<32> UnalignedReadName("__tsan_unaligned_read" + ByteSizeStr);
306+
TsanUnalignedRead[i] = M.getOrInsertFunction(
307+
UnalignedReadName, Attr, IRB.getVoidTy(), IntptrTy, IRB.getInt32Ty(),
308+
Int8PtrTy, IRB.getInt32Ty(), Int8PtrTy);
309+
310+
SmallString<32> UnalignedWriteName("__tsan_unaligned_write" + ByteSizeStr);
311+
TsanUnalignedWrite[i] = M.getOrInsertFunction(
312+
UnalignedWriteName, Attr, IRB.getVoidTy(), IntptrTy, IRB.getInt32Ty(),
313+
Int8PtrTy, IRB.getInt32Ty(), Int8PtrTy);
302314
}
303315
}
304316

@@ -406,9 +418,9 @@ bool ThreadSanitizerOnSpirv::isUnsupportedDeviceGlobal(
406418
// TODO: Will support global variable with local address space later.
407419
if (G.getAddressSpace() == kSpirOffloadLocalAS)
408420
return true;
409-
// Global variables have constant value or constant address space will not
410-
// trigger race condition.
411-
if (G.isConstant() || G.getAddressSpace() == kSpirOffloadConstantAS)
421+
// Global variables have constant address space will not trigger race
422+
// condition.
423+
if (G.getAddressSpace() == kSpirOffloadConstantAS)
412424
return true;
413425
return false;
414426
}
@@ -691,7 +703,7 @@ static bool shouldInstrumentReadWriteFromAddress(const Module *M, Value *Addr) {
691703
if (Triple(M->getTargetTriple()).isSPIROrSPIRV()) {
692704
auto *OrigValue = getUnderlyingObject(Addr);
693705
if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
694-
return true;
706+
return false;
695707

696708
auto AddrAS = cast<PointerType>(Addr->getType()->getScalarType())
697709
->getPointerAddressSpace();
@@ -1016,6 +1028,9 @@ bool ThreadSanitizer::instrumentLoadOrStore(const InstructionInfo &II,
10161028
else if (IsVolatile)
10171029
OnAccessFunc = IsWrite ? TsanUnalignedVolatileWrite[Idx]
10181030
: TsanUnalignedVolatileRead[Idx];
1031+
else if (Spirv)
1032+
OnAccessFunc = IsWrite ? Spirv->TsanUnalignedWrite[Idx]
1033+
: Spirv->TsanUnalignedRead[Idx];
10191034
else
10201035
OnAccessFunc = IsWrite ? TsanUnalignedWrite[Idx] : TsanUnalignedRead[Idx];
10211036
}

‎llvm/test/Instrumentation/ThreadSanitizer/SPIRV/basic.ll

Lines changed: 103 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,16 @@ entry:
5050
ret void
5151
}
5252

53+
; Function Attrs: sanitize_thread
54+
define linkonce_odr dso_local spir_func void @write_16_bytes(ptr addrspace(4) %a) #0 {
55+
; CHECK-LABEL: void @write_16_bytes
56+
entry:
57+
store <4 x i32> <i32 0, i32 0, i32 0, i32 0>, ptr addrspace(4) %a, align 16
58+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
59+
; CHECK: call void @__tsan_write16
60+
ret void
61+
}
62+
5363
define linkonce_odr dso_local spir_func i8 @read_1_byte(ptr addrspace(4) %a) #0 {
5464
; CHECK-LABEL: i8 @read_1_byte
5565
entry:
@@ -86,4 +96,97 @@ entry:
8696
ret i64 %tmp1
8797
}
8898

99+
; Function Attrs: sanitize_thread
100+
define linkonce_odr dso_local spir_func void @read_16_bytes(ptr addrspace(4) %a) #0 {
101+
; CHECK-LABEL: void @read_16_bytes
102+
entry:
103+
%temp1 = load <4 x i32>, ptr addrspace(4) %a, align 16
104+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
105+
; CHECK: call void @__tsan_read16
106+
ret void
107+
}
108+
109+
; Function Attrs: sanitize_thread
110+
define linkonce_odr dso_local spir_func void @unaligned_write_2_bytes(ptr addrspace(4) %a) #0 {
111+
; CHECK-LABEL: void @unaligned_write_2_bytes
112+
entry:
113+
%tmp1 = load i16, ptr addrspace(4) %a, align 2
114+
%inc = add i16 %tmp1, 1
115+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
116+
; CHECK: call void @__tsan_unaligned_write2
117+
store i16 %inc, ptr addrspace(4) %a, align 1
118+
ret void
119+
}
120+
121+
; Function Attrs: sanitize_thread
122+
define linkonce_odr dso_local spir_func void @unaligned_write_4_bytes(ptr addrspace(4) %a) #0 {
123+
; CHECK-LABEL: void @unaligned_write_4_bytes
124+
entry:
125+
%tmp1 = load i32, ptr addrspace(4) %a, align 4
126+
%inc = add i32 %tmp1, 1
127+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
128+
; CHECK: call void @__tsan_unaligned_write4
129+
store i32 %inc, ptr addrspace(4) %a, align 1
130+
ret void
131+
}
132+
133+
; Function Attrs: sanitize_thread
134+
define linkonce_odr dso_local spir_func void @unaligned_write_8_bytes(ptr addrspace(4) %a) #0 {
135+
; CHECK-LABEL: void @unaligned_write_8_bytes
136+
entry:
137+
%tmp1 = load i64, ptr addrspace(4) %a, align 8
138+
%inc = add i64 %tmp1, 1
139+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
140+
; CHECK: call void @__tsan_unaligned_write8
141+
store i64 %inc, ptr addrspace(4) %a, align 1
142+
ret void
143+
}
144+
145+
; Function Attrs: sanitize_thread
146+
define linkonce_odr dso_local spir_func void @unaligned_write_16_bytes(ptr addrspace(4) %a) #0 {
147+
; CHECK-LABEL: void @unaligned_write_16_bytes
148+
entry:
149+
store <4 x i32> <i32 0, i32 0, i32 0, i32 0>, ptr addrspace(4) %a, align 1
150+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
151+
; CHECK: call void @__tsan_unaligned_write16
152+
ret void
153+
}
154+
155+
define linkonce_odr dso_local spir_func i16 @unaligned_read_2_bytes(ptr addrspace(4) %a) #0 {
156+
; CHECK-LABEL: i16 @unaligned_read_2_bytes
157+
entry:
158+
%tmp1 = load i16, ptr addrspace(4) %a, align 1
159+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
160+
; CHECK: call void @__tsan_unaligned_read2
161+
ret i16 %tmp1
162+
}
163+
164+
define linkonce_odr dso_local spir_func i32 @unaligned_read_4_bytes(ptr addrspace(4) %a) #0 {
165+
; CHECK-LABEL: i32 @unaligned_read_4_bytes
166+
entry:
167+
%tmp1 = load i32, ptr addrspace(4) %a, align 1
168+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
169+
; CHECK: call void @__tsan_unaligned_read4
170+
ret i32 %tmp1
171+
}
172+
173+
define linkonce_odr dso_local spir_func i64 @unaligned_read_8_bytes(ptr addrspace(4) %a) #0 {
174+
; CHECK-LABEL: i64 @unaligned_read_8_bytes
175+
entry:
176+
%tmp1 = load i64, ptr addrspace(4) %a, align 1
177+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
178+
; CHECK: call void @__tsan_unaligned_read8
179+
ret i64 %tmp1
180+
}
181+
182+
; Function Attrs: sanitize_thread
183+
define linkonce_odr dso_local spir_func void @unaligned_read_16_bytes(ptr addrspace(4) %a) #0 {
184+
; CHECK-LABEL: void @unaligned_read_16_bytes
185+
entry:
186+
%temp1 = load <4 x i32>, ptr addrspace(4) %a, align 1
187+
; CHECK: ptrtoint ptr addrspace(4) %a to i64
188+
; CHECK: call void @__tsan_unaligned_read16
189+
ret void
190+
}
191+
89192
attributes #0 = { sanitize_thread }

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

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4,12 +4,10 @@ target triple = "spir64-unknown-unknown"
44

55
@dev_global = external addrspace(1) global { [4 x i32] }
66
@dev_global_no_users = dso_local addrspace(1) global { [4 x i32] } zeroinitializer
7-
@.str = external addrspace(1) constant [59 x i8]
87
@__spirv_BuiltInGlobalInvocationId = external addrspace(1) constant <3 x i64>
98

109
; CHECK: @__TsanDeviceGlobalMetadata
1110
; CHECK-NOT: @dev_global_no_users
12-
; CHECK-NOT: @.str
1311
; CHECK-NOT: @__spirv_BuiltInGlobalInvocationId
1412
; CHECK-SAME: @dev_global
1513

‎sycl/test-e2e/ThreadSanitizer/check_access16.cpp

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
// REQUIRES: linux, cpu || (gpu && level_zero)
2+
// ALLOW_RETRIES: 10
3+
// RUN: %{build} %device_tsan_flags -O2 -g -o %t1.out
4+
// RUN: %{run} %t1.out 2>&1 | FileCheck %s
5+
#include "sycl/detail/core.hpp"
6+
#include "sycl/usm.hpp"
7+
#include "sycl/vector.hpp"
8+
9+
int main() {
10+
sycl::queue Q;
11+
auto *array = sycl::malloc_device<sycl::int3>(1, Q);
12+
13+
Q.submit([&](sycl::handler &h) {
14+
h.parallel_for<class Test>(sycl::nd_range<1>(32, 8),
15+
[=](sycl::nd_item<1>) {
16+
sycl::int3 vec1 = {1, 1, 1};
17+
sycl::int3 vec2 = {2, 2, 2};
18+
array[0] = vec1 / vec2;
19+
});
20+
}).wait();
21+
// CHECK: WARNING: DeviceSanitizer: data race
22+
// CHECK-NEXT: When write of size 8 at 0x{{.*}} in kernel <{{.*}}Test>
23+
// CHECK-NEXT: #0 {{.*}}check_access16.cpp:[[@LINE-5]]
24+
25+
return 0;
26+
}

‎sycl/test-e2e/ThreadSanitizer/check_unaligned_access.cpp

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
// REQUIRES: linux, cpu || (gpu && level_zero)
2+
// ALLOW_RETRIES: 10
3+
// RUN: %{build} %device_tsan_flags -O2 -g -o %t1.out
4+
// RUN: %{run} %t1.out 2>&1 | FileCheck %s
5+
#include "sycl/detail/core.hpp"
6+
#include "sycl/usm.hpp"
7+
8+
struct __attribute__((packed)) S {
9+
char c;
10+
int x;
11+
};
12+
13+
int main() {
14+
sycl::queue Q;
15+
auto *array = sycl::malloc_device<S>(1, Q);
16+
17+
Q.submit([&](sycl::handler &h) {
18+
h.parallel_for<class Test>(sycl::nd_range<1>(32, 8),
19+
[=](sycl::nd_item<1>) { array[0].x++; });
20+
}).wait();
21+
// CHECK: WARNING: DeviceSanitizer: data race
22+
// CHECK-NEXT: When write of size 4 at 0x{{.*}} in kernel <{{.*}}Test>
23+
// CHECK-NEXT: #0 {{.*}}check_unaligned_access.cpp:[[@LINE-4]]
24+
25+
sycl::free(array, Q);
26+
return 0;
27+
}

0 commit comments

Comments
 (0)