80   if (impl->diag) {  in CeedOperatorDestroy_Sycl()
81     CeedCallBackend(CeedFree(&impl->diag->h_eval_mode_in));  in CeedOperatorDestroy_Sycl()
82     CeedCallBackend(CeedFree(&impl->diag->h_eval_mode_out));  in CeedOperatorDestroy_Sycl()
85     CeedCallSycl(ceed, sycl::free(impl->diag->d_eval_mode_in, sycl_data->sycl_context));  in CeedOperatorDestroy_Sycl()
86     CeedCallSycl(ceed, sycl::free(impl->diag->d_eval_mode_out, sycl_data->sycl_context));  in CeedOperatorDestroy_Sycl()
87     CeedCallSycl(ceed, sycl::free(impl->diag->d_identity, sycl_data->sycl_context));  in CeedOperatorDestroy_Sycl()
88     CeedCallSycl(ceed, sycl::free(impl->diag->d_interp_in, sycl_data->sycl_context));  in CeedOperatorDestroy_Sycl()
89     CeedCallSycl(ceed, sycl::free(impl->diag->d_interp_out, sycl_data->sycl_context));  in CeedOperatorDestroy_Sycl()
90     CeedCallSycl(ceed, sycl::free(impl->diag->d_grad_in, sycl_data->sycl_context));  in CeedOperatorDestroy_Sycl()
91     CeedCallSycl(ceed, sycl::free(impl->diag->d_grad_out, sycl_data->sycl_context));  in CeedOperatorDestroy_Sycl()
93     CeedCallBackend(CeedVectorDestroy(&impl->diag->elem_diag));  in CeedOperatorDestroy_Sycl()
94     CeedCallBackend(CeedVectorDestroy(&impl->diag->point_block_elem_diag));  in CeedOperatorDestroy_Sycl()
95     CeedCallBackend(CeedElemRestrictionDestroy(&impl->diag->diag_rstr));  in CeedOperatorDestroy_Sycl()
96     CeedCallBackend(CeedElemRestrictionDestroy(&impl->diag->point_block_diag_rstr));  in CeedOperatorDestroy_Sycl()
97     CeedCallBackend(CeedBasisDestroy(&impl->diag->basis_in));  in CeedOperatorDestroy_Sycl()
98     CeedCallBackend(CeedBasisDestroy(&impl->diag->basis_out));  in CeedOperatorDestroy_Sycl()
100   CeedCallBackend(CeedFree(&impl->diag));  in CeedOperatorDestroy_Sycl()
739   CeedCallBackend(CeedCalloc(1, &impl->diag));  in CeedOperatorAssembleDiagonalSetup_Sycl()
740   CeedOperatorDiag_Sycl *diag = impl->diag;  in CeedOperatorAssembleDiagonalSetup_Sycl()  local
742   CeedCallBackend(CeedBasisReferenceCopy(basis_in, &diag->basis_in));  in CeedOperatorAssembleDiagonalSetup_Sycl()
743   CeedCallBackend(CeedBasisReferenceCopy(basis_out, &diag->basis_out));  in CeedOperatorAssembleDiagonalSetup_Sycl()
744   diag->h_eval_mode_in    = eval_mode_in;  in CeedOperatorAssembleDiagonalSetup_Sycl()
745   diag->h_eval_mode_out   = eval_mode_out;  in CeedOperatorAssembleDiagonalSetup_Sycl()
746   diag->num_eval_mode_in  = num_eval_mode_in;  in CeedOperatorAssembleDiagonalSetup_Sycl()
747   diag->num_eval_mode_out = num_eval_mode_out;  in CeedOperatorAssembleDiagonalSetup_Sycl()
754   diag->num_nodes = num_nodes;  in CeedOperatorAssembleDiagonalSetup_Sycl()
755   diag->num_qpts  = num_qpts;  in CeedOperatorAssembleDiagonalSetup_Sycl()
756   diag->num_comp  = num_comp;  in CeedOperatorAssembleDiagonalSetup_Sycl()
778 …CeedCallSycl(ceed, diag->d_identity = sycl::malloc_device<CeedScalar>(i_len, sycl_data->sycl_devic…  in CeedOperatorAssembleDiagonalSetup_Sycl()
779 …sycl::event identity_copy = sycl_data->sycl_queue.copy<CeedScalar>(identity, diag->d_identity, i_l…  in CeedOperatorAssembleDiagonalSetup_Sycl()
785 …CeedCallSycl(ceed, diag->d_interp_in = sycl::malloc_device<CeedScalar>(i_len, sycl_data->sycl_devi…  in CeedOperatorAssembleDiagonalSetup_Sycl()
786 …sycl::event interp_in_copy = sycl_data->sycl_queue.copy<CeedScalar>(interp_in, diag->d_interp_in, …  in CeedOperatorAssembleDiagonalSetup_Sycl()
790 …CeedCallSycl(ceed, diag->d_interp_out = sycl::malloc_device<CeedScalar>(i_len, sycl_data->sycl_dev…  in CeedOperatorAssembleDiagonalSetup_Sycl()
791 …sycl::event interp_out_copy = sycl_data->sycl_queue.copy<CeedScalar>(interp_out, diag->d_interp_ou…  in CeedOperatorAssembleDiagonalSetup_Sycl()
796 …CeedCallSycl(ceed, diag->d_grad_in = sycl::malloc_device<CeedScalar>(g_len, sycl_data->sycl_device…  in CeedOperatorAssembleDiagonalSetup_Sycl()
797 …sycl::event grad_in_copy = sycl_data->sycl_queue.copy<CeedScalar>(grad_in, diag->d_grad_in, g_len,…  in CeedOperatorAssembleDiagonalSetup_Sycl()
801 …CeedCallSycl(ceed, diag->d_grad_out = sycl::malloc_device<CeedScalar>(g_len, sycl_data->sycl_devic…  in CeedOperatorAssembleDiagonalSetup_Sycl()
802 …sycl::event grad_out_copy = sycl_data->sycl_queue.copy<CeedScalar>(grad_out, diag->d_grad_out, g_l…  in CeedOperatorAssembleDiagonalSetup_Sycl()
806 …CeedCallSycl(ceed, diag->d_eval_mode_in = sycl::malloc_device<CeedEvalMode>(num_eval_mode_in, sycl…  in CeedOperatorAssembleDiagonalSetup_Sycl()
807 …sycl::event eval_mode_in_copy = sycl_data->sycl_queue.copy<CeedEvalMode>(eval_mode_in, diag->d_eva…  in CeedOperatorAssembleDiagonalSetup_Sycl()
810 …CeedCallSycl(ceed, diag->d_eval_mode_out = sycl::malloc_device<CeedEvalMode>(num_eval_mode_out, sy…  in CeedOperatorAssembleDiagonalSetup_Sycl()
811 …sycl::event eval_mode_out_copy = sycl_data->sycl_queue.copy<CeedEvalMode>(eval_mode_out, diag->d_e…  in CeedOperatorAssembleDiagonalSetup_Sycl()
815   CeedCallBackend(CeedElemRestrictionReferenceCopy(rstr_out, &diag->diag_rstr));  in CeedOperatorAssembleDiagonalSetup_Sycl()
833 …const CeedOperatorDiag_Sycl *diag, const CeedScalar *assembled_qf_array, CeedScalar *elem_diag_arr…  in CeedOperatorLinearDiagonal_Sycl()  argument
834   const CeedSize      num_nodes         = diag->num_nodes;  in CeedOperatorLinearDiagonal_Sycl()
835   const CeedSize      num_qpts          = diag->num_qpts;  in CeedOperatorLinearDiagonal_Sycl()
836   const CeedSize      num_comp          = diag->num_comp;  in CeedOperatorLinearDiagonal_Sycl()
837   const CeedSize      num_eval_mode_in  = diag->num_eval_mode_in;  in CeedOperatorLinearDiagonal_Sycl()
838   const CeedSize      num_eval_mode_out = diag->num_eval_mode_out;  in CeedOperatorLinearDiagonal_Sycl()
839   const CeedScalar   *identity          = diag->d_identity;  in CeedOperatorLinearDiagonal_Sycl()
840   const CeedScalar   *interp_in         = diag->d_interp_in;  in CeedOperatorLinearDiagonal_Sycl()
841   const CeedScalar   *grad_in           = diag->d_grad_in;  in CeedOperatorLinearDiagonal_Sycl()
842   const CeedScalar   *interp_out        = diag->d_interp_out;  in CeedOperatorLinearDiagonal_Sycl()
843   const CeedScalar   *grad_out          = diag->d_grad_out;  in CeedOperatorLinearDiagonal_Sycl()
844   const CeedEvalMode *eval_mode_in      = diag->d_eval_mode_in;  in CeedOperatorLinearDiagonal_Sycl()
845   const CeedEvalMode *eval_mode_out     = diag->d_eval_mode_out;  in CeedOperatorLinearDiagonal_Sycl()
935   if (!impl->diag) {  in CeedOperatorAssembleDiagonalCore_Sycl()
938   CeedOperatorDiag_Sycl *diag = impl->diag;  in CeedOperatorAssembleDiagonalCore_Sycl()  local
940   assert(diag != NULL);  in CeedOperatorAssembleDiagonalCore_Sycl()
943   if (is_point_block && !diag->point_block_diag_rstr) {  in CeedOperatorAssembleDiagonalCore_Sycl()
944 …CeedCallBackend(CeedOperatorCreateActivePointBlockRestriction(diag->diag_rstr, &diag->point_block_…  in CeedOperatorAssembleDiagonalCore_Sycl()
946   CeedElemRestriction diag_rstr = is_point_block ? diag->point_block_diag_rstr : diag->diag_rstr;  in CeedOperatorAssembleDiagonalCore_Sycl()
949   CeedVector elem_diag = is_point_block ? diag->point_block_elem_diag : diag->elem_diag;  in CeedOperatorAssembleDiagonalCore_Sycl()
953     if (is_point_block) diag->point_block_elem_diag = elem_diag;  in CeedOperatorAssembleDiagonalCore_Sycl()
954     else diag->elem_diag = elem_diag;  in CeedOperatorAssembleDiagonalCore_Sycl()
964 …ratorLinearDiagonal_Sycl(sycl_data->sycl_queue, is_point_block, num_elem, diag, assembled_qf_array…  in CeedOperatorAssembleDiagonalCore_Sycl()