From 6b85dcf227a7e494c331d02d5307c1d2c51bda40 Mon Sep 17 00:00:00 2001
From: Satyanvesh Dittakavi <53337087+satyanveshd@users.noreply.github.com>
Date: Fri, 10 Oct 2025 00:21:44 +0530
Subject: [PATCH] SWDEV-557093 - Add nested tiled partition in HIP cooperative
 groups (#1166)

---
 .../amd_detail/amd_hip_cooperative_groups.h   | 19 +++++++++++++++++++
 1 file changed, 19 insertions(+)

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 d499a93b47..322cab668b 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
@@ -921,6 +921,15 @@ template <unsigned int tileSize, class ParentCGTy> class thread_block_tile_type
     coalesced_info.tiled_info.is_tiled = true;
   }
 
+  __CG_QUALIFIER__ thread_block_tile_type(unsigned int meta_group_rank,
+                                          unsigned int meta_group_size)
+      : tiled_group(numThreads) {
+    coalesced_info.tiled_info.num_threads = numThreads;
+    coalesced_info.tiled_info.is_tiled = true;
+    coalesced_info.tiled_info.meta_group_rank = meta_group_rank;
+    coalesced_info.tiled_info.meta_group_size = meta_group_size;
+  }
+
  public:
   using tbtBase::num_threads;
   using tbtBase::size;
@@ -1189,6 +1198,16 @@ template <unsigned int size> struct tiled_partition_internal<size, thread_block>
       : thread_block_tile<size, thread_block>(g) {}
 };
 
+// ParentCGTy = thread_block_tile<ParentSize, GrandParentCGTy> specialization
+template <unsigned int size, unsigned int ParentSize, class GrandParentCGTy>
+struct tiled_partition_internal<size, thread_block_tile<ParentSize, GrandParentCGTy> >
+    : public thread_block_tile<size, thread_block_tile<ParentSize, GrandParentCGTy> > {
+  static_assert(size <= ParentSize, "Sub tile size must be <= parent tile size in tiled_partition");
+
+  __CG_QUALIFIER__ tiled_partition_internal(const thread_block_tile<ParentSize, GrandParentCGTy>& g)
+      : thread_block_tile<size, thread_block_tile<ParentSize, GrandParentCGTy> >(g) {}
+};
+
 }  // namespace impl
 
 /** \ingroup CooperativeGConstruct