[DevTSAN] Fix missing symbols __tsan[_unaligned]_[read|write]16

libdevice/sanitizer/tsan_rtl.cpp

@@ -58,6 +58,8 @@ inline constexpr uptr RoundDownTo(uptr x, uptr boundary) {
   return x & ~(boundary - 1);
 }
 
+inline constexpr uptr Min(uptr a, uptr b) { return a < b ? a : b; }
+
 inline void ConvertGenericPointer(uptr &addr, uint32_t &as) {
   auto old = addr;
   if ((addr = (uptr)ToPrivate((void *)old))) {
@@ -341,6 +343,89 @@ TSAN_CHECK(write, true, 2)
 TSAN_CHECK(write, true, 4)
 TSAN_CHECK(write, true, 8)
 
+DEVICE_EXTERN_C_NOINLINE void
+__tsan_write16(uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,
+               uint32_t line, const char __SYCL_CONSTANT__ *func) {
+  __tsan_write8(addr, as, file, line, func);
+  __tsan_write8(addr + 8, as, file, line, func);
+}
+
+DEVICE_EXTERN_C_NOINLINE void
+__tsan_read16(uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,
+              uint32_t line, const char __SYCL_CONSTANT__ *func) {
+  __tsan_read8(addr, as, file, line, func);
+  __tsan_read8(addr + 8, as, file, line, func);
+}
+
+#define TSAN_UNALIGNED_CHECK(type, is_write, size)                             \
+  DEVICE_EXTERN_C_NOINLINE void __tsan_unaligned_##type##size(                 \
+      uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,              \
+      uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
+    __SYCL_GLOBAL__ RawShadow *shadow_mem = MemToShadow(addr, as);             \
+    if (!shadow_mem)                                                           \
+      return;                                                                  \
+    Sid sid = GetCurrentSid();                                                 \
+    uint16_t current_clock = IncrementEpoch(sid) + 1;                          \
+    AccessType type = is_write ? kAccessWrite : kAccessRead;                   \
+    uptr size1 = Min(size, RoundUpTo(addr + 1, kShadowCell) - addr);           \
+    {                                                                          \
+      TSAN_DEBUG(__spirv_ocl_printf(                                           \
+          __tsan_print_raw_shadow, (void *)addr, as, (void *)shadow_mem,       \
+          shadow_mem[0], shadow_mem[1], shadow_mem[2], shadow_mem[3]));        \
+      Shadow cur(addr, size1, current_clock, sid, type);                       \
+      TSAN_DEBUG(__spirv_ocl_printf(__tsan_print_shadow_value, (void *)addr,   \
+                                    as, size1, cur.access(), cur.sid(),        \
+                                    cur.clock(), is_write));                   \
+      if (ContainsSameAccess(shadow_mem, cur, type))                           \
+        goto SECOND;                                                           \
+      if (CheckRace(shadow_mem, cur, type, addr, size1, file, line, func))     \
+        return;                                                                \
+    }                                                                          \
+  SECOND:                                                                      \
+    uptr size2 = size - size1;                                                 \
+    if (size2 == 0)                                                            \
+      return;                                                                  \
+    shadow_mem += kShadowCnt;                                                  \
+    {                                                                          \
+      TSAN_DEBUG(                                                              \
+          __spirv_ocl_printf(__tsan_print_raw_shadow, (void *)(addr + size1),  \
+                             as, (void *)shadow_mem, shadow_mem[0],            \
+                             shadow_mem[1], shadow_mem[2], shadow_mem[3]));    \
+      Shadow cur(0, size2, current_clock, sid, type);                          \
+      TSAN_DEBUG(__spirv_ocl_printf(                                           \
+          __tsan_print_shadow_value, (void *)(addr + size1), as, size2,        \
+          cur.access(), cur.sid(), cur.clock(), is_write));                    \
+      if (ContainsSameAccess(shadow_mem, cur, type))                           \
+        return;                                                                \
+      CheckRace(shadow_mem, cur, type, addr + size1, size2, file, line, func); \
+    }                                                                          \
+  }
+
+TSAN_UNALIGNED_CHECK(read, false, 1)
+TSAN_UNALIGNED_CHECK(read, false, 2)
+TSAN_UNALIGNED_CHECK(read, false, 4)
+TSAN_UNALIGNED_CHECK(read, false, 8)
+TSAN_UNALIGNED_CHECK(write, true, 1)
+TSAN_UNALIGNED_CHECK(write, true, 2)
+TSAN_UNALIGNED_CHECK(write, true, 4)
+TSAN_UNALIGNED_CHECK(write, true, 8)
+
+DEVICE_EXTERN_C_NOINLINE void
+__tsan_unaligned_write16(uptr addr, uint32_t as,
+                         const char __SYCL_CONSTANT__ *file, uint32_t line,
+                         const char __SYCL_CONSTANT__ *func) {
+  __tsan_unaligned_write8(addr, as, file, line, func);
+  __tsan_unaligned_write8(addr + 8, as, file, line, func);
+}
+
+DEVICE_EXTERN_C_NOINLINE void
+__tsan_unaligned_read16(uptr addr, uint32_t as,
+                        const char __SYCL_CONSTANT__ *file, uint32_t line,
+                        const char __SYCL_CONSTANT__ *func) {
+  __tsan_unaligned_read8(addr, as, file, line, func);
+  __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;

llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp

@@ -154,6 +154,8 @@ struct ThreadSanitizerOnSpirv {
   FunctionCallee TsanGroupBarrier;
   FunctionCallee TsanRead[kNumberOfAccessSizes];
   FunctionCallee TsanWrite[kNumberOfAccessSizes];
+  FunctionCallee TsanUnalignedRead[kNumberOfAccessSizes];
+  FunctionCallee TsanUnalignedWrite[kNumberOfAccessSizes];
 
   friend struct ThreadSanitizer;
 };
@@ -299,6 +301,16 @@ void ThreadSanitizerOnSpirv::initialize() {
     TsanWrite[i] = M.getOrInsertFunction(WriteName, Attr, IRB.getVoidTy(),
                                          IntptrTy, IRB.getInt32Ty(), Int8PtrTy,
                                          IRB.getInt32Ty(), Int8PtrTy);
+
+    SmallString<32> UnalignedReadName("__tsan_unaligned_read" + ByteSizeStr);
+    TsanUnalignedRead[i] = M.getOrInsertFunction(
+        UnalignedReadName, Attr, IRB.getVoidTy(), IntptrTy, IRB.getInt32Ty(),
+        Int8PtrTy, IRB.getInt32Ty(), Int8PtrTy);
+
+    SmallString<32> UnalignedWriteName("__tsan_unaligned_write" + ByteSizeStr);
+    TsanUnalignedWrite[i] = M.getOrInsertFunction(
+        UnalignedWriteName, Attr, IRB.getVoidTy(), IntptrTy, IRB.getInt32Ty(),
+        Int8PtrTy, IRB.getInt32Ty(), Int8PtrTy);
   }
 }
 
@@ -406,9 +418,9 @@ bool ThreadSanitizerOnSpirv::isUnsupportedDeviceGlobal(
   // 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)
+  // Global variables have constant address space will not trigger race
+  // condition.
+  if (G.getAddressSpace() == kSpirOffloadConstantAS)
     return true;
   return false;
 }
@@ -691,7 +703,7 @@ static bool shouldInstrumentReadWriteFromAddress(const Module *M, Value *Addr) {
   if (Triple(M->getTargetTriple()).isSPIROrSPIRV()) {
     auto *OrigValue = getUnderlyingObject(Addr);
     if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
-      return true;
+      return false;
 
     auto AddrAS = cast<PointerType>(Addr->getType()->getScalarType())
                       ->getPointerAddressSpace();
@@ -1016,6 +1028,9 @@ bool ThreadSanitizer::instrumentLoadOrStore(const InstructionInfo &II,
     else if (IsVolatile)
       OnAccessFunc = IsWrite ? TsanUnalignedVolatileWrite[Idx]
                              : TsanUnalignedVolatileRead[Idx];
+    else if (Spirv)
+      OnAccessFunc = IsWrite ? Spirv->TsanUnalignedWrite[Idx]
+                             : Spirv->TsanUnalignedRead[Idx];
     else
       OnAccessFunc = IsWrite ? TsanUnalignedWrite[Idx] : TsanUnalignedRead[Idx];
   }

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

@@ -50,6 +50,16 @@ entry:
   ret void
 }
 
+; Function Attrs: sanitize_thread
+define linkonce_odr dso_local spir_func void @write_16_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: void @write_16_bytes
+entry:
+  store <4 x i32> <i32 0, i32 0, i32 0, i32 0>, ptr addrspace(4) %a, align 16
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_write16
+  ret void
+}
+
 define linkonce_odr dso_local spir_func i8 @read_1_byte(ptr addrspace(4) %a) #0 {
 ; CHECK-LABEL: i8 @read_1_byte
 entry:
@@ -86,4 +96,97 @@ entry:
   ret i64 %tmp1
 }
 
