From c00cd118d3aff893845715c577d6fc31a598c867 Mon Sep 17 00:00:00 2001 From: Sourabh Betigeri Date: Thu, 1 Jun 2023 22:55:49 -0700 Subject: [PATCH] SWDEV-402274 - Adds CG API group_dim() in class thread_block Change-Id: I1c2569fffbc6fc730df9ba80786459c71140d327 [ROCm/clr commit: 422d3ea03559d6ec9e8d5132efd12a3a48c13281] --- .../include/hip/amd_detail/amd_hip_cooperative_groups.h | 1 + .../include/hip/amd_detail/hip_cooperative_groups_helper.h | 6 +++++- 2 files changed, 6 insertions(+), 1 deletion(-) 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(blockDim.x), static_cast(blockDim.y), + static_cast(blockDim.z))); +} + } // namespace workgroup namespace tiled_group {