Pass sycl::queue by reference to avoid needless copy

The copy operation would involved atomics and might get expensive.

Author: oleksandr-pavlyk

common_ext/cy_kde.pyx

@@ -24,7 +24,7 @@ from dpctl.sycl cimport unwrap_queue
 
 cdef extern from "kde.hpp" namespace "example":
     void kernel_density_estimate[T](
-        dpcpp_queue,   # execution queue
+        dpcpp_queue &, # execution queue
         uint16_t,      # datadimensionality
         const T *,     # evaluation points
         T *,           # output for kde values

common_ext/py_kde.cpp

@@ -28,7 +28,7 @@ namespace py = pybind11;
 template <typename T>
 py::array_t<T>
 py_kde_eval_t(
-    sycl::queue q,
+    sycl::queue &q,
     py::array_t<T, py::array::c_style|py::array::forcecast> x,
     py::array_t<T, py::array::c_style|py::array::forcecast> data,
     T h)
@@ -72,7 +72,7 @@ py_kde_eval_t(
 
 py::array
 py_kde_eval(
-    sycl::queue exec_q,
+    sycl::queue &exec_q,
     py::array x,
     py::array data,
     py::object h)

common_src/kde.hpp

@@ -14,7 +14,7 @@
 
 #pragma once
 
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 #include <cstdint>
 #include <iostream>
 
@@ -51,7 +51,7 @@ T1 upper_quotient_of(T1 n, T2 wg) {
 template <typename T>
 void
 kernel_density_estimate(
-    sycl::queue q,
+    sycl::queue &q,
     std::uint16_t dim, // dimensionality of the data
     const T* x,        // points at which KDE is evaluated
     T *f,              // where values of kde(x, h) are written to