From 152f34312494fd64ffd4412964ebf7dd548212cf Mon Sep 17 00:00:00 2001 From: Marko Arandjelovic Date: Mon, 17 Jun 2024 14:01:16 +0200 Subject: [PATCH] SWDEV-472345 - Fix coalesced group size In case when the tile size is greater than the number of active threads, the coalesced group size should be equal to the number of active threads. Change-Id: I1d41322f2428a07862a590cb5d34b01243383b7c --- hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h index 6c038d33ec..418f96949d 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h +++ b/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h @@ -104,7 +104,7 @@ class thread_group { * @ingroup API * @{ * This section describes the cooperative groups functions of HIP runtime API. - * + * * The cooperative groups provides flexible thread parallel programming algorithms, threads * cooperate and share data to perform collective computations. * @@ -218,7 +218,7 @@ class thread_block : public thread_group { if (!tile_size || (tile_size > __AMDGCN_WAVEFRONT_SIZE) || !pow2) { __hip_assert(false && "invalid tile size"); } - + auto block_size = size(); auto rank = thread_rank(); auto partitions = (block_size + tile_size - 1) / tile_size; @@ -330,7 +330,7 @@ class coalesced_group : public thread_group { __CG_QUALIFIER__ coalesced_group new_tiled_group(unsigned int tile_size) const { const bool pow2 = ((tile_size & (tile_size - 1)) == 0); - if (!tile_size || (tile_size > size()) || !pow2) { + if (!tile_size || !pow2) { return coalesced_group(0); }