Add synchronization points in example

oleksandr-pavlyk · oleksandr-pavlyk · commit f31fb287b0d5 · 2024-06-05T14:22:00.000-05:00
diff --git a/examples/cython/usm_memory/blackscholes/blackscholes.pyx b/examples/cython/usm_memory/blackscholes/blackscholes.pyx
@@ -106,11 +106,15 @@ def black_scholes_price(c_dpt.usm_ndarray option_params_arr):
         call_put_prices = dpt.empty((n_opts, 2), dtype='d', sycl_queue=q)
         dp1 = <double *>option_params_arr.get_data()
         dp2 = <double *>call_put_prices.get_data()
+        # ensure content of dp1 and dp2 is no longer worked on
+        exec_q_ptr[0].wait()
         cpp_blackscholes[double](exec_q_ptr[0], n_opts, dp1, dp2)
     elif (typenum_ == c_dpt.UAR_FLOAT):
         call_put_prices = dpt.empty((n_opts, 2), dtype='f', sycl_queue=q)
         fp1 = <float *>option_params_arr.get_data()
         fp2 = <float *>call_put_prices.get_data()
+        # ensure content of fp1 and fp2 is no longer worked on
+        exec_q_ptr[0].wait()
         cpp_blackscholes[float](exec_q_ptr[0], n_opts, fp1, fp2)
     else:
         raise ValueError("Unsupported data-type")
@@ -196,11 +200,13 @@ def populate_params(
 
     if (typenum_ == c_dpt.UAR_DOUBLE):
         dp = <double *>option_params_arr.get_data()
+        exec_q_ptr[0].wait()
         cpp_populate_params[double](
             exec_q_ptr[0], n_opts, dp, pl, ph, sl, sh, tl, th, rl, rh, vl, vh, seed
         )
     elif (typenum_ == c_dpt.UAR_FLOAT):
         fp = <float *>option_params_arr.get_data()
+        exec_q_ptr[0].wait()
         cpp_populate_params[float](
             exec_q_ptr[0], n_opts, fp, pl, ph, sl, sh, tl, th, rl, rh, vl, vh, seed
         )
diff --git a/examples/cython/usm_memory/src/sycl_blackscholes.hpp b/examples/cython/usm_memory/src/sycl_blackscholes.hpp
@@ -47,7 +47,7 @@ constexpr int CALL = 0;
 constexpr int PUT = 1;
 
 template <typename T>
-void cpp_blackscholes(sycl::queue q, size_t n_opts, T *params, T *callput)
+void cpp_blackscholes(sycl::queue &q, size_t n_opts, T *params, T *callput)
 {
     using data_t = T;
 
@@ -57,56 +57,54 @@ void cpp_blackscholes(sycl::queue q, size_t n_opts, T *params, T *callput)
         data_t half = one / two;
 
         cgh.parallel_for<class black_scholes_kernel<T>>(
-		sycl::range<1>(n_opts),
-		[=](sycl::id<1> idx)
-        {
-            const size_t i = n_params * idx[0];
-            const data_t opt_price = params[i + PRICE];
-            const data_t opt_strike = params[i + STRIKE];
-            const data_t opt_maturity = params[i + MATURITY];
-            const data_t opt_rate = params[i + RATE];
-            const data_t opt_volatility = params[i + VOLATILITY];
-            data_t a, b, c, y, z, e, d1, d1c, d2, d2c, w1, w2;
-            data_t mr = -opt_rate,
-                   sig_sig_two = two * opt_volatility * opt_volatility;
-
-            a = sycl::log(opt_price / opt_strike);
-            b = opt_maturity * mr;
-            z = opt_maturity * sig_sig_two;
-
-            c = quarter * z;
-            e = sycl::exp(b);
-            y = sycl::rsqrt(z);
-
-            a = b - a;
-            w1 = (a - c) * y;
-            w2 = (a + c) * y;
-
-            if (w1 < zero) {
-                d1 = sycl::erfc(w1) * half;
-                d1c = one - d1;
-            }
-            else {
-                d1c = sycl::erfc(-w1) * half;
-                d1 = one - d1c;
-            }
-            if (w2 < zero) {
-                d2 = sycl::erfc(w2) * half;
-                d2c = one - d2;
-            }
-            else {
-                d2c = sycl::erfc(-w2) * half;
-                d2 = one - d2c;
-            }
-
-            e *= opt_strike;
-            data_t call_price = opt_price * d1 - e * d2;
-            data_t put_price = e * d2c - opt_price * d1c;
-
-            const size_t callput_i = n_prices * idx[0];
-            callput[callput_i + CALL] = call_price;
-            callput[callput_i + PUT] = put_price;
-        });
+            sycl::range<1>(n_opts), [=](sycl::id<1> idx) {
+                const size_t i = n_params * idx[0];
+                const data_t opt_price = params[i + PRICE];
+                const data_t opt_strike = params[i + STRIKE];
+                const data_t opt_maturity = params[i + MATURITY];
+                const data_t opt_rate = params[i + RATE];
+                const data_t opt_volatility = params[i + VOLATILITY];
+                data_t a, b, c, y, z, e, d1, d1c, d2, d2c, w1, w2;
+                data_t mr = -opt_rate,
+                       sig_sig_two = two * opt_volatility * opt_volatility;
+
+                a = sycl::log(opt_price / opt_strike);
+                b = opt_maturity * mr;
+                z = opt_maturity * sig_sig_two;
+
+                c = quarter * z;
+                e = sycl::exp(b);
+                y = sycl::rsqrt(z);
+
+                a = b - a;
+                w1 = (a - c) * y;
+                w2 = (a + c) * y;
+
+                if (w1 < zero) {
+                    d1 = sycl::erfc(w1) * half;
+                    d1c = one - d1;
+                }
+                else {
+                    d1c = sycl::erfc(-w1) * half;
+                    d1 = one - d1c;
+                }
+                if (w2 < zero) {
+                    d2 = sycl::erfc(w2) * half;
+                    d2c = one - d2;
+                }
+                else {
+                    d2c = sycl::erfc(-w2) * half;
+                    d2 = one - d2c;
+                }
+
+                e *= opt_strike;
+                data_t call_price = opt_price * d1 - e * d2;
+                data_t put_price = e * d2c - opt_price * d1c;
+
+                const size_t callput_i = n_prices * idx[0];
+                callput[callput_i + CALL] = call_price;
+                callput[callput_i + PUT] = put_price;
+            });
     });
 
     e.wait_and_throw();
