diff --git a/libdevice/sanitizer/tsan_rtl.cpp b/libdevice/sanitizer/tsan_rtl.cpp
index 6c7b6f26bb9bd..fa467b70901bf 100644
--- a/libdevice/sanitizer/tsan_rtl.cpp
+++ b/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() {
diff --git a/sycl/test-e2e/ThreadSanitizer/check_access16.cpp b/sycl/test-e2e/ThreadSanitizer/check_access16.cpp
index cd7940f700e4d..f09ba66dbd36c 100644
--- a/sycl/test-e2e/ThreadSanitizer/check_access16.cpp
+++ b/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};
diff --git a/sycl/test-e2e/ThreadSanitizer/check_buffer.cpp b/sycl/test-e2e/ThreadSanitizer/check_buffer.cpp
index 86b96f9b60fdd..7f1c32f20dea7 100644
--- a/sycl/test-e2e/ThreadSanitizer/check_buffer.cpp
+++ b/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;
diff --git a/sycl/test-e2e/ThreadSanitizer/check_device_global.cpp b/sycl/test-e2e/ThreadSanitizer/check_device_global.cpp
index d0cc151424ac6..288832f98589c 100644
--- a/sycl/test-e2e/ThreadSanitizer/check_device_global.cpp
+++ b/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
diff --git a/sycl/test-e2e/ThreadSanitizer/check_device_usm.cpp b/sycl/test-e2e/ThreadSanitizer/check_device_usm.cpp
new file mode 100644
index 0000000000000..4134ae2ee7f27
--- /dev/null
+++ b/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;
+}
diff --git a/sycl/test-e2e/ThreadSanitizer/check_host_usm.cpp b/sycl/test-e2e/ThreadSanitizer/check_host_usm.cpp
new file mode 100644
index 0000000000000..668ee9e3befc3
--- /dev/null
+++ b/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;
+}
diff --git a/sycl/test-e2e/ThreadSanitizer/check_shared_usm.cpp b/sycl/test-e2e/ThreadSanitizer/check_shared_usm.cpp
new file mode 100644
index 0000000000000..39745dbf08f0f
--- /dev/null
+++ b/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;
+}
diff --git a/sycl/test-e2e/ThreadSanitizer/check_unaligned_access.cpp b/sycl/test-e2e/ThreadSanitizer/check_unaligned_access.cpp
index b1cd5572b722d..4101e4189f643 100644
--- a/sycl/test-e2e/ThreadSanitizer/check_unaligned_access.cpp
+++ b/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
diff --git a/sycl/test-e2e/ThreadSanitizer/check_usm.cpp b/sycl/test-e2e/ThreadSanitizer/check_usm.cpp
deleted file mode 100644
index 456844b57af55..0000000000000
--- a/sycl/test-e2e/ThreadSanitizer/check_usm.cpp
+++ /dev/null
@@ -1,34 +0,0 @@
-// REQUIRES: linux, cpu || (gpu && level_zero)
-// ALLOW_RETRIES: 10
-// RUN: %{build} %device_tsan_flags -DMALLOC_DEVICE -O0 -g -o %t1.out
-// RUN: %{run} %t1.out 2>&1 | FileCheck %s
-// RUN: %{build} %device_tsan_flags -DMALLOC_DEVICE -O2 -g -o %t2.out
-// RUN: %{run} %t2.out 2>&1 | FileCheck %s
-// RUN: %{build} %device_tsan_flags -DMALLOC_HOST -O2 -g -o %t3.out
-// RUN: %{run} %t3.out 2>&1 | FileCheck %s
-// RUN: %{build} %device_tsan_flags -DMALLOC_SHARED -O2 -g -o %t4.out
-// RUN: %{run} %t4.out 2>&1 | FileCheck %s
-#include "sycl/detail/core.hpp"
-#include "sycl/usm.hpp"
-
-int main() {
-  sycl::queue Q;
-#if defined(MALLOC_DEVICE)
-  auto *array = sycl::malloc_device<char>(1, Q);
-#elif defined(MALLOC_HOST)
-  auto *array = sycl::malloc_host<char>(1, Q);
-#else // defined(MALLOC_SHARED)
-  auto *array = sycl::malloc_shared<char>(1, Q);
-#endif
-
-  Q.submit([&](sycl::handler &h) {
-     h.parallel_for<class Test>(sycl::nd_range<1>(32, 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_usm.cpp:[[@LINE-4]]
-
-  sycl::free(array, Q);
-  return 0;
-}