[SYCL][COMPAT] Added constant memory_region for read_only memory (#11024)

Alcpz · web-flow · commit ab578107b6f8 · 2023-09-07T14:48:49.000+01:00
This PR extends the memory wrapper of SYCLcompat to support allocating
read only memory.

To allow the allocation of `device_memory` directly for source to source
translation:
- Extended `device_memory` to support a queue parameter in its
constructor, which defaults to `syclcompat::get_default_queue`
- Added device_read_only property to device memory if usm device read
only extension is available for `memory_region::constant`

Documentation was already updated, except for the usm device read only
extension
diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md
@@ -42,6 +42,7 @@ Specifically, this library depends on the following SYCL extensions:
     ../extensions/supported/sycl_ext_oneapi_assert.asciidoc)
 * [sycl_ext_oneapi_enqueue_barrier](
     ../extensions/supported/sycl_ext_oneapi_enqueue_barrier.asciidoc)
+* [sycl_ext_oneapi_usm_device_read_only](../extensions/supported/sycl_ext_oneapi_usm_device_read_only.asciidoc)
 
 ## Usage
 
@@ -265,7 +266,7 @@ void vectorAdd(const float *A, const float *B, float *C, int n,
 Then, `vectorAdd` can be launched like this:
 
 ``` c++
-syclcompat::launch<vectorAdd>(blocksPerGrid, threadsPerBlock, mem_size, d_A, 
+syclcompat::launch<vectorAdd>(blocksPerGrid, threadsPerBlock, mem_size, d_A,
                               d_B, d_C, n);
 ```
 
diff --git a/sycl/include/syclcompat/memory.hpp b/sycl/include/syclcompat/memory.hpp
@@ -876,12 +876,14 @@ template <class T, memory_region Memory, size_t Dimension> class device_memory {
   using value_t = typename detail::memory_traits<Memory, T>::value_t;
   using compat_accessor_t = syclcompat::accessor<T, Memory, Dimension>;
 
-  device_memory() : device_memory(sycl::range<Dimension>(1)) {}
+  device_memory(sycl::queue q = get_default_queue())
+      : device_memory(sycl::range<Dimension>(1), q) {}
 
   /// Constructor of 1-D array with initializer list
   device_memory(const sycl::range<Dimension> &in_range,
-                std::initializer_list<value_t> &&init_list)
-      : device_memory(in_range) {
+                std::initializer_list<value_t> &&init_list,
+                sycl::queue q = get_default_queue())
+      : device_memory(in_range, q) {
     assert(init_list.size() <= in_range.size());
     _host_ptr = (value_t *)std::malloc(_size);
     std::memset(_host_ptr, 0, _size);
@@ -892,8 +894,9 @@ template <class T, memory_region Memory, size_t Dimension> class device_memory {
   template <size_t D = Dimension>
   device_memory(
       const typename std::enable_if<D == 2, sycl::range<2>>::type &in_range,
-      std::initializer_list<std::initializer_list<value_t>> &&init_list)
-      : device_memory(in_range) {
+      std::initializer_list<std::initializer_list<value_t>> &&init_list,
+      sycl::queue q = get_default_queue())
+      : device_memory(in_range, q) {
     assert(init_list.size() <= in_range[0]);
     _host_ptr = (value_t *)std::malloc(_size);
     std::memset(_host_ptr, 0, _size);
@@ -906,9 +909,10 @@ template <class T, memory_region Memory, size_t Dimension> class device_memory {
   }
 
   /// Constructor with range
-  device_memory(const sycl::range<Dimension> &range_in)
+  device_memory(const sycl::range<Dimension> &range_in,
+                sycl::queue q = get_default_queue())
       : _size(range_in.size() * sizeof(T)), _range(range_in), _reference(false),
-        _host_ptr(nullptr), _device_ptr(nullptr) {
+        _host_ptr(nullptr), _device_ptr(nullptr), _q(q) {
     static_assert((Memory == memory_region::global) ||
                       (Memory == memory_region::constant) ||
                       (Memory == memory_region::usm_shared),
@@ -918,19 +922,28 @@ template <class T, memory_region Memory, size_t Dimension> class device_memory {
   }
 
   /// Constructor with range
-  template <class... Args>
+  // enable_if_t SFINAE to avoid ambiguity with
+  // device_memory(Args... Arguments, sycl::queue q)
+  template <class... Args, size_t D = Dimension,
+            typename = std::enable_if_t<sizeof...(Args) == D>>
   device_memory(Args... Arguments)
-      : device_memory(sycl::range<Dimension>(Arguments...)) {}
+      : device_memory(sycl::range<Dimension>(Arguments...),
+                      get_default_queue()) {}
+
+  /// Constructor with range and queue
+  template <class... Args>
+  device_memory(Args... Arguments, sycl::queue q)
+      : device_memory(sycl::range<Dimension>(Arguments...), q) {}
 
   ~device_memory() {
     if (_device_ptr && !_reference)
-      free(_device_ptr);
+      syclcompat::free(_device_ptr, _q);
     if (_host_ptr)
       std::free(_host_ptr);
   }
 
   /// Allocate memory with default queue, and init memory if has initial value.
-  void init() { init(get_default_queue()); }
+  void init() { init(_q); }
   /// Allocate memory with specified queue, and init memory if has initial
   /// value.
   void init(sycl::queue q) {
@@ -946,12 +959,12 @@ template <class T, memory_region Memory, size_t Dimension> class device_memory {
   /// The variable is assigned to a device pointer.
   void assign(value_t *src, size_t size) {
     this->~device_memory();
-    new (this) device_memory(src, size);
+    new (this) device_memory(src, size, _q);
   }
 
   /// Get memory pointer of the memory object, which is virtual pointer when
   /// usm is not used, and device pointer when usm is used.
-  value_t *get_ptr() { return get_ptr(get_default_queue()); }
+  value_t *get_ptr() { return get_ptr(_q); }
   /// Get memory pointer of the memory object, which is virtual pointer when
   /// usm is not used, and device pointer when usm is used.
   value_t *get_ptr(sycl::queue q) {
@@ -977,16 +990,25 @@ template <class T, memory_region Memory, size_t Dimension> class device_memory {
   }
 
 private:
-  device_memory(value_t *memory_ptr, size_t size)
+  device_memory(value_t *memory_ptr, size_t size,
+                sycl::queue q = get_default_queue())
       : _size(size), _range(size / sizeof(T)), _reference(true),
-        _device_ptr(memory_ptr) {}
+        _device_ptr(memory_ptr), _q(q) {}
 
   void allocate_device(sycl::queue q) {
     if (Memory == memory_region::usm_shared) {
       _device_ptr = (value_t *)sycl::malloc_shared(_size, q.get_device(),
                                                    q.get_context());
       return;
     }
+#ifdef SYCL_EXT_ONEAPI_USM_DEVICE_READ_ONLY
+    if (Memory == memory_region::constant) {
+      _device_ptr = (value_t *)sycl::malloc_device(
+          _size, q.get_device(), q.get_context(),
+          sycl::ext::oneapi::property::usm::device_read_only());
+      return;
+    }
+#endif
     _device_ptr = (value_t *)detail::malloc(_size, q);
   }
 
@@ -995,6 +1017,7 @@ template <class T, memory_region Memory, size_t Dimension> class device_memory {
   bool _reference;
   value_t *_host_ptr;
   value_t *_device_ptr;
+  sycl::queue _q;
 };
 template <class T, memory_region Memory>
 class device_memory<T, Memory, 0> : public device_memory<T, Memory, 1> {
diff --git a/sycl/test-e2e/syclcompat/memory/memory_management_test1.cpp b/sycl/test-e2e/syclcompat/memory/memory_management_test1.cpp
@@ -368,11 +368,120 @@ template <typename T> void test_fill_q() {
   free(h_A);
 }
 
+void test_constant_memcpy() {
+  std::cout << __PRETTY_FUNCTION__ << std::endl;
+
+  constexpr size_t size = 2000;
+  constexpr size_t offset = 1000;
+
+  syclcompat::constant_memory<float, 1> d_A(size);
+  syclcompat::constant_memory<float, 1> d_B(size);
+
+  float *h_A = (float *)malloc(size / 2 * sizeof(float));
+  float *h_B = (float *)malloc(size / 2 * sizeof(float));
+  float *h_C = (float *)malloc(size * sizeof(float));
+  float *h_D = (float *)malloc(size * sizeof(float));
+
+  for (int i = 0; i < size / 2; i++) {
+    h_A[i] = 1.0f;
+    h_B[i] = 2.0f;
+  }
+
+  // hostA[0..999] -> deviceA[0..999]
+  // hostB[0..999] -> deviceA[1000..1999]
+  // deviceA[0..1999] -> hostC[0..1999]
+  // deviceA[0..999] -> deviceB[0..999]
+  // deviceA[1000..1999] -> deviceB[1000..1999]
+  // deviceB[0..1999] -> hostD[0..1999]
+
+  syclcompat::memcpy(d_A.get_ptr(), h_A, offset * sizeof(float));
+  syclcompat::memcpy((char *)d_A.get_ptr() + offset * sizeof(float), h_B,
+                     (size - offset) * sizeof(float));
+  syclcompat::memcpy(h_C, d_A.get_ptr(), size * sizeof(float));
+  syclcompat::memcpy(d_B.get_ptr(), d_A.get_ptr(), offset * sizeof(float));
+  syclcompat::memcpy((char *)d_B.get_ptr() + offset * sizeof(float),
+                     (void *)((size_t)d_A.get_ptr() + offset * sizeof(float)),
+                     (size - offset) * sizeof(float));
+  syclcompat::memcpy(h_D, d_B.get_ptr(), size * sizeof(float));
+
+  // verify hostD
+  for (int i = 0; i < offset; i++) {
+    assert(fabs(h_A[i] - h_D[i]) <= 1e-5);
+  }
+
+  for (int i = offset; i < size; i++) {
+    assert(fabs(h_B[i - offset] - h_D[i]) <= 1e-5);
+  }
+
+  free(h_A);
+  free(h_B);
+  free(h_C);
+  free(h_D);
+}
+
+void test_constant_memcpy_q() {
+  std::cout << __PRETTY_FUNCTION__ << std::endl;
+
+  sycl::queue q{{sycl::property::queue::in_order()}};
+
+  constexpr size_t size = 2000;
+  constexpr size_t offset = 1000;
+  syclcompat::constant_memory<float, 1> d_A(size, q);
+  syclcompat::constant_memory<float, 1> d_B(size, q);
+
+  float *h_A = (float *)malloc(size / 2 * sizeof(float));
+  float *h_B = (float *)malloc(size / 2 * sizeof(float));
+  float *h_C = (float *)malloc(size * sizeof(float));
+  float *h_D = (float *)malloc(size * sizeof(float));
+
+  for (int i = 0; i < size / 2; i++) {
+    h_A[i] = 1.0f;
+    h_B[i] = 2.0f;
+  }
+
+  // hostA[0..999] -> deviceA[0..999]
+  // hostB[0..999] -> deviceA[1000..1999]
+  // deviceA[0..1999] -> hostC[0..1999]
+  // deviceA[0..999] -> deviceB[0..999]
+  // deviceA[1000..1999] -> deviceB[1000..1999]
+  // deviceB[0..1999] -> hostD[0..1999]
+
+  syclcompat::memcpy(d_A.get_ptr(), h_A, offset * sizeof(float), q);
+
+  syclcompat::memcpy((char *)d_A.get_ptr() + offset * sizeof(float), h_B,
+                     (size - offset) * sizeof(float), q);
+  syclcompat::memcpy(h_C, d_A.get_ptr(), size * sizeof(float), q);
+
+  syclcompat::memcpy(d_B.get_ptr(), d_A.get_ptr(), offset * sizeof(float), q);
+
+  syclcompat::memcpy((char *)d_B.get_ptr() + offset * sizeof(float),
+                     (void *)((size_t)d_A.get_ptr() + offset * sizeof(float)),
+                     (size - offset) * sizeof(float), q);
+
+  syclcompat::memcpy(h_D, d_B.get_ptr(), size * sizeof(float), q);
+
+  // verify hostD
+  for (int i = 0; i < offset; i++) {
+    assert(fabs(h_A[i] - h_D[i]) <= 1e-5);
+  }
+
+  for (int i = offset; i < size; i++) {
+    assert(fabs(h_B[i - offset] - h_D[i]) <= 1e-5);
+  }
+
+  free(h_A);
+  free(h_B);
+  free(h_C);
+  free(h_D);
+}
+
 int main() {
   test_memcpy();
   test_memcpy_q();
   test_memset();
   test_memset_q();
+  test_constant_memcpy();
+  test_constant_memcpy_q();
 
   INSTANTIATE_ALL_TYPES(value_type_list, test_memcpy_t);
   INSTANTIATE_ALL_TYPES(value_type_list, test_memcpy_t_q);
diff --git a/sycl/test-e2e/syclcompat/memory/memory_management_test2.cpp b/sycl/test-e2e/syclcompat/memory/memory_management_test2.cpp
@@ -39,6 +39,9 @@
 
 #include "memory_common.hpp"
 
+constexpr size_t DataW = 100;
+constexpr size_t DataH = 100;
+
 void test_memcpy_pitched() {
   std::cout << __PRETTY_FUNCTION__ << std::endl;
 
@@ -151,13 +154,6 @@ void test_memcpy_kernel() {
   free(h_C);
 }
 
-#define DataW 100
-#define DataH 100
-
-syclcompat::global_memory<float, 2> g_A(DataW, DataH);
-syclcompat::global_memory<float, 2> g_B(DataW, DataH);
-syclcompat::global_memory<float, 2> g_C(DataW, DataH);
-
 void test_global_memory() {
   std::cout << __PRETTY_FUNCTION__ << std::endl;
 
@@ -172,6 +168,10 @@ void test_global_memory() {
     }
   }
 
+  syclcompat::global_memory<float, 2> g_A(DataW, DataH);
+  syclcompat::global_memory<float, 2> g_B(DataW, DataH);
+  syclcompat::global_memory<float, 2> g_C(DataW, DataH);
+
   g_A.init();
   g_B.init();
   g_C.init();
@@ -216,13 +216,13 @@ void test_global_memory() {
   }
 }
 
-syclcompat::shared_memory<float, 1> s_A(DataW);
-syclcompat::shared_memory<float, 1> s_B(DataW);
-syclcompat::shared_memory<float, 1> s_C(DataW);
-
 void test_shared_memory() {
   std::cout << __PRETTY_FUNCTION__ << std::endl;
 
+  syclcompat::shared_memory<float, 1> s_A(DataW);
+  syclcompat::shared_memory<float, 1> s_B(DataW);
+  syclcompat::shared_memory<float, 1> s_C(DataW);
+
   s_A.init();
   s_B.init();
   s_C.init();
@@ -256,6 +256,60 @@ void test_shared_memory() {
   }
 }
 
+void test_constant_memory() {
+  std::cout << __PRETTY_FUNCTION__ << std::endl;
+
+  float h_A[DataW][DataH];
+  float h_B[DataW][DataH];
+  float h_C[DataW][DataH];
+
+  for (int i = 0; i < DataW; i++) {
+    for (int j = 0; j < DataH; j++) {
+      h_A[i][j] = 1.0f;
+      h_B[i][j] = 2.0f;
+    }
+  }
+
+  syclcompat::constant_memory<float, 2> c_A(DataW, DataH);
+  syclcompat::constant_memory<float, 2> c_B(DataW, DataH);
+  syclcompat::global_memory<float, 2> g_C(DataW, DataH);
+
+  c_A.init();
+  c_B.init();
+  g_C.init();
+  syclcompat::memcpy((void *)c_A.get_ptr(), (void *)&h_A[0][0],
+                     DataW * DataH * sizeof(float));
+  syclcompat::memcpy((void *)c_B.get_ptr(), (void *)&h_B[0][0],
+                     DataW * DataH * sizeof(float));
+
+  {
+    syclcompat::get_default_queue().submit([&](sycl::handler &cgh) {
+      auto c_A_acc = c_A.get_access(cgh);
+      auto c_B_acc = c_B.get_access(cgh);
+      auto g_C_acc = g_C.get_access(cgh);
+      cgh.parallel_for(sycl::range<2>(DataW, DataH), [=](sycl::id<2> id) {
+        syclcompat::accessor<float, syclcompat::memory_region::constant, 2> A(
+            c_A_acc);
+        syclcompat::accessor<float, syclcompat::memory_region::constant, 2> B(
+            c_B_acc);
+        syclcompat::accessor<float, syclcompat::memory_region::global, 2> C(
+            g_C_acc);
+        int i = id[0], j = id[1];
+        C[i][j] = A[i][j] + B[i][j];
+      });
+    });
+    syclcompat::get_default_queue().wait_and_throw();
+  }
+  syclcompat::memcpy((void *)&h_C[0][0], (void *)g_C.get_ptr(),
+                     DataW * DataH * sizeof(float));
+  // verify hostD
+  for (int i = 0; i < DataW; i++) {
+    for (int j = 0; j < DataH; j++) {
+      assert(fabs(h_C[i][j] - h_A[i][j] - h_B[i][j]) <= 1e-5);
+    }
+  }
+}
+
 void test_memcpy_pitched_q() {
   std::cout << __PRETTY_FUNCTION__ << std::endl;
 
@@ -312,6 +366,6 @@ int main() {
 
   test_global_memory();
   test_shared_memory();
-
+  test_constant_memory();
   return 0;
 }
diff --git a/sycl/test-e2e/syclcompat/memory/memory_management_test3.cpp b/sycl/test-e2e/syclcompat/memory/memory_management_test3.cpp
