Update dpctl sycl kernel example

Author: ndgrigorian

examples/pybind11/use_dpctl_sycl_kernel/example.py

@@ -16,11 +16,12 @@
 
 # coding: utf-8
 
+import numpy as np
 import use_kernel as eg
 
 import dpctl
+import dpctl.memory as dpmem
 import dpctl.program as dppr
-import dpctl.tensor as dpt
 
 # create execution queue, targeting default selected device
 q = dpctl.SyclQueue()
@@ -38,10 +39,18 @@
 assert krn.num_args == 2
 
 # Construct the argument, and allocate memory for the result
-x = dpt.arange(0, stop=13, step=1, dtype="i4", sycl_queue=q)
-y = dpt.empty_like(x)
+x = np.arange(0, stop=13, step=1, dtype="i4")
+y = np.empty_like(x)
+x_dev = dpmem.MemoryUSMDevice(x.nbytes, queue=q)
+y_dev = dpmem.MemoryUSMDevice(y.nbytes, queue=q)
 
-eg.submit_custom_kernel(q, krn, src=x, dst=y)
+# Copy input data to the device
+q.memcpy(dest=x_dev, src=x, count=x.nbytes)
+
+eg.submit_custom_kernel(q, krn, src=x_dev, dst=y_dev)
+
+# Copy result data back to host
+q.memcpy(dest=y, src=y_dev, count=y.nbytes)
 
 # output the result
-print(dpt.asnumpy(y))
+print(y)

examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py

@@ -23,8 +23,8 @@
 import use_kernel as uk
 
 import dpctl
-import dpctl.program as dpm
-import dpctl.tensor as dpt
+import dpctl.memory as dpmem
+import dpctl.program as dppr
 
 
 def _get_spv_path():
@@ -45,7 +45,7 @@ def test_kernel_can_be_found():
         q = dpctl.SyclQueue()
     except dpctl.SyclQueueCreationError:
         pytest.skip("Could not create default queue")
-    pr = dpm.create_program_from_spirv(q, il, "")
+    pr = dppr.create_program_from_spirv(q, il, "")
     assert pr.has_sycl_kernel("double_it")
 
 
@@ -57,14 +57,20 @@ def test_kernel_submit_through_extension():
         q = dpctl.SyclQueue()
     except dpctl.SyclQueueCreationError:
         pytest.skip("Could not create default queue")
-    pr = dpm.create_program_from_spirv(q, il, "")
+    pr = dppr.create_program_from_spirv(q, il, "")
     krn = pr.get_sycl_kernel("double_it")
     assert krn.num_args == 2
 
-    x = dpt.arange(0, stop=13, step=1, dtype="i4", sycl_queue=q)
-    y = dpt.zeros_like(x)
+    x = np.arange(0, stop=13, step=1, dtype="i4")
+    y = np.empty_like(x)
 
-    q.wait()
-    uk.submit_custom_kernel(q, krn, x, y, [])
+    x_usm = dpmem.MemoryUSMDevice(x.nbytes, queue=q)
+    y_usm = dpmem.MemoryUSMDevice(y.nbytes, queue=q)
 
-    assert np.array_equal(dpt.asnumpy(y), np.arange(0, 26, step=2, dtype="i4"))
+    ev = q.memcpy_async(dest=x_usm, src=x, count=x_usm.nbytes)
+
+    uk.submit_custom_kernel(q, krn, x_usm, y_usm, [ev])
+
+    q.memcpy(dest=y, src=y_usm, count=y.nbytes)
+
+    assert np.array_equal(y, np.arange(0, 26, step=2, dtype="i4"))

examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp

@@ -37,45 +37,37 @@ namespace py = pybind11;
 
 void submit_custom_kernel(sycl::queue &q,
                           sycl::kernel &krn,
-                          dpctl::tensor::usm_ndarray x,
-                          dpctl::tensor::usm_ndarray y,
+                          dpctl::memory::usm_memory x,
+                          dpctl::memory::usm_memory y,
                           const std::vector<sycl::event> &depends = {})
 {
-    if (x.get_ndim() != 1 || !x.is_c_contiguous() || y.get_ndim() != 1 ||
-        !y.is_c_contiguous())
-    {
-        throw py::value_error(
-            "src and dst arguments must be 1D and contiguous.");
-    }
+    const std::size_t nbytes_x = x.get_nbytes();
+    const std::size_t nbytes_y = y.get_nbytes();
 
-    auto const &api = dpctl::detail::dpctl_capi::get();
-    if (x.get_typenum() != api.UAR_INT32_ || y.get_typenum() != api.UAR_INT32_)
-    {
-        throw py::value_error(
-            "src and dst arguments must have int32 element data types.");
+    if (nbytes_x != nbytes_y) {
+        throw py::value_error("src and dst arguments must have equal nbytes.");
+    }
+    if (nbytes_x % sizeof(std::int32_t) != 0) {
+        throw py::value_error("src and dst must be interpretable as int32 "
+                              "(nbytes must be a multiple of 4).");
     }
 
-    size_t n_x = x.get_size();
-    size_t n_y = y.get_size();
+    auto *x_data = reinterpret_cast<std::int32_t *>(x.get_pointer());
+    auto *y_data = reinterpret_cast<std::int32_t *>(y.get_pointer());
 
-    if (n_x != n_y) {
-        throw py::value_error("src and dst arguments must have equal size.");
-    }
+    const std::size_t n_elems = nbytes_x / sizeof(std::int32_t);
 
     if (!dpctl::utils::queues_are_compatible(q, {x.get_queue(), y.get_queue()}))
     {
         throw std::runtime_error(
             "Execution queue is not compatible with allocation queues");
     }
 
-    void *x_data = x.get_data<void>();
-    void *y_data = y.get_data<void>();
-
     sycl::event e = q.submit([&](sycl::handler &cgh) {
         cgh.depends_on(depends);
         cgh.set_arg(0, x_data);
         cgh.set_arg(1, y_data);
-        cgh.parallel_for(sycl::range<1>(n_x), krn);
+        cgh.parallel_for(sycl::range<1>(n_elems), krn);
     });
 
     e.wait();