Implementation of DPCTLQueue_Fill8/16/32/64/128 function to SyclInterface library (#831)

* Implementation of DPCTLQueue_Fill8/16/32/64/128 function

Also added tests and docstrings.

* genhtml uses --no-source to work around issue with non-existing integration headers

* gen_coverage script runs llvm-cov rather than lcov-genhtml target, since HTMLs are not used for CI coverage collection

* Added commended out use of -save-temps=obj

Whenever C++ source code contains device code, DPC++ would perform
two compilation passes (for host and for device portions) and combine
these two into fat object file. In the process of doing this it generates
temporary integration headers. So functions from these headers get
instrumented and .lcov file would contain references to deleted header
files, in the form of /tmp/dpctl_sycl_queue_integrface-header-hex.h

Added -save-temps instructs DPC++ to keep all the temporaries around and
makes these headers available for llvm-profdata and llvm-cov to use.

The side-effect of using -save-temps=obj is that preprocessor output
files are also kept around and referenced by instrumentation tools.
These files are very large. So much so that generated .json file can
not be uploaded to coverall service.

Therefore the -save-temps=obj is commented out. May still be useful
for manual profile collection.

References to deleted header files are still present (and wreck
llvm-lcov calls), but that can be addressed using llvm-cov option to
ignore certain files (like /tmp/dpctl_*.h)

* increase compute intensity of kern3 in the test

* Fixed logic for collection of profiling binaries

Since addition of --save-temps=obj, a slew of host/device object
files were added and llvm-cov can not work with some of them.

Since binaries can actually be shared objects, use that instead.

Process traces for libtensor and libsyclinterface shared objects

* llvm-cov target must export to dpctl.lcov for coveralls to work

* added diagnostic to the workflow

* Replaced use of numba-dpex in example with use of dpnp

This is until we have GEMM implemented in dpctl

* Use -ignore-filename-regex option of llvm-cov to exclude references to deleted integration headers

Co-authored-by: Vladislav Perevezentsev <[email protected]>

Author: oleksandr-pavlyk

.github/workflows/conda-package.yml

@@ -380,7 +380,7 @@ jobs:
           source $CONDA/etc/profile.d/conda.sh
           conda activate
           CHANNELS="-c $GITHUB_WORKSPACE/channel -c dppy/label/dev -c intel --override-channels"
-          conda install -n examples -y $CHANNELS numba-dppy numpy dpctl || exit 1
+          conda install -n examples -y $CHANNELS numpy dpctl dpnp || exit 1
       - name: Build and run examples with native extensions
         shell: bash -l {0}
         run: |

.github/workflows/generate-coverage.yaml

@@ -110,6 +110,9 @@ jobs:
               fh3 = open('combined-dpctl-c-api-coverage.json', 'w'); \
               json.dump(f3, fh3); fh3.close()" || exit 1
           # merge combined file with coverage data and upload
+          ls -lh dpctl-c-api-coverage.json pytest-dpctl-c-api-coverage.json \
+               combined-dpctl-c-api-coverage.json \
+               $(find _skbuild -name dpctl.lcov) $(find . -name dpctl_pytest.lcov)
           coveralls --service=github --merge=combined-dpctl-c-api-coverage.json
         env:
           GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}

dpctl/tests/test_sycl_kernel_submit.py

@@ -135,6 +135,11 @@ def test_async_submit():
         "   global unsigned int *arg2)"
         "{"
         "  size_t index = get_global_id(0);"
+        "  size_t i = 0; "
+        "  size_t unused_sum = 0;"
+        "  for (i = 0; i < 4000; i++) { "
+        "       unused_sum += i;"
+        "  } "
         "  res[index] = "
         "      (arg1[index] < arg2[index]) ? arg1[index] : arg2[index];"
         "}"

examples/python/dppy_kernel.py renamed to examples/python/sycl_timer.py

@@ -15,49 +15,50 @@
 # limitations under the License.
 
 
-import numba_dppy
+import dpnp
 import numpy as np
 
 import dpctl
+import dpctl.tensor as dpt
 from dpctl import SyclTimer
 
-
-@numba_dppy.kernel
-def dppy_gemm(a, b, c):
-    i = numba_dppy.get_global_id(0)
-    j = numba_dppy.get_global_id(1)
-    if i >= c.shape[0] or j >= c.shape[1]:
-        return
-    c[i, j] = 0
-    for k in range(c.shape[0]):
-        c[i, j] += a[i, k] * b[k, j]
-
-
-X = 1024
-Y = 16
-global_size = X, X
-
-griddim = X, X
-blockdim = Y, Y
-
-a = np.arange(X * X, dtype=np.float32).reshape(X, X)
-b = np.array(np.random.random(X * X), dtype=np.float32).reshape(X, X)
-c = np.ones_like(a).reshape(X, X)
+n = 4000
 
 try:
