From 6159b0eba0cd0d13141c1967097611f6f67a4e73 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 [ROCm/clr commit: 152f34312494fd64ffd4412964ebf7dd548212cf] --- .../include/hip/amd_detail/amd_hip_cooperative_groups.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h index 6c038d33ec..418f96949d 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h +++ b/projects/clr/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); }