Merge remote-tracking branch 'origin/release-2021' into master

[alexxy/gromacs.git] / src / gromacs / ewald / pme_spread.clh

diff --git a/src/gromacs/ewald/pme_spread.clh b/src/gromacs/ewald/pme_spread.clh
index 7f6f4beccdf4b224100d49c9953ec88721ca0d08..de624b5f95a1cf72ad0414b28cfeced18892d380 100644 (file)
--- a/src/gromacs/ewald/pme_spread.clh
+++ b/src/gromacs/ewald/pme_spread.clh
@@ -455,9 +455,18 @@ __attribute__((reqd_work_group_size(order, order, atomsPerBlock))) __kernel void
         pme_gpu_stage_atom_data(sm_coordinates, gm_coordinates, DIM);
 
         barrier(CLK_LOCAL_MEM_FENCE);
-        calculate_splines(kernelParams, atomIndexOffset, sm_coordinates, sm_coefficients, sm_theta,
-                          sm_gridlineIndices, sm_fractCoords, gm_theta, gm_dtheta,
-                          gm_gridlineIndices, gm_fractShiftsTable, gm_gridlineIndicesTable);
+        calculate_splines(kernelParams,
+                          atomIndexOffset,
+                          sm_coordinates,
+                          sm_coefficients,
+                          sm_theta,
+                          sm_gridlineIndices,
+                          sm_fractCoords,
+                          gm_theta,
+                          gm_dtheta,
+                          gm_gridlineIndices,
+                          gm_fractShiftsTable,
+                          gm_gridlineIndicesTable);
 #if !defined(_AMD_SOURCE_) && !defined(_NVIDIA_SOURCE_)
         /* This is only here for execution of e.g. 32-sized warps on 16-wide hardware; this was
          * __syncwarp() in CUDA. #2519
@@ -474,8 +483,8 @@ __attribute__((reqd_work_group_size(order, order, atomsPerBlock))) __kernel void
         /* Spline data - only thetas (dthetas will only be needed in gather) */
         pme_gpu_stage_atom_data(sm_theta, gm_theta, DIM * order);
         /* Gridline indices - they're actually int and not float, but C99 is angry about overloads */
-        pme_gpu_stage_atom_data((__local float*)sm_gridlineIndices,
-                                (__global const float*)gm_gridlineIndices, DIM);
+        pme_gpu_stage_atom_data(
+                (__local float*)sm_gridlineIndices, (__global const float*)gm_gridlineIndices, DIM);
 
         barrier(CLK_LOCAL_MEM_FENCE);
     }