diff --git a/.github/workflows/release-notes.yml b/.github/workflows/release-notes.yml index 7b297a7433..a8aa9d1972 100644 --- a/.github/workflows/release-notes.yml +++ b/.github/workflows/release-notes.yml @@ -20,7 +20,8 @@ jobs: - name: Check release notes run: | git fetch origin main - if git diff origin/main --exit-code include/ceed/*; then + if git diff origin/main --exit-code include/*.h && \ + git diff origin/main --exit-code include/ceed/*.h; then echo "No public interface changes detected" elif git diff origin/main --exit-code CHANGELOG.md; then echo "API changes detected, but release notes not updated" diff --git a/include/ceed/jit-source/cuda/cuda-shared-basis-tensor-flattened-templates.h b/include/ceed/jit-source/cuda/cuda-shared-basis-tensor-flattened-templates.h index 54594b3af4..b9bb0ddaa4 100644 --- a/include/ceed/jit-source/cuda/cuda-shared-basis-tensor-flattened-templates.h +++ b/include/ceed/jit-source/cuda/cuda-shared-basis-tensor-flattened-templates.h @@ -570,15 +570,15 @@ inline __device__ void GradTransposeTensor3dFlattened(SharedData_Cuda &data, Cee if (Q_1D != T_1D) QUnpack3d(data, t_id_x, t_id_y, t_id_z, r_U); for (CeedInt comp = 0; comp < NUM_COMP; comp++) { - ContractTransposeZ3dFlattened(data, &r_U[comp + 0 * NUM_COMP], c_B, r_t1); - ContractTransposeY3dFlattened(data, r_t1, c_B, r_t2); - ContractTransposeX3dFlattened(data, r_t2, c_G, &r_V[comp]); - ContractTransposeZ3dFlattened(data, &r_U[comp + 1 * NUM_COMP], c_B, r_t1); - ContractTransposeY3dFlattened(data, r_t1, c_G, r_t2); - ContractTransposeAddX3dFlattened(data, r_t2, c_B, &r_V[comp]); - ContractTransposeZ3dFlattened(data, &r_U[comp + 2 * NUM_COMP], c_G, r_t1); - ContractTransposeY3dFlattened(data, r_t1, c_B, r_t2); - ContractTransposeAddX3dFlattened(data, r_t2, c_B, &r_V[comp]); + ContractTransposeZ3dFlattened(data, t_id_x, t_id_y, t_id_z, &r_U[comp + 0 * NUM_COMP], c_B, r_t1); + ContractTransposeY3dFlattened(data, t_id_x, t_id_y, t_id_z, r_t1, c_B, r_t2); + ContractTransposeX3dFlattened(data, t_id_x, t_id_y, t_id_z, r_t2, c_G, &r_V[comp]); + ContractTransposeZ3dFlattened(data, t_id_x, t_id_y, t_id_z, &r_U[comp + 1 * NUM_COMP], c_B, r_t1); + ContractTransposeY3dFlattened(data, t_id_x, t_id_y, t_id_z, r_t1, c_G, r_t2); + ContractTransposeAddX3dFlattened(data, t_id_x, t_id_y, t_id_z, r_t2, c_B, &r_V[comp]); + ContractTransposeZ3dFlattened(data, t_id_x, t_id_y, t_id_z, &r_U[comp + 2 * NUM_COMP], c_G, r_t1); + ContractTransposeY3dFlattened(data, t_id_x, t_id_y, t_id_z, r_t1, c_B, r_t2); + ContractTransposeAddX3dFlattened(data, t_id_x, t_id_y, t_id_z, r_t2, c_B, &r_V[comp]); } __syncthreads(); if (P_1D != T_1D) QPack3d(data, t_id_x, t_id_y, t_id_z, r_V); @@ -611,22 +611,19 @@ inline __device__ void GradTensorCollocated3dFlattened(SharedData_Cuda &data, Ce // 3D derivatives transpose //------------------------------------------------------------------------------ template -inline __device__ void GradTransposeTensor3dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, - const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { +inline __device__ void GradTransposeTensorCollocated3dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *c_B, + const CeedScalar *c_G, CeedScalar *__restrict__ r_V) { const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_id_x / (T_1D * T_1D); CeedScalar r_t1[1], r_t2[1]; if (Q_1D != T_1D) QUnpack3d(data, t_id_x, t_id_y, t_id_z, r_U); for (CeedInt comp = 0; comp < NUM_COMP; comp++) { - ContractTransposeZ3dFlattened(data, &r_U[comp + 0 * NUM_COMP], c_B, r_t1); - ContractTransposeY3dFlattened(data, r_t1, c_B, r_t2); - ContractTransposeX3dFlattened(data, r_t2, c_G, &r_V[comp]); - ContractTransposeZ3dFlattened(data, &r_U[comp + 1 * NUM_COMP], c_B, r_t1); - ContractTransposeY3dFlattened(data, r_t1, c_G, r_t2); - ContractTransposeAddX3dFlattened(data, r_t2, c_B, &r_V[comp]); - ContractTransposeZ3dFlattened(data, &r_U[comp + 2 * NUM_COMP], c_G, r_t1); - ContractTransposeY3dFlattened(data, r_t1, c_B, r_t2); - ContractTransposeAddX3dFlattened(data, r_t2, c_B, &r_V[comp]); + ContractTransposeZ3dFlattened(data, t_id_x, t_id_y, t_id_z, &r_U[comp + 2 * NUM_COMP], c_G, r_t2); + ContractTransposeAddY3dFlattened(data, t_id_x, t_id_y, t_id_z, &r_U[comp + 1 * NUM_COMP], c_G, r_t2); + ContractTransposeAddX3dFlattened(data, t_id_x, t_id_y, t_id_z, &r_U[comp + 0 * NUM_COMP], c_G, r_t2); + ContractTransposeZ3dFlattened(data, t_id_x, t_id_y, t_id_z, r_t2, c_B, r_t1); + ContractTransposeY3dFlattened(data, t_id_x, t_id_y, t_id_z, r_t1, c_B, r_t2); + ContractTransposeX3dFlattened(data, t_id_x, t_id_y, t_id_z, r_t2, c_B, &r_V[comp]); } __syncthreads(); if (P_1D != T_1D) QPack3d(data, t_id_x, t_id_y, t_id_z, r_V);