44 inline __device__ void InterpAtPoints1d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__…  in InterpAtPoints1d()  argument
53     if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp];  in InterpAtPoints1d()
57       r_V[comp] += chebyshev_x[i] * data.slice[i];  in InterpAtPoints1d()
66 inline __device__ void InterpTransposeAtPoints1d(SharedData_Hip &data, const CeedInt p, const CeedS…  in InterpTransposeAtPoints1d()  argument
73     if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0;  in InterpTransposeAtPoints1d()
78 …atomicAdd(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1…  in InterpTransposeAtPoints1d()
83     if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x];  in InterpTransposeAtPoints1d()
91 inline __device__ void GradAtPoints1d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__re…  in GradAtPoints1d()  argument
100     if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp];  in GradAtPoints1d()
104       r_V[comp] += chebyshev_x[i] * data.slice[i];  in GradAtPoints1d()
113 inline __device__ void GradTransposeAtPoints1d(SharedData_Hip &data, const CeedInt p, const CeedSca…  in GradTransposeAtPoints1d()  argument
120     if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0;  in GradTransposeAtPoints1d()
125 …atomicAdd(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1…  in GradTransposeAtPoints1d()
130     if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x];  in GradTransposeAtPoints1d()
142 inline __device__ void InterpAtPoints2d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__…  in InterpAtPoints2d()  argument
151 …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()
158         buffer[i] += chebyshev_x[j] * data.slice[j + i * Q_1D];  in InterpAtPoints2d()
173 inline __device__ void InterpTransposeAtPoints2d(SharedData_Hip &data, const CeedInt p, const CeedS…  in InterpTransposeAtPoints2d()  argument
180 …  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()
193       const CeedInt ii = (i + data.t_id_y) % Q_1D;  in InterpTransposeAtPoints2d()
196         const CeedInt jj = (j + data.t_id_x) % Q_1D;  in InterpTransposeAtPoints2d()
198 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd(&data.slice[jj + ii * Q_1D], chebyshev_x[j…  in InterpTransposeAtPoints2d()
203 …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()
211 inline __device__ void GradAtPoints2d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__re…  in GradAtPoints2d()  argument
220 …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()
229           buffer[i] += chebyshev_x[j] * data.slice[j + i * Q_1D];  in GradAtPoints2d()
246 inline __device__ void GradTransposeAtPoints2d(SharedData_Hip &data, const CeedInt p, const CeedSca…  in GradTransposeAtPoints2d()  argument
253 …  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()
269         const CeedInt ii = (i + data.t_id_y) % Q_1D;  in GradTransposeAtPoints2d()
272           const CeedInt jj = (j + data.t_id_x) % Q_1D;  in GradTransposeAtPoints2d()
274 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd(&data.slice[jj + ii * Q_1D], chebyshev_x[j…  in GradTransposeAtPoints2d()
280 …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()
292 inline __device__ void InterpAtPoints3d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__…  in InterpAtPoints3d()  argument
306 …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()
313           buffer[i] += chebyshev_x[j] * data.slice[j + i * Q_1D];  in InterpAtPoints3d()
329 inline __device__ void InterpTransposeAtPoints3d(SharedData_Hip &data, const CeedInt p, const CeedS…  in InterpTransposeAtPoints3d()  argument
341 …  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()
354         const CeedInt ii = (i + data.t_id_y) % Q_1D;  in InterpTransposeAtPoints3d()
357           const CeedInt jj = (j + data.t_id_x) % Q_1D;  in InterpTransposeAtPoints3d()
359 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd(&data.slice[jj + ii * Q_1D], chebyshev_x[j…  in InterpTransposeAtPoints3d()
364 …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()
373 inline __device__ void GradAtPoints3d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__re…  in GradAtPoints3d()  argument
390 …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()
400             buffer[i] += chebyshev_x[j] * data.slice[j + i * Q_1D];  in GradAtPoints3d()
420 inline __device__ void GradTransposeAtPoints3d(SharedData_Hip &data, const CeedInt p, const CeedSca…  in GradTransposeAtPoints3d()  argument
435 …  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(&data.slice[jj + ii * Q_1D], chebyshev_x[j…  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()