Fix SYCL PME Solve kernel

- Avoid infinite recursion in sycl_2020::isfinite.
- Make one of the accessors optional.

diff --git a/src/gromacs/ewald/pme_solve_sycl.cpp b/src/gromacs/ewald/pme_solve_sycl.cpp
index 633cf31e8e690b5cfacadcdb8726dcb5699a5156..46883060e9b1367d8f34e48a145b530d250b1309 100644 (file)
--- a/src/gromacs/ewald/pme_solve_sycl.cpp
+++ b/src/gromacs/ewald/pme_solve_sycl.cpp
@@ -65,12 +65,15 @@ template<GridOrdering gridOrdering, bool computeEnergyAndVirial, int subGroupSiz
 auto makeSolveKernel(cl::sycl::handler&                            cgh,
                      DeviceAccessor<float, mode::read>             a_splineModuli,
                      DeviceAccessor<SolveKernelParams, mode::read> a_solveKernelParams,
-                     DeviceAccessor<float, mode::read_write>       a_virialAndEnergy,
-                     DeviceAccessor<float, mode::read_write>       a_fourierGrid)
+                     OptionalAccessor<float, mode::read_write, computeEnergyAndVirial> a_virialAndEnergy,
+                     DeviceAccessor<float, mode::read_write> a_fourierGrid)
 {
     cgh.require(a_splineModuli);
     cgh.require(a_solveKernelParams);
-    cgh.require(a_virialAndEnergy);
+    if constexpr (computeEnergyAndVirial)
+    {
+        cgh.require(a_virialAndEnergy);
+    }
     cgh.require(a_fourierGrid);
 
     /* Reduce 7 outputs per warp in the shared memory */
@@ -281,7 +284,7 @@ auto makeSolveKernel(cl::sycl::handler&                            cgh,
         }
 
         /* Optional energy/virial reduction */
-        if (computeEnergyAndVirial)
+        if constexpr (computeEnergyAndVirial)
         {
             /* A tricky shuffle reduction inspired by reduce_force_j_warp_shfl.
              * The idea is to reduce 7 energy/virial components into a single variable (aligned by

diff --git a/src/gromacs/gpu_utils/sycl_kernel_utils.h b/src/gromacs/gpu_utils/sycl_kernel_utils.h
index aba31db05351b8ca1d83937c98d64fd1761f0af1..1f44ca24d01ce39ccb2c829b3b67a4b4cabe9d2c 100644 (file)
--- a/src/gromacs/gpu_utils/sycl_kernel_utils.h
+++ b/src/gromacs/gpu_utils/sycl_kernel_utils.h
@@ -157,9 +157,9 @@ __device__ __host__ static inline bool isfinite(Real value)
     // https://github.com/illuhad/hipSYCL/issues/636
 #    ifdef SYCL_DEVICE_ONLY
 #        if defined(HIPSYCL_PLATFORM_CUDA) && defined(__HIPSYCL_ENABLE_CUDA_TARGET__)
-    return isfinite(value);
+    return ::isfinite(value);
 #        elif defined(HIPSYCL_PLATFORM_ROCM) && defined(__HIPSYCL_ENABLE_HIP_TARGET__)
-    return isfinite(value);
+    return ::isfinite(value);
 #        else
 #            error "Unsupported hipSYCL target"
 #        endif