From 417fcb11f88b79fe64f092abd4ff8927e2368995 Mon Sep 17 00:00:00 2001
From: "Zhao, Yang2" <yang2.zhao@intel.com>
Date: Mon, 23 Sep 2024 08:48:04 +0200
Subject: [PATCH 1/3] fix wrong test

---
 .../AddressSanitizer/nullpointer/global_nullptr.cpp      | 9 ++-------
 1 file changed, 2 insertions(+), 7 deletions(-)

diff --git a/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
index b280662c9548f..260da1f6fb524 100644
--- a/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
+++ b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
@@ -15,17 +15,12 @@
 
 int main() {
   sycl::queue Q;
-  constexpr std::size_t N = 4;
+  constexpr std::size_t N = 2;
   int *array = nullptr;
 
   Q.submit([&](sycl::handler &h) {
     h.parallel_for<class MyKernel>(
-        sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) {
-          auto private_array =
-              sycl::ext::oneapi::experimental::static_address_cast<
-                  sycl::access::address_space::private_space>(array);
-          private_array[0] = 0;
-        });
+        sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) { array[0] = 0; });
     Q.wait();
   });
   // CHECK: ERROR: DeviceSanitizer: null-pointer-access on Unknown Memory

From 6ba0490e65ab682c5eb2d5ed8b6ae8f8c5ac2992 Mon Sep 17 00:00:00 2001
From: "Zhao, Yang2" <yang2.zhao@intel.com>
Date: Mon, 23 Sep 2024 08:49:06 +0200
Subject: [PATCH 2/3] fix

---
 sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
index 260da1f6fb524..66e16bdc011c6 100644
--- a/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
+++ b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
@@ -15,7 +15,7 @@
 
 int main() {
   sycl::queue Q;
-  constexpr std::size_t N = 2;
+  constexpr std::size_t N = 4;
   int *array = nullptr;
 
   Q.submit([&](sycl::handler &h) {

From 5fd4ac170469268853e6053d860b0c7db05cf1fe Mon Sep 17 00:00:00 2001
From: "Zhao, Yang2" <yang2.zhao@intel.com>
Date: Mon, 23 Sep 2024 09:39:13 +0200
Subject: [PATCH 3/3] fix O0 case

---
 libdevice/sanitizer_utils.cpp                 | 20 +++++++------------
 .../nullpointer/global_nullptr.cpp            |  2 --
 2 files changed, 7 insertions(+), 15 deletions(-)

diff --git a/libdevice/sanitizer_utils.cpp b/libdevice/sanitizer_utils.cpp
index 1074b6ac7e3cd..2e70e14b8168a 100644
--- a/libdevice/sanitizer_utils.cpp
+++ b/libdevice/sanitizer_utils.cpp
@@ -120,22 +120,20 @@ __SYCL_PRIVATE__ void *ToPrivate(void *ptr) {
   return __spirv_GenericCastToPtrExplicit_ToPrivate(ptr, 7);
 }
 
-inline bool ConvertGenericPointer(uptr &addr, uint32_t &as) {
+inline void ConvertGenericPointer(uptr &addr, uint32_t &as) {
   auto old = addr;
   if ((addr = (uptr)ToPrivate((void *)old))) {
     as = ADDRESS_SPACE_PRIVATE;
   } else if ((addr = (uptr)ToLocal((void *)old))) {
     as = ADDRESS_SPACE_LOCAL;
-  } else if ((addr = (uptr)ToGlobal((void *)old))) {
-    as = ADDRESS_SPACE_GLOBAL;
   } else {
-    if (__AsanDebug)
-      __spirv_ocl_printf(__generic_to_fail, old);
-    return false;
+    // FIXME: I'm not sure if we need to check ADDRESS_SPACE_CONSTANT,
+    // but this can really simplify the generic pointer conversion logic
+    as = ADDRESS_SPACE_GLOBAL;
+    addr = old;
   }
   if (__AsanDebug)
     __spirv_ocl_printf(__generic_to, old, addr, as);
-  return true;
 }
 
 inline uptr MemToShadow_CPU(uptr addr) {
@@ -144,9 +142,7 @@ inline uptr MemToShadow_CPU(uptr addr) {
 
 inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
   if (as == ADDRESS_SPACE_GENERIC) {
-    if (!ConvertGenericPointer(addr, as)) {
-      return 0;
-    }
+    ConvertGenericPointer(addr, as);
   }
 
   if (as == ADDRESS_SPACE_GLOBAL) {     // global
@@ -233,9 +229,7 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
 
 inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
   if (as == ADDRESS_SPACE_GENERIC) {
-    if (!ConvertGenericPointer(addr, as)) {
-      return 0;
-    }
+    ConvertGenericPointer(addr, as);
   }
 
   if (as == ADDRESS_SPACE_GLOBAL) { // global
diff --git a/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
index 66e16bdc011c6..6d436b58bfcee 100644
--- a/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
+++ b/sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
@@ -11,8 +11,6 @@
 
 #include <sycl/detail/core.hpp>
 
-#include <sycl/ext/oneapi/experimental/address_cast.hpp>
-
 int main() {
   sycl::queue Q;
   constexpr std::size_t N = 4;