[SYCL] Adjust multi_ptr deduction guides (#9751)

This commit adjusts the multi_ptr deduction guides following the changes
to the SYCL 2020 specification in
KhronosGroup/SYCL-Docs#405.

Fixes #9692.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/include/sycl/multi_ptr.hpp

@@ -1274,10 +1274,21 @@ class multi_ptr<const void, Space, access::decorated::legacy> {
 };
 
 #ifdef __cpp_deduction_guides
-template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
+template <int dimensions, access::placeholder isPlaceholder,
           typename PropertyListT, class T>
-multi_ptr(accessor<T, dimensions, Mode, access::target::device, isPlaceholder,
-                   PropertyListT>)
+multi_ptr(accessor<T, dimensions, access::mode::read, access::target::device,
+                   isPlaceholder, PropertyListT>)
+    -> multi_ptr<const T, access::address_space::global_space,
+                 access::decorated::no>;
+template <int dimensions, access::placeholder isPlaceholder,
+          typename PropertyListT, class T>
+multi_ptr(accessor<T, dimensions, access::mode::write, access::target::device,
+                   isPlaceholder, PropertyListT>)
+    -> multi_ptr<T, access::address_space::global_space, access::decorated::no>;
+template <int dimensions, access::placeholder isPlaceholder,
+          typename PropertyListT, class T>
+multi_ptr(accessor<T, dimensions, access::mode::read_write,
+                   access::target::device, isPlaceholder, PropertyListT>)
     -> multi_ptr<T, access::address_space::global_space, access::decorated::no>;
 template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
           typename PropertyListT, class T>

sycl/test/multi_ptr/ctad.cpp

@@ -19,27 +19,47 @@ int main() {
   using sycl::access::address_space;
   using sycl::access::mode;
   using sycl::access::target;
-  using deviceAcc = sycl::accessor<int, 1, mode::read, target::device>;
-  using globlAcc = sycl::accessor<int, 1, mode::read, target::global_buffer>;
+  using rDeviceAcc = sycl::accessor<int, 1, mode::read, target::device>;
+  using rGloblAcc = sycl::accessor<int, 1, mode::read, target::global_buffer>;
+  using wDeviceAcc = sycl::accessor<int, 1, mode::write, target::device>;
+  using wGloblAcc = sycl::accessor<int, 1, mode::write, target::global_buffer>;
+  using rwDeviceAcc = sycl::accessor<int, 1, mode::read_write, target::device>;
+  using rwGloblAcc =
+      sycl::accessor<int, 1, mode::read_write, target::global_buffer>;
   using constAcc = sycl::accessor<int, 1, mode::read, target::constant_buffer>;
   using localAcc = sycl::local_accessor<int, 1>;
-  using localAccDep = sycl::accessor<int, 1, mode::read, target::local>;
-  using deviceCTAD = decltype(sycl::multi_ptr(std::declval<deviceAcc>()));
-  using globlCTAD = decltype(sycl::multi_ptr(std::declval<globlAcc>()));
+  using localAccDep = sycl::local_accessor<int, 1>;
+  using rDeviceCTAD = decltype(sycl::multi_ptr(std::declval<rDeviceAcc>()));
+  using rGloblCTAD = decltype(sycl::multi_ptr(std::declval<rGloblAcc>()));
+  using wDeviceCTAD = decltype(sycl::multi_ptr(std::declval<wDeviceAcc>()));
+  using wGloblCTAD = decltype(sycl::multi_ptr(std::declval<wGloblAcc>()));
+  using rwDeviceCTAD = decltype(sycl::multi_ptr(std::declval<rwDeviceAcc>()));
+  using rwGloblCTAD = decltype(sycl::multi_ptr(std::declval<rwGloblAcc>()));
   using constCTAD = decltype(sycl::multi_ptr(std::declval<constAcc>()));
   using localCTAD = decltype(sycl::multi_ptr(std::declval<localAcc>()));
   using localCTADDep = decltype(sycl::multi_ptr(std::declval<localAccDep>()));
   using deviceMPtr = sycl::multi_ptr<int, address_space::global_space,
                                      sycl::access::decorated::no>;
   using globlMPtr = sycl::multi_ptr<int, address_space::global_space,
                                     sycl::access::decorated::no>;
+  using deviceConstMPtr =
+      sycl::multi_ptr<const int, address_space::global_space,
+                      sycl::access::decorated::no>;
+  using globlConstMPtr = sycl::multi_ptr<const int, address_space::global_space,
+                                         sycl::access::decorated::no>;
   using constMPtr = sycl::multi_ptr<int, address_space::constant_space,
                                     sycl::access::decorated::legacy>;
   using localMPtr = sycl::multi_ptr<int, address_space::local_space,
                                     sycl::access::decorated::no>;
-  static_assert(std::is_same<deviceCTAD, deviceMPtr>::value);
-  static_assert(std::is_same<deviceCTAD, globlMPtr>::value);
-  static_assert(std::is_same<globlCTAD, globlMPtr>::value);
+  static_assert(std::is_same<rwDeviceCTAD, deviceMPtr>::value);
+  static_assert(std::is_same<rwDeviceCTAD, globlMPtr>::value);
+  static_assert(std::is_same<rwGloblCTAD, globlMPtr>::value);
+  static_assert(std::is_same<wDeviceCTAD, deviceMPtr>::value);
+  static_assert(std::is_same<wDeviceCTAD, globlMPtr>::value);
+  static_assert(std::is_same<wGloblCTAD, globlMPtr>::value);
+  static_assert(std::is_same<rDeviceCTAD, deviceConstMPtr>::value);
+  static_assert(std::is_same<rDeviceCTAD, globlConstMPtr>::value);
+  static_assert(std::is_same<rGloblCTAD, globlConstMPtr>::value);
   static_assert(std::is_same<constCTAD, constMPtr>::value);
   static_assert(std::is_same<localCTAD, localMPtr>::value);
   static_assert(std::is_same<localCTADDep, localMPtr>::value);

sycl/test/regression/read_only_accessor_multi_ptr.cpp

@@ -0,0 +1,19 @@
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning
+// expected-no-diagnostics
+
+// Tests that read-only accessors can be used in multi_ptr deduction guides.
+
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::queue Q;
+  int Data = 0;
+  sycl::buffer<int, 1> Buf{&Data, {1}};
+
+  Q.submit([&](sycl::handler &CGH) {
+     sycl::accessor<int, 1, sycl::access_mode::read, sycl::target::device> Acc(
+         Buf, CGH);
+     CGH.single_task([=] { auto MPtr = sycl::multi_ptr(Acc); });
+   }).wait_and_throw();
+  return 0;
+}