29   const CeedInt  elem_size    = impl->elem_size;  in CeedElemRestrictionStridedNoTranspose_Sycl()  local
35   sycl::range<1> kernel_range(num_elem * elem_size);  in CeedElemRestrictionStridedNoTranspose_Sycl()
41     const CeedInt loc_node = node % elem_size;  in CeedElemRestrictionStridedNoTranspose_Sycl()
42     const CeedInt elem     = node / elem_size;  in CeedElemRestrictionStridedNoTranspose_Sycl()
45 …v[loc_node + comp * elem_size * num_elem + elem * elem_size] = u[loc_node * stride_nodes + comp * …  in CeedElemRestrictionStridedNoTranspose_Sycl()
56   const CeedInt  elem_size   = impl->elem_size;  in CeedElemRestrictionOffsetNoTranspose_Sycl()  local
62   sycl::range<1> kernel_range(num_elem * elem_size);  in CeedElemRestrictionOffsetNoTranspose_Sycl()
69     const CeedInt loc_node = node % elem_size;  in CeedElemRestrictionOffsetNoTranspose_Sycl()
70     const CeedInt elem     = node / elem_size;  in CeedElemRestrictionOffsetNoTranspose_Sycl()
73       v[loc_node + comp * elem_size * num_elem + elem * elem_size] = u[ind + comp * comp_stride];  in CeedElemRestrictionOffsetNoTranspose_Sycl()
84   const CeedInt elem_size    = impl->elem_size;  in CeedElemRestrictionStridedTranspose_Sycl()  local
91   sycl::range<1> kernel_range(num_elem * elem_size);  in CeedElemRestrictionStridedTranspose_Sycl()
97     const CeedInt loc_node = node % elem_size;  in CeedElemRestrictionStridedTranspose_Sycl()
98     const CeedInt elem     = node / elem_size;  in CeedElemRestrictionStridedTranspose_Sycl()
101 … stride_comp + elem * stride_elem] += u[loc_node + comp * elem_size * num_elem + elem * elem_size];  in CeedElemRestrictionStridedTranspose_Sycl()
113   const CeedInt  elem_size     = impl->elem_size;  in CeedElemRestrictionOffsetTranspose_Sycl()  local
136       CeedInt       loc_node = t_ind % elem_size;  in CeedElemRestrictionOffsetTranspose_Sycl()
137       CeedInt       elem     = t_ind / elem_size;  in CeedElemRestrictionOffsetTranspose_Sycl()
139       value += u[loc_node + comp * elem_size * num_elem + elem * elem_size];  in CeedElemRestrictionOffsetTranspose_Sycl()
254 …CeedInt                   num_elem, elem_size, num_comp, num_nodes = 0, *ind_to_offset, *l_vec_ind…  in CeedElemRestrictionOffset_Sycl()  local
260   CeedCallBackend(CeedElemRestrictionGetElementSize(rstr, &elem_size));  in CeedElemRestrictionOffset_Sycl()
266   const CeedInt size_indices = num_elem * elem_size;  in CeedElemRestrictionOffset_Sycl()
290 …for (CeedInt i = 0; i < elem_size; ++i) ++t_offsets[ind_to_offset[indices[elem_size * e + i]] + 1];  in CeedElemRestrictionOffset_Sycl()
296     for (CeedInt i = 0; i < elem_size; ++i) {  in CeedElemRestrictionOffset_Sycl()
297       const CeedInt lid                          = elem_size * e + i;  in CeedElemRestrictionOffset_Sycl()
343   CeedInt                   num_elem, num_comp, elem_size, comp_stride = 1;  in CeedElemRestrictionCreate_Sycl()  local
351   CeedCallBackend(CeedElemRestrictionGetElementSize(rstr, &elem_size));  in CeedElemRestrictionCreate_Sycl()
352   const CeedInt size       = num_elem * elem_size;  in CeedElemRestrictionCreate_Sycl()
353   CeedInt       strides[3] = {1, size, elem_size};  in CeedElemRestrictionCreate_Sycl()
376   impl->elem_size   = elem_size;  in CeedElemRestrictionCreate_Sycl()
386     CeedInt layout[3] = {1, size, elem_size};  in CeedElemRestrictionCreate_Sycl()
403             CeedCallBackend(CeedMalloc(elem_size * num_elem, &impl->h_offsets_owned));  in CeedElemRestrictionCreate_Sycl()
404             memcpy(impl->h_offsets_owned, offsets, elem_size * num_elem * sizeof(CeedInt));  in CeedElemRestrictionCreate_Sycl()
457         CeedCallBackend(CeedMalloc(elem_size * num_elem, &impl->h_offsets_owned));  in CeedElemRestrictionCreate_Sycl()
461 … data->sycl_queue.copy<CeedInt>(impl->d_offsets, impl->h_offsets_owned, elem_size * num_elem, {e});  in CeedElemRestrictionCreate_Sycl()