diff --git a/backends/magma/ceed-magma-basis.c b/backends/magma/ceed-magma-basis.c index 3f0cd2f640..2122342be8 100644 --- a/backends/magma/ceed-magma-basis.c +++ b/backends/magma/ceed-magma-basis.c @@ -54,19 +54,6 @@ static int CeedBasisApply_Magma(CeedBasis basis, CeedInt num_elem, CeedTranspose else CeedCheck(e_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is required for this CeedEvalMode"); CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); - // Clear v for transpose operation - if (t_mode == CEED_TRANSPOSE) { - CeedSize length; - - CeedCallBackend(CeedVectorGetLength(v, &length)); - if (CEED_SCALAR_TYPE == CEED_SCALAR_FP32) { - magmablas_slaset(MagmaFull, length, 1, 0.0, 0.0, (float *)d_v, length, data->queue); - } else { - magmablas_dlaset(MagmaFull, length, 1, 0.0, 0.0, (double *)d_v, length, data->queue); - } - ceed_magma_queue_sync(data->queue); - } - // Apply basis operation switch (e_mode) { case CEED_EVAL_INTERP: { @@ -289,19 +276,6 @@ static int CeedBasisApplyNonTensor_Magma(CeedBasis basis, CeedInt num_elem, Ceed else CeedCheck(e_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is required for this CeedEvalMode"); CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); - // Clear v for transpose operation - if (t_mode == CEED_TRANSPOSE) { - CeedSize length; - - CeedCallBackend(CeedVectorGetLength(v, &length)); - if (CEED_SCALAR_TYPE == CEED_SCALAR_FP32) { - magmablas_slaset(MagmaFull, length, 1, 0.0, 0.0, (float *)d_v, length, data->queue); - } else { - magmablas_dlaset(MagmaFull, length, 1, 0.0, 0.0, (double *)d_v, length, data->queue); - } - ceed_magma_queue_sync(data->queue); - } - // Apply basis operation if (e_mode != CEED_EVAL_WEIGHT) { const CeedScalar *d_b = NULL; diff --git a/include/ceed/jit-source/magma/magma-basis-grad-1d.h b/include/ceed/jit-source/magma/magma-basis-grad-1d.h index a338be2055..21572c5c3a 100644 --- a/include/ceed/jit-source/magma/magma-basis-grad-1d.h +++ b/include/ceed/jit-source/magma/magma-basis-grad-1d.h @@ -18,7 +18,7 @@ //////////////////////////////////////////////////////////////////////////////// // grad basis action (1D) template -static __device__ __inline__ void magma_grad_1d_device(const T *sT, magma_trans_t transT, T *sU[NUM_COMP], T *sV[NUM_COMP], const int tx) { +static __device__ __inline__ void magma_grad_1d_device(const T *sT, T *sU[NUM_COMP], T *sV[NUM_COMP], const int tx) { // Assumptions // 1. 1D threads of size max(P,Q) // 2. sU[i] is 1xP: in shared memory @@ -28,10 +28,9 @@ static __device__ __inline__ void magma_grad_1d_device(const T *sT, magma_trans_ // 6. Must sync before and after call // 7. Note that the layout for U and V is different from 2D/3D problem - T rv; if (tx < Q) { for (int comp = 0; comp < NUM_COMP; comp++) { - rv = (transT == MagmaTrans) ? sV[comp][tx] : 0.0; + T rv = 0.0; for (int i = 0; i < P; i++) { rv += sU[comp][i] * sT(i, tx); } @@ -46,10 +45,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ const int dstrdU, CeedScalar *dV, const int estrdV, const int cstrdV, const int dstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaNoTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -79,7 +77,7 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ read_1d(dU, cstrdU, sU, tx); __syncthreads(); - magma_grad_1d_device(sT, transT, sU, sV, tx); + magma_grad_1d_device(sT, sU, sV, tx); __syncthreads(); // write V @@ -92,10 +90,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ const int dstrdU, CeedScalar *dV, const int estrdV, const int cstrdV, const int dstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -124,11 +121,8 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ // read U read_1d(dU, cstrdU, sU, tx); - // read V - read_1d(dV, cstrdV, sV, tx); - __syncthreads(); - magma_grad_1d_device(sT, transT, sU, sV, tx); + magma_grad_1d_device(sT, sU, sV, tx); __syncthreads(); // write V diff --git a/include/ceed/jit-source/magma/magma-basis-grad-2d.h b/include/ceed/jit-source/magma/magma-basis-grad-2d.h index 64397f431b..8ed41718c0 100644 --- a/include/ceed/jit-source/magma/magma-basis-grad-2d.h +++ b/include/ceed/jit-source/magma/magma-basis-grad-2d.h @@ -16,6 +16,21 @@ #define sT(i, j) sT[(j)*P + (i)] #define sTmp(i, j, ldw) sTmp[(j) * (ldw) + (i)] +//////////////////////////////////////////////////////////////////////////////// +// Helper function to add or set into V +template +struct magma_grad_2d_device_accumulate; + +template +struct magma_grad_2d_device_accumulate { + static __device__ __inline__ void op(T &rV, const T &rTmp) { rV += rTmp; } +}; + +template +struct magma_grad_2d_device_accumulate { + static __device__ __inline__ void op(T &rV, const T &rTmp) { rV = rTmp; } +}; + //////////////////////////////////////////////////////////////////////////////// // grad basis action (2D) // This function is called two times at a higher level for 2D @@ -24,10 +39,9 @@ // i_DIM -- the index of the outermost loop over dimensions in grad // i_DIM_U -- which dim index of rU is accessed (always 0 for notrans, 0 or 1 for trans) // i_DIM_V -- which dim index of rV is accessed (0 or 1 for notrans, always 0 for trans) -// the scalar beta is used to specify whether to accumulate to rV, or overwrite it -template +template static __device__ __inline__ void magma_grad_2d_device(const T *sTinterp, const T *sTgrad, T rU[DIM_U][NUM_COMP][rU_SIZE], - T rV[DIM_V][NUM_COMP][rV_SIZE], T beta, const int tx, T rTmp, T *swork) { + T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx, T rTmp, T *swork) { // Assumptions // 0. This device routine applies grad for one dim only (i_DIM), so it should be called twice for 2D // 1. 1D threads of size max(P,Q) @@ -68,8 +82,7 @@ static __device__ __inline__ void magma_grad_2d_device(const T *sTinterp, const for (int i = 0; i < P; i++) { rTmp += sTmp(tx, i, sld) * sT(i, j); } - rV[i_DIM_V][comp][j] *= beta; - rV[i_DIM_V][comp][j] += rTmp; + magma_grad_2d_device_accumulate::op(rV[i_DIM_V][comp][j], rTmp); } } __syncthreads(); @@ -82,10 +95,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ const int dstrdU, CeedScalar *dV, const int estrdV, const int cstrdV, const int dstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaNoTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -109,22 +121,21 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ read_T_notrans_gm2sm(tx, dgrad1d, sTgrad); } - // No need to read V ( required only in transposed grad ) - const CeedScalar beta = 0.0; - /* read U (idim = 0 for dU, i_DIM = 0 for rU) -- there is a sync at the end of this function */ read_U_2d(dU + (0 * dstrdU), cstrdU, rU, sTmp, tx); /* first call (i_DIM = 0, i_DIM_U = 0, i_DIM_V = 0) -- output from rV[0][][] into dV (idim = 0) */ - magma_grad_2d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_2d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, + sTmp); /* there is a sync at the end of magma_grad_2d_device */ write_V_2d(dV + (0 * dstrdV), cstrdV, rV, tx); /* second call (i_DIM = 1, i_DIM_U = 0, i_DIM_V = 0) -- output from rV[0][][] into dV (idim = 1) */ - magma_grad_2d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_2d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, + sTmp); /* there is a sync at the end of magma_grad_2d_device */ write_V_2d(dV + (1 * dstrdV), cstrdV, rV, tx); } @@ -135,10 +146,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ const int dstrdU, CeedScalar *dV, const int estrdV, const int cstrdV, const int dstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -163,23 +173,18 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ } __syncthreads(); - /* read V (since this is transposed mode -- - idim = 0 for dV, i_DIM = 0 for rV) */ - const CeedScalar beta = 1.0; - read_V_2d(dV + (0 * dstrdV), cstrdV, rV, tx); - /* read U (idim = 0 for dU, i_DIM = 0 for rU) -- there is a sync at the end of this function */ read_U_2d(dU + (0 * dstrdU), cstrdU, rU, sTmp, tx); /* first call (i_DIM = 0, i_DIM_U = 0, i_DIM_V = 0) */ - magma_grad_2d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_2d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, sTmp); /* there is a sync at the end of magma_grad_2d_device */ /* read U (idim = 1 for dU, i_DIM = 0 for rU) -- there is a sync at the end of this function */ read_U_2d(dU + (1 * dstrdU), cstrdU, rU, sTmp, tx); /* second call (i_DIM = 1, i_DIM_U = 0, i_DIM_V = 0) */ - magma_grad_2d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_2d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, sTmp); /* there is a sync at the end of magma_grad_2d_device */ // write V diff --git a/include/ceed/jit-source/magma/magma-basis-grad-3d.h b/include/ceed/jit-source/magma/magma-basis-grad-3d.h index 6a56d47c66..632fe1e5c9 100644 --- a/include/ceed/jit-source/magma/magma-basis-grad-3d.h +++ b/include/ceed/jit-source/magma/magma-basis-grad-3d.h @@ -17,6 +17,21 @@ #define sTmp(i, j, ldw) sTmp[(j) * (ldw) + (i)] #define sTmp2(i, j, ldw) sTmp2[(j) * (ldw) + (i)] +//////////////////////////////////////////////////////////////////////////////// +// Helper function to add or set into V +template +struct magma_grad_3d_device_accumulate; + +template +struct magma_grad_3d_device_accumulate { + static __device__ __inline__ void op(T &rV, const T &rTmp) { rV += rTmp; } +}; + +template +struct magma_grad_3d_device_accumulate { + static __device__ __inline__ void op(T &rV, const T &rTmp) { rV = rTmp; } +}; + //////////////////////////////////////////////////////////////////////////////// // grad basis action (3D) // This function is called three times at a higher level for 3D @@ -25,10 +40,9 @@ // i_DIM -- the index of the outermost loop over dimensions in grad // i_DIM_U -- which dim index of rU is accessed (always 0 for notrans, 0, 1, or 2 for trans) // i_DIM_V -- which dim index of rV is accessed (0, 1, or 2 for notrans, always 0 for trans) -// the scalar beta is used to specify whether to accumulate to rV, or overwrite it -template +template static __device__ __inline__ void magma_grad_3d_device(const T *sTinterp, const T *sTgrad, T rU[DIM_U][NUM_COMP][rU_SIZE], - T rV[DIM_V][NUM_COMP][rV_SIZE], T beta, const int tx, T rTmp, T *swork) { + T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx, T rTmp, T *swork) { // Assumptions // 0. This device routine applies grad for one dim only (i_DIM), so it should be thrice for 3D // 1. 1D threads of size max(P,Q)^2 @@ -90,8 +104,7 @@ static __device__ __inline__ void magma_grad_3d_device(const T *sTinterp, const for (int i = 0; i < P; i++) { rTmp += sTmp(tx, i, sld) * sT(i, j); } - rV[i_DIM_V][comp][j] *= beta; - rV[i_DIM_V][comp][j] += rTmp; + magma_grad_3d_device_accumulate::op(rV[i_DIM_V][comp][j], rTmp); } } __syncthreads(); @@ -104,10 +117,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q *BASIS_MAX_P_Q, MA const int dstrdU, CeedScalar *dV, const int estrdV, const int cstrdV, const int dstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaNoTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -132,28 +144,28 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q *BASIS_MAX_P_Q, MA } __syncthreads(); - // No need to read V ( required only in transposed grad ) - const CeedScalar beta = 0.0; - /* read U (idim = 0 for dU, i_DIM = 0 for rU) -- there is a sync at the end of this function */ read_U_3d(dU + (0 * dstrdU), cstrdU, rU, sTmp, tx); /* first call (i_DIM = 0, i_DIM_U = 0, i_DIM_V = 0) -- output from rV[0][][] into dV (idim = 0) */ - magma_grad_3d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_3d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, + sTmp); /* there is a sync at the end of magma_grad_3d_device */ write_V_3d(dV + (0 * dstrdV), cstrdV, rV, tx); /* second call (i_DIM = 1, i_DIM_U = 0, i_DIM_V = 0) -- output from rV[0][][] into dV (idim = 1) */ - magma_grad_3d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_3d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, + sTmp); /* there is a sync at the end of magma_grad_3d_device */ write_V_3d(dV + (1 * dstrdV), cstrdV, rV, tx); /* third call (i_DIM = 2, i_DIM_U = 0, i_DIM_V = 0) -- output from rV[0][][] into dV (idim = 2) */ - magma_grad_3d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_3d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, + sTmp); /* there is a sync at the end of magma_grad_3d_device */ write_V_3d(dV + (2 * dstrdV), cstrdV, rV, tx); } @@ -164,10 +176,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q *BASIS_MAX_P_Q, MA const int dstrdU, CeedScalar *dV, const int estrdV, const int cstrdV, const int dstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -192,29 +203,25 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q *BASIS_MAX_P_Q, MA } __syncthreads(); - // read V (since this is transposed mode) - const CeedScalar beta = 1.0; - read_V_3d(dV + (0 * dstrdV), cstrdV, rV, tx); - /* read U (idim = 0 for dU, i_DIM = 0 for rU) -- there is a sync at the end of this function */ read_U_3d(dU + (0 * dstrdU), cstrdU, rU, sTmp, tx); /* then first call (i_DIM = 0, i_DIM_U = 0, i_DIM_V = 0) */ - magma_grad_3d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_3d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, sTmp); /* there is a sync at the end of magma_grad_3d_device */ /* read U (idim = 1 for dU, i_DIM = 0 for rU) -- there is a sync at the end of this function */ read_U_3d(dU + (1 * dstrdU), cstrdU, rU, sTmp, tx); /* then second call (i_DIM = 1, i_DIM_U = 0, i_DIM_V = 0) */ - magma_grad_3d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_3d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, sTmp); /* there is a sync at the end of magma_grad_3d_device */ /* read U (idim = 2 for dU, i_DIM = 0 for rU) -- there is a sync at the end of this function */ read_U_3d(dU + (2 * dstrdU), cstrdU, rU, sTmp, tx); /* then third call (i_DIM = 2, i_DIM_U = 0, i_DIM_V = 0) */ - magma_grad_3d_device(sTinterp, sTgrad, rU, rV, beta, tx, rTmp, sTmp); + magma_grad_3d_device(sTinterp, sTgrad, rU, rV, tx, rTmp, sTmp); /* there is a sync at the end of magma_grad_3d_device */ // write V diff --git a/include/ceed/jit-source/magma/magma-basis-interp-1d.h b/include/ceed/jit-source/magma/magma-basis-interp-1d.h index 0c49fa3696..9b6787d24d 100644 --- a/include/ceed/jit-source/magma/magma-basis-interp-1d.h +++ b/include/ceed/jit-source/magma/magma-basis-interp-1d.h @@ -18,7 +18,7 @@ //////////////////////////////////////////////////////////////////////////////// // interp basis action (1D) template -static __device__ __inline__ void magma_interp_1d_device(const T *sT, magma_trans_t transT, T *sU[NUM_COMP], T *sV[NUM_COMP], const int tx) { +static __device__ __inline__ void magma_interp_1d_device(const T *sT, T *sU[NUM_COMP], T *sV[NUM_COMP], const int tx) { // Assumptions // 1. 1D threads of size max(P,Q) // 2. sU[i] is 1xP: in shared memory @@ -28,10 +28,9 @@ static __device__ __inline__ void magma_interp_1d_device(const T *sT, magma_tran // 6. Must sync before and after call // 7. Note that the layout for U and V is different from 2D/3D problem - T rv; if (tx < Q) { for (int comp = 0; comp < NUM_COMP; comp++) { - rv = (transT == MagmaTrans) ? sV[comp][tx] : 0.0; + T rv = 0.0; for (int i = 0; i < P; i++) { rv += sU[comp][i] * sT(i, tx); // sT[tx * P + i]; } @@ -46,10 +45,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ const int cstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaNoTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -79,7 +77,7 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ read_1d(dU, cstrdU, sU, tx); __syncthreads(); - magma_interp_1d_device(sT, transT, sU, sV, tx); + magma_interp_1d_device(sT, sU, sV, tx); __syncthreads(); // write V @@ -92,10 +90,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ const int cstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -124,11 +121,8 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ // read U read_1d(dU, cstrdU, sU, tx); - // read V - read_1d(dV, cstrdV, sV, tx); - __syncthreads(); - magma_interp_1d_device(sT, transT, sU, sV, tx); + magma_interp_1d_device(sT, sU, sV, tx); __syncthreads(); // write V diff --git a/include/ceed/jit-source/magma/magma-basis-interp-2d.h b/include/ceed/jit-source/magma/magma-basis-interp-2d.h index ece85a26fe..206f5cc32f 100644 --- a/include/ceed/jit-source/magma/magma-basis-interp-2d.h +++ b/include/ceed/jit-source/magma/magma-basis-interp-2d.h @@ -19,8 +19,8 @@ //////////////////////////////////////////////////////////////////////////////// // interp basis action (2D) template -static __device__ __inline__ void magma_interp_2d_device(const T *sT, magma_trans_t transT, T rU[DIM_U][NUM_COMP][rU_SIZE], - T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx, T rTmp, T *swork) { +static __device__ __inline__ void magma_interp_2d_device(const T *sT, T rU[DIM_U][NUM_COMP][rU_SIZE], T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx, + T rTmp, T *swork) { // Assumptions // 1. 1D threads of size max(P,Q) // 2. input: rU[DIM_U x NUM_COMP x rU_SIZE] in registers (per thread) @@ -71,10 +71,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ const int cstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaNoTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -100,7 +99,7 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ read_U_2d(dU, cstrdU, rU, sTmp, tx); // no sync needed here -- read_U_2d already syncs at the end - magma_interp_2d_device(sT, transT, rU, rV, tx, rTmp, sTmp); + magma_interp_2d_device(sT, rU, rV, tx, rTmp, sTmp); __syncthreads(); // write V @@ -113,10 +112,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ const int cstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -138,14 +136,11 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q, MAGMA_MAXTHREADS_ read_T_trans_gm2sm(tx, dT, sT); } - // read V - read_V_2d(dV, cstrdV, rV, tx); - // read U -- there is a sync at the end of this function read_U_2d(dU, cstrdU, rU, sTmp, tx); // no sync needed here -- read_U_2d already syncs at the end - magma_interp_2d_device(sT, transT, rU, rV, tx, rTmp, sTmp); + magma_interp_2d_device(sT, rU, rV, tx, rTmp, sTmp); __syncthreads(); // write V diff --git a/include/ceed/jit-source/magma/magma-basis-interp-3d.h b/include/ceed/jit-source/magma/magma-basis-interp-3d.h index 9c5020c3ff..2d6fa106ca 100644 --- a/include/ceed/jit-source/magma/magma-basis-interp-3d.h +++ b/include/ceed/jit-source/magma/magma-basis-interp-3d.h @@ -19,8 +19,8 @@ //////////////////////////////////////////////////////////////////////////////// // interp basis action (3D) template -static __device__ __inline__ void magma_interp_3d_device(const T *sT, magma_trans_t transT, T rU[DIM_U][NUM_COMP][rU_SIZE], - T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx, T rTmp[Q], T *swork) { +static __device__ __inline__ void magma_interp_3d_device(const T *sT, T rU[DIM_U][NUM_COMP][rU_SIZE], T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx, + T rTmp[Q], T *swork) { // Assumptions // 1. 1D threads of size max(P,Q)^2 // 2. input: rU[DIM_U x NUM_COMP x rU_SIZE] in registers (per thread) @@ -99,10 +99,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q *BASIS_MAX_P_Q, MA const int cstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaNoTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -128,7 +127,7 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q *BASIS_MAX_P_Q, MA read_U_3d(dU, cstrdU, rU, sTmp, tx); // there is a sync at the end of this function - magma_interp_3d_device(sT, transT, rU, rV, tx, rTmp, sTmp); + magma_interp_3d_device(sT, rU, rV, tx, rTmp, sTmp); __syncthreads(); // write V @@ -141,10 +140,9 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q *BASIS_MAX_P_Q, MA const int cstrdV, const int nelem) { MAGMA_DEVICE_SHARED(CeedScalar, shared_data) - const int tx = threadIdx.x; - const int ty = threadIdx.y; - const int elem_id = (blockIdx.x * blockDim.y) + ty; - magma_trans_t transT = MagmaTrans; + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int elem_id = (blockIdx.x * blockDim.y) + ty; if (elem_id >= nelem) return; @@ -166,14 +164,11 @@ extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_MAX_P_Q *BASIS_MAX_P_Q, MA read_T_trans_gm2sm(tx, dT, sT); } - // read V - read_V_3d(dV, cstrdV, rV, tx); - // read U (idim = 0 for dU, i_DIM = 0 for rU, u_dimstride is always 0) read_U_3d(dU, cstrdU, rU, sTmp, tx); // there is a sync at the end of this function - magma_interp_3d_device(sT, transT, rU, rV, tx, rTmp, sTmp); + magma_interp_3d_device(sT, rU, rV, tx, rTmp, sTmp); __syncthreads(); // write V diff --git a/include/ceed/jit-source/magma/magma-common-defs.h b/include/ceed/jit-source/magma/magma-common-defs.h index 30c68dfb1e..0cb7c400cf 100644 --- a/include/ceed/jit-source/magma/magma-common-defs.h +++ b/include/ceed/jit-source/magma/magma-common-defs.h @@ -22,6 +22,4 @@ // Define macro for computing the total threads in a block for use with __launch_bounds__() #define MAGMA_BASIS_BOUNDS(x, maxt) (x * MAGMA_BASIS_NTCOL(x, maxt)) -typedef enum { MagmaNoTrans = 111, MagmaTrans = 112, MagmaConjTrans = 113, Magma_ConjTrans = MagmaConjTrans } magma_trans_t; - #endif // CEED_MAGMA_COMMON_DEFS_H