| /libCEED/include/ceed/jit-source/sycl/ |
| H A D | sycl-shared-basis-tensor.h | 28 CeedScalar r_V[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in Interp() local
38 Interp1d(BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, r_U, s_B, r_V, elem_scratch); in Interp()
39 …ed1d(BASIS_NUM_COMP, BASIS_Q_1D, num_elem, 1, BASIS_NUM_QPTS * num_elem, BASIS_NUM_QPTS, r_V, d_V); in Interp()
43 InterpTensor2d(BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, r_U, s_B, r_V, elem_scratch); in Interp()
44 …ed2d(BASIS_NUM_COMP, BASIS_Q_1D, num_elem, 1, BASIS_NUM_QPTS * num_elem, BASIS_NUM_QPTS, r_V, d_V); in Interp()
48 InterpTensor3d(BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, r_U, s_B, r_V, elem_scratch); in Interp()
49 …ed3d(BASIS_NUM_COMP, BASIS_Q_1D, num_elem, 1, BASIS_NUM_QPTS * num_elem, BASIS_NUM_QPTS, r_V, d_V); in Interp()
62 CeedScalar r_V[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in InterpTranspose() local
72 InterpTranspose1d(BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, r_U, s_B, r_V, elem_scratch); in InterpTranspose()
73 …1d(BASIS_NUM_COMP, BASIS_P_1D, num_elem, 1, BASIS_NUM_NODES * num_elem, BASIS_NUM_NODES, r_V, d_V); in InterpTranspose()
[all …]
|
| H A D | sycl-shared-basis-tensor-templates.h | 58 …local const CeedScalar *restrict s_B, private CeedScalar *restrict r_V, local CeedScalar *restrict… in Interp1d() argument
60 ContractX1d(P_1D, Q_1D, r_U + comp, s_B, r_V + comp, scratch); in Interp1d()
68 …local const CeedScalar *restrict s_B, private CeedScalar *restrict r_V, local CeedScalar *restrict… in InterpTranspose1d() argument
70 ContractTransposeX1d(P_1D, Q_1D, r_U + comp, s_B, r_V + comp, scratch); in InterpTranspose1d()
78 …local const CeedScalar *restrict s_G, private CeedScalar *restrict r_V, local CeedScalar *restrict… in Grad1d() argument
80 ContractX1d(P_1D, Q_1D, r_U + comp, s_G, r_V + comp, scratch); in Grad1d()
88 …local const CeedScalar *restrict s_G, private CeedScalar *restrict r_V, local CeedScalar *restrict… in GradTranspose1d() argument
90 ContractTransposeX1d(P_1D, Q_1D, r_U + comp, s_G, r_V + comp, scratch); in GradTranspose1d()
209 …local const CeedScalar *restrict s_B, private CeedScalar *restrict r_V, local CeedScalar *restrict… in InterpTensor2d() argument
214 ContractY2d(P_1D, Q_1D, r_t, s_B, r_V + comp, scratch); in InterpTensor2d()
[all …]
|
| H A D | sycl-gen-templates.h | 286 …const local CeedScalar *s_G, private CeedScalar *restrict r_V, local CeedScalar *restrict scratch)… in gradCollo3d() argument
298 r_V[comp + 0 * num_comp] = 0.0; in gradCollo3d()
300 …r_V[comp + 0 * num_comp] += s_G[i + item_id_x * Q_1D] * scratch[i + item_id_y * T_1D]; // Contrac… in gradCollo3d()
303 r_V[comp + 1 * num_comp] = 0.0; in gradCollo3d()
305 …r_V[comp + 1 * num_comp] += s_G[i + item_id_y * Q_1D] * scratch[item_id_x + i * T_1D]; // Contrac… in gradCollo3d()
308 r_V[comp + 2 * num_comp] = 0.0; in gradCollo3d()
309 …for (CeedInt i = 0; i < Q_1D; ++i) r_V[comp + 2 * num_comp] += s_G[i + q * Q_1D] * r_U[i + comp * … in gradCollo3d()
320 …const local CeedScalar *restrict s_G, private CeedScalar *restrict r_V, local CeedScalar *restrict… in gradColloTranspose3d() argument
333 …r_V[q + comp * Q_1D] += s_G[item_id_x + i * Q_1D] * scratch[i + item_id_y * T_1D]; // Contract x … in gradColloTranspose3d()
345 …r_V[q + comp * Q_1D] += s_G[item_id_y + i * Q_1D] * scratch[item_id_x + i * T_1D]; // Contract y … in gradColloTranspose3d()
[all …]
|
| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-shared-basis-tensor-templates.h | 52 …a &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, CeedScalar *__restrict__ r_V) { in Interp1d() argument
54 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in Interp1d()
63 CeedScalar *__restrict__ r_V) { in InterpTranspose1d() argument
65 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTranspose1d()
74 CeedScalar *__restrict__ r_V) { in InterpCollocatedNodes1d() argument
76 r_V[comp] = r_U[comp]; in InterpCollocatedNodes1d()
85 CeedScalar *__restrict__ r_V) { in InterpTransposeCollocatedNodes1d() argument
87 r_V[comp] = r_U[comp]; in InterpTransposeCollocatedNodes1d()
96 CeedScalar *__restrict__ r_V) { in Grad1d() argument
98 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_G, &r_V[comp]); in Grad1d()
[all …]
|
| H A D | cuda-shared-basis-tensor-at-points.h | 37 CeedScalar r_V[BASIS_NUM_COMP]; in InterpAtPoints() local
67 … InterpAtPoints1d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_C, r_X, r_V); in InterpAtPoints()
69 … InterpAtPoints2d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_C, r_X, r_V); in InterpAtPoints()
71 … InterpAtPoints3d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_C, r_X, r_V); in InterpAtPoints()
73 …BASIS_NUM_PTS>(data, elem, p, BASIS_NUM_PTS, 1, num_elem * BASIS_NUM_PTS, BASIS_NUM_PTS, r_V, d_V); in InterpAtPoints()
93 CeedScalar r_V[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in InterpTransposeAtPoints() local
106 for (CeedInt i = 0; i < BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1); i++) r_V[i] = 0.0; in InterpTransposeAtPoints()
108 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, r_V, d_V); in InterpTransposeAtPoints()
110 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, r_V, d_V); in InterpTransposeAtPoints()
113 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, r_V, d_V); in InterpTransposeAtPoints()
[all …]
|
| H A D | cuda-shared-basis-tensor.h | 29 CeedScalar r_V[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in Interp() local
40 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
41 …tStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D, r_V, d_V); in Interp()
44 InterpTensor2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
45 … BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * BASIS_Q_1D * num_elem, BASIS_Q_1D * BASIS_Q_1D, r_V, d_V); in Interp()
49 InterpTensor3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
51 … BASIS_Q_1D * BASIS_Q_1D * BASIS_Q_1D, r_V, d_V); in Interp()
98 CeedScalar r_V[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in InterpTranspose() local
109 InterpTranspose1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in InterpTranspose()
110 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, r_V, d_V); in InterpTranspose()
[all …]
|
| H A D | cuda-ref-basis-nontensor-templates.h | 21 CeedScalar r_V[Q_COMP]; in Contract() local
27 for (CeedInt d = 0; d < Q_COMP; d++) r_V[d] = 0.0; in Contract()
31 for (CeedInt d = 0; d < Q_COMP; d++) r_V[d] += d_B[i + t_id * P + d * P * Q] * val; in Contract()
34 d_V[elem * strides_elem_V + comp * strides_comp_V + d * strides_q_comp_V + t_id] = r_V[d]; in Contract()
48 CeedScalar r_V; in ContractTranspose() local
53 r_V = 0.0; in ContractTranspose()
56 for (CeedInt i = 0; i < Q; i++) r_V += d_B[t_id + i * P + d * P * Q] * U[i]; in ContractTranspose()
58 d_V[elem * strides_elem_V + comp * strides_comp_V + t_id] += r_V; in ContractTranspose()
|
| H A D | cuda-shared-basis-nontensor-templates.h | 48 CeedScalar *__restrict__ r_V) { in InterpNonTensor() argument
50 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpNonTensor()
59 CeedScalar *__restrict__ r_V) { in InterpTransposeNonTensor() argument
61 r_V[comp] = 0.0; in InterpTransposeNonTensor()
62 ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTransposeNonTensor()
70 …a &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { in GradNonTensor() argument
73 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], &c_G[dim * P * Q], &r_V[comp + dim * NUM_COMP]); in GradNonTensor()
83 CeedScalar *__restrict__ r_V) { in GradTransposeNonTensor() argument
84 for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_V[comp] = 0.0; in GradTransposeNonTensor()
87 …tractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp + dim * NUM_COMP], &c_G[dim * P * Q], &r_V[comp]); in GradTransposeNonTensor()
|
| H A D | cuda-shared-basis-tensor-flattened-templates.h | 132 CeedScalar *__restrict__ r_V) { in InterpTensor2dFlattened() argument
139 ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[comp]); in InterpTensor2dFlattened()
143 if (Q_1D != T_1D) QPack2d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTensor2dFlattened()
151 CeedScalar *__restrict__ r_V) { in InterpTransposeTensor2dFlattened() argument
158 …ContractTransposeX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[com… in InterpTransposeTensor2dFlattened()
161 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTransposeTensor2dFlattened()
169 CeedScalar *__restrict__ r_V) { in InterpTensorCollocatedNodes2dFlattened() argument
174 r_V[comp] = r_U[comp]; in InterpTensorCollocatedNodes2dFlattened()
178 if (Q_1D != T_1D) QPack2d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTensorCollocatedNodes2dFlattened()
186 … CeedScalar *__restrict__ r_V) { in InterpTransposeTensorCollocatedNodes2dFlattened() argument
[all …]
|
| H A D | cuda-shared-basis-nontensor.h | 29 CeedScalar r_V[BASIS_NUM_COMP]; in Interp() local
39 InterpNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
40 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, r_V, d_… in Interp()
56 CeedScalar r_V[BASIS_NUM_COMP]; in InterpTranspose() local
66 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in InterpTranspose()
67 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_… in InterpTranspose()
83 CeedScalar r_V[BASIS_NUM_COMP]; in InterpTransposeAdd() local
93 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in InterpTransposeAdd()
94 …SumElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_V); in InterpTransposeAdd()
112 CeedScalar r_V[BASIS_NUM_COMP * BASIS_DIM]; in Grad() local
[all …]
|
| H A D | cuda-shared-basis-tensor-at-points-templates.h | 45 CeedScalar *__restrict__ r_V) { in InterpAtPoints1d() argument
48 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in InterpAtPoints1d()
56 r_V[comp] += chebyshev_x[i] * data.slice[i]; in InterpAtPoints1d()
91 CeedScalar *__restrict__ r_V) { in GradAtPoints1d() argument
95 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in GradAtPoints1d()
103 r_V[comp] += chebyshev_x[i] * data.slice[i]; in GradAtPoints1d()
142 CeedScalar *__restrict__ r_V) { in InterpAtPoints2d() argument
143 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in InterpAtPoints2d()
163 r_V[comp] += chebyshev_x[i] * buffer[i]; in InterpAtPoints2d()
211 CeedScalar *__restrict__ r_V) { in GradAtPoints2d() argument
[all …]
|
| H A D | cuda-gen-templates.h | 480 CeedScalar *__restrict__ r_V) { in GradColloSlice3d() argument
487 r_V[comp + 0 * NUM_COMP] = 0.0; in GradColloSlice3d()
489 … r_V[comp + 0 * NUM_COMP] += c_G[i + data.t_id_x * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSlice3d()
492 r_V[comp + 1 * NUM_COMP] = 0.0; in GradColloSlice3d()
494 … r_V[comp + 1 * NUM_COMP] += c_G[i + data.t_id_y * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSlice3d()
497 r_V[comp + 2 * NUM_COMP] = 0.0; in GradColloSlice3d()
499 r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1D] * r_U[i + comp * Q_1D]; in GradColloSlice3d()
510 CeedScalar *__restrict__ r_V) { in GradColloSliceTranspose3d() argument
518 r_V[q + comp * Q_1D] += c_G[data.t_id_x + i * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSliceTranspose3d()
525 r_V[q + comp * Q_1D] += c_G[data.t_id_y + i * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSliceTranspose3d()
[all …]
|
| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-shared-basis-tensor-templates.h | 52 …p &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, CeedScalar *__restrict__ r_V) { in Interp1d() argument
54 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in Interp1d()
63 CeedScalar *__restrict__ r_V) { in InterpTranspose1d() argument
65 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTranspose1d()
74 CeedScalar *__restrict__ r_V) { in InterpCollocatedNodes1d() argument
76 r_V[comp] = r_U[comp]; in InterpCollocatedNodes1d()
85 CeedScalar *__restrict__ r_V) { in InterpTransposeCollocatedNodes1d() argument
87 r_V[comp] = r_U[comp]; in InterpTransposeCollocatedNodes1d()
96 CeedScalar *__restrict__ r_V) { in Grad1d() argument
98 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_G, &r_V[comp]); in Grad1d()
[all …]
|
| H A D | hip-shared-basis-tensor-at-points.h | 38 CeedScalar r_V[BASIS_NUM_COMP]; in __launch_bounds__() local
68 … InterpAtPoints1d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_C, r_X, r_V); in __launch_bounds__()
70 … InterpAtPoints2d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_C, r_X, r_V); in __launch_bounds__()
72 … InterpAtPoints3d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_C, r_X, r_V); in __launch_bounds__()
74 …BASIS_NUM_PTS>(data, elem, p, BASIS_NUM_PTS, 1, num_elem * BASIS_NUM_PTS, BASIS_NUM_PTS, r_V, d_V); in __launch_bounds__()
94 CeedScalar r_V[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in __launch_bounds__() local
107 for (CeedInt i = 0; i < BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1); i++) r_V[i] = 0.0; in __launch_bounds__()
109 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, r_V, d_V); in __launch_bounds__()
111 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, r_V, d_V); in __launch_bounds__()
114 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, r_V, d_V); in __launch_bounds__()
[all …]
|
| H A D | hip-shared-basis-tensor.h | 30 CeedScalar r_V[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in __launch_bounds__() local
41 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
42 …tStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D, r_V, d_V); in __launch_bounds__()
45 InterpTensor2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
46 … BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * BASIS_Q_1D * num_elem, BASIS_Q_1D * BASIS_Q_1D, r_V, d_V); in __launch_bounds__()
50 InterpTensor3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
52 … BASIS_Q_1D * BASIS_Q_1D * BASIS_Q_1D, r_V, d_V); in __launch_bounds__()
99 CeedScalar r_V[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in __launch_bounds__() local
110 InterpTranspose1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
111 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, r_V, d_V); in __launch_bounds__()
[all …]
|
| H A D | hip-ref-basis-nontensor-templates.h | 21 CeedScalar r_V[Q_COMP]; in Contract() local
27 for (CeedInt d = 0; d < Q_COMP; d++) r_V[d] = 0.0; in Contract()
31 for (CeedInt d = 0; d < Q_COMP; d++) r_V[d] += d_B[i + t_id * P + d * P * Q] * val; in Contract()
34 d_V[elem * strides_elem_V + comp * strides_comp_V + d * strides_q_comp_V + t_id] = r_V[d]; in Contract()
48 CeedScalar r_V; in ContractTranspose() local
53 r_V = 0.0; in ContractTranspose()
56 for (CeedInt i = 0; i < Q; i++) r_V += d_B[t_id + i * P + d * P * Q] * U[i]; in ContractTranspose()
58 d_V[elem * strides_elem_V + comp * strides_comp_V + t_id] += r_V; in ContractTranspose()
|
| H A D | hip-shared-basis-nontensor-templates.h | 48 CeedScalar *__restrict__ r_V) { in InterpNonTensor() argument
50 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpNonTensor()
59 CeedScalar *__restrict__ r_V) { in InterpTransposeNonTensor() argument
61 r_V[comp] = 0.0; in InterpTransposeNonTensor()
62 ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTransposeNonTensor()
70 …p &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { in GradNonTensor() argument
73 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], &c_G[dim * P * Q], &r_V[comp + dim * NUM_COMP]); in GradNonTensor()
83 CeedScalar *__restrict__ r_V) { in GradTransposeNonTensor() argument
84 for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_V[comp] = 0.0; in GradTransposeNonTensor()
87 …tractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp + dim * NUM_COMP], &c_G[dim * P * Q], &r_V[comp]); in GradTransposeNonTensor()
|
| H A D | hip-shared-basis-tensor-flattened-templates.h | 132 CeedScalar *__restrict__ r_V) { in InterpTensor2dFlattened() argument
139 ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[comp]); in InterpTensor2dFlattened()
143 if (Q_1D != T_1D) QPack2d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTensor2dFlattened()
151 CeedScalar *__restrict__ r_V) { in InterpTransposeTensor2dFlattened() argument
158 …ContractTransposeX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[com… in InterpTransposeTensor2dFlattened()
161 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTransposeTensor2dFlattened()
169 CeedScalar *__restrict__ r_V) { in InterpTensorCollocatedNodes2dFlattened() argument
174 r_V[comp] = r_U[comp]; in InterpTensorCollocatedNodes2dFlattened()
178 if (Q_1D != T_1D) QPack2d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTensorCollocatedNodes2dFlattened()
186 … CeedScalar *__restrict__ r_V) { in InterpTransposeTensorCollocatedNodes2dFlattened() argument
[all …]
|
| H A D | hip-shared-basis-nontensor.h | 30 CeedScalar r_V[BASIS_NUM_COMP]; in __launch_bounds__() local
40 InterpNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
41 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, r_V, d_… in __launch_bounds__()
57 CeedScalar r_V[BASIS_NUM_COMP]; in __launch_bounds__() local
67 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
68 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_… in __launch_bounds__()
84 CeedScalar r_V[BASIS_NUM_COMP]; in __launch_bounds__() local
94 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
95 …SumElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_V); in __launch_bounds__()
114 CeedScalar r_V[BASIS_NUM_COMP * BASIS_DIM]; in __launch_bounds__() local
[all …]
|
| H A D | hip-shared-basis-tensor-at-points-templates.h | 45 CeedScalar *__restrict__ r_V) { in InterpAtPoints1d() argument
48 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in InterpAtPoints1d()
57 r_V[comp] += chebyshev_x[i] * data.slice[i]; in InterpAtPoints1d()
92 CeedScalar *__restrict__ r_V) { in GradAtPoints1d() argument
96 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in GradAtPoints1d()
104 r_V[comp] += chebyshev_x[i] * data.slice[i]; in GradAtPoints1d()
143 CeedScalar *__restrict__ r_V) { in InterpAtPoints2d() argument
144 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in InterpAtPoints2d()
164 r_V[comp] += chebyshev_x[i] * buffer[i]; in InterpAtPoints2d()
212 CeedScalar *__restrict__ r_V) { in GradAtPoints2d() argument
[all …]
|
| H A D | hip-gen-templates.h | 477 CeedScalar *__restrict__ r_V) { in GradColloSlice3d() argument
484 r_V[comp + 0 * NUM_COMP] = 0.0; in GradColloSlice3d()
486 … r_V[comp + 0 * NUM_COMP] += c_G[i + data.t_id_x * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSlice3d()
489 r_V[comp + 1 * NUM_COMP] = 0.0; in GradColloSlice3d()
491 … r_V[comp + 1 * NUM_COMP] += c_G[i + data.t_id_y * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSlice3d()
494 r_V[comp + 2 * NUM_COMP] = 0.0; in GradColloSlice3d()
496 r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1D] * r_U[i + comp * Q_1D]; in GradColloSlice3d()
507 CeedScalar *__restrict__ r_V) { in GradColloSliceTranspose3d() argument
515 r_V[q + comp * Q_1D] += c_G[data.t_id_x + i * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSliceTranspose3d()
522 r_V[q + comp * Q_1D] += c_G[data.t_id_y + i * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSliceTranspose3d()
[all …]
|