| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-shared-basis-tensor-at-points.h | 36 CeedScalar r_C[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in InterpAtPoints() local
49 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in InterpAtPoints()
52 InterpTensor2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in InterpAtPoints()
56 InterpTensor3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in InterpAtPoints()
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()
92 CeedScalar r_C[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in InterpTransposeAtPoints() local
103 for (CeedInt i = 0; i < BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1); i++) r_C[i] = 0.0; in InterpTransposeAtPoints()
125 …TransposeAtPoints1d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_U, r_X, r_C); in InterpTransposeAtPoints()
[all …]
|
| H A D | cuda-shared-basis-tensor-at-points-templates.h | 44 …s1d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__restrict__ r_C, const CeedScalar *… in InterpAtPoints1d() argument
52 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in InterpAtPoints1d()
66 CeedScalar *__restrict__ r_C) { in InterpTransposeAtPoints1d() argument
82 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in InterpTransposeAtPoints1d()
90 …s1d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__restrict__ r_C, const CeedScalar *… in GradAtPoints1d() argument
99 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in GradAtPoints1d()
113 CeedScalar *__restrict__ r_C) { in GradTransposeAtPoints1d() argument
129 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in GradTransposeAtPoints1d()
141 …s2d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__restrict__ r_C, const CeedScalar *… in InterpAtPoints2d() argument
150 …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[comp]; in InterpAtPoints2d()
[all …]
|
| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-shared-basis-tensor-at-points.h | 37 CeedScalar r_C[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in __launch_bounds__() local
50 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in __launch_bounds__()
53 InterpTensor2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in __launch_bounds__()
57 InterpTensor3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in __launch_bounds__()
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__()
93 CeedScalar r_C[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in __launch_bounds__() local
104 for (CeedInt i = 0; i < BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1); i++) r_C[i] = 0.0; in __launch_bounds__()
126 …TransposeAtPoints1d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_U, r_X, r_C); in __launch_bounds__()
[all …]
|
| H A D | hip-shared-basis-tensor-at-points-templates.h | 44 …ts1d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__restrict__ r_C, const CeedScalar *… in InterpAtPoints1d() argument
53 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in InterpAtPoints1d()
67 … const CeedScalar *__restrict__ r_X, CeedScalar *__restrict__ r_C) { in InterpTransposeAtPoints1d() argument
83 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in InterpTransposeAtPoints1d()
91 …ts1d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__restrict__ r_C, const CeedScalar *… in GradAtPoints1d() argument
100 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in GradAtPoints1d()
114 … const CeedScalar *__restrict__ r_X, CeedScalar *__restrict__ r_C) { in GradTransposeAtPoints1d() argument
130 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in GradTransposeAtPoints1d()
142 …ts2d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__restrict__ r_C, const CeedScalar *… in InterpAtPoints2d() argument
151 …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[comp]; in InterpAtPoints2d()
[all …]
|