+; Function Attrs: sanitize_thread
+define linkonce_odr dso_local spir_func void @read_16_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: void @read_16_bytes
+entry:
+  %temp1 = load <4 x i32>, ptr addrspace(4) %a, align 16
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_read16
+  ret void
+}
+
+; Function Attrs: sanitize_thread
+define linkonce_odr dso_local spir_func void @unaligned_write_2_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: void @unaligned_write_2_bytes
+entry:
+  %tmp1 = load i16, ptr addrspace(4) %a, align 2
+  %inc = add i16 %tmp1, 1
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_unaligned_write2
+  store i16 %inc, ptr addrspace(4) %a, align 1
+  ret void
+}
+
+; Function Attrs: sanitize_thread
+define linkonce_odr dso_local spir_func void @unaligned_write_4_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: void @unaligned_write_4_bytes
+entry:
+  %tmp1 = load i32, ptr addrspace(4) %a, align 4
+  %inc = add i32 %tmp1, 1
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_unaligned_write4
+  store i32 %inc, ptr addrspace(4) %a, align 1
+  ret void
+}
+
+; Function Attrs: sanitize_thread
+define linkonce_odr dso_local spir_func void @unaligned_write_8_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: void @unaligned_write_8_bytes
+entry:
+  %tmp1 = load i64, ptr addrspace(4) %a, align 8
+  %inc = add i64 %tmp1, 1
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_unaligned_write8
+  store i64 %inc, ptr addrspace(4) %a, align 1
+  ret void
+}
+
+; Function Attrs: sanitize_thread
+define linkonce_odr dso_local spir_func void @unaligned_write_16_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: void @unaligned_write_16_bytes
+entry:
+  store <4 x i32> <i32 0, i32 0, i32 0, i32 0>, ptr addrspace(4) %a, align 1
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_unaligned_write16
+  ret void
+}
+
+define linkonce_odr dso_local spir_func i16 @unaligned_read_2_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: i16 @unaligned_read_2_bytes
+entry:
+  %tmp1 = load i16, ptr addrspace(4) %a, align 1
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_unaligned_read2
+  ret i16 %tmp1
+}
+
+define linkonce_odr dso_local spir_func i32 @unaligned_read_4_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: i32 @unaligned_read_4_bytes
+entry:
+  %tmp1 = load i32, ptr addrspace(4) %a, align 1
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_unaligned_read4
+  ret i32 %tmp1
+}
+
+define linkonce_odr dso_local spir_func i64 @unaligned_read_8_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: i64 @unaligned_read_8_bytes
+entry:
+  %tmp1 = load i64, ptr addrspace(4) %a, align 1
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_unaligned_read8
+  ret i64 %tmp1
+}
+
+; Function Attrs: sanitize_thread
+define linkonce_odr dso_local spir_func void @unaligned_read_16_bytes(ptr addrspace(4) %a) #0 {
+; CHECK-LABEL: void @unaligned_read_16_bytes
+entry:
+  %temp1 = load <4 x i32>, ptr addrspace(4) %a, align 1
+  ; CHECK: ptrtoint ptr addrspace(4) %a to i64
+  ; CHECK: call void @__tsan_unaligned_read16
+  ret void
+}
+
 attributes #0 = { sanitize_thread }

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

