Home
last modified time | relevance | path

Searched refs:c_G (Results 1 – 12 of 12) sorted by relevance

/libCEED/include/ceed/jit-source/cuda/
H A Dcuda-shared-basis-tensor-templates.h95 …dData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in Grad1d() argument
98 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_G, &r_V[comp]); in Grad1d()
106 …dData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in GradTranspose1d() argument
109 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_G, &r_V[comp]); in GradTranspose1d()
256 …dData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in GradTensor2d() argument
260 ContractX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_G, r_t); in GradTensor2d()
263 ContractY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_G, &r_V[comp + 1 * NUM_COMP]); in GradTensor2d()
271 …dData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in GradTransposeTensor2d() argument
276 ContractTransposeX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_G, &r_V[comp]); in GradTransposeTensor2d()
277 ContractTransposeY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp + 1 * NUM_COMP], c_G, r_t); in GradTransposeTensor2d()
[all …]
H A Dcuda-shared-basis-tensor-flattened-templates.h201 …(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in GradTensor2dFlattened() argument
208 ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_G, r_t); in GradTensor2dFlattened()
211 …ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_G, &r_V[comp + 1 * N… in GradTensor2dFlattened()
223 … const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { in GradTransposeTensor2dFlattened() argument
230 …ContractTransposeX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_G, &r_V[com… in GradTransposeTensor2dFlattened()
231 …2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp + 1 * NUM_COMP], c_G, r_t); in GradTransposeTensor2dFlattened()
243 … const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { in GradTensorCollocatedNodes2dFlattened() argument
248 …ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_G, &r_V[comp … in GradTensorCollocatedNodes2dFlattened()
249 …ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_G, &r_V[comp … in GradTensorCollocatedNodes2dFlattened()
261 … const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { in GradTransposeTensorCollocatedNodes2dFlattened() argument
[all …]
H A Dcuda-shared-basis-nontensor-templates.h70 …aredData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, CeedScalar *__rest… 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()
82 …ransposeNonTensor(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, in GradTransposeNonTensor() argument
87 …ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp + dim * NUM_COMP], &c_G[dim * P * Q], &r_V[com… in GradTransposeNonTensor()
H A Dcuda-shared-basis-nontensor.h101 extern "C" __global__ void Grad(const CeedInt num_elem, const CeedScalar *c_G, const CeedScalar *__… in Grad() argument
116 LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G); in Grad()
127 extern "C" __global__ void GradTranspose(const CeedInt num_elem, const CeedScalar *c_G, const CeedS… in GradTranspose() argument
143 LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G); in GradTranspose()
154 extern "C" __global__ void GradTransposeAdd(const CeedInt num_elem, const CeedScalar *c_G, const Ce… in GradTransposeAdd() argument
170 LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G); in GradTransposeAdd()
H A Dcuda-shared-basis-tensor.h227 …void Grad(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, const CeedScalar *… in Grad() argument
245 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in Grad()
270 …ollocated(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, const CeedScalar *… in GradCollocated() argument
286 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in GradCollocated()
310 …Transpose(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, const CeedScalar *… in GradTranspose() argument
328 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in GradTranspose()
353 … void GradCollocatedTranspose(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, in GradCollocatedTranspose() argument
369 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in GradCollocatedTranspose()
393 …nsposeAdd(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, const CeedScalar *… in GradTransposeAdd() argument
411 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in GradTransposeAdd()
[all …]
H A Dcuda-gen-templates.h479 …(SharedData_Cuda &data, const CeedInt q, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, in GradColloSlice3d() argument
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()
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()
499 r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1D] * r_U[i + comp * Q_1D]; in GradColloSlice3d()
509 …(SharedData_Cuda &data, const CeedInt q, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, 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()
529 r_V[i + comp * Q_1D] += c_G[i + q * Q_1D] * r_U[comp + 2 * NUM_COMP]; in GradColloSliceTranspose3d()
/libCEED/include/ceed/jit-source/hip/
H A Dhip-shared-basis-tensor-templates.h95 …edData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in Grad1d() argument
98 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_G, &r_V[comp]); in Grad1d()
106 …edData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in GradTranspose1d() argument
109 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_G, &r_V[comp]); in GradTranspose1d()
255 …edData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in GradTensor2d() argument
259 ContractX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_G, r_t); in GradTensor2d()
262 ContractY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_G, &r_V[comp + 1 * NUM_COMP]); in GradTensor2d()
270 …edData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in GradTransposeTensor2d() argument
275 ContractTransposeX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_G, &r_V[comp]); in GradTransposeTensor2d()
276 ContractTransposeY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp + 1 * NUM_COMP], c_G, r_t); in GradTransposeTensor2d()
[all …]
H A Dhip-shared-basis-tensor-flattened-templates.h201 …d(SharedData_Hip &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *c_G, in GradTensor2dFlattened() argument
208 ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_G, r_t); in GradTensor2dFlattened()
211 …ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_G, &r_V[comp + 1 * N… in GradTensor2dFlattened()
223 … const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { in GradTransposeTensor2dFlattened() argument
230 …ContractTransposeX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_G, &r_V[com… in GradTransposeTensor2dFlattened()
231 …2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp + 1 * NUM_COMP], c_G, r_t); in GradTransposeTensor2dFlattened()
243 … const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { in GradTensorCollocatedNodes2dFlattened() argument
248 …ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_G, &r_V[comp … in GradTensorCollocatedNodes2dFlattened()
249 …ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_G, &r_V[comp … in GradTensorCollocatedNodes2dFlattened()
261 … const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { in GradTransposeTensorCollocatedNodes2dFlattened() argument
[all …]
H A Dhip-shared-basis-nontensor-templates.h70 …haredData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, CeedScalar *__rest… 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()
82 …TransposeNonTensor(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, in GradTransposeNonTensor() argument
87 …ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp + dim * NUM_COMP], &c_G[dim * P * Q], &r_V[com… in GradTransposeNonTensor()
H A Dhip-shared-basis-nontensor.h103 …void Grad(const CeedInt num_elem, const CeedScalar *c_G, const CeedScalar *__restrict__ d_U, CeedS… in __launch_bounds__()
118 LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G); in __launch_bounds__()
130 …void GradTranspose(const CeedInt num_elem, const CeedScalar *c_G, const CeedScalar *__restrict__ d… in __launch_bounds__()
145 LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G); in __launch_bounds__()
157 …void GradTransposeAdd(const CeedInt num_elem, const CeedScalar *c_G, const CeedScalar *__restrict_… in __launch_bounds__()
172 LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G); in __launch_bounds__()
H A Dhip-shared-basis-tensor.h229 …CK_SIZE) __global__ void Grad(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, in __launch_bounds__()
247 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in __launch_bounds__()
273 …void GradCollocated(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, const Ce… in __launch_bounds__()
289 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in __launch_bounds__()
314 …void GradTranspose(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, const Cee… in __launch_bounds__()
332 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in __launch_bounds__()
358 …Transpose(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, const CeedScalar *… in __launch_bounds__()
374 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in __launch_bounds__()
399 …void GradTransposeAdd(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_G, const … in __launch_bounds__()
417 LoadMatrix<BASIS_Q_1D, BASIS_HAS_COLLOCATED_GRAD ? BASIS_Q_1D : BASIS_P_1D>(data, c_G, s_G); in __launch_bounds__()
[all …]
H A Dhip-gen-templates.h476 …d(SharedData_Hip &data, const CeedInt q, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, in GradColloSlice3d() argument
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()
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()
496 r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1D] * r_U[i + comp * Q_1D]; in GradColloSlice3d()
506 …d(SharedData_Hip &data, const CeedInt q, const CeedScalar *__restrict__ r_U, const CeedScalar *c_G, 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()
526 r_V[i + comp * Q_1D] += c_G[i + q * Q_1D] * r_U[comp + 2 * NUM_COMP]; in GradColloSliceTranspose3d()