Skip to content

Commit

Permalink
Adds missing changes from previous git-cherry-pick
Browse files Browse the repository at this point in the history
  • Loading branch information
@uumesh
uumesh committed Dec 19, 2023
1 parent 10c7a55 commit 2552f7c
Show file tree
Hide file tree
Showing 4 changed files with 19 additions and 22 deletions.
25 changes: 13 additions & 12 deletions backends/sycl-ref/ceed-sycl-ref-operator.sycl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<CeedScalar>(i_len, sycl_data->sycl_device, sycl_data->sycl_context));
sycl::event identity_copy = sycl_data->sycl_queue.copy<CeedScalar>(identity, diag->d_identity, i_len, {e});
sycl::event identity_copy = sycl_data->sycl_queue.copy<CeedScalar>(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<CeedScalar>(i_len, sycl_data->sycl_device, sycl_data->sycl_context));
sycl::event interp_in_copy = sycl_data->sycl_queue.copy<CeedScalar>(interp_in, diag->d_interp_in, i_len, {e});
sycl::event interp_in_copy = sycl_data->sycl_queue.copy<CeedScalar>(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<CeedScalar>(i_len, sycl_data->sycl_device, sycl_data->sycl_context));
sycl::event interp_out_copy = sycl_data->sycl_queue.copy<CeedScalar>(interp_out, diag->d_interp_out, i_len, {e});
sycl::event interp_out_copy = sycl_data->sycl_queue.copy<CeedScalar>(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<CeedScalar>(g_len, sycl_data->sycl_device, sycl_data->sycl_context));
sycl::event grad_in_copy = sycl_data->sycl_queue.copy<CeedScalar>(grad_in, diag->d_grad_in, g_len, {e});
sycl::event grad_in_copy = sycl_data->sycl_queue.copy<CeedScalar>(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<CeedScalar>(g_len, sycl_data->sycl_device, sycl_data->sycl_context));
sycl::event grad_out_copy = sycl_data->sycl_queue.copy<CeedScalar>(grad_out, diag->d_grad_out, g_len, {e});
sycl::event grad_out_copy = sycl_data->sycl_queue.copy<CeedScalar>(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<CeedEvalMode>(num_e_mode_in, sycl_data->sycl_device, sycl_data->sycl_context));
sycl::event e_mode_in_copy = sycl_data->sycl_queue.copy<CeedEvalMode>(e_mode_in, diag->d_e_mode_in, num_e_mode_in, {e});
sycl::event e_mode_in_copy = sycl_data->sycl_queue.copy<CeedEvalMode>(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<CeedEvalMode>(num_e_mode_out, sycl_data->sycl_device, sycl_data->sycl_context));
sycl::event e_mode_out_copy = sycl_data->sycl_queue.copy<CeedEvalMode>(e_mode_out, diag->d_e_mode_out, num_e_mode_out, {e});
sycl::event e_mode_out_copy = sycl_data->sycl_queue.copy<CeedEvalMode>(e_mode_out, diag->d_e_mode_out, num_e_mode_out, e);
copy_events.push_back(e_mode_out_copy);

// Restriction
Expand Down Expand Up @@ -1109,11 +1109,12 @@ static int CeedSingleOperatorAssembleSetup_Sycl(CeedOperator op) {
if (eval_mode == CEED_EVAL_INTERP) {
std::vector<sycl::event> e;
if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()};
sycl_data->sycl_queue.copy<CeedScalar>(interp_in, &asmb->d_B_in[mat_start], elem_size * num_qpts, {e});
sycl_data->sycl_queue.copy<CeedScalar>(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<sycl::event> e;
std::vector<sycl::event> e;
if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()};
sycl_data->sycl_queue.copy<CeedScalar>(grad_in, &asmb->d_B_in[mat_start], dim * elem_size * num_qpts, e);
mat_start += dim * elem_size * num_qpts;
}
}
Expand All @@ -1138,12 +1139,12 @@ static int CeedSingleOperatorAssembleSetup_Sycl(CeedOperator op) {
if (eval_mode == CEED_EVAL_INTERP) {
std::vector<sycl::event> e;
if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()};
sycl_data->sycl_queue.copy<CeedScalar>(interp_out, &asmb->d_B_out[mat_start], elem_size * num_qpts, {e});
sycl_data->sycl_queue.copy<CeedScalar>(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<sycl::event> e;
if (!sycl_data->sycl_queue.is_in_order()) e = {sycl_data->sycl_queue.ext_oneapi_submit_barrier()};
sycl_data->sycl_queue.copy<CeedScalar>(grad_out, &asmb->d_B_out[mat_start], dim * elem_size * num_qpts, {e});
sycl_data->sycl_queue.copy<CeedScalar>(grad_out, &asmb->d_B_out[mat_start], dim * elem_size * num_qpts, e);
mat_start += dim * elem_size * num_qpts;
}
}
Expand Down Expand Up @@ -1187,7 +1188,7 @@ static int CeedOperatorLinearAssemble_Sycl(sycl::queue &sycl_queue, const CeedOp

std::vector<sycl::event> e;
if (!sycl_queue.is_in_order()) e = {sycl_queue.ext_oneapi_submit_barrier()};
sycl_queue.parallel_for<CeedOperatorSyclLinearAssemble>(kernel_range, {e}, [=](sycl::id<3> idx) {
sycl_queue.parallel_for<CeedOperatorSyclLinearAssemble>(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
Expand Down
6 changes: 3 additions & 3 deletions backends/sycl-ref/ceed-sycl-ref-qfunctioncontext.sycl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -44,7 +44,7 @@ static inline int CeedQFunctionContextSyncH2D_Sycl(const CeedQFunctionContext ct
}
std::vector<sycl::event> 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;
}
Expand Down Expand Up @@ -81,7 +81,7 @@ static inline int CeedQFunctionContextSyncD2H_Sycl(const CeedQFunctionContext ct

std::vector<sycl::event> 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;
}
Expand Down Expand Up @@ -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: {
Expand Down
8 changes: 2 additions & 6 deletions backends/sycl-ref/ceed-sycl-vector.sycl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -66,9 +66,7 @@ static inline int CeedVectorSyncH2D_Sycl(const CeedVector vec) {
// Copy from host to device
std::vector<sycl::event> e;
if (!data->sycl_queue.is_in_order()) e = {data->sycl_queue.ext_oneapi_submit_barrier()};
sycl::event copy_event = data->sycl_queue.copy<CeedScalar>(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<CeedScalar>(impl->h_array, impl->d_array, length, e).wait_and_throw());
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -100,9 +98,7 @@ static inline int CeedVectorSyncD2H_Sycl(const CeedVector vec) {
// Copy from device to host
std::vector<sycl::event> e;
if (!data->sycl_queue.is_in_order()) e = {data->sycl_queue.ext_oneapi_submit_barrier()};
sycl::event copy_event = data->sycl_queue.copy<CeedScalar>(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<CeedScalar>(impl->d_array, impl->h_array, length, e).wait_and_throw());
return CEED_ERROR_SUCCESS;
}

Expand Down
2 changes: 1 addition & 1 deletion backends/sycl-shared/ceed-sycl-shared-basis.sycl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -66,7 +66,7 @@ int CeedBasisApplyTensor_Sycl_shared(CeedBasis basis, const CeedInt num_elem, Ce

std::vector<sycl::event> 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);
Expand Down

0 comments on commit 2552f7c

Please sign in to comment.