diff --git a/include/ceed/jit-source/cuda/cuda-gen-templates.h b/include/ceed/jit-source/cuda/cuda-gen-templates.h index b5373320c5..17b9658cb5 100644 --- a/include/ceed/jit-source/cuda/cuda-gen-templates.h +++ b/include/ceed/jit-source/cuda/cuda-gen-templates.h @@ -179,13 +179,14 @@ inline __device__ void WriteLVecStrided2d(SharedData_Cuda &data, const CeedInt e template inline __device__ void ReadLVecStandard3d(SharedData_Cuda &data, const CeedInt num_nodes, const CeedInt elem, const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ d_u, CeedScalar *__restrict__ r_u) { - if (data.t_id_x < P_1d && data.t_id_y < P_1d) + if (data.t_id_x < P_1d && data.t_id_y < P_1d) { for (CeedInt z = 0; z < P_1d; z++) { const CeedInt node = data.t_id_x + data.t_id_y * P_1d + z * P_1d * P_1d; const CeedInt ind = indices[node + elem * P_1d * P_1d * P_1d]; for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[z + comp * P_1d] = d_u[ind + COMP_STRIDE * comp]; } + } } //------------------------------------------------------------------------------ @@ -194,13 +195,14 @@ inline __device__ void ReadLVecStandard3d(SharedData_Cuda &data, const CeedInt n template inline __device__ void ReadLVecStrided3d(SharedData_Cuda &data, const CeedInt elem, const CeedScalar *__restrict__ d_u, CeedScalar *__restrict__ r_u) { - if (data.t_id_x < P_1d && data.t_id_y < P_1d) + if (data.t_id_x < P_1d && data.t_id_y < P_1d) { for (CeedInt z = 0; z < P_1d; z++) { const CeedInt node = data.t_id_x + data.t_id_y * P_1d + z * P_1d * P_1d; const CeedInt ind = node * STRIDES_NODE + elem * STRIDES_ELEM; for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[z + comp * P_1d] = d_u[ind + comp * STRIDES_COMP]; } + } } //------------------------------------------------------------------------------ @@ -238,13 +240,14 @@ inline __device__ void ReadEVecSliceStrided3d(SharedData_Cuda &data, const CeedI template inline __device__ void WriteLVecStandard3d(SharedData_Cuda &data, const CeedInt num_nodes, const CeedInt elem, const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ r_v, CeedScalar *__restrict__ d_v) { - if (data.t_id_x < P_1d && data.t_id_y < P_1d) + if (data.t_id_x < P_1d && data.t_id_y < P_1d) { for (CeedInt z = 0; z < P_1d; z++) { const CeedInt node = data.t_id_x + data.t_id_y * P_1d + z * P_1d * P_1d; const CeedInt ind = indices[node + elem * P_1d * P_1d * P_1d]; for (CeedInt comp = 0; comp < NUM_COMP; comp++) atomicAdd(&d_v[ind + COMP_STRIDE * comp], r_v[z + comp * P_1d]); } + } } //------------------------------------------------------------------------------ @@ -253,13 +256,14 @@ inline __device__ void WriteLVecStandard3d(SharedData_Cuda &data, const CeedInt template inline __device__ void WriteLVecStrided3d(SharedData_Cuda &data, const CeedInt elem, const CeedScalar *__restrict__ r_v, CeedScalar *__restrict__ d_v) { - if (data.t_id_x < P_1d && data.t_id_y < P_1d) + if (data.t_id_x < P_1d && data.t_id_y < P_1d) { for (CeedInt z = 0; z < P_1d; z++) { const CeedInt node = data.t_id_x + data.t_id_y * P_1d + z * P_1d * P_1d; const CeedInt ind = node * STRIDES_NODE + elem * STRIDES_ELEM; for (CeedInt comp = 0; comp < NUM_COMP; comp++) d_v[ind + comp * STRIDES_COMP] += r_v[z + comp * P_1d]; } + } } //------------------------------------------------------------------------------ @@ -274,15 +278,19 @@ inline __device__ void GradColloSlice3d(SharedData_Cuda &data, const CeedInt q, __syncthreads(); // X derivative r_V[comp + 0 * NUM_COMP] = 0.0; - for (CeedInt i = 0; i < Q_1d; i++) - r_V[comp + 0 * NUM_COMP] += c_G[i + data.t_id_x * Q_1d] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction (X derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[comp + 0 * NUM_COMP] += c_G[i + data.t_id_x * Q_1d] * data.slice[i + data.t_id_y * T_1D]; + } // Y derivative r_V[comp + 1 * NUM_COMP] = 0.0; - for (CeedInt i = 0; i < Q_1d; i++) - r_V[comp + 1 * NUM_COMP] += c_G[i + data.t_id_y * Q_1d] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction (Y derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[comp + 1 * NUM_COMP] += c_G[i + data.t_id_y * Q_1d] * data.slice[data.t_id_x + i * T_1D]; + } // Z derivative r_V[comp + 2 * NUM_COMP] = 0.0; - for (CeedInt i = 0; i < Q_1d; i++) r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1d] * r_U[i + comp * Q_1d]; // Contract z direction (Z derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1d] * r_U[i + comp * Q_1d]; + } __syncthreads(); } } @@ -296,21 +304,24 @@ inline __device__ void GradColloSliceTranspose3d(SharedData_Cuda &data, const Ce CeedScalar *__restrict__ r_V) { if (data.t_id_x < Q_1d && data.t_id_y < Q_1d) { for (CeedInt comp = 0; comp < NUM_COMP; comp++) { - // X derivative data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 0 * NUM_COMP]; __syncthreads(); - for (CeedInt i = 0; i < Q_1d; i++) - r_V[q + comp * Q_1d] += c_G[data.t_id_x + i * Q_1d] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction (X derivative) + // X derivative + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[q + comp * Q_1d] += c_G[data.t_id_x + i * Q_1d] * data.slice[i + data.t_id_y * T_1D]; + } __syncthreads(); // Y derivative data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 1 * NUM_COMP]; __syncthreads(); - for (CeedInt i = 0; i < Q_1d; i++) - r_V[q + comp * Q_1d] += c_G[data.t_id_y + i * Q_1d] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction (Y derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[q + comp * Q_1d] += c_G[data.t_id_y + i * Q_1d] * data.slice[data.t_id_x + i * T_1D]; + } __syncthreads(); // Z derivative - for (CeedInt i = 0; i < Q_1d; i++) - r_V[i + comp * Q_1d] += c_G[i + q * Q_1d] * r_U[comp + 2 * NUM_COMP]; // PARTIAL contract z direction (Z derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[i + comp * Q_1d] += c_G[i + q * Q_1d] * r_U[comp + 2 * NUM_COMP]; + } } } } diff --git a/include/ceed/jit-source/hip/hip-gen-templates.h b/include/ceed/jit-source/hip/hip-gen-templates.h index 7059d8dafe..56521a3685 100644 --- a/include/ceed/jit-source/hip/hip-gen-templates.h +++ b/include/ceed/jit-source/hip/hip-gen-templates.h @@ -145,13 +145,14 @@ inline __device__ void WriteLVecStrided2d(SharedData_Hip &data, const CeedInt el template inline __device__ void ReadLVecStandard3d(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ d_u, CeedScalar *__restrict__ r_u) { - if (data.t_id_x < P_1d && data.t_id_y < P_1d) + if (data.t_id_x < P_1d && data.t_id_y < P_1d) { for (CeedInt z = 0; z < P_1d; z++) { const CeedInt node = data.t_id_x + data.t_id_y * P_1d + z * P_1d * P_1d; const CeedInt ind = indices[node + elem * P_1d * P_1d * P_1d]; for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[z + comp * P_1d] = d_u[ind + COMP_STRIDE * comp]; } + } } //------------------------------------------------------------------------------ @@ -159,13 +160,14 @@ inline __device__ void ReadLVecStandard3d(SharedData_Hip &data, const CeedInt nu //------------------------------------------------------------------------------ template inline __device__ void ReadLVecStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedScalar *__restrict__ d_u, CeedScalar *__restrict__ r_u) { - if (data.t_id_x < P_1d && data.t_id_y < P_1d) + if (data.t_id_x < P_1d && data.t_id_y < P_1d) { for (CeedInt z = 0; z < P_1d; z++) { const CeedInt node = data.t_id_x + data.t_id_y * P_1d + z * P_1d * P_1d; const CeedInt ind = node * STRIDES_NODE + elem * STRIDES_ELEM; for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[z + comp * P_1d] = d_u[ind + comp * STRIDES_COMP]; } + } } //------------------------------------------------------------------------------ @@ -203,13 +205,14 @@ inline __device__ void ReadEVecSliceStrided3d(SharedData_Hip &data, const CeedIn template inline __device__ void WriteLVecStandard3d(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ r_v, CeedScalar *__restrict__ d_v) { - if (data.t_id_x < P_1d && data.t_id_y < P_1d) + if (data.t_id_x < P_1d && data.t_id_y < P_1d) { for (CeedInt z = 0; z < P_1d; z++) { const CeedInt node = data.t_id_x + data.t_id_y * P_1d + z * P_1d * P_1d; const CeedInt ind = indices[node + elem * P_1d * P_1d * P_1d]; for (CeedInt comp = 0; comp < NUM_COMP; comp++) atomicAdd(&d_v[ind + COMP_STRIDE * comp], r_v[z + comp * P_1d]); } + } } //------------------------------------------------------------------------------ @@ -218,13 +221,14 @@ inline __device__ void WriteLVecStandard3d(SharedData_Hip &data, const CeedInt n template inline __device__ void WriteLVecStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedScalar *__restrict__ r_v, CeedScalar *__restrict__ d_v) { - if (data.t_id_x < P_1d && data.t_id_y < P_1d) + if (data.t_id_x < P_1d && data.t_id_y < P_1d) { for (CeedInt z = 0; z < P_1d; z++) { const CeedInt node = data.t_id_x + data.t_id_y * P_1d + z * P_1d * P_1d; const CeedInt ind = node * STRIDES_NODE + elem * STRIDES_ELEM; for (CeedInt comp = 0; comp < NUM_COMP; comp++) d_v[ind + comp * STRIDES_COMP] += r_v[z + comp * P_1d]; } + } } //------------------------------------------------------------------------------ @@ -239,15 +243,19 @@ inline __device__ void GradColloSlice3d(SharedData_Hip &data, const CeedInt q, c __syncthreads(); // X derivative r_V[comp + 0 * NUM_COMP] = 0.0; - for (CeedInt i = 0; i < Q_1d; i++) - r_V[comp + 0 * NUM_COMP] += c_G[i + data.t_id_x * Q_1d] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction (X derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[comp + 0 * NUM_COMP] += c_G[i + data.t_id_x * Q_1d] * data.slice[i + data.t_id_y * T_1D]; + } // Y derivative r_V[comp + 1 * NUM_COMP] = 0.0; - for (CeedInt i = 0; i < Q_1d; i++) - r_V[comp + 1 * NUM_COMP] += c_G[i + data.t_id_y * Q_1d] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction (Y derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[comp + 1 * NUM_COMP] += c_G[i + data.t_id_y * Q_1d] * data.slice[data.t_id_x + i * T_1D]; + } // Z derivative r_V[comp + 2 * NUM_COMP] = 0.0; - for (CeedInt i = 0; i < Q_1d; i++) r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1d] * r_U[i + comp * Q_1d]; // Contract z direction (Z derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1d] * r_U[i + comp * Q_1d]; + } __syncthreads(); } } @@ -264,18 +272,21 @@ inline __device__ void GradColloSliceTranspose3d(SharedData_Hip &data, const Cee // X derivative data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 0 * NUM_COMP]; __syncthreads(); - for (CeedInt i = 0; i < Q_1d; i++) - r_V[q + comp * Q_1d] += c_G[data.t_id_x + i * Q_1d] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction (X derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[q + comp * Q_1d] += c_G[data.t_id_x + i * Q_1d] * data.slice[i + data.t_id_y * T_1D]; + } __syncthreads(); // Y derivative data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 1 * NUM_COMP]; __syncthreads(); - for (CeedInt i = 0; i < Q_1d; i++) - r_V[q + comp * Q_1d] += c_G[data.t_id_y + i * Q_1d] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction (Y derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[q + comp * Q_1d] += c_G[data.t_id_y + i * Q_1d] * data.slice[data.t_id_x + i * T_1D]; + } __syncthreads(); // Z derivative - for (CeedInt i = 0; i < Q_1d; i++) - r_V[i + comp * Q_1d] += c_G[i + q * Q_1d] * r_U[comp + 2 * NUM_COMP]; // PARTIAL contract z direction (Z derivative) + for (CeedInt i = 0; i < Q_1d; i++) { + r_V[i + comp * Q_1d] += c_G[i + q * Q_1d] * r_U[comp + 2 * NUM_COMP]; + } } } }