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 5baf21db53..ca9dae4437 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
@@ -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
diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
index 041f4fd1b8..391c6b9201 100644
--- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
+++ b/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 {