From ee52a0500914945ae0c01b9ad83a0cce5ff80a27 Mon Sep 17 00:00:00 2001 From: Rahul Manocha Date: Fri, 1 Mar 2024 01:40:44 +0000 Subject: [PATCH] SWDEV-448281 - Fix for CG Tile Partition Dynamic Size Change-Id: Ibdccec45470caf55a8c00194e6da0dde73e4b180 [ROCm/clr commit: e5e7687fd86c5b19d6ab51b7664b673f0116c661] --- .../hip/amd_detail/amd_hip_cooperative_groups.h | 12 +++++++++--- 1 file changed, 9 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 8b1a0c067d..c01039a7e1 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 @@ -216,12 +216,18 @@ 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; + auto tail = (partitions * tile_size) - block_size; + auto partition_size = tile_size - tail * (rank >= (partitions - 1) * tile_size); + thread_group tiledGroup = thread_group(internal::cg_tiled_group, partition_size); - thread_group tiledGroup = thread_group(internal::cg_tiled_group, tile_size); tiledGroup.coalesced_info.tiled_info.size = tile_size; tiledGroup.coalesced_info.tiled_info.is_tiled = true; - tiledGroup.coalesced_info.tiled_info.meta_group_rank = thread_rank() / tile_size; - tiledGroup.coalesced_info.tiled_info.meta_group_size = (size() + tile_size - 1) / tile_size; + tiledGroup.coalesced_info.tiled_info.meta_group_rank = rank / tile_size; + tiledGroup.coalesced_info.tiled_info.meta_group_size = partitions; return tiledGroup; }