From 2552f7c6298859b3c1db9051a444573bcdd52f26 Mon Sep 17 00:00:00 2001 From: Umesh Unnikrishnan Date: Tue, 19 Dec 2023 20:37:03 +0000 Subject: [PATCH] Adds missing changes from previous git-cherry-pick --- .../sycl-ref/ceed-sycl-ref-operator.sycl.cpp | 25 ++++++++++--------- .../ceed-sycl-ref-qfunctioncontext.sycl.cpp | 6 ++--- backends/sycl-ref/ceed-sycl-vector.sycl.cpp | 8 ++---- .../ceed-sycl-shared-basis.sycl.cpp | 2 +- 4 files changed, 19 insertions(+), 22 deletions(-) diff --git a/backends/sycl-ref/ceed-sycl-ref-operator.sycl.cpp b/backends/sycl-ref/ceed-sycl-ref-operator.sycl.cpp index b06e622bad..81883b326f 100644 --- a/backends/sycl-ref/ceed-sycl-ref-operator.sycl.cpp +++ b/backends/sycl-ref/ceed-sycl-ref-operator.sycl.cpp @@ -762,39 +762,39 @@ static inline int CeedOperatorAssembleDiagonalSetup_Sycl(CeedOperator op) { CeedCallBackend(CeedCalloc(num_qpts * num_nodes, &identity)); for (CeedSize i = 0; i < (num_nodes < num_qpts ? num_nodes : num_qpts); i++) identity[i * num_nodes + i] = 1.0; CeedCallSycl(ceed, diag->d_identity = sycl::malloc_device(i_len, sycl_data->sycl_device, sycl_data->sycl_context)); - sycl::event identity_copy = sycl_data->sycl_queue.copy(identity, diag->d_identity, i_len, {e}); + sycl::event identity_copy = sycl_data->sycl_queue.copy(identity, diag->d_identity, i_len, e); copy_events.push_back(identity_copy); } // CEED_EVAL_INTERP CeedCallBackend(CeedBasisGetInterp(basis_in, &interp_in)); CeedCallSycl(ceed, diag->d_interp_in = sycl::malloc_device(i_len, sycl_data->sycl_device, sycl_data->sycl_context)); - sycl::event interp_in_copy = sycl_data->sycl_queue.copy(interp_in, diag->d_interp_in, i_len, {e}); + sycl::event interp_in_copy = sycl_data->sycl_queue.copy(interp_in, diag->d_interp_in, i_len, e); copy_events.push_back(interp_in_copy); CeedCallBackend(CeedBasisGetInterp(basis_out, &interp_out)); CeedCallSycl(ceed, diag->d_interp_out = sycl::malloc_device(i_len, sycl_data->sycl_device, sycl_data->sycl_context)); - sycl::event interp_out_copy = sycl_data->sycl_queue.copy(interp_out, diag->d_interp_out, i_len, {e}); + sycl::event interp_out_copy = sycl_data->sycl_queue.copy(interp_out, diag->d_interp_out, i_len, e); copy_events.push_back(interp_out_copy); // CEED_EVAL_GRAD CeedCallBackend(CeedBasisGetGrad(basis_in, &grad_in)); CeedCallSycl(ceed, diag->d_grad_in = sycl::malloc_device(g_len, sycl_data->sycl_device, sycl_data->sycl_context)); - sycl::event grad_in_copy = sycl_data->sycl_queue.copy(grad_in, diag->d_grad_in, g_len, {e}); + sycl::event grad_in_copy = sycl_data->sycl_queue.copy(grad_in, diag->d_grad_in, g_len, e); copy_events.push_back(grad_in_copy); CeedCallBackend(CeedBasisGetGrad(basis_out, &grad_out)); CeedCallSycl(ceed, diag->d_grad_out = sycl::malloc_device(g_len, sycl_data->sycl_device, sycl_data->sycl_context)); - sycl::event grad_out_copy = sycl_data->sycl_queue.copy(grad_out, diag->d_grad_out, g_len, {e}); + sycl::event grad_out_copy = sycl_data->sycl_queue.copy(grad_out, diag->d_grad_out, g_len, e); copy_events.push_back(grad_out_copy); // Arrays of e_modes CeedCallSycl(ceed, diag->d_e_mode_in = sycl::malloc_device(num_e_mode_in, sycl_data->sycl_device, sycl_data->sycl_context)); - sycl::event e_mode_in_copy = sycl_data->sycl_queue.copy(e_mode_in, diag->d_e_mode_in, num_e_mode_in, {e}); + sycl::event e_mode_in_copy = sycl_data->sycl_queue.copy(e_mode_in, diag->d_e_mode_in, num_e_mode_in, e); copy_events.push_back(e_mode_in_copy); CeedCallSycl(ceed, diag->d_e_mode_out = sycl::malloc_device(num_e_mode_out, sycl_data->sycl_device, sycl_data->sycl_context)); - sycl::event e_mode_out_copy = sycl_data->sycl_queue.copy(e_mode_out, diag->d_e_mode_out, num_e_mode_out, {e}); + sycl::event e_mode_out_copy = sycl_data->sycl_queue.copy(e_mode_out, diag->d_e_mode_out, num_e_mode_out, e); copy_events.push_back(e_mode_out_copy); // Restriction @@ -1109,11 +1109,12 @@ static int CeedSingleOperatorAssembleSetup_Sycl(CeedOperator op) { if (eval_mode == CEED_EVAL_INTERP) { std::vector e; if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()}; - sycl_data->sycl_queue.copy(interp_in, &asmb->d_B_in[mat_start], elem_size * num_qpts, {e}); + sycl_data->sycl_queue.copy(interp_in, &asmb->d_B_in[mat_start], elem_size * num_qpts, e); mat_start += elem_size * num_qpts; } else if (eval_mode == CEED_EVAL_GRAD) { - td::vector e; + std::vector e; if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()}; + sycl_data->sycl_queue.copy(grad_in, &asmb->d_B_in[mat_start], dim * elem_size * num_qpts, e); mat_start += dim * elem_size * num_qpts; } } @@ -1138,12 +1139,12 @@ static int CeedSingleOperatorAssembleSetup_Sycl(CeedOperator op) { if (eval_mode == CEED_EVAL_INTERP) { std::vector e; if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()}; - sycl_data->sycl_queue.copy(interp_out, &asmb->d_B_out[mat_start], elem_size * num_qpts, {e}); + sycl_data->sycl_queue.copy(interp_out, &asmb->d_B_out[mat_start], elem_size * num_qpts, e); mat_start += elem_size * num_qpts; } else if (eval_mode == CEED_EVAL_GRAD) { std::vector e; if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()}; - sycl_data->sycl_queue.copy(grad_out, &asmb->d_B_out[mat_start], dim * elem_size * num_qpts, {e}); + sycl_data->sycl_queue.copy(grad_out, &asmb->d_B_out[mat_start], dim * elem_size * num_qpts, e); mat_start += dim * elem_size * num_qpts; } } @@ -1187,7 +1188,7 @@ static int CeedOperatorLinearAssemble_Sycl(sycl::queue &sycl_queue, const CeedOp std::vector e; if (!sycl_queue.is_in_order()) e = {sycl_queue.ext_oneapi_submit_barrier()}; - sycl_queue.parallel_for(kernel_range, {e}, [=](sycl::id<3> idx) { + sycl_queue.parallel_for(kernel_range, e, [=](sycl::id<3> idx) { const int e = idx.get(0); // Element index const int l = idx.get(1); // The output column index of each B^TDB operation const int i = idx.get(2); // The output row index of each B^TDB operation diff --git a/backends/sycl-ref/ceed-sycl-ref-qfunctioncontext.sycl.cpp b/backends/sycl-ref/ceed-sycl-ref-qfunctioncontext.sycl.cpp index 0451aab364..fca40fe21e 100644 --- a/backends/sycl-ref/ceed-sycl-ref-qfunctioncontext.sycl.cpp +++ b/backends/sycl-ref/ceed-sycl-ref-qfunctioncontext.sycl.cpp @@ -44,7 +44,7 @@ static inline int CeedQFunctionContextSyncH2D_Sycl(const CeedQFunctionContext ct } std::vector e; if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()}; - sycl::event copy_event = sycl_data->sycl_queue.memcpy(impl->d_data, impl->h_data, ctx_size, {e}); + sycl::event copy_event = sycl_data->sycl_queue.memcpy(impl->d_data, impl->h_data, ctx_size, e); CeedCallSycl(ceed, copy_event.wait_and_throw()); return CEED_ERROR_SUCCESS; } @@ -81,7 +81,7 @@ static inline int CeedQFunctionContextSyncD2H_Sycl(const CeedQFunctionContext ct std::vector e; if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()}; - sycl::event copy_event = sycl_data->sycl_queue.memcpy(impl->h_data, impl->d_data, ctx_size, {e}); + sycl::event copy_event = sycl_data->sycl_queue.memcpy(impl->h_data, impl->d_data, ctx_size, e); CeedCallSycl(ceed, copy_event.wait_and_throw()); return CEED_ERROR_SUCCESS; } @@ -222,7 +222,7 @@ static int CeedQFunctionContextSetDataDevice_Sycl(const CeedQFunctionContext ctx CeedCallSycl(ceed, impl->d_data_owned = sycl::malloc_device(ctx_size, sycl_data->sycl_device, sycl_data->sycl_context)); impl->d_data_borrowed = NULL; impl->d_data = impl->d_data_owned; - sycl::event copy_event = sycl_data->sycl_queue.memcpy(impl->d_data, data, ctx_size, {e}); + sycl::event copy_event = sycl_data->sycl_queue.memcpy(impl->d_data, data, ctx_size, e); CeedCallSycl(ceed, copy_event.wait_and_throw()); } break; case CEED_OWN_POINTER: { diff --git a/backends/sycl-ref/ceed-sycl-vector.sycl.cpp b/backends/sycl-ref/ceed-sycl-vector.sycl.cpp index 53a8e86e8b..1e21ca8bca 100644 --- a/backends/sycl-ref/ceed-sycl-vector.sycl.cpp +++ b/backends/sycl-ref/ceed-sycl-vector.sycl.cpp @@ -66,9 +66,7 @@ static inline int CeedVectorSyncH2D_Sycl(const CeedVector vec) { // Copy from host to device std::vector e; if (!data->sycl_queue.is_in_order()) e = {data->sycl_queue.ext_oneapi_submit_barrier()}; - sycl::event copy_event = data->sycl_queue.copy(impl->h_array, impl->d_array, length, {e}); - // Wait for copy to finish and handle exceptions. - CeedCallSycl(ceed, copy_event.wait_and_throw()); + CeedCallSycl(ceed, data->sycl_queue.copy(impl->h_array, impl->d_array, length, e).wait_and_throw()); return CEED_ERROR_SUCCESS; } @@ -100,9 +98,7 @@ static inline int CeedVectorSyncD2H_Sycl(const CeedVector vec) { // Copy from device to host std::vector e; if (!data->sycl_queue.is_in_order()) e = {data->sycl_queue.ext_oneapi_submit_barrier()}; - sycl::event copy_event = data->sycl_queue.copy(impl->d_array, impl->h_array, length, {e}); - // Wait for copy to finish and handle exceptions. - CeedCallSycl(ceed, copy_event.wait_and_throw()); + CeedCallSycl(ceed,data->sycl_queue.copy(impl->d_array, impl->h_array, length, e).wait_and_throw()); return CEED_ERROR_SUCCESS; } diff --git a/backends/sycl-shared/ceed-sycl-shared-basis.sycl.cpp b/backends/sycl-shared/ceed-sycl-shared-basis.sycl.cpp index 6e623e3d70..7dfe940116 100644 --- a/backends/sycl-shared/ceed-sycl-shared-basis.sycl.cpp +++ b/backends/sycl-shared/ceed-sycl-shared-basis.sycl.cpp @@ -66,7 +66,7 @@ int CeedBasisApplyTensor_Sycl_shared(CeedBasis basis, const CeedInt num_elem, Ce std::vector e; if (!ceed_Sycl->sycl_queue.is_in_order()) e = {ceed_Sycl->sycl_queue.ext_oneapi_submit_barrier()}; - + ceed_Sycl->sycl_queue.submit([&](sycl::handler &cgh) { cgh.depends_on(e); cgh.set_args(num_elem, impl->d_interp_1d, d_u, d_v);