-    q = dpctl.SyclQueue("opencl:gpu", property="enable_profiling")
+    q = dpctl.SyclQueue(property="enable_profiling")
 except dpctl.SyclQueueCreationError:
     print(
         "Skipping the example, as dpctl.SyclQueue targeting "
-        "opencl:gpu device could not be created"
+        "default device could not be created"
     )
     exit(0)
+
+a = dpt.reshape(dpt.arange(n * n, dtype=np.float32, sycl_queue=q), (n, n))
+b = dpt.reshape(
+    dpt.asarray(np.random.random(n * n), dtype=np.float32, sycl_queue=q), (n, n)
+)
+
 timer = SyclTimer(time_scale=1)
-with dpctl.device_context(q):
+
+wall_times = []
+device_times = []
+print(
+    f"Performing matrix multiplication of two {n} by {n} matrices "
+    f"on {q.sycl_device.name}, repeating 5 times."
+)
+for _ in range(5):
     with timer(q):
-        dppy_gemm[griddim, blockdim](a, b, c)
-        cc = np.dot(a, b)
+        a_matmul_b = dpnp.matmul(a, b)
     host_time, device_time = timer.dt
+    wall_times.append(host_time)
+    device_times.append(device_time)
+
+c = dpnp.asnumpy(a_matmul_b)
+cc = np.dot(dpnp.asnumpy(a), dpnp.asnumpy(b))
 
-print("Wall time: ", host_time, "\nDevice time: ", device_time)
-print(np.allclose(c, cc))
+print("Wall time: ", wall_times, "\nDevice time: ", device_times)
+print(
+    "Accuracy test: passed."
+    if np.allclose(c, cc)
+    else (f"Accuracy test: failed. Discrepancy {np.max(np.abs(c-cc))}")
+)

libsyclinterface/cmake/modules/SetupCoverage.cmake

@@ -11,6 +11,7 @@ function(setup_coverage_generation)
         "-fprofile-instr-generate "
         "-fcoverage-mapping "
         "-fno-sycl-use-footer "
+#        "-save-temps=obj "
         "-DDPCTL_COVERAGE "
     )
 

libsyclinterface/include/dpctl_sycl_queue_interface.h

@@ -205,7 +205,7 @@ DPCTLQueue_GetDevice(__dpctl_keep const DPCTLSyclQueueRef QRef);
  * @ingroup QueueInterface
  */
 DPCTL_API
-DPCTLSyclEventRef
+__dpctl_give DPCTLSyclEventRef
 DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef,
                        __dpctl_keep const DPCTLSyclQueueRef QRef,
                        __dpctl_keep void **Args,
@@ -253,7 +253,7 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef,
  * @ingroup QueueInterface
  */
 DPCTL_API
-DPCTLSyclEventRef
+__dpctl_give DPCTLSyclEventRef
 DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef,
                          __dpctl_keep const DPCTLSyclQueueRef QRef,
                          __dpctl_keep void **Args,
@@ -287,10 +287,11 @@ void DPCTLQueue_Wait(__dpctl_keep const DPCTLSyclQueueRef QRef);
  * @ingroup QueueInterface
  */
 DPCTL_API
-DPCTLSyclEventRef DPCTLQueue_Memcpy(__dpctl_keep const DPCTLSyclQueueRef QRef,
-                                    void *Dest,
-                                    const void *Src,
-                                    size_t Count);
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_Memcpy(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                  void *Dest,
+                  const void *Src,
+                  size_t Count);
 
 /*!
  * @brief C-API wrapper for ``sycl::queue::prefetch``.
@@ -303,9 +304,10 @@ DPCTLSyclEventRef DPCTLQueue_Memcpy(__dpctl_keep const DPCTLSyclQueueRef QRef,
  * @ingroup QueueInterface
  */
 DPCTL_API
-DPCTLSyclEventRef DPCTLQueue_Prefetch(__dpctl_keep DPCTLSyclQueueRef QRef,
-                                      const void *Ptr,
-                                      size_t Count);
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_Prefetch(__dpctl_keep DPCTLSyclQueueRef QRef,
+                    const void *Ptr,
+                    size_t Count);
 
 /*!
  * @brief C-API wrapper for ``sycl::queue::mem_advise``.
@@ -321,10 +323,11 @@ DPCTLSyclEventRef DPCTLQueue_Prefetch(__dpctl_keep DPCTLSyclQueueRef QRef,
  * @ingroup QueueInterface
  */
 DPCTL_API
