From 2fdbade459a61eee49f1716c2d090f04e4195917 Mon Sep 17 00:00:00 2001 From: John Afaganis Date: Tue, 10 Jun 2025 08:34:54 -0600 Subject: [PATCH] Remove usage of 'warpSize' variable as it has been deprecated (#2295) * SWDEV-535598 - remove usage of 'warpSize' variable as it has been deprecated. Ideally get_warp_size() should not be constexpr but this is just a workaround * SWDEV-535598 - remove comment from get_warp_size as constexpr is required for this repo --------- Co-authored-by: Gerardo Hernandez [ROCm/composable_kernel commit: 6635d1bb888e3f51ec1125ff7d6f54a2ec054a10] --- include/ck/utility/get_id.hpp | 7 +++++-- include/ck_tile/core/arch/arch.hpp | 7 +++++-- 2 files changed, 10 insertions(+), 4 deletions(-) diff --git a/include/ck/utility/get_id.hpp b/include/ck/utility/get_id.hpp index 77564c6130..fd0d1024b2 100644 --- a/include/ck/utility/get_id.hpp +++ b/include/ck/utility/get_id.hpp @@ -9,8 +9,11 @@ namespace ck { __host__ __device__ constexpr index_t get_warp_size() { - // warpSize is defined by HIP - return warpSize; +#if defined(__GFX9__) || !defined(__HIP_DEVICE_COMPILE__) + return 64; +#else + return 32; +#endif } __device__ index_t get_thread_local_1d_id() { return threadIdx.x; } diff --git a/include/ck_tile/core/arch/arch.hpp b/include/ck_tile/core/arch/arch.hpp index 1d3cf5c010..3dd9604b01 100644 --- a/include/ck_tile/core/arch/arch.hpp +++ b/include/ck_tile/core/arch/arch.hpp @@ -50,8 +50,11 @@ enum struct memory_operation_enum : std::uint16_t CK_TILE_HOST_DEVICE constexpr index_t get_warp_size() { - // warpSize is defined by HIP - return warpSize; +#if defined(__GFX9__) || !defined(__HIP_DEVICE_COMPILE__) + return 64; +#else + return 32; +#endif } CK_TILE_DEVICE index_t get_grid_size() { return gridDim.x; }