[DeviceSanitizer] Fix wrong test for nullptr detection (#15467)

The content of "global_nullptr.cpp" was same with "private_nullptr.cpp".
This patch fix this.

Author: AllanZyne

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

sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp

@@ -11,21 +11,14 @@
 
 #include <sycl/detail/core.hpp>
 
-#include <sycl/ext/oneapi/experimental/address_cast.hpp>
-
 int main() {
   sycl::queue Q;
   constexpr std::size_t N = 4;
   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