Fix shift_left/right function on hipSYCL

Having separate __host__ and __device__ definitions confused the
compiler when called from a non-annotated function. As a result,
a __host__ version, containing `assert(false)`, was always called,
leading to UB.

Refs #3965

diff --git a/src/gromacs/gpu_utils/sycl_kernel_utils.h b/src/gromacs/gpu_utils/sycl_kernel_utils.h
index b261d9d42d66956531641269988aefd4bec0f654..6c80300adcbaf3c90af6beaf313c8881bbf3683d 100644 (file)
--- a/src/gromacs/gpu_utils/sycl_kernel_utils.h
+++ b/src/gromacs/gpu_utils/sycl_kernel_utils.h
@@ -87,7 +87,9 @@ static inline void atomicFetchAdd(DeviceAccessor<float, mode_atomic> acc, const
 namespace sycl_2020
 {
 #if GMX_SYCL_HIPSYCL
-__device__ static inline float shift_left(sycl_2020::sub_group, float var, sycl_2020::sub_group::linear_id_type delta)
+__device__ __host__ static inline float shift_left(sycl_2020::sub_group,
+                                                   float                                var,
+                                                   sycl_2020::sub_group::linear_id_type delta)
 {
     // No sycl::sub_group::shift_left / shuffle_down in hipSYCL yet
 #    ifdef SYCL_DEVICE_ONLY
@@ -108,12 +110,6 @@ __device__ static inline float shift_left(sycl_2020::sub_group, float var, sycl_
     return NAN;
 #    endif
 }
-__host__ static inline float shift_left(sycl_2020::sub_group, float, sycl_2020::sub_group::linear_id_type)
-{
-    // Should never be called
-    assert(false);
-    return NAN;
-}
 #elif GMX_SYCL_DPCPP
 static inline float shift_left(sycl_2020::sub_group sg, float var, sycl_2020::sub_group::linear_id_type delta)
 {
@@ -122,9 +118,9 @@ static inline float shift_left(sycl_2020::sub_group sg, float var, sycl_2020::su
 #endif
 
 #if GMX_SYCL_HIPSYCL
-__device__ static inline float shift_right(sycl_2020::sub_group,
-                                           float                                var,
-                                           sycl_2020::sub_group::linear_id_type delta)
+__device__ __host__ static inline float shift_right(sycl_2020::sub_group,
+                                                    float                                var,
+                                                    sycl_2020::sub_group::linear_id_type delta)
 {
     // No sycl::sub_group::shift_right / shuffle_up in hipSYCL yet
 #    ifdef SYCL_DEVICE_ONLY
@@ -145,12 +141,6 @@ __device__ static inline float shift_right(sycl_2020::sub_group,
     return NAN;
 #    endif
 }
-__host__ static inline float shift_right(sycl_2020::sub_group, float, sycl_2020::sub_group::linear_id_type)
-{
-    // Should never be called
-    assert(false);
-    return NAN;
-}
 #elif GMX_SYCL_DPCPP
 static inline float shift_right(sycl_2020::sub_group sg, float var, sycl_2020::sub_group::linear_id_type delta)
 {