| /libCEED/include/ceed/jit-source/magma/ |
| H A D | magma-basis-grad-3d.h | 124 …CeedScalar rV[1][BASIS_NUM_COMP][BASIS_Q] = {0.0}; // here DIM_V = 1, but might be different for …
133 CeedScalar *sTgrad = sTinterp + BASIS_P * BASIS_Q;
134 CeedScalar *sTmp = sTgrad + BASIS_P * BASIS_Q;
135 …* (max(BASIS_P * BASIS_P * BASIS_P, (BASIS_P * BASIS_P * BASIS_Q) + (BASIS_P * BASIS_Q * BASIS_Q))…
139 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dinterp1d, sTinterp);
140 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dgrad1d, sTgrad);
150 …magma_grad_3d_device<CeedScalar, 1, 1, BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_P, BASIS_Q, 0, 0, 0…
153 write_V_3d<CeedScalar, BASIS_Q, 1, BASIS_NUM_COMP, BASIS_Q, 0>(dV + (0 * dstrdV), cstrdV, rV, tx);
157 …magma_grad_3d_device<CeedScalar, 1, 1, BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_P, BASIS_Q, 1, 0, 0…
160 write_V_3d<CeedScalar, BASIS_Q, 1, BASIS_NUM_COMP, BASIS_Q, 0>(dV + (1 * dstrdV), cstrdV, rV, tx);
[all …]
|
| H A D | magma-basis-grad-2d.h | 102 …CeedScalar rV[1][BASIS_NUM_COMP][BASIS_Q] = {0.0}; // here DIM_V = 1, but might be different for …
111 CeedScalar *sTgrad = sTinterp + BASIS_P * BASIS_Q;
112 CeedScalar *sTmp = sTgrad + BASIS_P * BASIS_Q;
117 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dinterp1d, sTinterp);
118 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dgrad1d, sTgrad);
127 …magma_grad_2d_device<CeedScalar, 1, 1, BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_P, BASIS_Q, 0, 0, 0…
130 write_V_2d<CeedScalar, BASIS_Q, 1, BASIS_NUM_COMP, BASIS_Q, 0>(dV + (0 * dstrdV), cstrdV, rV, tx);
134 …magma_grad_2d_device<CeedScalar, 1, 1, BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_P, BASIS_Q, 1, 0, 0…
137 write_V_2d<CeedScalar, BASIS_Q, 1, BASIS_NUM_COMP, BASIS_Q, 0>(dV + (1 * dstrdV), cstrdV, rV, tx);
152 …CeedScalar rU[1][BASIS_NUM_COMP][BASIS_Q] = {0.0}; // here DIM_U = 1, but might be different for …
[all …]
|
| H A D | magma-basis-interp-3d.h | 106 …CeedScalar rV[1][BASIS_NUM_COMP][BASIS_Q] = {0.0}; // for a non-fused operator BASIS_DIM is alway… in __launch_bounds__()
107 CeedScalar rTmp[BASIS_Q] = {0.0}; in __launch_bounds__()
115 CeedScalar *sTmp = sT + BASIS_P * BASIS_Q; in __launch_bounds__()
116 sTmp += ty * (max(BASIS_P * BASIS_P * BASIS_MAX_P_Q, BASIS_P * BASIS_Q * BASIS_Q)); in __launch_bounds__()
120 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dT, sT); in __launch_bounds__()
127 …magma_interp_3d_device<CeedScalar, 1, 1, BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_P, BASIS_Q>(sT, r… in __launch_bounds__()
131 write_V_3d<CeedScalar, BASIS_Q, 1, BASIS_NUM_COMP, BASIS_Q, 0>(dV, cstrdV, rV, tx); in __launch_bounds__()
146 …CeedScalar rU[1][BASIS_NUM_COMP][BASIS_Q] = {0.0}; // for a non-fused operator BASIS_DIM is alway… in __launch_bounds__()
156 CeedScalar *sTmp = sT + BASIS_Q * BASIS_P; in __launch_bounds__()
157 sTmp += ty * (max(BASIS_Q * BASIS_Q * BASIS_MAX_P_Q, BASIS_Q * BASIS_P * BASIS_P)); in __launch_bounds__()
[all …]
|
| H A D | magma-basis-interp-2d.h | 78 …CeedScalar rV[1][BASIS_NUM_COMP][BASIS_Q] = {0.0}; // for a non-fused operator BASIS_DIM is alway… in __launch_bounds__()
87 CeedScalar *sTmp = sT + BASIS_P * BASIS_Q; in __launch_bounds__()
92 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dT, sT); in __launch_bounds__()
99 …magma_interp_2d_device<CeedScalar, 1, 1, BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_P, BASIS_Q>(sT, r… in __launch_bounds__()
103 write_V_2d<CeedScalar, BASIS_Q, 1, BASIS_NUM_COMP, BASIS_Q, 0>(dV, cstrdV, rV, tx); in __launch_bounds__()
118 …CeedScalar rU[1][BASIS_NUM_COMP][BASIS_Q] = {0.0}; // for a non-fused operator BASIS_DIM is alway… in __launch_bounds__()
128 CeedScalar *sTmp = sT + BASIS_Q * BASIS_P; in __launch_bounds__()
129 sTmp += ty * (BASIS_Q * BASIS_MAX_P_Q); in __launch_bounds__()
133 read_T_trans_gm2sm<BASIS_Q, BASIS_P>(tx, dT, sT); in __launch_bounds__()
137 read_U_2d<CeedScalar, BASIS_Q, 1, BASIS_NUM_COMP, BASIS_Q, 0>(dU, cstrdU, rU, sTmp, tx); in __launch_bounds__()
[all …]
|
| H A D | magma-basis-weight-3d.h | 34 extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q *BASIS_Q, MAGMA_MAXTHREADS_3D)) __global__ in __launch_bounds__() argument
44 …CeedScalar rV[1][1][BASIS_Q]; // allocate with BASIS_DIM=BASIS_NUM_COMP=1, but sizes may differ f… in __launch_bounds__()
52 if (tx < BASIS_Q) { in __launch_bounds__()
57 magma_weight_3d_device<CeedScalar, 1, 1, BASIS_Q, 0, 0>(sTweight, rV, tx); in __launch_bounds__()
60 if (tx < (BASIS_Q * BASIS_Q)) { in __launch_bounds__()
61 for (int j = 0; j < BASIS_Q; j++) { in __launch_bounds__()
62 dV[j * (BASIS_Q * BASIS_Q) + tx] = rV[0][0][j]; in __launch_bounds__()
|
| H A D | magma-basis-grad-1d.h | 60 CeedScalar *sW = sT + BASIS_P * BASIS_Q; in __launch_bounds__()
61 sU[0] = sW + ty * BASIS_NUM_COMP * (BASIS_P + BASIS_Q); in __launch_bounds__()
65 sV[comp] = sV[comp - 1] + (1 * BASIS_Q); in __launch_bounds__()
70 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dTgrad, sT); in __launch_bounds__()
77 magma_grad_1d_device<CeedScalar, BASIS_DIM, BASIS_NUM_COMP, BASIS_P, BASIS_Q>(sT, sU, sV, tx); in __launch_bounds__()
81 write_1d<CeedScalar, BASIS_Q, BASIS_NUM_COMP>(sV, dV, cstrdV, tx); in __launch_bounds__()
105 CeedScalar *sW = sT + BASIS_Q * BASIS_P; in __launch_bounds__()
106 sU[0] = sW + ty * BASIS_NUM_COMP * (BASIS_Q + BASIS_P); in __launch_bounds__()
107 sV[0] = sU[0] + (BASIS_NUM_COMP * 1 * BASIS_Q); in __launch_bounds__()
109 sU[comp] = sU[comp - 1] + (1 * BASIS_Q); in __launch_bounds__()
[all …]
|
| H A D | magma-basis-interp-1d.h | 60 CeedScalar *sW = sT + BASIS_P * BASIS_Q; in __launch_bounds__()
61 sU[0] = sW + ty * BASIS_NUM_COMP * (BASIS_P + BASIS_Q); in __launch_bounds__()
65 sV[comp] = sV[comp - 1] + (1 * BASIS_Q); in __launch_bounds__()
70 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dT, sT); in __launch_bounds__()
77 magma_interp_1d_device<CeedScalar, BASIS_DIM, BASIS_NUM_COMP, BASIS_P, BASIS_Q>(sT, sU, sV, tx); in __launch_bounds__()
81 write_1d<CeedScalar, BASIS_Q, BASIS_NUM_COMP>(sV, dV, cstrdV, tx); in __launch_bounds__()
105 CeedScalar *sW = sT + BASIS_Q * BASIS_P; in __launch_bounds__()
106 sU[0] = sW + ty * BASIS_NUM_COMP * (BASIS_Q + BASIS_P); in __launch_bounds__()
107 sV[0] = sU[0] + (BASIS_NUM_COMP * 1 * BASIS_Q); in __launch_bounds__()
109 sU[comp] = sU[comp - 1] + (1 * BASIS_Q); in __launch_bounds__()
[all …]
|
| H A D | magma-basis-weight-nontensor.h | 13 extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q, MAGMA_MAXTHREADS_1D)) __global__ in __launch_bounds__() argument
24 dV += id * BASIS_Q; in __launch_bounds__()
28 CeedScalar *sV = sqweight + BASIS_Q; in __launch_bounds__()
29 sV += ty * BASIS_Q; in __launch_bounds__()
32 if (ty == 0 && tx < BASIS_Q) { in __launch_bounds__()
37 if (tx < BASIS_Q) { in __launch_bounds__()
|
| H A D | magma-basis-weight-2d.h | 33 extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q, MAGMA_MAXTHREADS_2D)) __global__ in __launch_bounds__() argument
43 …CeedScalar rV[1][1][BASIS_Q]; // allocate with BASIS_DIM=BASIS_NUM_COMP=1, but sizes may differ f… in __launch_bounds__()
51 if (ty == 0 && tx < BASIS_Q) { in __launch_bounds__()
56 magma_weight_2d_device<CeedScalar, 1, 1, BASIS_Q, 0, 0>(sTweight, rV, tx); in __launch_bounds__()
59 if (tx < BASIS_Q) { in __launch_bounds__()
60 for (int j = 0; j < BASIS_Q; j++) { in __launch_bounds__()
61 dV[j * BASIS_Q + tx] = rV[0][0][j]; in __launch_bounds__()
|
| H A D | magma-basis-weight-1d.h | 25 extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q, MAGMA_MAXTHREADS_1D)) __global__ in __launch_bounds__() argument
40 CeedScalar *sV = sTweight + BASIS_Q; in __launch_bounds__()
41 sV += ty * BASIS_Q; in __launch_bounds__()
44 if (ty == 0 && tx < BASIS_Q) { in __launch_bounds__()
49 magma_weight_1d_device<CeedScalar, BASIS_Q>(sTweight, sV, tx); in __launch_bounds__()
|
| H A D | magma-basis-interp-deriv-nontensor.h | 256 extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q, MAGMA_MAXTHREADS_1D)) __global__ in __launch_bounds__() argument
261 …magma_basis_nontensor_device_n1<CeedScalar, BASIS_P, BASIS_Q, BASIS_NB_INTERP_N>(n, dA, dB, dC, (C… in __launch_bounds__()
263 …magma_basis_nontensor_device_n<CeedScalar, BASIS_Q_COMP_INTERP, BASIS_P, BASIS_Q, BASIS_NB_INTERP_… in __launch_bounds__()
273 …magma_basis_nontensor_device_t1<CeedScalar, BASIS_P, BASIS_Q, BASIS_NB_INTERP_T>(n, dA, dB, dC, (C… in __launch_bounds__()
275 …magma_basis_nontensor_device_t<CeedScalar, BASIS_Q_COMP_INTERP, BASIS_P, BASIS_Q, BASIS_NB_INTERP_… in __launch_bounds__()
285 …magma_basis_nontensor_device_ta1<CeedScalar, BASIS_P, BASIS_Q, BASIS_NB_INTERP_T>(n, dA, dB, dC, (… in __launch_bounds__()
287 …magma_basis_nontensor_device_ta<CeedScalar, BASIS_Q_COMP_INTERP, BASIS_P, BASIS_Q, BASIS_NB_INTERP… in __launch_bounds__()
292 extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q, MAGMA_MAXTHREADS_1D)) __global__ in __launch_bounds__() argument
297 …magma_basis_nontensor_device_n1<CeedScalar, BASIS_P, BASIS_Q, BASIS_NB_DERIV_N>(n, dA, dB, dC, (Ce… in __launch_bounds__()
299 …magma_basis_nontensor_device_n<CeedScalar, BASIS_Q_COMP_DERIV, BASIS_P, BASIS_Q, BASIS_NB_DERIV_N>… in __launch_bounds__()
[all …]
|
| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-ref-basis-nontensor.h | 24 …BASIS_NUM_COMP, BASIS_Q_COMP_INTERP, BASIS_P, BASIS_Q>(elem, BASIS_P, BASIS_Q, BASIS_P * num_elem,… in Interp()
25 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in Interp()
32 …Transpose<BASIS_NUM_COMP, BASIS_Q_COMP_INTERP, BASIS_P, BASIS_Q>(elem, BASIS_Q, BASIS_P, BASIS_Q *… in InterpTranspose()
33 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in InterpTranspose()
43 …<BASIS_NUM_COMP, BASIS_Q_COMP_DERIV, BASIS_P, BASIS_Q>(elem, BASIS_P, BASIS_Q, BASIS_P * num_elem,… in Deriv()
44 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in Deriv()
51 …tTranspose<BASIS_NUM_COMP, BASIS_Q_COMP_DERIV, BASIS_P, BASIS_Q>(elem, BASIS_Q, BASIS_P, BASIS_Q *… in DerivTranspose()
52 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in DerivTranspose()
64 d_V[elem * BASIS_Q + t_id] = q_weight[t_id]; in Weight()
|
| H A D | hip-shared-basis-nontensor.h | 33 __shared__ CeedScalar s_B[BASIS_P * BASIS_Q]; in __launch_bounds__()
34 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in __launch_bounds__()
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__()
60 __shared__ CeedScalar s_B[BASIS_P * BASIS_Q]; in __launch_bounds__()
61 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in __launch_bounds__()
66 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U… in __launch_bounds__()
67 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
87 __shared__ CeedScalar s_B[BASIS_P * BASIS_Q]; in __launch_bounds__()
88 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in __launch_bounds__()
[all …]
|
| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-ref-basis-nontensor.h | 24 …BASIS_NUM_COMP, BASIS_Q_COMP_INTERP, BASIS_P, BASIS_Q>(elem, BASIS_P, BASIS_Q, BASIS_P * num_elem,… in Interp()
25 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in Interp()
32 …Transpose<BASIS_NUM_COMP, BASIS_Q_COMP_INTERP, BASIS_P, BASIS_Q>(elem, BASIS_Q, BASIS_P, BASIS_Q *… in InterpTranspose()
33 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in InterpTranspose()
43 …<BASIS_NUM_COMP, BASIS_Q_COMP_DERIV, BASIS_P, BASIS_Q>(elem, BASIS_P, BASIS_Q, BASIS_P * num_elem,… in Deriv()
44 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in Deriv()
51 …tTranspose<BASIS_NUM_COMP, BASIS_Q_COMP_DERIV, BASIS_P, BASIS_Q>(elem, BASIS_Q, BASIS_P, BASIS_Q *… in DerivTranspose()
52 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in DerivTranspose()
64 d_V[elem * BASIS_Q + t_id] = q_weight[t_id]; in Weight()
|
| H A D | cuda-shared-basis-nontensor.h | 32 __shared__ CeedScalar s_B[BASIS_P * BASIS_Q]; in Interp()
33 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in Interp()
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()
59 __shared__ CeedScalar s_B[BASIS_P * BASIS_Q]; in InterpTranspose()
60 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in InterpTranspose()
65 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U… in InterpTranspose()
66 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in InterpTranspose()
86 __shared__ CeedScalar s_B[BASIS_P * BASIS_Q]; in InterpTransposeAdd()
87 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in InterpTransposeAdd()
[all …]
|