| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-shared-basis-tensor-templates.h | 52 …aredData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, CeedScalar *__rest… in Interp1d() argument 54 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in Interp1d() 62 …InterpTranspose1d(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTranspose1d() argument 65 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTranspose1d() 73 …CollocatedNodes1d(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpCollocatedNodes1d() argument 84 …CollocatedNodes1d(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTransposeCollocatedNodes1d() argument 95 …aredData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *… in Grad1d() argument 106 …aredData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *… in GradTranspose1d() argument 208 …id InterpTensor2d(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTensor2d() argument 212 ContractX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_B, r_t); in InterpTensor2d() [all …]
|
| H A D | cuda-shared-basis-tensor-flattened-templates.h | 131 …InterpTensor2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTensor2dFlattened() argument 138 ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_B, r_t); in InterpTensor2dFlattened() 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() 150 …nsposeTensor2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTransposeTensor2dFlattened() argument 157 …ContractTransposeY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_B, r… in InterpTransposeTensor2dFlattened() 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() 168 …locatedNodes2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTensorCollocatedNodes2dFlattened() argument 185 …locatedNodes2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTransposeTensorCollocatedNodes2dFlattened() argument 201 …ned(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *… in GradTensor2dFlattened() argument 209 …ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[comp + 0 * N… in GradTensor2dFlattened() [all …]
|
| H A D | cuda-shared-basis-nontensor-templates.h | 47 …d InterpNonTensor(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpNonTensor() argument 50 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpNonTensor() 58 …ransposeNonTensor(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTransposeNonTensor() argument 62 ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTransposeNonTensor()
|
| H A D | cuda-shared-basis-tensor.h | 18 extern "C" __global__ void Interp(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *… in Interp() argument 33 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in Interp() 56 extern "C" __global__ void InterpCollocated(const CeedInt num_elem, const CeedScalar *c_B, const Ce… in InterpCollocated() argument 86 extern "C" __global__ void InterpTranspose(const CeedInt num_elem, const CeedScalar *c_B, const Cee… in InterpTranspose() argument 102 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in InterpTranspose() 125 extern "C" __global__ void InterpCollocatedTranspose(const CeedInt num_elem, const CeedScalar *c_B,… in InterpCollocatedTranspose() argument 155 extern "C" __global__ void InterpTransposeAdd(const CeedInt num_elem, const CeedScalar *c_B, const … in InterpTransposeAdd() argument 171 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in InterpTransposeAdd() 194 …oid InterpCollocatedTransposeAdd(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *… in InterpCollocatedTransposeAdd() argument 227 extern "C" __global__ void Grad(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *c_… in Grad() argument [all …]
|
| H A D | cuda-shared-basis-tensor-at-points.h | 23 …void InterpAtPoints(const CeedInt num_elem, const CeedScalar *__restrict__ c_B, const CeedInt *__r… in InterpAtPoints() argument 41 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in InterpAtPoints() 78 …__global__ void InterpTransposeAtPoints(const CeedInt num_elem, const CeedScalar *__restrict__ c_B, in InterpTransposeAtPoints() argument 97 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in InterpTransposeAtPoints() 148 …lobal__ void InterpTransposeAddAtPoints(const CeedInt num_elem, const CeedScalar *__restrict__ c_B, in InterpTransposeAddAtPoints() argument 167 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in InterpTransposeAddAtPoints() 210 extern "C" __global__ void GradAtPoints(const CeedInt num_elem, const CeedScalar *__restrict__ c_B,… in GradAtPoints() argument 228 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in GradAtPoints() 265 …" __global__ void GradTransposeAtPoints(const CeedInt num_elem, const CeedScalar *__restrict__ c_B, in GradTransposeAtPoints() argument 284 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in GradTransposeAtPoints() [all …]
|
| H A D | cuda-shared-basis-nontensor.h | 18 extern "C" __global__ void Interp(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *… in Interp() argument 33 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in Interp() 44 extern "C" __global__ void InterpTranspose(const CeedInt num_elem, const CeedScalar *c_B, const Cee… in InterpTranspose() argument 60 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in InterpTranspose() 71 extern "C" __global__ void InterpTransposeAdd(const CeedInt num_elem, const CeedScalar *c_B, const … in InterpTransposeAdd() argument 87 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in InterpTransposeAdd()
|
| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-shared-basis-tensor-templates.h | 52 …haredData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, CeedScalar *__rest… in Interp1d() argument 54 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in Interp1d() 62 … InterpTranspose1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTranspose1d() argument 65 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTranspose1d() 73 …pCollocatedNodes1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpCollocatedNodes1d() argument 84 …eCollocatedNodes1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTransposeCollocatedNodes1d() argument 95 …haredData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *… in Grad1d() argument 106 …haredData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *… in GradTranspose1d() argument 208 …haredData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, CeedScalar *__rest… in InterpTensor2d() argument 211 ContractX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_B, r_t); in InterpTensor2d() [all …]
|
| H A D | hip-shared-basis-tensor-flattened-templates.h | 131 … InterpTensor2dFlattened(SharedData_Hip &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTensor2dFlattened() argument 138 ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_B, r_t); in InterpTensor2dFlattened() 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() 150 …ansposeTensor2dFlattened(SharedData_Hip &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTransposeTensor2dFlattened() argument 157 …ContractTransposeY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_B, r… in InterpTransposeTensor2dFlattened() 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() 168 …llocatedNodes2dFlattened(SharedData_Hip &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTensorCollocatedNodes2dFlattened() argument 185 …llocatedNodes2dFlattened(SharedData_Hip &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTransposeTensorCollocatedNodes2dFlattened() argument 201 …ened(SharedData_Hip &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, const CeedScalar *… in GradTensor2dFlattened() argument 209 …ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[comp + 0 * N… in GradTensor2dFlattened() [all …]
|
| H A D | hip-shared-basis-nontensor-templates.h | 47 …id InterpNonTensor(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpNonTensor() argument 50 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpNonTensor() 58 …TransposeNonTensor(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *c_B, in InterpTransposeNonTensor() argument 62 ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTransposeNonTensor()
|
| H A D | hip-shared-basis-tensor.h | 19 …void Interp(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, Cee… in __launch_bounds__() 34 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__() 58 …void InterpCollocated(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict_… in __launch_bounds__() 88 …void InterpTranspose(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__… in __launch_bounds__() 103 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__() 127 …void InterpCollocatedTranspose(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__… in __launch_bounds__() 157 …void InterpTransposeAdd(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restric… in __launch_bounds__() 172 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__() 196 …void InterpCollocatedTransposeAdd(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar … in __launch_bounds__() 229 …BLOCK_SIZE) __global__ void Grad(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *… in __launch_bounds__() [all …]
|
| H A D | hip-shared-basis-tensor-at-points.h | 24 …void InterpAtPoints(const CeedInt num_elem, const CeedScalar *c_B, const CeedInt *points_per_elem,… in __launch_bounds__() 42 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__() 80 …void InterpTransposeAtPoints(const CeedInt num_elem, const CeedScalar *c_B, const CeedInt *points_… in __launch_bounds__() 98 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__() 150 …void InterpTransposeAddAtPoints(const CeedInt num_elem, const CeedScalar *c_B, const CeedInt *poin… in __launch_bounds__() 168 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__() 212 …void GradAtPoints(const CeedInt num_elem, const CeedScalar *c_B, const CeedInt *points_per_elem, c… in __launch_bounds__() 230 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__() 268 …void GradTransposeAtPoints(const CeedInt num_elem, const CeedScalar *c_B, const CeedInt *points_pe… in __launch_bounds__() 286 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__() [all …]
|
| H A D | hip-shared-basis-nontensor.h | 19 …void Interp(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, Cee… in __launch_bounds__() 34 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in __launch_bounds__() 46 …void InterpTranspose(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__… in __launch_bounds__() 61 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in __launch_bounds__() 73 …void InterpTransposeAdd(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restric… in __launch_bounds__() 88 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in __launch_bounds__()
|