[SYCL] Fix multi_ptr relational operators customised for nullptr. (#13201)

`multi_ptr` relational operators taking a `std::nullptr_t` are written
in a way that assume it is the lowest possible value
(example
https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/multi_ptr.hpp#L1575).
However the C++ specs states there is no ordering requirement
(https://eel.is/c++draft/expr.rel#4.3).

In practice, this is causing issues in the CUDA and AMDGPU backend. For
instance, in CUDA the nullptr in the local address space is a non `0`
value and `0` in this address space is the root of the allocated local
memory.

Author: mmoadeli

sycl/include/sycl/multi_ptr.hpp

@@ -1560,56 +1560,64 @@ template <typename ElementType, access::address_space Space,
           access::decorated DecorateAddress>
 bool operator>(const multi_ptr<ElementType, Space, DecorateAddress> &lhs,
                std::nullptr_t) {
-  return lhs.get() != nullptr;
+  return lhs.get() >
+         multi_ptr<ElementType, Space, DecorateAddress>(nullptr).get();
 }
 
 template <typename ElementType, access::address_space Space,
           access::decorated DecorateAddress>
 bool operator>(std::nullptr_t,
-               const multi_ptr<ElementType, Space, DecorateAddress> &) {
-  return false;
+               const multi_ptr<ElementType, Space, DecorateAddress> &rhs) {
+  return multi_ptr<ElementType, Space, DecorateAddress>(nullptr).get() >
+         rhs.get();
 }
 
 template <typename ElementType, access::address_space Space,
           access::decorated DecorateAddress>
-bool operator<(const multi_ptr<ElementType, Space, DecorateAddress> &,
+bool operator<(const multi_ptr<ElementType, Space, DecorateAddress> &lhs,
                std::nullptr_t) {
-  return false;
+  return lhs.get() <
+         multi_ptr<ElementType, Space, DecorateAddress>(nullptr).get();
 }
 
 template <typename ElementType, access::address_space Space,
           access::decorated DecorateAddress>
 bool operator<(std::nullptr_t,
                const multi_ptr<ElementType, Space, DecorateAddress> &rhs) {
-  return rhs.get() != nullptr;
+  return multi_ptr<ElementType, Space, DecorateAddress>(nullptr).get() <
+         rhs.get();
 }
 
 template <typename ElementType, access::address_space Space,
           access::decorated DecorateAddress>
-bool operator>=(const multi_ptr<ElementType, Space, DecorateAddress> &,
+bool operator>=(const multi_ptr<ElementType, Space, DecorateAddress> &lhs,
                 std::nullptr_t) {
-  return true;
+  return lhs.get() >=
+         multi_ptr<ElementType, Space, DecorateAddress>(nullptr).get();
 }
 
 template <typename ElementType, access::address_space Space,
           access::decorated DecorateAddress>
 bool operator>=(std::nullptr_t,
                 const multi_ptr<ElementType, Space, DecorateAddress> &rhs) {
-  return rhs.get() == nullptr;
+  return multi_ptr<ElementType, Space, DecorateAddress>(nullptr).get() >=
+         rhs.get();
 }
 
 template <typename ElementType, access::address_space Space,
           access::decorated DecorateAddress>
 bool operator<=(const multi_ptr<ElementType, Space, DecorateAddress> &lhs,
                 std::nullptr_t) {
-  return lhs.get() == nullptr;
+  return lhs.get() <=
+         multi_ptr<ElementType, Space, DecorateAddress>(nullptr).get();
 }
 
 template <typename ElementType, access::address_space Space,
           access::decorated DecorateAddress>
 bool operator<=(std::nullptr_t,
-                const multi_ptr<ElementType, Space, DecorateAddress> &) {
-  return true;
+                const multi_ptr<ElementType, Space, DecorateAddress> &rhs) {
+  return multi_ptr<ElementType, Space, DecorateAddress>(nullptr).get() <=
+         rhs.get();
 }
 
 } // namespace _V1

sycl/test-e2e/Basic/multi_ptr_null_relational_operators.cpp

@@ -0,0 +1,85 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/multi_ptr.hpp>
+
+template <typename T, typename AccessorType,
+          sycl::access::address_space address_space,
+          sycl::access::decorated decorated>
+void check(sycl::multi_ptr<T, address_space, decorated> &mp,
+           AccessorType &dev_acc) {
+  using multi_ptr_t = sycl::multi_ptr<T, address_space, decorated>;
+  multi_ptr_t null_mp;
+  dev_acc[0] = nullptr == null_mp;
+  dev_acc[0] += nullptr != mp;
+  dev_acc[0] += std::less<multi_ptr_t>()(nullptr, mp) == nullptr < mp;
+  dev_acc[0] += std::less<multi_ptr_t>()(mp, nullptr) == mp < nullptr;
+  dev_acc[0] += std::less_equal<multi_ptr_t>()(nullptr, mp) == nullptr <= mp;
+  dev_acc[0] += std::less_equal<multi_ptr_t>()(mp, nullptr) == mp <= nullptr;
+  dev_acc[0] += std::greater<multi_ptr_t>()(nullptr, mp) == nullptr > mp;
+  dev_acc[0] += std::greater<multi_ptr_t>()(mp, nullptr) == mp > nullptr;
+  dev_acc[0] += std::greater_equal<multi_ptr_t>()(nullptr, mp) == nullptr >= mp;
+  dev_acc[0] += std::greater_equal<multi_ptr_t>()(mp, nullptr) == mp >= nullptr;
+}
+
+template <typename T, sycl::access::address_space address_space,
+          sycl::access::decorated decorated>
+void nullptrRelationalOperatorTest() {
+  using multi_ptr_t = sycl::multi_ptr<int, address_space, decorated>;
+  try {
+    sycl::queue queue;
+    sycl::buffer<int, 1> buf(1);
+    queue
+        .submit([&](sycl::handler &cgh) {
+          auto dev_acc = buf.get_access<sycl::access::mode::write>(cgh);
+          if constexpr (address_space ==
+                        sycl::access::address_space::local_space) {
+            sycl::local_accessor<int, 1> locAcc(1, cgh);
+            cgh.parallel_for(sycl::nd_range<1>{1, 1}, [=](sycl::id<1>) {
+              locAcc[0] = 1;
+              multi_ptr_t mp(locAcc);
+              check(mp, dev_acc);
+            });
+          } else if constexpr (address_space ==
+                               sycl::access::address_space::private_space) {
+            cgh.single_task([=] {
+              T priv_arr[1];
+              sycl::multi_ptr<T, address_space, decorated> mp =
+                  sycl::address_space_cast<address_space, decorated>(priv_arr);
+              check(mp, dev_acc);
+            });
+          } else {
+            cgh.single_task([=] {
+              multi_ptr_t mp(dev_acc);
+              check(mp, dev_acc);
+            });
+          }
+        })
+        .wait_and_throw();
+    assert(sycl::host_accessor{buf}[0] == 10);
+  } catch (sycl::exception e) {
+    std::cout << "SYCL exception caught: " << e.what();
+    return;
+  }
+}
+
+int main() {
+  nullptrRelationalOperatorTest<int, sycl::access::address_space::local_space,
+                                sycl::access::decorated::yes>();
+  nullptrRelationalOperatorTest<int, sycl::access::address_space::local_space,
+                                sycl::access::decorated::no>();
+  nullptrRelationalOperatorTest<int, sycl::access::address_space::global_space,
+                                sycl::access::decorated::yes>();
+  nullptrRelationalOperatorTest<int, sycl::access::address_space::global_space,
+                                sycl::access::decorated::no>();
+  nullptrRelationalOperatorTest<int, sycl::access::address_space::generic_space,
+                                sycl::access::decorated::yes>();
+  nullptrRelationalOperatorTest<int, sycl::access::address_space::generic_space,
+                                sycl::access::decorated::no>();
+  nullptrRelationalOperatorTest<int, sycl::access::address_space::private_space,
+                                sycl::access::decorated::yes>();
+  nullptrRelationalOperatorTest<int, sycl::access::address_space::private_space,
+                                sycl::access::decorated::no>();
+  return 0;
+}