From 77f47f0591f097c22410d0f62da10ed11a93cb69 Mon Sep 17 00:00:00 2001
From: =?utf8?q?Szil=C3=A1rd=20P=C3=A1ll?= <pall.szilard@gmail.com>
Date: Wed, 21 Apr 2021 15:55:30 +0000
Subject: [PATCH] Remove isnormal() check around the atomic increment

---
 src/gromacs/gpu_utils/sycl_kernel_utils.h | 26 ++++++++++-------------
 1 file changed, 11 insertions(+), 15 deletions(-)

diff --git a/src/gromacs/gpu_utils/sycl_kernel_utils.h b/src/gromacs/gpu_utils/sycl_kernel_utils.h
index 9896a51d3e..45831f715b 100644
--- a/src/gromacs/gpu_utils/sycl_kernel_utils.h
+++ b/src/gromacs/gpu_utils/sycl_kernel_utils.h
@@ -68,25 +68,21 @@ template<class IndexType>
 static inline void atomicFetchAdd(DeviceAccessor<float, mode_atomic> acc, const IndexType idx, const float val)
 {
 #if GMX_SYCL_DPCPP
-    if (cl::sycl::isnormal(val))
-    {
-        sycl_2020::atomic_ref<float, sycl_2020::memory_order::relaxed, sycl_2020::memory_scope::device, cl::sycl::access::address_space::global_space>
-                fout_atomic(acc[idx]);
-        fout_atomic.fetch_add(val);
-    }
+    sycl_2020::atomic_ref<float, sycl_2020::memory_order::relaxed, sycl_2020::memory_scope::device, cl::sycl::access::address_space::global_space>
+            fout_atomic(acc[idx]);
+    fout_atomic.fetch_add(val);
 #elif GMX_SYCL_HIPSYCL
-    if (std::isnormal(val)) // No sycl::isnormal in hipSYCL
-    {
 #    ifdef SYCL_DEVICE_ONLY
-        /* While there is support for float atomics on device, the host implementation uses
-         * Clang's __atomic_fetch_add intrinsic, that, at least in Clang 11, does not support
-         * floats. Luckily, we don't want to run on host. */
-        acc[idx].fetch_add(val);
+    /* While there is support for float atomics on device, the host implementation uses
+     * Clang's __atomic_fetch_add intrinsic, that, at least in Clang 11, does not support
+     * floats. Luckily, we don't want to run on host. */
+    acc[idx].fetch_add(val);
 #    else
-        GMX_UNUSED_VALUE(acc);
-        GMX_UNUSED_VALUE(idx);
+    GMX_ASSERT(false, "hipSYCL host codepath not supported");
+    GMX_UNUSED_VALUE(val);
+    GMX_UNUSED_VALUE(acc);
+    GMX_UNUSED_VALUE(idx);
 #    endif
-    }
 #endif
 }
 
-- 
2.22.0