From f3663d440ca7b4d85d60dc5931d30a4b085c2a2c Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Fri, 2 Oct 2020 22:45:37 -0400 Subject: [PATCH] Fix warpSize for gfx10 Change-Id: I8029ebeb91db1efa8e166ad349aaae6364e8069d [ROCm/clr commit: 71586b02e3a2272eaf4acdb8b72bb83e17967c9e] --- .../hipamd/include/hip/hcc_detail/device_functions.h | 12 +++++++++++- 1 file changed, 11 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index a5fe9425f4..822c626fa2 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -306,7 +306,17 @@ __device__ static inline int __hip_move_dpp_N(int src) { bound_ctrl); } -static constexpr int warpSize = 64; +// FIXME: Remove the following workaround once the clang change is released. +// This is for backward compatibility with older clang which does not define +// __AMDGCN_WAVEFRONT_SIZE. It does not consider -mwavefrontsize64. +#ifndef __AMDGCN_WAVEFRONT_SIZE +#if __gfx1010__ || __gfx1011__ || __gfx1012__ || __gfx1030__ || __gfx1031__ +#define __AMDGCN_WAVEFRONT_SIZE 32 +#else +#define __AMDGCN_WAVEFRONT_SIZE 64 +#endif +#endif +static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE; __device__ inline