@@ -4,12 +4,10 @@ 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
 

sycl/test-e2e/ThreadSanitizer/check_access16.cpp

@@ -0,0 +1,26 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// ALLOW_RETRIES: 10
+// RUN: %{build} %device_tsan_flags -O2 -g -o %t1.out
+// RUN: %{run} %t1.out 2>&1 | FileCheck %s
+#include "sycl/detail/core.hpp"
+#include "sycl/usm.hpp"
+#include "sycl/vector.hpp"
+
+int main() {
+  sycl::queue Q;
+  auto *array = sycl::malloc_device<sycl::int3>(1, Q);
+
+  Q.submit([&](sycl::handler &h) {
+     h.parallel_for<class Test>(sycl::nd_range<1>(32, 8),
+                                [=](sycl::nd_item<1>) {
+                                  sycl::int3 vec1 = {1, 1, 1};
+                                  sycl::int3 vec2 = {2, 2, 2};
+                                  array[0] = vec1 / vec2;
+                                });
+   }).wait();
+  // CHECK: WARNING: DeviceSanitizer: data race
+  // CHECK-NEXT: When write of size 8 at 0x{{.*}} in kernel <{{.*}}Test>
+  // CHECK-NEXT: #0 {{.*}}check_access16.cpp:[[@LINE-5]]
+
+  return 0;
+}

sycl/test-e2e/ThreadSanitizer/check_unaligned_access.cpp

@@ -0,0 +1,27 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// ALLOW_RETRIES: 10
+// RUN: %{build} %device_tsan_flags -O2 -g -o %t1.out
+// RUN: %{run} %t1.out 2>&1 | FileCheck %s
+#include "sycl/detail/core.hpp"
+#include "sycl/usm.hpp"
+
+struct __attribute__((packed)) S {
+  char c;
+  int x;
+};
+
+int main() {
+  sycl::queue Q;
+  auto *array = sycl::malloc_device<S>(1, Q);
+
+  Q.submit([&](sycl::handler &h) {
+     h.parallel_for<class Test>(sycl::nd_range<1>(32, 8),
+                                [=](sycl::nd_item<1>) { array[0].x++; });
+   }).wait();
+  // CHECK: WARNING: DeviceSanitizer: data race
+  // CHECK-NEXT: When write of size 4 at 0x{{.*}} in kernel <{{.*}}Test>
+  // CHECK-NEXT: #0 {{.*}}check_unaligned_access.cpp:[[@LINE-4]]
+
+  sycl::free(array, Q);
+  return 0;
+}