44 inline __device__ void InterpAtPoints1d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *_…  in InterpAtPoints1d()  argument
52     if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp];  in InterpAtPoints1d()
56       r_V[comp] += chebyshev_x[i] * data.slice[i];  in InterpAtPoints1d()
65 inline __device__ void InterpTransposeAtPoints1d(SharedData_Cuda &data, const CeedInt p, const Ceed…  in InterpTransposeAtPoints1d()  argument
72     if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0;  in InterpTransposeAtPoints1d()
77 …atomicAdd_block(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x)…  in InterpTransposeAtPoints1d()
82     if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x];  in InterpTransposeAtPoints1d()
90 inline __device__ void GradAtPoints1d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__r…  in GradAtPoints1d()  argument
99     if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp];  in GradAtPoints1d()
103       r_V[comp] += chebyshev_x[i] * data.slice[i];  in GradAtPoints1d()
112 inline __device__ void GradTransposeAtPoints1d(SharedData_Cuda &data, const CeedInt p, const CeedSc…  in GradTransposeAtPoints1d()  argument
119     if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0;  in GradTransposeAtPoints1d()
124 …atomicAdd_block(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x)…  in GradTransposeAtPoints1d()
129     if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x];  in GradTransposeAtPoints1d()
141 inline __device__ void InterpAtPoints2d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *_…  in InterpAtPoints2d()  argument
150 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = r_C[c…  in InterpAtPoints2d()
157         buffer[i] += chebyshev_x[j] * data.slice[j + i * Q_1D];  in InterpAtPoints2d()
172 inline __device__ void InterpTransposeAtPoints2d(SharedData_Cuda &data, const CeedInt p, const Ceed…  in InterpTransposeAtPoints2d()  argument
179 …  if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = 0.0;  in InterpTransposeAtPoints2d()
192       const CeedInt ii = (i + data.t_id_y) % Q_1D;  in InterpTransposeAtPoints2d()
195         const CeedInt jj = (j + data.t_id_x) % Q_1D;  in InterpTransposeAtPoints2d()
197 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd_block(&data.slice[jj + ii * Q_1D], chebysh…  in InterpTransposeAtPoints2d()
202 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) r_C[comp] += data.slice[data.t_id_x + data.t_id_y * …  in InterpTransposeAtPoints2d()
210 inline __device__ void GradAtPoints2d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__r…  in GradAtPoints2d()  argument
219 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = r_C[c…  in GradAtPoints2d()
228           buffer[i] += chebyshev_x[j] * data.slice[j + i * Q_1D];  in GradAtPoints2d()
245 inline __device__ void GradTransposeAtPoints2d(SharedData_Cuda &data, const CeedInt p, const CeedSc…  in GradTransposeAtPoints2d()  argument
252 …  if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = 0.0;  in GradTransposeAtPoints2d()
268         const CeedInt ii = (i + data.t_id_y) % Q_1D;  in GradTransposeAtPoints2d()
271           const CeedInt jj = (j + data.t_id_x) % Q_1D;  in GradTransposeAtPoints2d()
273 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd_block(&data.slice[jj + ii * Q_1D], chebysh…  in GradTransposeAtPoints2d()
279 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) r_C[comp] += data.slice[data.t_id_x + data.t_id_y * …  in GradTransposeAtPoints2d()
291 inline __device__ void InterpAtPoints3d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *_…  in InterpAtPoints3d()  argument
305 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = r_C[k…  in InterpAtPoints3d()
312           buffer[i] += chebyshev_x[j] * data.slice[j + i * Q_1D];  in InterpAtPoints3d()
328 inline __device__ void InterpTransposeAtPoints3d(SharedData_Cuda &data, const CeedInt p, const Ceed…  in InterpTransposeAtPoints3d()  argument
340 …  if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = 0.0;  in InterpTransposeAtPoints3d()
353         const CeedInt ii = (i + data.t_id_y) % Q_1D;  in InterpTransposeAtPoints3d()
356           const CeedInt jj = (j + data.t_id_x) % Q_1D;  in InterpTransposeAtPoints3d()
358 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd_block(&data.slice[jj + ii * Q_1D], chebysh…  in InterpTransposeAtPoints3d()
363 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) r_C[k + comp * Q_1D] += data.slice[data.t_id_x + dat…  in InterpTransposeAtPoints3d()
372 inline __device__ void GradAtPoints3d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__r…  in GradAtPoints3d()  argument
389 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = r_C[k…  in GradAtPoints3d()
399             buffer[i] += chebyshev_x[j] * data.slice[j + i * Q_1D];  in GradAtPoints3d()
419 inline __device__ void GradTransposeAtPoints3d(SharedData_Cuda &data, const CeedInt p, const CeedSc…  in GradTransposeAtPoints3d()  argument
434 …  if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = 0.0;  in GradTransposeAtPoints3d()
453           const CeedInt ii = (i + data.t_id_y) % Q_1D;  in GradTransposeAtPoints3d()
456             const CeedInt jj = (j + data.t_id_x) % Q_1D;  in GradTransposeAtPoints3d()
458 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd_block(&data.slice[jj + ii * Q_1D], chebysh…  in GradTransposeAtPoints3d()
464 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) r_C[k + comp * Q_1D] += data.slice[data.t_id_x + dat…  in GradTransposeAtPoints3d()