From 90578d5e07a70dfddd347c6482df6a6bef8c7036 Mon Sep 17 00:00:00 2001 From: Rupam Chetia Date: Wed, 20 Sep 2023 17:08:55 +0530 Subject: [PATCH] SWDEV-421132 - [catch2][dtest] Adding basic tests for Cooperative Groups methods meta_group_size() and meta_group_rank() Change-Id: I7b9df6a3d656e0303e357e2610da9303e93f5600 [ROCm/hip-tests commit: ae8090bde868f4798794131253dfae6c72886c62] --- .../catch/unit/cooperativeGrps/CMakeLists.txt | 1 + .../coalesced_tiled_groups_metagrp.cc | 139 ++++++++++++++++++ 2 files changed, 140 insertions(+) create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/coalesced_tiled_groups_metagrp.cc diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt index e546df0398..829758b622 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt @@ -11,6 +11,7 @@ set(TEST_SRC coalesced_groups_shfl_up.cc hipCGTiledPartition.cc hipCGCoalescedGroups.cc + coalesced_tiled_groups_metagrp.cc ) if(HIP_PLATFORM STREQUAL "nvidia") set_source_files_properties(hipCGMultiGridGroupType.cc PROPERTIES COMPILE_FLAGS "-rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/coalesced_tiled_groups_metagrp.cc b/projects/hip-tests/catch/unit/cooperativeGrps/coalesced_tiled_groups_metagrp.cc new file mode 100644 index 0000000000..a7f9ddc7e7 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/coalesced_tiled_groups_metagrp.cc @@ -0,0 +1,139 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include + +/** + * @addtogroup coalesced_group thread_block_tile + * @{ + * @ingroup CooperativeGroupTest + * meta_group_size() + * meta_group_rank() + */ + +using namespace cooperative_groups; + +constexpr auto total_elem = 1 << 16; +constexpr auto block_size = 256; +constexpr auto test_size = 32; + +static __global__ void kernel_coalesced_grp(int *mgrpSize, int *mgrpRank) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + if (id % 2 == 0) { + coalesced_group threadBlockCGTy = coalesced_threads(); + mgrpSize[id] = threadBlockCGTy.meta_group_size(); + mgrpRank[id] = threadBlockCGTy.meta_group_rank(); + } else { + coalesced_group threadBlockCGTx = coalesced_threads(); + mgrpSize[id] = threadBlockCGTx.meta_group_size(); + mgrpRank[id] = threadBlockCGTx.meta_group_rank(); + } +} + +static __global__ void kernel_tiledgrp_threadblk(int *mgrpSize, + int *mgrpRank) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + thread_block_tile tiledGr = + tiled_partition(this_thread_block()); + mgrpSize[id] = tiledGr.meta_group_size(); + mgrpRank[id] = tiledGr.meta_group_rank(); +} + +/** + * Test Description + * ------------------------ + * - Verify the values returned by meta_group_size() and + * meta_group_rank() for thread block tile with thread block + * group as parent. + * ------------------------ + * - catch\unit\cooperativeGrps\coalesced_tiled_groups_metagrp.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_tiled_groups_metagrp_basic") { + int *mgrpSize_d = nullptr, *mgrpRank_d = nullptr; + int *mgrpSize_h = nullptr, *mgrpRank_h = nullptr; + mgrpSize_h = new int[total_elem]; + REQUIRE(mgrpSize_h != nullptr); + mgrpRank_h = new int[total_elem]; + REQUIRE(mgrpRank_h != nullptr); + + HIP_CHECK(hipMalloc(&mgrpSize_d, total_elem*sizeof(int))); + HIP_CHECK(hipMalloc(&mgrpRank_d, total_elem*sizeof(int))); + SECTION("Parent Group = thread block group") { + hipLaunchKernelGGL(kernel_tiledgrp_threadblk, total_elem/block_size, + block_size, 0, 0, mgrpSize_d, mgrpRank_d); + HIP_CHECK(hipMemcpy(mgrpRank_h, mgrpRank_d, total_elem*sizeof(int), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(mgrpSize_h, mgrpSize_d, total_elem*sizeof(int), + hipMemcpyDeviceToHost)); + for (int i = 0; i < total_elem; i++) { + REQUIRE(mgrpRank_h[i] >= 0); + REQUIRE(mgrpRank_h[i] < (block_size/test_size)); + REQUIRE(mgrpSize_h[i] == (block_size/test_size)); + } + } + HIP_CHECK(hipFree(mgrpSize_d)); + HIP_CHECK(hipFree(mgrpRank_d)); + delete[] mgrpSize_h; + delete[] mgrpRank_h; +} + +/** + * Test Description + * ------------------------ + * - Verify the values returned by meta_group_size() and + * meta_group_rank() for Coalesced Groups. + * ------------------------ + * - catch\unit\cooperativeGrps\coalesced_tiled_groups_metagrp.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_coalesced_groups_metagrp_basic") { + int *mgrpSize_d = nullptr, *mgrpRank_d = nullptr; + int *mgrpSize_h = nullptr, *mgrpRank_h = nullptr; + mgrpSize_h = new int[total_elem]; + REQUIRE(mgrpSize_h != nullptr); + mgrpRank_h = new int[total_elem]; + REQUIRE(mgrpRank_h != nullptr); + + HIP_CHECK(hipMalloc(&mgrpSize_d, total_elem*sizeof(int))); + HIP_CHECK(hipMalloc(&mgrpRank_d, total_elem*sizeof(int))); + + hipLaunchKernelGGL(kernel_coalesced_grp, total_elem/block_size, + block_size, 0, 0, mgrpSize_d, mgrpRank_d); + HIP_CHECK(hipMemcpy(mgrpRank_h, mgrpRank_d, total_elem*sizeof(int), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(mgrpSize_h, mgrpSize_d, total_elem*sizeof(int), + hipMemcpyDeviceToHost)); + for (int i = 0; i < total_elem; i++) { + REQUIRE(mgrpRank_h[i] == 0); + REQUIRE(mgrpSize_h[i] == 1); + } + HIP_CHECK(hipFree(mgrpSize_d)); + HIP_CHECK(hipFree(mgrpRank_d)); + delete[] mgrpSize_h; + delete[] mgrpRank_h; +}