[DevTSAN] Treat each work item as a thread for GPU device

libdevice/sanitizer/tsan_rtl.cpp

@@ -76,7 +76,7 @@ inline void ConvertGenericPointer(uptr &addr, uint32_t &as) {
 }
 
 inline Epoch IncrementEpoch(Sid sid) {
-  return atomicAdd(&TsanLaunchInfo->Clock[sid].clk_[sid], 1);
+  return TsanLaunchInfo->Clock[sid].clk_[sid]++;
 }
 
 inline __SYCL_GLOBAL__ RawShadow *MemToShadow_CPU(uptr addr, uint32_t) {
@@ -145,12 +145,11 @@ inline Sid GetCurrentSid_CPU() {
   return wg_lid;
 }
 
-// For GPU device, each sub group is a thread
+// For GPU device, each work item is a thread
 inline Sid GetCurrentSid_GPU() {
   // sub-group linear id
-  const auto sg_lid =
-      __spirv_BuiltInGlobalLinearId / __spirv_BuiltInSubgroupSize;
-  return sg_lid;
+  const auto lid = __spirv_BuiltInGlobalLinearId;
+  return lid;
 }
 
 inline Sid GetCurrentSid() {

sycl/test-e2e/ThreadSanitizer/check_access16.cpp

@@ -11,7 +11,7 @@ int main() {
   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),
+     h.parallel_for<class Test>(sycl::nd_range<1>(128, 8),
                                 [=](sycl::nd_item<1>) {
                                   sycl::int3 vec1 = {1, 1, 1};
                                   sycl::int3 vec2 = {2, 2, 2};

sycl/test-e2e/ThreadSanitizer/check_buffer.cpp

@@ -7,7 +7,7 @@
 
 #include <sycl/detail/core.hpp>
 
-static const int N = 16;
+static const int N = 128;
 
 int main() {
   sycl::queue q;

sycl/test-e2e/ThreadSanitizer/check_device_global.cpp

@@ -19,7 +19,7 @@ int main() {
   sycl::queue Q;
 
   Q.submit([&](sycl::handler &h) {
-     h.parallel_for<class Test>(sycl::nd_range<1>(32, 8),
+     h.parallel_for<class Test>(sycl::nd_range<1>(128, 8),
                                 [=](sycl::nd_item<1>) { dev_global[0]++; });
    }).wait();
   // CHECK: WARNING: DeviceSanitizer: data race

sycl/test-e2e/ThreadSanitizer/check_device_usm.cpp

@@ -0,0 +1,22 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// ALLOW_RETRIES: 10
+// RUN: %{build} %device_tsan_flags -O0 -g -o %t.out
+// RUN: %{run} %t.out 2>&1 | FileCheck %s
+#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>(128, 8),
+                                [=](sycl::nd_item<1>) { array[0]++; });
+   }).wait();
+  // CHECK: WARNING: DeviceSanitizer: data race
+  // CHECK-NEXT: When write of size 1 at 0x{{.*}} in kernel <{{.*}}Test>
+  // CHECK-NEXT: #0 {{.*}}check_device_usm.cpp:[[@LINE-4]]
+
+  sycl::free(array, Q);
+  return 0;
+}

sycl/test-e2e/ThreadSanitizer/check_host_usm.cpp

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

sycl/test-e2e/ThreadSanitizer/check_shared_usm.cpp

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

sycl/test-e2e/ThreadSanitizer/check_unaligned_access.cpp

@@ -15,7 +15,7 @@ int main() {
   auto *array = sycl::malloc_device<S>(1, Q);
 
   Q.submit([&](sycl::handler &h) {
-     h.parallel_for<class Test>(sycl::nd_range<1>(32, 8),
+     h.parallel_for<class Test>(sycl::nd_range<1>(128, 8),
                                 [=](sycl::nd_item<1>) { array[0].x++; });
    }).wait();
   // CHECK: WARNING: DeviceSanitizer: data race