diff --git a/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp b/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp
@@ -52,27 +52,12 @@ struct DMatrix
     DMatrix(const DMatrix &) = default;
     DMatrix(DMatrix &&) = default;
 
-    size_t get_n() const
-    {
-        return n_;
-    }
-    size_t get_m() const
-    {
-        return m_;
-    }
-    vec_t &get_vector()
-    {
-        return vec_;
-    }
-    sycl::queue get_queue() const
-    {
-        return q_;
-    }
+    size_t get_n() const { return n_; }
+    size_t get_m() const { return m_; }
+    vec_t &get_vector() { return vec_; }
+    sycl::queue get_queue() const { return q_; }
 
-    double get_element(size_t i, size_t j)
-    {
-        return vec_.at(i * m_ + j);
-    }
+    double get_element(size_t i, size_t j) { return vec_.at(i * m_ + j); }
 
 private:
     size_t n_;
@@ -114,6 +99,9 @@ py::dict construct_sua_iface(DMatrix &m)
     iface["typestr"] = "|f8";
     iface["syclobj"] = syclobj;
 
+    // ensure that content of array is flushed out
+    m.get_queue().wait();
+
     return iface;
 }
 
diff --git a/examples/pybind11/onemkl_gemv/solve.py b/examples/pybind11/onemkl_gemv/solve.py
@@ -127,12 +127,14 @@ def cg_solve(A, b):
     converged is False if solver has not converged, or the iteration number
     """
     exec_queue = A.sycl_queue
+    exec_queue.wait()
+
     x = dpt.zeros_like(b)
     Ap = dpt.empty_like(x)
 
     all_host_tasks = []
-    r = dpt.copy(b)  # synchronous copy
-    p = dpt.copy(b)  # synchronous copy
+    r = dpt.copy(b)
+    p = dpt.copy(b)
 
     rsold = sycl_gemm.norm_squared_blocking(exec_queue, r)
     if rsold < 1e-20:
diff --git a/examples/pybind11/onemkl_gemv/sycl_gemm/_onemkl.cpp b/examples/pybind11/onemkl_gemv/sycl_gemm/_onemkl.cpp
@@ -41,14 +41,15 @@ namespace py = pybind11;
 using dpctl::utils::keep_args_alive;
 
 std::pair<sycl::event, sycl::event>
-py_gemv(sycl::queue q,
+py_gemv(sycl::queue &q,
         dpctl::tensor::usm_ndarray matrix,
         dpctl::tensor::usm_ndarray vector,
         dpctl::tensor::usm_ndarray result,
         const std::vector<sycl::event> &depends = {})
 {
     if (matrix.get_ndim() != 2 || vector.get_ndim() != 1 ||
-        result.get_ndim() != 1) {
+        result.get_ndim() != 1)
+    {
         throw std::runtime_error(
             "Inconsistent dimensions, expecting matrix and a vector");
     }
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py b/examples/pybind11/use_dpctl_sycl_kernel/tests/test_user_kernel.py
@@ -64,6 +64,7 @@ def test_kernel_submit_through_extension():
     x = dpt.arange(0, stop=13, step=1, dtype="i4", sycl_queue=q)
     y = dpt.zeros_like(x)
 
+    q.wait()
     uk.submit_custom_kernel(q, krn, x, y, [])
 
     assert np.array_equal(dpt.asnumpy(y), np.arange(0, 26, step=2, dtype="i4"))
diff --git a/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp b/examples/pybind11/use_dpctl_sycl_kernel/use_kernel/_example.cpp
@@ -35,8 +35,8 @@
 
 namespace py = pybind11;
 
-void submit_custom_kernel(sycl::queue q,
-                          sycl::kernel krn,
+void submit_custom_kernel(sycl::queue &q,
+                          sycl::kernel &krn,
                           dpctl::tensor::usm_ndarray x,
                           dpctl::tensor::usm_ndarray y,
                           const std::vector<sycl::event> &depends = {})
