From fe0b3c8bd13326c1c5fd607a527f4dfb9bb1f5ff Mon Sep 17 00:00:00 2001 From: Andrey Alekseenko Date: Sat, 8 May 2021 14:23:23 +0300 Subject: [PATCH] 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 --- src/gromacs/gpu_utils/sycl_kernel_utils.h | 22 ++++++---------------- 1 file changed, 6 insertions(+), 16 deletions(-) diff --git a/src/gromacs/gpu_utils/sycl_kernel_utils.h b/src/gromacs/gpu_utils/sycl_kernel_utils.h index b261d9d42d..6c80300adc 100644 --- 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 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) { -- 2.22.0