42   CeedBasis_Sycl_shared *impl;  in CeedBasisApplyTensor_Sycl_shared()  local
46   CeedCallBackend(CeedBasisGetData(basis, &impl));  in CeedBasisApplyTensor_Sycl_shared()
56       CeedInt       *lrange         = impl->interp_local_range;  in CeedBasisApplyTensor_Sycl_shared()
64 …ycl::kernel *interp_kernel = (t_mode == CEED_TRANSPOSE) ? impl->interp_transpose_kernel : impl->in…  in CeedBasisApplyTensor_Sycl_shared()
71         cgh.set_args(num_elem, impl->d_interp_1d, d_u, d_v);  in CeedBasisApplyTensor_Sycl_shared()
77       CeedInt       *lrange         = impl->grad_local_range;  in CeedBasisApplyTensor_Sycl_shared()
85 …cl::kernel     *grad_kernel = (t_mode == CEED_TRANSPOSE) ? impl->grad_transpose_kernel : impl->gra…  in CeedBasisApplyTensor_Sycl_shared()
86 … const CeedScalar *d_grad_1d   = (impl->d_collo_grad_1d) ? impl->d_collo_grad_1d : impl->d_grad_1d;  in CeedBasisApplyTensor_Sycl_shared()
94         cgh.set_args(num_elem, impl->d_interp_1d, d_grad_1d, d_u, d_v);  in CeedBasisApplyTensor_Sycl_shared()
99       CeedInt       *lrange         = impl->weight_local_range;  in CeedBasisApplyTensor_Sycl_shared()
109 …CeedCheck(impl->d_q_weight_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weight_1d not set", …  in CeedBasisApplyTensor_Sycl_shared()
114         cgh.set_args(num_elem, impl->d_q_weight_1d, d_v);  in CeedBasisApplyTensor_Sycl_shared()
115         cgh.parallel_for(kernel_range, *(impl->weight_kernel));  in CeedBasisApplyTensor_Sycl_shared()
141   CeedBasis_Sycl_shared *impl;  in CeedBasisDestroy_Sycl_shared()  local
144   CeedCallBackend(CeedBasisGetData(basis, &impl));  in CeedBasisDestroy_Sycl_shared()
147   if (impl->d_q_weight_1d) CeedCallSycl(ceed, sycl::free(impl->d_q_weight_1d, data->sycl_context));  in CeedBasisDestroy_Sycl_shared()
148   CeedCallSycl(ceed, sycl::free(impl->d_interp_1d, data->sycl_context));  in CeedBasisDestroy_Sycl_shared()
149   CeedCallSycl(ceed, sycl::free(impl->d_grad_1d, data->sycl_context));  in CeedBasisDestroy_Sycl_shared()
150   CeedCallSycl(ceed, sycl::free(impl->d_collo_grad_1d, data->sycl_context));  in CeedBasisDestroy_Sycl_shared()
152   delete impl->interp_kernel;  in CeedBasisDestroy_Sycl_shared()
153   delete impl->interp_transpose_kernel;  in CeedBasisDestroy_Sycl_shared()
154   delete impl->grad_kernel;  in CeedBasisDestroy_Sycl_shared()
155   delete impl->grad_transpose_kernel;  in CeedBasisDestroy_Sycl_shared()
156   delete impl->weight_kernel;  in CeedBasisDestroy_Sycl_shared()
157   delete impl->sycl_module;  in CeedBasisDestroy_Sycl_shared()
159   CeedCallBackend(CeedFree(&impl));  in CeedBasisDestroy_Sycl_shared()
175   CeedBasis_Sycl_shared *impl;  in CeedBasisCreateTensorH1_Sycl_shared()  local
178   CeedCallBackend(CeedCalloc(1, &impl));  in CeedBasisCreateTensorH1_Sycl_shared()
186   CeedInt *interp_lrange = impl->interp_local_range;  in CeedBasisCreateTensorH1_Sycl_shared()
191   CeedInt *grad_lrange = impl->grad_local_range;  in CeedBasisCreateTensorH1_Sycl_shared()
196   CeedCallBackend(ComputeLocalRange(ceed, dim, Q_1d, impl->weight_local_range));  in CeedBasisCreateTensorH1_Sycl_shared()
205 …CeedCallSycl(ceed, impl->d_q_weight_1d = sycl::malloc_device<CeedScalar>(Q_1d, data->sycl_device, …  in CeedBasisCreateTensorH1_Sycl_shared()
206 …sycl::event copy_weight = data->sycl_queue.copy<CeedScalar>(q_weight_1d, impl->d_q_weight_1d, Q_1d…  in CeedBasisCreateTensorH1_Sycl_shared()
211 …CeedCallSycl(ceed, impl->d_interp_1d = sycl::malloc_device<CeedScalar>(interp_length, data->sycl_d…  in CeedBasisCreateTensorH1_Sycl_shared()
212 …sycl::event copy_interp = data->sycl_queue.copy<CeedScalar>(interp_1d, impl->d_interp_1d, interp_l…  in CeedBasisCreateTensorH1_Sycl_shared()
215 …CeedCallSycl(ceed, impl->d_grad_1d = sycl::malloc_device<CeedScalar>(interp_length, data->sycl_dev…  in CeedBasisCreateTensorH1_Sycl_shared()
216 …sycl::event copy_grad = data->sycl_queue.copy<CeedScalar>(grad_1d, impl->d_grad_1d, interp_length,…  in CeedBasisCreateTensorH1_Sycl_shared()
222   impl->d_collo_grad_1d          = NULL;  in CeedBasisCreateTensorH1_Sycl_shared()
231 …CeedCallSycl(ceed, impl->d_collo_grad_1d = sycl::malloc_device<CeedScalar>(cgrad_length, data->syc…  in CeedBasisCreateTensorH1_Sycl_shared()
232 …CeedCallSycl(ceed, data->sycl_queue.copy<CeedScalar>(collo_grad_1d, impl->d_collo_grad_1d, cgrad_l…  in CeedBasisCreateTensorH1_Sycl_shared()
262 …CeedCallBackend(CeedBuildModule_Sycl(ceed, basis_kernel_source, &impl->sycl_module, jit_constants)…  in CeedBasisCreateTensorH1_Sycl_shared()
265   CeedCallBackend(CeedGetKernel_Sycl(ceed, impl->sycl_module, "Interp", &impl->interp_kernel));  in CeedBasisCreateTensorH1_Sycl_shared()
266 …CeedCallBackend(CeedGetKernel_Sycl(ceed, impl->sycl_module, "InterpTranspose", &impl->interp_trans…  in CeedBasisCreateTensorH1_Sycl_shared()
267   CeedCallBackend(CeedGetKernel_Sycl(ceed, impl->sycl_module, "Grad", &impl->grad_kernel));  in CeedBasisCreateTensorH1_Sycl_shared()
268 …CeedCallBackend(CeedGetKernel_Sycl(ceed, impl->sycl_module, "GradTranspose", &impl->grad_transpose…  in CeedBasisCreateTensorH1_Sycl_shared()
269   CeedCallBackend(CeedGetKernel_Sycl(ceed, impl->sycl_module, "Weight", &impl->weight_kernel));  in CeedBasisCreateTensorH1_Sycl_shared()
276   CeedCallBackend(CeedBasisSetData(basis, impl));  in CeedBasisCreateTensorH1_Sycl_shared()