Merge branch 'jm-gromacs-2021.4-patched' into '2022'

Include patch to fix GPU synchronization bug in GROMACS.

See merge request hps-public/easybuild-repository!1330

Golden_Repo/g/GROMACS/GROMACS-2021.4-gpsmkl-2021b-plumed.eb

@@ -38,7 +38,16 @@ toolchainopts = {'openmp': True, 'usempi': True}
 
 source_urls = ['ftp://ftp.gromacs.org/pub/gromacs/']
 sources = [SOURCELOWER_TAR_GZ]
-checksums = ['cb708a3e3e83abef5ba475fdb62ef8d42ce8868d68f52dafdb6702dc9742ba1d']
+
+# Fix missing synchronization in GPU code.
+patches = [
+    'GROMACS-2021_add-missing-sync.patch',
+]
+
+checksums = [
+    'cb708a3e3e83abef5ba475fdb62ef8d42ce8868d68f52dafdb6702dc9742ba1d',
+    '52ee257309ff7761c2dd5b26de7dbc63f8ba698082adb88e2843f90e3f9168bf',  # GROMACS-2021_add-missing-sync.patch
+]
 
 builddependencies = [
     ('CMake', '3.21.1', '', SYSTEM),

Golden_Repo/g/GROMACS/GROMACS-2021.4-gpsmkl-2021b.eb

@@ -35,7 +35,16 @@ toolchainopts = {'openmp': True, 'usempi': True}
 
 source_urls = ['ftp://ftp.gromacs.org/pub/gromacs/']
 sources = [SOURCELOWER_TAR_GZ]
-checksums = ['cb708a3e3e83abef5ba475fdb62ef8d42ce8868d68f52dafdb6702dc9742ba1d']
+
+# Fix missing synchronization in GPU code.
+patches = [
+    'GROMACS-2021_add-missing-sync.patch',
+]
+
+checksums = [
+    'cb708a3e3e83abef5ba475fdb62ef8d42ce8868d68f52dafdb6702dc9742ba1d',
+    '52ee257309ff7761c2dd5b26de7dbc63f8ba698082adb88e2843f90e3f9168bf',  # GROMACS-2021_add-missing-sync.patch
+]
 
 builddependencies = [
     ('CMake', '3.21.1', '', SYSTEM),

Golden_Repo/g/GROMACS/GROMACS-2021_add-missing-sync.patch

+Add missing sync in LINCS and SETTLE CUDA kernels
+https://gitlab.com/gromacs/gromacs/-/merge_requests/2499
+Patch added to EasyBuild by Simon Branford (University of Birmingham)
+diff --git a/src/gromacs/mdlib/lincs_gpu.cu b/src/gromacs/mdlib/lincs_gpu.cu
+index 27ad86f17f..4964a3d040 100644
+--- a/src/gromacs/mdlib/lincs_gpu.cu
++++ b/src/gromacs/mdlib/lincs_gpu.cu
+@@ -382,11 +382,15 @@ __launch_bounds__(c_maxThreadsPerBlock) __global__
+                             sm_threadVirial[d * blockDim.x + (threadIdx.x + dividedAt)];
+                 }
+             }
+-            // Syncronize if not within one warp
++            // Synchronize if not within one warp
+             if (dividedAt > warpSize / 2)
+             {
+                 __syncthreads();
+             }
++            else
++            {
++                __syncwarp();
++            }
+         }
+         // First 6 threads in the block add the results of 6 tensor components to the global memory address.
+         if (threadIdx.x < 6)
+diff --git a/src/gromacs/mdlib/settle_gpu.cu b/src/gromacs/mdlib/settle_gpu.cu
+index c02f0d5884..c6ca1e2452 100644
+--- a/src/gromacs/mdlib/settle_gpu.cu
++++ b/src/gromacs/mdlib/settle_gpu.cu
+@@ -352,6 +352,10 @@ __launch_bounds__(c_maxThreadsPerBlock) __global__
+             {
+                 __syncthreads();
+             }
++            else
++            {
++                __syncwarp();
++            }
+         }
+         // First 6 threads in the block add the 6 components of virial to the global memory address
+         if (tib < 6)