From cbd1e7a8d76fac6ce3b2c7f824fe27c37bb54b27 Mon Sep 17 00:00:00 2001
From: Andrey Alekseenko <al42and@gmail.com>
Date: Mon, 18 Oct 2021 18:06:49 +0200
Subject: [PATCH] Fix SYCL PME Solve kernel

- Avoid infinite recursion in sycl_2020::isfinite.
- Make one of the accessors optional.
---
 src/gromacs/ewald/pme_solve_sycl.cpp      | 11 +++++++----
 src/gromacs/gpu_utils/sycl_kernel_utils.h |  4 ++--
 2 files changed, 9 insertions(+), 6 deletions(-)

diff --git a/src/gromacs/ewald/pme_solve_sycl.cpp b/src/gromacs/ewald/pme_solve_sycl.cpp
index 633cf31e8e..46883060e9 100644
--- 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 aba31db053..1f44ca24d0 100644
--- 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
-- 
2.22.0