From 6f750d2361435985d03f5346fe80cc0adbee6f3c Mon Sep 17 00:00:00 2001 From: Jaydeep Patel Date: Thu, 22 Feb 2024 14:36:50 +0000 Subject: [PATCH] SWDEV-444257 - Less # of CUs i.e. 2 tends to .. - make blockdim to become 0 and hence need validation. - run cooperative kernel with param lesser # of threads than actual # of elements and hence need correction in validation part. Change-Id: Ie48a038fe6f8cf25d7f6460703648a360f511f35 [ROCm/hip-tests commit: ad602b0011295d58be08934e71412f951591dea6] --- .../catch/unit/cooperativeGrps/cooperative_groups_common.hh | 4 +++- .../catch/unit/streamperthread/hipStreamPerThrdTsts.cc | 1 + 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh b/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh index 19ad0dd092..d12128187f 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh +++ b/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh @@ -70,7 +70,9 @@ template bool CheckDimensions(unsigned int device, T kernel, dim3 bloc HIP_CHECK(hipGetDeviceProperties(&props, device)); num_sm = props.multiProcessorCount; - if ((blocks.x * blocks.y * blocks.z) > max_blocks_per_sm * num_sm) { + if ((blocks.x * blocks.y * blocks.z) > max_blocks_per_sm * num_sm || + blocks.x <= 0 || blocks.y <= 0 || blocks.z <= 0 || + threads.x <= 0 || threads.y <= 0 || threads.z <= 0) { return false; } diff --git a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc index c14b38ce19..0b432ee1b5 100644 --- a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc +++ b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc @@ -505,6 +505,7 @@ TEST_CASE("Unit_hipStreamPerThread_CoopLaunch") { warp_size, 0)); int max_active_blocks = max_blocks_per_sm * num_sms; int *Ad = nullptr, *Ah = nullptr, *DNumElms = nullptr, NumElms = 4096; + NumElms = max_active_blocks * warp_size > NumElms ? NumElms : max_active_blocks * warp_size; int Const = 123; Ah = new int[NumElms]; for (int i = 0; i < NumElms; ++i) {