-DPCTLSyclEventRef DPCTLQueue_MemAdvise(__dpctl_keep DPCTLSyclQueueRef QRef,
-                                       const void *Ptr,
-                                       size_t Count,
-                                       int Advice);
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_MemAdvise(__dpctl_keep DPCTLSyclQueueRef QRef,
+                     const void *Ptr,
+                     size_t Count,
+                     int Advice);
 
 /*!
  * @brief C-API wrapper for sycl::queue::is_in_order that indicates whether
@@ -365,7 +368,7 @@ size_t DPCTLQueue_Hash(__dpctl_keep const DPCTLSyclQueueRef QRef);
  *           ``sycl::queue::submit_barrier()`` function.
  */
 DPCTL_API
-DPCTLSyclEventRef
+__dpctl_give DPCTLSyclEventRef
 DPCTLQueue_SubmitBarrier(__dpctl_keep const DPCTLSyclQueueRef QRef);
 
 /*!
@@ -380,7 +383,7 @@ DPCTLQueue_SubmitBarrier(__dpctl_keep const DPCTLSyclQueueRef QRef);
  *           ``sycl::queue::submit_barrier()`` function.
  */
 DPCTL_API
-DPCTLSyclEventRef DPCTLQueue_SubmitBarrierForEvents(
+__dpctl_give DPCTLSyclEventRef DPCTLQueue_SubmitBarrierForEvents(
     __dpctl_keep const DPCTLSyclQueueRef QRef,
     __dpctl_keep const DPCTLSyclEventRef *DepEvents,
     size_t NDepEvents);
@@ -397,9 +400,101 @@ DPCTLSyclEventRef DPCTLQueue_SubmitBarrierForEvents(
  * @ingroup QueueInterface
  */
 DPCTL_API
-DPCTLSyclEventRef DPCTLQueue_Memset(__dpctl_keep const DPCTLSyclQueueRef QRef,
-                                    void *USMRef,
-                                    uint8_t Value,
-                                    size_t Count);
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_Memset(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                  void *USMRef,
+                  uint8_t Value,
+                  size_t Count);
+
+/*!
+ * @brief C-API wrapper for ``sycl::queue::fill``.
+ *
+ * @param    QRef           An opaque pointer to the ``sycl::queue``.
+ * @param    USMRef         An USM pointer to the memory to fill.
+ * @param    Value          A uint8_t value to fill.
+ * @param    Count          A number of uint8_t elements to fill.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::fill`` function.
+ * @ingroup QueueInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_Fill8(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                 void *USMRef,
+                 uint8_t Value,
+                 size_t Count);
+
+/*!
+ * @brief C-API wrapper for ``sycl::queue::fill``.
+ *
+ * @param    QRef           An opaque pointer to the ``sycl::queue``.
+ * @param    USMRef         An USM pointer to the memory to fill.
+ * @param    Value          A uint16_t value to fill.
+ * @param    Count          A number of uint16_t elements to fill.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::fill`` function.
+ * @ingroup QueueInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_Fill16(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                  void *USMRef,
+                  uint16_t Value,
+                  size_t Count);
+
+/*!
+ * @brief C-API wrapper for ``sycl::queue::fill``.
+ *
+ * @param    QRef           An opaque pointer to the ``sycl::queue``.
+ * @param    USMRef         An USM pointer to the memory to fill.
+ * @param    Value          A uint32_t value to fill.
+ * @param    Count          A number of uint32_t elements to fill.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::fill`` function.
+ * @ingroup QueueInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_Fill32(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                  void *USMRef,
+                  uint32_t Value,
+                  size_t Count);
+
+/*!
+ * @brief C-API wrapper for ``sycl::queue::fill``.
+ *
+ * @param    QRef           An opaque pointer to the ``sycl::queue``.
+ * @param    USMRef         An USM pointer to the memory to fill.
+ * @param    Value          A uint64_t value to fill.
+ * @param    Count          A number of uint64_t elements to fill.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::fill`` function.
+ * @ingroup QueueInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_Fill64(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                  void *USMRef,
+                  uint64_t Value,
+                  size_t Count);
+
+/*!
+ * @brief C-API wrapper for ``sycl::queue::fill``.
+ *
+ * @param    QRef           An opaque pointer to the ``sycl::queue``.
+ * @param    USMRef         An USM pointer to the memory to fill.
+ * @param    Value          A pointer to uint64_t array of 2 elements with value
+ * to fill.
+ * @param    Count          A number of 128-bit elements to fill.
+ * @return   An opaque pointer to the ``sycl::event`` returned by the
+ *           ``sycl::queue::fill`` function.
+ * @ingroup QueueInterface
+ */
+DPCTL_API
+__dpctl_give DPCTLSyclEventRef
+DPCTLQueue_Fill128(__dpctl_keep const DPCTLSyclQueueRef QRef,
+                   void *USMRef,
+                   uint64_t *Value,
+                   size_t Count);
 
 DPCTL_C_EXTERN_C_END