From 819e537dc5ac984c8d4ea22160c19e915556a316 Mon Sep 17 00:00:00 2001 From: Fabian Ritter Date: Fri, 10 May 2024 07:32:03 -0400 Subject: [PATCH] Add __device__ specifier to HIP's warpSize constexpr variable The warpSize variable is set to the value of the __AMDGCN_WAVEFRONT_SIZE macro, which is a meaningless default in host code. The resolution for SWDEV-449015 will introduce diagnostics for uses of this macro in host code, which includes the current definition of the warpSize variable. With the __device__ specifier, the definition of the warpSize variable will not cause these diagnostics. This change does not stop the variable from being used in host code since clang intentionally does not diagnose uses of __device__ constexpr variables in host code. Change-Id: I0317217affe94fdf2dfd9ad0f134e68f5173245f --- hipamd/include/hip/amd_detail/amd_warp_functions.h | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/include/hip/amd_detail/amd_warp_functions.h b/hipamd/include/hip/amd_detail/amd_warp_functions.h index 98f8896cd9..46e3fbf8ee 100644 --- a/hipamd/include/hip/amd_detail/amd_warp_functions.h +++ b/hipamd/include/hip/amd_detail/amd_warp_functions.h @@ -73,6 +73,7 @@ __device__ static inline int __hip_move_dpp_N(int src) { bound_ctrl); } +__device__ static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE; // warp vote function __all __any __ballot