SWDEV-402274 - Adds CG API group_dim() in class thread_block

Change-Id: I1c2569fffbc6fc730df9ba80786459c71140d327


[ROCm/clr commit: 422d3ea035]

projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h

@@ -252,6 +252,7 @@ class thread_block : public thread_group {
   __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::workgroup::thread_rank(); }
   __CG_QUALIFIER__ bool is_valid() const { return internal::workgroup::is_valid(); }
   __CG_QUALIFIER__ void sync() const { internal::workgroup::sync(); }
+  __CG_QUALIFIER__ dim3 group_dim() { return internal::workgroup::block_dim(); }
 };
 
 /** \brief   User exposed API interface to construct workgroup cooperative

projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h

@@ -195,12 +195,16 @@ __CG_STATIC_QUALIFIER__ uint32_t thread_rank() {
 }
 
 __CG_STATIC_QUALIFIER__ bool is_valid() {
-  // TODO(mahesha) any functionality need to be added here? I believe not
   return true;
 }
 
 __CG_STATIC_QUALIFIER__ void sync() { __syncthreads(); }
 
+__CG_STATIC_QUALIFIER__ dim3 block_dim() {
+  return (dim3(static_cast<uint32_t>(blockDim.x), static_cast<uint32_t>(blockDim.y),
+          static_cast<uint32_t>(blockDim.z)));
+}
+
 }  // namespace workgroup
 
 namespace tiled_group {