From d537bb385998ce2308be63eb6ae992e83785bddf Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 19 Aug 2020 10:47:32 -0400 Subject: [PATCH 1/4] Fix hipcc for extracing clang version On teamcity clang --version returns something like clang version 12.0.0 Target: x86_64-unknown-linux-gnu currently hipcc does not handle it correctly since [^ ] causes new line and Targe: matched. \S will make sure only version string matched. Change-Id: I3659d6daf89b709d2b001e6b8024d69d8b79efc0 --- bin/hipcc | 2 +- bin/hipconfig | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index 22fd249bfd..0e3027e03a 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -201,7 +201,7 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") { } $HIP_CLANG_VERSION = `$HIPCC --version`; - $HIP_CLANG_VERSION=~/.*clang version ([^ ]+).*/; + $HIP_CLANG_VERSION=~/.*clang version (\S+).*/; $HIP_CLANG_VERSION=$1; if (! defined $HIP_CLANG_INCLUDE_PATH) { diff --git a/bin/hipconfig b/bin/hipconfig index 8659669f5f..b0c01796cc 100755 --- a/bin/hipconfig +++ b/bin/hipconfig @@ -139,7 +139,7 @@ if ($HIP_COMPILER eq "clang") { $HIP_CLANG_PATH = "$HIP_ROCCLR_HOME/bin"; } $HIP_CLANG_VERSION = `$HIP_CLANG_PATH/clang++ --version`; - $HIP_CLANG_VERSION=~/.*clang version ([^ ]+).*/; + $HIP_CLANG_VERSION=~/.*clang version (\S+).*/; $HIP_CLANG_VERSION=$1; $CPP_CONFIG = " -D__HIP_PLATFORM_HCC__= -I$HIP_PATH/include -I$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION -I$HSA_PATH/include"; From 9c75797d9f667ef0463808db9480ebfb8296972e Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Mon, 24 Aug 2020 17:25:21 -0400 Subject: [PATCH 2/4] SWDEV-249499 - Add message to Guarantee assert code in hip runtime. Change-Id: I1850463675bbc3e6ff5e3cef8211b1fd40e55f96 --- rocclr/hip_code_object.cpp | 2 +- rocclr/hip_global.cpp | 41 +++++++++++++++++++++----------------- rocclr/hip_platform.cpp | 2 +- 3 files changed, 25 insertions(+), 20 deletions(-) diff --git a/rocclr/hip_code_object.cpp b/rocclr/hip_code_object.cpp index dff006d335..8b89647e89 100755 --- a/rocclr/hip_code_object.cpp +++ b/rocclr/hip_code_object.cpp @@ -124,7 +124,7 @@ hipError_t DynCO::loadCodeObject(const char* fname, const void* image) { const void *mmap_ptr = nullptr; size_t mmap_size = 0; - guarantee(fname || image); + guarantee((fname || image) && "Both filename or image are nullptr"); if (fname != nullptr) { /* We are given file name */ diff --git a/rocclr/hip_global.cpp b/rocclr/hip_global.cpp index 680d16b03f..1211d99556 100755 --- a/rocclr/hip_global.cpp +++ b/rocclr/hip_global.cpp @@ -14,20 +14,20 @@ DeviceVar::DeviceVar(std::string name, hipModule_t hmod) : shadowVptr(nullptr), amd::Program* program = as_amd(reinterpret_cast(hmod)); device::Program* dev_program = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]); if (dev_program == nullptr) { - DevLogPrintfError("Cannot get Device Function for module: 0x%x \n", hmod); - guarantee(false); + DevLogPrintfError("Cannot get Device Program for module: 0x%x \n", hmod); + guarantee(false && "Cannot get Device Program"); } if(!dev_program->createGlobalVarObj(&amd_mem_obj_, &device_ptr_, &size_, name.c_str())) { DevLogPrintfError("Cannot create Global Var obj for symbol: %s \n", name.c_str()); - guarantee(false); + guarantee(false && "Cannot create GlobalVar Obj"); } // Handle size 0 symbols if (size_ != 0) { if (amd_mem_obj_ == nullptr || device_ptr_ == nullptr) { DevLogPrintfError("Cannot get memory for creating device Var: %s", name.c_str()); - guarantee(false); + guarantee(false && "Cannot get memory for creating device var"); } amd::MemObjMap::AddMemObj(device_ptr_, amd_mem_obj_); } @@ -57,13 +57,13 @@ DeviceFunc::DeviceFunc(std::string name, hipModule_t hmod) : dflock_("function l const amd::Symbol *symbol = program->findSymbol(name.c_str()); if (symbol == nullptr) { DevLogPrintfError("Cannot find Symbol with name: %s \n", name.c_str()); - guarantee(false); + guarantee(false && "Cannot find Symbol"); } kernel_ = new amd::Kernel(*program, *symbol, name); if (kernel_ == nullptr) { DevLogPrintfError("Cannot create kernel with name: %s \n", name.c_str()); - guarantee(false); + guarantee(false && "Cannot Create kernel"); } } @@ -88,7 +88,7 @@ Function::~Function() { } hipError_t Function::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod) { - guarantee(dFunc_.size() == g_devices.size()); + guarantee((dFunc_.size() == g_devices.size()) && "dFunc Size mismatch"); if (dFunc_[ihipGetDevice()] == nullptr) { dFunc_[ihipGetDevice()] = new DeviceFunc(name_, hmod); } @@ -98,9 +98,10 @@ hipError_t Function::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod) { } hipError_t Function::getStatFunc(hipFunction_t* hfunc, int deviceId) { - guarantee(modules_ != nullptr); - guarantee(deviceId >= 0); - guarantee(static_cast(deviceId) < modules_->size()); + guarantee(modules_ != nullptr && "Module not initialized"); + guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero"); + guarantee((static_cast(deviceId) < modules_->size()) + && "Invalid DeviceId, greater than no of code objects"); hipModule_t module = (*modules_)[deviceId].first; FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second; @@ -121,9 +122,10 @@ hipError_t Function::getStatFunc(hipFunction_t* hfunc, int deviceId) { } hipError_t Function::getStatFuncAttr(hipFuncAttributes* func_attr, int deviceId) { - guarantee(modules_ != nullptr); - guarantee(deviceId >= 0); - guarantee(static_cast(deviceId) < modules_->size()); + guarantee((modules_ != nullptr) && "Module not initialized"); + guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero"); + guarantee((static_cast(deviceId) < modules_->size()) + && "Invalid DeviceId, greater than no of code objects"); hipModule_t module = (*modules_)[deviceId].first; FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second; @@ -175,9 +177,11 @@ Var::~Var() { } hipError_t Var::getDeviceVar(DeviceVar** dvar, int deviceId, hipModule_t hmod) { - guarantee(deviceId >= 0); - guarantee(static_cast(deviceId) < g_devices.size()); - guarantee(dVar_.size() == g_devices.size()); + guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero"); + guarantee((static_cast(deviceId) < g_devices.size()) + && "Invalid DeviceId, greater than no of code objects"); + guarantee((dVar_.size() == g_devices.size()) + && "Device Var not initialized to size"); if (dVar_[deviceId] == nullptr) { dVar_[deviceId] = new DeviceVar(name_, hmod); @@ -188,8 +192,9 @@ hipError_t Var::getDeviceVar(DeviceVar** dvar, int deviceId, hipModule_t hmod) { } hipError_t Var::getStatDeviceVar(DeviceVar** dvar, int deviceId) { - guarantee(deviceId >= 0); - guarantee(static_cast(deviceId) < g_devices.size()); + guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero"); + guarantee((static_cast(deviceId) < g_devices.size()) + && "Invalid DeviceId, greater than no of code objects"); hipModule_t module = (*modules_)[deviceId].first; FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second; diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 3935e4d5f8..2d74d8f440 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -127,7 +127,7 @@ extern "C" void __hipRegisterFunction( hipError_t hip_error = hipSuccess; for (size_t dev_idx = 0; dev_idx < g_devices.size(); ++dev_idx) { hip_error = PlatformState::instance().getStatFunc(&hfunc, hostFunction, dev_idx); - guarantee(hip_error == hipSuccess); + guarantee((hip_error == hipSuccess) && "Cannot Retrieve Static function"); } } } From 2380cf7805e16b06ed8853fe0ebfa805cdce1a77 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 24 Aug 2020 21:45:34 +0000 Subject: [PATCH 3/4] Bump version to 3.9 Change-Id: I5dfc997336fac42f3b6b0998a4bb212462de38f9 --- bin/hipconfig | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipconfig b/bin/hipconfig index b0c01796cc..d7e3cc5b5a 100755 --- a/bin/hipconfig +++ b/bin/hipconfig @@ -1,7 +1,7 @@ #!/usr/bin/perl -w $HIP_BASE_VERSION_MAJOR = "3"; -$HIP_BASE_VERSION_MINOR = "8"; +$HIP_BASE_VERSION_MINOR = "9"; # Need perl > 5.10 to use logic-defined or use 5.006; use v5.10.1; From 15aff89ba1aef1976d5f3b6c385e3695be0e74bd Mon Sep 17 00:00:00 2001 From: mshivama Date: Sun, 9 Aug 2020 13:30:08 +0530 Subject: [PATCH 4/4] Support public apis for cooperative group types. Change-Id: I01346afde485e82c34b7868b9241b34c542d0cf9 --- .../hip/hcc_detail/hip_cooperative_groups.h | 31 +- tests/src/cg/hipCGGridGroupType.cpp | 157 +++++++++++ .../src/cg/hipCGGridGroupTypeViaBaseType.cpp | 157 +++++++++++ .../src/cg/hipCGGridGroupTypeViaPublicApi.cpp | 157 +++++++++++ tests/src/cg/hipCGMultiGridGroupType.cpp | 264 ++++++++++++++++++ .../cg/hipCGMultiGridGroupTypeViaBaseType.cpp | 238 ++++++++++++++++ .../hipCGMultiGridGroupTypeViaPublicApi.cpp | 238 ++++++++++++++++ tests/src/cg/hipCGThreadBlockType.cpp | 216 +++++++------- .../cg/hipCGThreadBlockTypeViaBaseType.cpp | 155 ++++++++++ .../cg/hipCGThreadBlockTypeViaPublicApi.cpp | 155 ++++++++++ 10 files changed, 1670 insertions(+), 98 deletions(-) create mode 100644 tests/src/cg/hipCGGridGroupType.cpp create mode 100644 tests/src/cg/hipCGGridGroupTypeViaBaseType.cpp create mode 100644 tests/src/cg/hipCGGridGroupTypeViaPublicApi.cpp create mode 100644 tests/src/cg/hipCGMultiGridGroupType.cpp create mode 100644 tests/src/cg/hipCGMultiGridGroupTypeViaBaseType.cpp create mode 100644 tests/src/cg/hipCGMultiGridGroupTypeViaPublicApi.cpp create mode 100644 tests/src/cg/hipCGThreadBlockTypeViaBaseType.cpp create mode 100644 tests/src/cg/hipCGThreadBlockTypeViaPublicApi.cpp diff --git a/include/hip/hcc_detail/hip_cooperative_groups.h b/include/hip/hcc_detail/hip_cooperative_groups.h index 73f6778a1c..353bdc5514 100644 --- a/include/hip/hcc_detail/hip_cooperative_groups.h +++ b/include/hip/hcc_detail/hip_cooperative_groups.h @@ -230,7 +230,8 @@ __CG_QUALIFIER__ uint32_t thread_group::thread_rank() const { return (static_cast(this)->thread_rank()); } default: { - return 0; //TODO(mahesha) + assert(false && "invalid cooperative group type"); + return -1; } } } @@ -247,6 +248,7 @@ __CG_QUALIFIER__ bool thread_group::is_valid() const { return (static_cast(this)->is_valid()); } default: { + assert(false && "invalid cooperative group type"); return false; } } @@ -266,9 +268,36 @@ __CG_QUALIFIER__ void thread_group::sync() const { static_cast(this)->sync(); break; } + default: { + assert(false && "invalid cooperative group type"); + } } } +/** + * Implemenation of publicly exposed `wrapper` APIs on top of basic cooperative + * group type APIs + */ +template +__CG_QUALIFIER__ uint32_t group_size(CGTy const &g) { + return g.size(); +} + +template +__CG_QUALIFIER__ uint32_t thread_rank(CGTy const &g) { + return g.thread_rank(); +} + +template +__CG_QUALIFIER__ bool is_valid(CGTy const &g) { + return g.is_valid(); +} + +template +__CG_QUALIFIER__ void sync(CGTy const &g) { + g.sync(); +} + } // namespace cooperative_groups #endif // __cplusplus diff --git a/tests/src/cg/hipCGGridGroupType.cpp b/tests/src/cg/hipCGGridGroupType.cpp new file mode 100644 index 0000000000..99fe5d88a2 --- /dev/null +++ b/tests/src/cg/hipCGGridGroupType.cpp @@ -0,0 +1,157 @@ +/* +Copyright (c) 2020 - present 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#include "hip/hip_cooperative_groups.h" +#include +#include + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +using namespace cooperative_groups; + +static __global__ +void kernel_cg_grid_group_type(int *sizeTestD, + int *thdRankTestD, + int *isValidTestD, + int *syncTestD) +{ + grid_group gg = this_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test size + sizeTestD[gIdx] = gg.size(); + + // Test thread_rank + thdRankTestD[gIdx] = gg.thread_rank(); + + // Test is_valid + isValidTestD[gIdx] = gg.is_valid(); + + // Test sync + __device__ int gm[2]; + if (blockIdx.x == 0 && threadIdx.x == 0) + gm[0] = 10; + else if (blockIdx.x == 1 && threadIdx.x == 0) + gm[1] = 20; + gg.sync(); + syncTestD[gIdx] = gm[1] * gm[0]; +} + +static void test_cg_grid_group_type(int blockSize) +{ + int nBytes = sizeof(int) * 2 * blockSize; + int *sizeTestD, *sizeTestH; + int *thdRankTestD, *thdRankTestH; + int *isValidTestD, *isValidTestH; + int *syncTestD, *syncTestH; + + // Allocate device memory + ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); + + // Allocate host memory + ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); + + // Launch Kernel + void *params[4]; + params[0] = &sizeTestD; + params[1] = &thdRankTestD; + params[2] = &isValidTestD; + params[3] = &syncTestD; + hipLaunchCooperativeKernel(kernel_cg_grid_group_type, + 2, + blockSize, + params, + 0, + 0); + + // Copy result from device to host + ASSERT_EQUAL(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + + // Validate results for both blocks together + for (int i = 0; i < 2 * blockSize; ++i) { + ASSERT_EQUAL(sizeTestH[i], 2 * blockSize); + ASSERT_EQUAL(thdRankTestH[i], i); + ASSERT_EQUAL(isValidTestH[i], 1); + ASSERT_EQUAL(syncTestH[i], 200); + } + + // Free device memory + ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); + ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); + ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess); + ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); + + //Free host memory + ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); +} + +int main() +{ + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + // Test block sizes which are powers of 2 + int i = 0; + while (true) { + int blockSize = pow(2, i); + if (blockSize > maxThreadsPerBlock) + break; + test_cg_grid_group_type(blockSize); + ++i; + } + + // Test some random block sizes + for(int j = 0; j < 10 ; ++j) { + int blockSize = rand() % maxThreadsPerBlock; + test_cg_grid_group_type(blockSize); + } + + passed(); +} diff --git a/tests/src/cg/hipCGGridGroupTypeViaBaseType.cpp b/tests/src/cg/hipCGGridGroupTypeViaBaseType.cpp new file mode 100644 index 0000000000..c1d3a66106 --- /dev/null +++ b/tests/src/cg/hipCGGridGroupTypeViaBaseType.cpp @@ -0,0 +1,157 @@ +/* +Copyright (c) 2020 - present 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#include "hip/hip_cooperative_groups.h" +#include +#include + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +using namespace cooperative_groups; + +static __global__ +void kernel_cg_grid_group_type_via_base_type(int *sizeTestD, + int *thdRankTestD, + int *isValidTestD, + int *syncTestD) +{ + thread_group tg = this_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test size + sizeTestD[gIdx] = tg.size(); + + // Test thread_rank + thdRankTestD[gIdx] = tg.thread_rank(); + + // Test is_valid + isValidTestD[gIdx] = tg.is_valid(); + + // Test sync + __device__ int gm[2]; + if (blockIdx.x == 0 && threadIdx.x == 0) + gm[0] = 10; + else if (blockIdx.x == 1 && threadIdx.x == 0) + gm[1] = 20; + tg.sync(); + syncTestD[gIdx] = gm[1] * gm[0]; +} + +static void test_cg_grid_group_type_via_base_type(int blockSize) +{ + int nBytes = sizeof(int) * 2 * blockSize; + int *sizeTestD, *sizeTestH; + int *thdRankTestD, *thdRankTestH; + int *isValidTestD, *isValidTestH; + int *syncTestD, *syncTestH; + + // Allocate device memory + ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); + + // Allocate host memory + ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); + + // Launch Kernel + void *params[4]; + params[0] = &sizeTestD; + params[1] = &thdRankTestD; + params[2] = &isValidTestD; + params[3] = &syncTestD; + hipLaunchCooperativeKernel(kernel_cg_grid_group_type_via_base_type, + 2, + blockSize, + params, + 0, + 0); + + // Copy result from device to host + ASSERT_EQUAL(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + + // Validate results for both blocks together + for (int i = 0; i < 2 * blockSize; ++i) { + ASSERT_EQUAL(sizeTestH[i], 2 * blockSize); + ASSERT_EQUAL(thdRankTestH[i], i); + ASSERT_EQUAL(isValidTestH[i], 1); + ASSERT_EQUAL(syncTestH[i], 200); + } + + // Free device memory + ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); + ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); + ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess); + ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); + + //Free host memory + ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); +} + +int main() +{ + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + // Test block sizes which are powers of 2 + int i = 0; + while (true) { + int blockSize = pow(2, i); + if (blockSize > maxThreadsPerBlock) + break; + test_cg_grid_group_type_via_base_type(blockSize); + ++i; + } + + // Test some random block sizes + for(int j = 0; j < 10 ; ++j) { + int blockSize = rand() % maxThreadsPerBlock; + test_cg_grid_group_type_via_base_type(blockSize); + } + + passed(); +} diff --git a/tests/src/cg/hipCGGridGroupTypeViaPublicApi.cpp b/tests/src/cg/hipCGGridGroupTypeViaPublicApi.cpp new file mode 100644 index 0000000000..50a2c2fd2e --- /dev/null +++ b/tests/src/cg/hipCGGridGroupTypeViaPublicApi.cpp @@ -0,0 +1,157 @@ +/* +Copyright (c) 2020 - present 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#include "hip/hip_cooperative_groups.h" +#include +#include + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +using namespace cooperative_groups; + +static __global__ +void kernel_cg_grid_group_type_via_public_api(int *sizeTestD, + int *thdRankTestD, + int *isValidTestD, + int *syncTestD) +{ + grid_group gg = this_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test group_size api + sizeTestD[gIdx] = group_size(gg); + + // Test thread_rank api + thdRankTestD[gIdx] = thread_rank(gg); + + // Test is_valid api + isValidTestD[gIdx] = is_valid(gg); + + // Test sync api + __device__ int gm[2]; + if (blockIdx.x == 0 && threadIdx.x == 0) + gm[0] = 10; + else if (blockIdx.x == 1 && threadIdx.x == 0) + gm[1] = 20; + sync(gg); + syncTestD[gIdx] = gm[1] * gm[0]; +} + +static void test_cg_grid_group_type_via_public_api(int blockSize) +{ + int nBytes = sizeof(int) * 2 * blockSize; + int *sizeTestD, *sizeTestH; + int *thdRankTestD, *thdRankTestH; + int *isValidTestD, *isValidTestH; + int *syncTestD, *syncTestH; + + // Allocate device memory + ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); + + // Allocate host memory + ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); + + // Launch Kernel + void *params[4]; + params[0] = &sizeTestD; + params[1] = &thdRankTestD; + params[2] = &isValidTestD; + params[3] = &syncTestD; + hipLaunchCooperativeKernel(kernel_cg_grid_group_type_via_public_api, + 2, + blockSize, + params, + 0, + 0); + + // Copy result from device to host + ASSERT_EQUAL(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + + // Validate results for both blocks together + for (int i = 0; i < 2 * blockSize; ++i) { + ASSERT_EQUAL(sizeTestH[i], 2 * blockSize); + ASSERT_EQUAL(thdRankTestH[i], i); + ASSERT_EQUAL(isValidTestH[i], 1); + ASSERT_EQUAL(syncTestH[i], 200); + } + + // Free device memory + ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); + ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); + ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess); + ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); + + //Free host memory + ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); +} + +int main() +{ + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + // Test block sizes which are powers of 2 + int i = 0; + while (true) { + int blockSize = pow(2, i); + if (blockSize > maxThreadsPerBlock) + break; + test_cg_grid_group_type_via_public_api(blockSize); + ++i; + } + + // Test some random block sizes + for(int j = 0; j < 10 ; ++j) { + int blockSize = rand() % maxThreadsPerBlock; + test_cg_grid_group_type_via_public_api(blockSize); + } + + passed(); +} diff --git a/tests/src/cg/hipCGMultiGridGroupType.cpp b/tests/src/cg/hipCGMultiGridGroupType.cpp new file mode 100644 index 0000000000..5a0529867a --- /dev/null +++ b/tests/src/cg/hipCGMultiGridGroupType.cpp @@ -0,0 +1,264 @@ +/* +Copyright (c) 2020 - present 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#include "hip/hip_cooperative_groups.h" +#include +#include +#include + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +using namespace cooperative_groups; + +static __global__ +void kernel_cg_multi_grid_group_type(int* numGridsTestD, + int* gridRankTestD, + int *sizeTestD, + int *thdRankTestD, + int *isValidTestD, + int *syncTestD, + int *syncResultD) +{ + multi_grid_group mg = this_multi_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test num_grids + numGridsTestD[gIdx] = mg.num_grids(); + + // Test grid_rank + gridRankTestD[gIdx] = mg.grid_rank(); + + // Test size + sizeTestD[gIdx] = mg.size(); + + // Test thread_rank + thdRankTestD[gIdx] = mg.thread_rank(); + + // Test is_valid + isValidTestD[gIdx] = mg.is_valid(); + + // Test sync + // + // Eech thread assign 1 to their respective location + syncTestD[gIdx] = 1; + // Grid level sync + this_grid().sync(); + // Thread 0 from work-group 0 of current grid (gpu) does grid level reduction + if (blockIdx.x == 0 && threadIdx.x == 0) { + for (int i = 1; i < gridDim.x * blockDim.x; ++i) { + syncTestD[0] += syncTestD[i]; + } + syncResultD[mg.grid_rank() + 1] = syncTestD[0]; + } + // multi-grid level sync + mg.sync(); + // grid (gpu) 0 does final reduction across all grids (gpus) + if (mg.grid_rank() == 0) { + syncResultD[0] = 0; + for (int i = 1; i <= mg.num_grids(); ++i) { + syncResultD[0] += syncResultD[i]; + } + } +} + +static void test_cg_multi_grid_group_type(int blockSize) +{ + // Get device count + constexpr int MaxGPUs = 8; + int nGpu = 0; + ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + + // Check if device suppurts multi gpu cooperative group support + hipDeviceProp_t deviceProp[MaxGPUs]; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + hipGetDeviceProperties(&deviceProp[i], 0); + if (!deviceProp[i].cooperativeMultiDeviceLaunch) { + printf("Device doesn't support multi gpu cooperative launch"); + return; + } + hipDeviceSynchronize(); + } + + // Create a stream each device + hipStream_t stream[MaxGPUs]; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + ASSERT_EQUAL(hipStreamCreate(&stream[i]), hipSuccess); + hipDeviceSynchronize(); + } + + // Allocate host and device memory + int nBytes = sizeof(int) * 2 * blockSize; + int *numGridsTestD[MaxGPUs], *numGridsTestH[MaxGPUs]; + int *gridRankTestD[MaxGPUs], *gridRankTestH[MaxGPUs]; + int *sizeTestD[MaxGPUs], *sizeTestH[MaxGPUs]; + int *thdRankTestD[MaxGPUs], *thdRankTestH[MaxGPUs]; + int *isValidTestD[MaxGPUs], *isValidTestH[MaxGPUs]; + int *syncTestD[MaxGPUs], *syncResultD; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + ASSERT_EQUAL(hipMalloc(&numGridsTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&gridRankTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&sizeTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdRankTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&isValidTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&syncTestD[i], nBytes), hipSuccess); + + ASSERT_EQUAL(hipHostMalloc(&numGridsTestH[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&gridRankTestH[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&sizeTestH[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdRankTestH[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&isValidTestH[i], nBytes), hipSuccess); + + if (i == 0) { + ASSERT_EQUAL( + hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), + hipSuccess); + } + + hipDeviceSynchronize(); + } + + // Launch Kernel + constexpr int NumKernelArgs = 7; + hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; + void* args[MaxGPUs * NumKernelArgs]; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + args[i * NumKernelArgs] = &numGridsTestD[i]; + args[i * NumKernelArgs + 1] = &gridRankTestD[i]; + args[i * NumKernelArgs + 2] = &sizeTestD[i]; + args[i * NumKernelArgs + 3] = &thdRankTestD[i]; + args[i * NumKernelArgs + 4] = &isValidTestD[i]; + args[i * NumKernelArgs + 5] = &syncTestD[i]; + args[i * NumKernelArgs + 6] = &syncResultD; + + launchParamsList[i].func = reinterpret_cast(kernel_cg_multi_grid_group_type); + launchParamsList[i].gridDim = 2; + launchParamsList[i].blockDim = blockSize; + launchParamsList[i].sharedMem = 0; + launchParamsList[i].stream = stream[i]; + launchParamsList[i].args = &args[i * NumKernelArgs]; + + hipDeviceSynchronize(); + } + hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + + // Copy result from device to host + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + ASSERT_EQUAL(hipMemcpy(numGridsTestH[i], numGridsTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(gridRankTestH[i], gridRankTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(sizeTestH[i], sizeTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdRankTestH[i], thdRankTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + + hipDeviceSynchronize(); + } + + // Validate results + for (int i = 0; i < nGpu; ++i) { + for (int j = 0; j < 2 * blockSize; ++j) { + //ASSERT_EQUAL(numGridsTestH[i][j], nGpu); + //ASSERT_EQUAL(gridRankTestH[i][j], i); + ASSERT_EQUAL(sizeTestH[i][j], nGpu * 2 * blockSize); + ASSERT_EQUAL(thdRankTestH[i][j], (i * 2 * blockSize) + j); + ASSERT_EQUAL(isValidTestH[i][j], 1); + } + ASSERT_EQUAL(syncResultD[i+1], 2 * blockSize); + } + ASSERT_EQUAL(syncResultD[0], nGpu * 2 * blockSize); + + // Free host and device memory + delete [] launchParamsList; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + ASSERT_EQUAL(hipFree(numGridsTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(gridRankTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(sizeTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(thdRankTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(isValidTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); + + if (i == 0) + ASSERT_EQUAL(hipFree(syncResultD), hipSuccess); + + ASSERT_EQUAL(hipHostFree(numGridsTestH[i]), hipSuccess); + ASSERT_EQUAL(hipHostFree(gridRankTestH[i]), hipSuccess); + ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); + ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); + + hipDeviceSynchronize(); + } +} + +int main() +{ + // Set `maxThreadsPerBlock` by taking minimum among all available devices + int nGpu = 0; + ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + int maxThreadsPerBlock = INT_MAX; + for (int i = 0; i < nGpu; i++) { + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); + hipDeviceSynchronize(); + } + + // Test block sizes which are powers of 2 + int i = 0; + while (true) { + int blockSize = pow(2, i); + if (blockSize > maxThreadsPerBlock) + break; + test_cg_multi_grid_group_type(blockSize); + ++i; + } + + // Test some random block sizes + for(int j = 0; j < 10 ; ++j) { + int blockSize = rand() % maxThreadsPerBlock; + test_cg_multi_grid_group_type(blockSize); + } + + passed(); +} diff --git a/tests/src/cg/hipCGMultiGridGroupTypeViaBaseType.cpp b/tests/src/cg/hipCGMultiGridGroupTypeViaBaseType.cpp new file mode 100644 index 0000000000..dae72f4cf8 --- /dev/null +++ b/tests/src/cg/hipCGMultiGridGroupTypeViaBaseType.cpp @@ -0,0 +1,238 @@ +/* +Copyright (c) 2020 - present 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#include "hip/hip_cooperative_groups.h" +#include +#include +#include + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +using namespace cooperative_groups; + +static __global__ +void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, + int *thdRankTestD, + int *isValidTestD, + int *syncTestD, + int *syncResultD) +{ + thread_group tg = this_multi_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test size + sizeTestD[gIdx] = tg.size(); + + // Test thread_rank + thdRankTestD[gIdx] = tg.thread_rank(); + + // Test is_valid + isValidTestD[gIdx] = tg.is_valid(); + + // Test sync + // + // Eech thread assign 1 to their respective location + syncTestD[gIdx] = 1; + // Grid level sync + this_grid().sync(); + // Thread 0 from work-group 0 of current grid (gpu) does grid level reduction + if (blockIdx.x == 0 && threadIdx.x == 0) { + for (int i = 1; i < gridDim.x * blockDim.x; ++i) { + syncTestD[0] += syncTestD[i]; + } + syncResultD[this_multi_grid().grid_rank() + 1] = syncTestD[0]; + } + // multi-grid level sync + tg.sync(); + // grid (gpu) 0 does final reduction across all grids (gpus) + if (this_multi_grid().grid_rank() == 0) { + syncResultD[0] = 0; + for (int i = 1; i <= this_multi_grid().num_grids(); ++i) { + syncResultD[0] += syncResultD[i]; + } + } +} + +static void test_cg_multi_grid_group_type_via_base_type(int blockSize) +{ + // Get device count + constexpr int MaxGPUs = 8; + int nGpu = 0; + ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + + // Check if device suppurts multi gpu cooperative group support + hipDeviceProp_t deviceProp[MaxGPUs]; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + hipGetDeviceProperties(&deviceProp[i], 0); + if (!deviceProp[i].cooperativeMultiDeviceLaunch) { + printf("Device doesn't support multi gpu cooperative launch"); + return; + } + hipDeviceSynchronize(); + } + + // Create a stream each device + hipStream_t stream[MaxGPUs]; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + ASSERT_EQUAL(hipStreamCreate(&stream[i]), hipSuccess); + hipDeviceSynchronize(); + } + + // Allocate host and device memory + int nBytes = sizeof(int) * 2 * blockSize; + int *sizeTestD[MaxGPUs], *sizeTestH[MaxGPUs]; + int *thdRankTestD[MaxGPUs], *thdRankTestH[MaxGPUs]; + int *isValidTestD[MaxGPUs], *isValidTestH[MaxGPUs]; + int *syncTestD[MaxGPUs], *syncResultD; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + ASSERT_EQUAL(hipMalloc(&sizeTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdRankTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&isValidTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&syncTestD[i], nBytes), hipSuccess); + + ASSERT_EQUAL(hipHostMalloc(&sizeTestH[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdRankTestH[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&isValidTestH[i], nBytes), hipSuccess); + + if (i == 0) { + ASSERT_EQUAL( + hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), + hipSuccess); + } + + hipDeviceSynchronize(); + } + + // Launch Kernel + constexpr int NumKernelArgs = 5; + hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; + void* args[MaxGPUs * NumKernelArgs]; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + args[i * NumKernelArgs ] = &sizeTestD[i]; + args[i * NumKernelArgs + 1] = &thdRankTestD[i]; + args[i * NumKernelArgs + 2] = &isValidTestD[i]; + args[i * NumKernelArgs + 3] = &syncTestD[i]; + args[i * NumKernelArgs + 4] = &syncResultD; + + launchParamsList[i].func = reinterpret_cast(kernel_cg_multi_grid_group_type_via_base_type); + launchParamsList[i].gridDim = 2; + launchParamsList[i].blockDim = blockSize; + launchParamsList[i].sharedMem = 0; + launchParamsList[i].stream = stream[i]; + launchParamsList[i].args = &args[i * NumKernelArgs]; + + hipDeviceSynchronize(); + } + hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + + // Copy result from device to host + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + ASSERT_EQUAL(hipMemcpy(sizeTestH[i], sizeTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdRankTestH[i], thdRankTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + + hipDeviceSynchronize(); + } + + // Validate results + for (int i = 0; i < nGpu; ++i) { + for (int j = 0; j < 2 * blockSize; ++j) { + ASSERT_EQUAL(sizeTestH[i][j], nGpu * 2 * blockSize); + ASSERT_EQUAL(thdRankTestH[i][j], (i * 2 * blockSize) + j); + ASSERT_EQUAL(isValidTestH[i][j], 1); + } + ASSERT_EQUAL(syncResultD[i+1], 2 * blockSize); + } + ASSERT_EQUAL(syncResultD[0], nGpu * 2 * blockSize); + + // Free host and device memory + delete [] launchParamsList; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + ASSERT_EQUAL(hipFree(sizeTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(thdRankTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(isValidTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); + + if (i == 0) + ASSERT_EQUAL(hipFree(syncResultD), hipSuccess); + + ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); + ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); + + hipDeviceSynchronize(); + } +} + +int main() +{ + // Set `maxThreadsPerBlock` by taking minimum among all available devices + int nGpu = 0; + ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + int maxThreadsPerBlock = INT_MAX; + for (int i = 0; i < nGpu; i++) { + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); + hipDeviceSynchronize(); + } + + // Test block sizes which are powers of 2 + int i = 0; + while (true) { + int blockSize = pow(2, i); + if (blockSize > maxThreadsPerBlock) + break; + test_cg_multi_grid_group_type_via_base_type(blockSize); + ++i; + } + + // Test some random block sizes + for(int j = 0; j < 10 ; ++j) { + int blockSize = rand() % maxThreadsPerBlock; + test_cg_multi_grid_group_type_via_base_type(blockSize); + } + + passed(); +} diff --git a/tests/src/cg/hipCGMultiGridGroupTypeViaPublicApi.cpp b/tests/src/cg/hipCGMultiGridGroupTypeViaPublicApi.cpp new file mode 100644 index 0000000000..2f2f378931 --- /dev/null +++ b/tests/src/cg/hipCGMultiGridGroupTypeViaPublicApi.cpp @@ -0,0 +1,238 @@ +/* +Copyright (c) 2020 - present 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#include "hip/hip_cooperative_groups.h" +#include +#include +#include + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +using namespace cooperative_groups; + +static __global__ +void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD, + int *thdRankTestD, + int *isValidTestD, + int *syncTestD, + int *syncResultD) +{ + multi_grid_group mg = this_multi_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test group_size api + sizeTestD[gIdx] = group_size(mg); + + // Test thread_rank api + thdRankTestD[gIdx] = thread_rank(mg); + + // Test is_valid api + isValidTestD[gIdx] = is_valid(mg); + + // Test sync api + // + // Eech thread assign 1 to their respective location + syncTestD[gIdx] = 1; + // Grid level sync + sync(this_grid()); + // Thread 0 from work-group 0 of current grid (gpu) does grid level reduction + if (blockIdx.x == 0 && threadIdx.x == 0) { + for (int i = 1; i < gridDim.x * blockDim.x; ++i) { + syncTestD[0] += syncTestD[i]; + } + syncResultD[this_multi_grid().grid_rank() + 1] = syncTestD[0]; + } + // multi-grid level sync via public api + sync(mg); + // grid (gpu) 0 does final reduction across all grids (gpus) + if (this_multi_grid().grid_rank() == 0) { + syncResultD[0] = 0; + for (int i = 1; i <= this_multi_grid().num_grids(); ++i) { + syncResultD[0] += syncResultD[i]; + } + } +} + +static void test_cg_multi_grid_group_type_via_public_api(int blockSize) +{ + // Get device count + constexpr int MaxGPUs = 8; + int nGpu = 0; + ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + + // Check if device suppurts multi gpu cooperative group support + hipDeviceProp_t deviceProp[MaxGPUs]; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + hipGetDeviceProperties(&deviceProp[i], 0); + if (!deviceProp[i].cooperativeMultiDeviceLaunch) { + printf("Device doesn't support multi gpu cooperative launch"); + return; + } + hipDeviceSynchronize(); + } + + // Create a stream each device + hipStream_t stream[MaxGPUs]; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + ASSERT_EQUAL(hipStreamCreate(&stream[i]), hipSuccess); + hipDeviceSynchronize(); + } + + // Allocate host and device memory + int nBytes = sizeof(int) * 2 * blockSize; + int *sizeTestD[MaxGPUs], *sizeTestH[MaxGPUs]; + int *thdRankTestD[MaxGPUs], *thdRankTestH[MaxGPUs]; + int *isValidTestD[MaxGPUs], *isValidTestH[MaxGPUs]; + int *syncTestD[MaxGPUs], *syncResultD; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + ASSERT_EQUAL(hipMalloc(&sizeTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdRankTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&isValidTestD[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&syncTestD[i], nBytes), hipSuccess); + + ASSERT_EQUAL(hipHostMalloc(&sizeTestH[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdRankTestH[i], nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&isValidTestH[i], nBytes), hipSuccess); + + if (i == 0) { + ASSERT_EQUAL( + hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), + hipSuccess); + } + + hipDeviceSynchronize(); + } + + // Launch Kernel + constexpr int NumKernelArgs = 5; + hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; + void* args[MaxGPUs * NumKernelArgs]; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + args[i * NumKernelArgs ] = &sizeTestD[i]; + args[i * NumKernelArgs + 1] = &thdRankTestD[i]; + args[i * NumKernelArgs + 2] = &isValidTestD[i]; + args[i * NumKernelArgs + 3] = &syncTestD[i]; + args[i * NumKernelArgs + 4] = &syncResultD; + + launchParamsList[i].func = reinterpret_cast(kernel_cg_multi_grid_group_type_via_public_api); + launchParamsList[i].gridDim = 2; + launchParamsList[i].blockDim = blockSize; + launchParamsList[i].sharedMem = 0; + launchParamsList[i].stream = stream[i]; + launchParamsList[i].args = &args[i * NumKernelArgs]; + + hipDeviceSynchronize(); + } + hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + + // Copy result from device to host + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + ASSERT_EQUAL(hipMemcpy(sizeTestH[i], sizeTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdRankTestH[i], thdRankTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), + hipSuccess); + + hipDeviceSynchronize(); + } + + // Validate results + for (int i = 0; i < nGpu; ++i) { + for (int j = 0; j < 2 * blockSize; ++j) { + ASSERT_EQUAL(sizeTestH[i][j], nGpu * 2 * blockSize); + ASSERT_EQUAL(thdRankTestH[i][j], (i * 2 * blockSize) + j); + ASSERT_EQUAL(isValidTestH[i][j], 1); + } + ASSERT_EQUAL(syncResultD[i+1], 2 * blockSize); + } + ASSERT_EQUAL(syncResultD[0], nGpu * 2 * blockSize); + + // Free host and device memory + delete [] launchParamsList; + for (int i = 0; i < nGpu; i++) { + ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + + ASSERT_EQUAL(hipFree(sizeTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(thdRankTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(isValidTestD[i]), hipSuccess); + ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); + + if (i == 0) + ASSERT_EQUAL(hipFree(syncResultD), hipSuccess); + + ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); + ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); + + hipDeviceSynchronize(); + } +} + +int main() +{ + // Set `maxThreadsPerBlock` by taking minimum among all available devices + int nGpu = 0; + ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + int maxThreadsPerBlock = INT_MAX; + for (int i = 0; i < nGpu; i++) { + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); + hipDeviceSynchronize(); + } + + // Test block sizes which are powers of 2 + int i = 0; + while (true) { + int blockSize = pow(2, i); + if (blockSize > maxThreadsPerBlock) + break; + test_cg_multi_grid_group_type_via_public_api(blockSize); + ++i; + } + + // Test some random block sizes + for(int j = 0; j < 10 ; ++j) { + int blockSize = rand() % maxThreadsPerBlock; + test_cg_multi_grid_group_type_via_public_api(blockSize); + } + + passed(); +} diff --git a/tests/src/cg/hipCGThreadBlockType.cpp b/tests/src/cg/hipCGThreadBlockType.cpp index 14c2e3ce2a..4e1de9e44a 100644 --- a/tests/src/cg/hipCGThreadBlockType.cpp +++ b/tests/src/cg/hipCGThreadBlockType.cpp @@ -22,143 +22,165 @@ THE SOFTWARE. /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s ../test_common.cpp * TEST: %t * HIT_END */ #include "test_common.h" #include "hip/hip_cooperative_groups.h" +#include +#include -#define HIP_ASSERT(lhs, rhs) assert(lhs == rhs) +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) using namespace cooperative_groups; static __global__ -void kernel_cg_thread_block_type(dim3 *groupIndexD, - dim3 *thdIndexD, - int *sizeD, - int *thdRankD, - int *isValidD, - int *syncValD) +void kernel_cg_thread_block_type(int *sizeTestD, + int *thdRankTestD, + int *isValidTestD, + int *syncTestD, + dim3 *groupIndexTestD, + dim3 *thdIndexTestD) { thread_block tb = this_thread_block(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; - // Consider the workgroup id (0, 1, 1) and thread id (0, 1, 2) for validation - // of the test - int isBlockIdx011 = - (hipBlockIdx_x == 0 && hipBlockIdx_y == 1 && hipBlockIdx_z == 1); - int isThreadIdx012 = - (hipThreadIdx_x == 0 && hipThreadIdx_y == 1 && hipThreadIdx_z == 2); + // Test size + sizeTestD[gIdx] = tb.size(); - if (isBlockIdx011 && isThreadIdx012) { - *groupIndexD = tb.group_index(); - *thdIndexD = tb.thread_index(); - *sizeD = tb.size(); - *thdRankD = tb.thread_rank(); - *isValidD = tb.is_valid(); - } + // Test thread_rank + thdRankTestD[gIdx] = tb.thread_rank(); - // Consider local thread id (0, 0, 0) and (0, 0, 1) for validation of sync() - // api - __shared__ int sVar[2]; - int isThreadIdx000 = - (hipThreadIdx_x == 0 && hipThreadIdx_y == 0 && hipThreadIdx_z == 0); - int isThreadIdx001 = - (hipThreadIdx_x == 0 && hipThreadIdx_y == 0 && hipThreadIdx_z == 1); + // Test is_valid + isValidTestD[gIdx] = tb.is_valid(); - if (isThreadIdx000) - sVar[0] = 10; - if (isThreadIdx001) - sVar[1] = 20; + // Test sync + __shared__ int sm[2]; + if (threadIdx.x == 0) + sm[0] = 10; + else if (threadIdx.x == 1) + sm[1] = 20; tb.sync(); - if (isBlockIdx011 && isThreadIdx012) - *syncValD = sVar[0] + sVar[1]; + syncTestD[gIdx] = sm[1] * sm[0]; + + // Test group_index + groupIndexTestD[gIdx] = tb.group_index(); + + // Test thread_index + thdIndexTestD[gIdx] = tb.thread_index(); } -static void test_cg_thread_block_type() +static void test_cg_thread_block_type(int blockSize) { - dim3 *groupIndexD, *groupIndexH; - dim3 *thdIndexD, *thdIndexH; - int *sizeD, *sizeH; - int *thdRankD, *thdRankH; - int *isValidD, *isValidH; - int *syncValD, *syncValH; + int nBytes = sizeof(int) * 2 * blockSize; + int nDim3Bytes = sizeof(dim3) * 2 * blockSize; + int *sizeTestD, *sizeTestH; + int *thdRankTestD, *thdRankTestH; + int *isValidTestD, *isValidTestH; + int *syncTestD, *syncTestH; + dim3 *groupIndexTestD, *groupIndexTestH; + dim3 *thdIndexTestD, *thdIndexTestH; // Allocate device memory - HIP_ASSERT(hipMalloc((void**)&groupIndexD, sizeof(dim3)), hipSuccess); - HIP_ASSERT(hipMalloc((void**)&thdIndexD, sizeof(dim3)), hipSuccess); - HIP_ASSERT(hipMalloc((void**)&sizeD, sizeof(int)), hipSuccess); - HIP_ASSERT(hipMalloc((void**)&thdRankD, sizeof(int)), hipSuccess); - HIP_ASSERT(hipMalloc((void**)&isValidD, sizeof(int)), hipSuccess); - HIP_ASSERT(hipMalloc((void**)&syncValD, sizeof(int)), hipSuccess); + ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&groupIndexTestD, nDim3Bytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdIndexTestD, nDim3Bytes), hipSuccess); // Allocate host memory - HIP_ASSERT(hipHostMalloc((void**)&groupIndexH, sizeof(dim3)), hipSuccess); - HIP_ASSERT(hipHostMalloc((void**)&thdIndexH, sizeof(dim3)), hipSuccess); - HIP_ASSERT(hipHostMalloc((void**)&sizeH, sizeof(int)), hipSuccess); - HIP_ASSERT(hipHostMalloc((void**)&thdRankH, sizeof(int)), hipSuccess); - HIP_ASSERT(hipHostMalloc((void**)&isValidH, sizeof(int)), hipSuccess); - HIP_ASSERT(hipHostMalloc((void**)&syncValH, sizeof(int)), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&groupIndexTestH, nDim3Bytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdIndexTestH, nDim3Bytes), hipSuccess); // Launch Kernel hipLaunchKernelGGL(kernel_cg_thread_block_type, - dim3(2, 2, 2), - dim3(4, 4, 4), + 2, + blockSize, 0, 0, - groupIndexD, - thdIndexD, - sizeD, - thdRankD, - isValidD, - syncValD); + sizeTestD, + thdRankTestD, + isValidTestD, + syncTestD, + groupIndexTestD, + thdIndexTestD); // Copy result from device to host - HIP_ASSERT(hipMemcpy(groupIndexH, groupIndexD, sizeof(dim3), hipMemcpyDeviceToHost), hipSuccess); - HIP_ASSERT(hipMemcpy(thdIndexH, thdIndexD, sizeof(dim3), hipMemcpyDeviceToHost), hipSuccess); - HIP_ASSERT(hipMemcpy(sizeH, sizeD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess); - HIP_ASSERT(hipMemcpy(thdRankH, thdRankD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess); - HIP_ASSERT(hipMemcpy(isValidH, isValidD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess); - HIP_ASSERT(hipMemcpy(syncValH, syncValD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess); + ASSERT_EQUAL(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(groupIndexTestH, groupIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdIndexTestH, thdIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost), + hipSuccess); - // Validate result - // Group index should be (0, 1, 1) - HIP_ASSERT(groupIndexH->x, 0); - HIP_ASSERT(groupIndexH->y, 1); - HIP_ASSERT(groupIndexH->z, 1); - // Thread index should be (0, 1, 2) - HIP_ASSERT(thdIndexH->x, 0); - HIP_ASSERT(thdIndexH->y, 1); - HIP_ASSERT(thdIndexH->z, 2); - // Workgroup size should be 64 - HIP_ASSERT(*sizeH, 64); - // Thread rank should be 36 - HIP_ASSERT(*thdRankH, 36); - // Call to is_valid() should return true - HIP_ASSERT(*isValidH, 1); - // syncVal should be equal to 30 - HIP_ASSERT(*syncValH, 30); + // Validate results for both blocks together + for (int i = 0; i < 2 * blockSize; ++i) { + ASSERT_EQUAL(sizeTestH[i], blockSize); + ASSERT_EQUAL(thdRankTestH[i], i % blockSize); + ASSERT_EQUAL(isValidTestH[i], 1); + ASSERT_EQUAL(syncTestH[i], 200); + ASSERT_EQUAL(groupIndexTestH[i].x, i / blockSize); + ASSERT_EQUAL(groupIndexTestH[i].y, 0); + ASSERT_EQUAL(groupIndexTestH[i].z, 0); + ASSERT_EQUAL(thdIndexTestH[i].x, i % blockSize); + ASSERT_EQUAL(thdIndexTestH[i].y, 0); + ASSERT_EQUAL(thdIndexTestH[i].z, 0); + } // Free device memory - HIP_ASSERT(hipFree(groupIndexD), hipSuccess); - HIP_ASSERT(hipFree(thdIndexD), hipSuccess); - HIP_ASSERT(hipFree(sizeD), hipSuccess); - HIP_ASSERT(hipFree(thdRankD), hipSuccess); - HIP_ASSERT(hipFree(isValidD), hipSuccess); - HIP_ASSERT(hipFree(syncValD), hipSuccess); + ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); + ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); + ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess); + ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); + ASSERT_EQUAL(hipFree(groupIndexTestD), hipSuccess); + ASSERT_EQUAL(hipFree(thdIndexTestD), hipSuccess); //Free host memory - HIP_ASSERT(hipHostFree(groupIndexH), hipSuccess); - HIP_ASSERT(hipHostFree(thdIndexH), hipSuccess); - HIP_ASSERT(hipHostFree(sizeH), hipSuccess); - HIP_ASSERT(hipHostFree(thdRankH), hipSuccess); - HIP_ASSERT(hipHostFree(isValidH), hipSuccess); - HIP_ASSERT(hipHostFree(syncValH), hipSuccess); + ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(groupIndexTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdIndexTestH), hipSuccess); } int main() { - test_cg_thread_block_type(); + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + // Test block sizes which are powers of 2 + int i = 0; + while (true) { + int blockSize = pow(2, i); + if (blockSize > maxThreadsPerBlock) + break; + test_cg_thread_block_type(blockSize); + ++i; + } + + // Test some random block sizes + for(int j = 0; j < 10 ; ++j) { + int blockSize = rand() % maxThreadsPerBlock; + test_cg_thread_block_type(blockSize); + } + passed(); } diff --git a/tests/src/cg/hipCGThreadBlockTypeViaBaseType.cpp b/tests/src/cg/hipCGThreadBlockTypeViaBaseType.cpp new file mode 100644 index 0000000000..d4c9402268 --- /dev/null +++ b/tests/src/cg/hipCGThreadBlockTypeViaBaseType.cpp @@ -0,0 +1,155 @@ +/* +Copyright (c) 2020 - present 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#include "hip/hip_cooperative_groups.h" +#include +#include + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +using namespace cooperative_groups; + +static __global__ +void kernel_cg_thread_block_type_via_base_type(int *sizeTestD, + int *thdRankTestD, + int *isValidTestD, + int *syncTestD) +{ + thread_group tg = this_thread_block(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test size + sizeTestD[gIdx] = tg.size(); + + // Test thread_rank + thdRankTestD[gIdx] = tg.thread_rank(); + + // Test is_valid + isValidTestD[gIdx] = tg.is_valid(); + + // Test sync + __shared__ int sm[2]; + if (threadIdx.x == 0) + sm[0] = 10; + else if (threadIdx.x == 1) + sm[1] = 20; + tg.sync(); + syncTestD[gIdx] = sm[1] * sm[0]; +} + +static void test_cg_thread_block_type_via_base_type(int blockSize) +{ + int nBytes = sizeof(int) * 2 * blockSize; + int *sizeTestD, *sizeTestH; + int *thdRankTestD, *thdRankTestH; + int *isValidTestD, *isValidTestH; + int *syncTestD, *syncTestH; + + // Allocate device memory + ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); + + // Allocate host memory + ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); + + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_thread_block_type_via_base_type, + 2, + blockSize, + 0, + 0, + sizeTestD, + thdRankTestD, + isValidTestD, + syncTestD); + + // Copy result from device to host + ASSERT_EQUAL(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + + // Validate results for both blocks together + for (int i = 0; i < 2 * blockSize; ++i) { + ASSERT_EQUAL(sizeTestH[i], blockSize); + ASSERT_EQUAL(thdRankTestH[i], i % blockSize); + ASSERT_EQUAL(isValidTestH[i], 1); + ASSERT_EQUAL(syncTestH[i], 200); + } + + // Free device memory + ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); + ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); + ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess); + ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); + + //Free host memory + ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); +} + +int main() +{ + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + // Test block sizes which are powers of 2 + int i = 0; + while (true) { + int blockSize = pow(2, i); + if (blockSize > maxThreadsPerBlock) + break; + test_cg_thread_block_type_via_base_type(blockSize); + ++i; + } + + // Test some random block sizes + for(int j = 0; j < 10 ; ++j) { + int blockSize = rand() % maxThreadsPerBlock; + test_cg_thread_block_type_via_base_type(blockSize); + } + + passed(); +} diff --git a/tests/src/cg/hipCGThreadBlockTypeViaPublicApi.cpp b/tests/src/cg/hipCGThreadBlockTypeViaPublicApi.cpp new file mode 100644 index 0000000000..d13e58b059 --- /dev/null +++ b/tests/src/cg/hipCGThreadBlockTypeViaPublicApi.cpp @@ -0,0 +1,155 @@ +/* +Copyright (c) 2020 - present 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#include "hip/hip_cooperative_groups.h" +#include +#include + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +using namespace cooperative_groups; + +static __global__ +void kernel_cg_thread_block_type_via_public_api(int *sizeTestD, + int *thdRankTestD, + int *isValidTestD, + int *syncTestD) +{ + thread_block tb = this_thread_block(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test group_size api + sizeTestD[gIdx] = group_size(tb); + + // Test thread_rank api + thdRankTestD[gIdx] = thread_rank(tb); + + // Test is_valid api + isValidTestD[gIdx] = is_valid(tb); + + // Test sync api + __shared__ int sm[2]; + if (threadIdx.x == 0) + sm[0] = 10; + else if (threadIdx.x == 1) + sm[1] = 20; + sync(tb); + syncTestD[gIdx] = sm[1] * sm[0]; +} + +static void test_cg_thread_block_type_via_public_api(int blockSize) +{ + int nBytes = sizeof(int) * 2 * blockSize; + int *sizeTestD, *sizeTestH; + int *thdRankTestD, *thdRankTestH; + int *isValidTestD, *isValidTestH; + int *syncTestD, *syncTestH; + + // Allocate device memory + ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess); + ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); + + // Allocate host memory + ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess); + ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); + + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_thread_block_type_via_public_api, + 2, + blockSize, + 0, + 0, + sizeTestD, + thdRankTestD, + isValidTestD, + syncTestD); + + // Copy result from device to host + ASSERT_EQUAL(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), + hipSuccess); + + // Validate results for both blocks together + for (int i = 0; i < 2 * blockSize; ++i) { + ASSERT_EQUAL(sizeTestH[i], blockSize); + ASSERT_EQUAL(thdRankTestH[i], i % blockSize); + ASSERT_EQUAL(isValidTestH[i], 1); + ASSERT_EQUAL(syncTestH[i], 200); + } + + // Free device memory + ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); + ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); + ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess); + ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); + + //Free host memory + ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); +} + +int main() +{ + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + // Test block sizes which are powers of 2 + int i = 0; + while (true) { + int blockSize = pow(2, i); + if (blockSize > maxThreadsPerBlock) + break; + test_cg_thread_block_type_via_public_api(blockSize); + ++i; + } + + // Test some random block sizes + for(int j = 0; j < 10 ; ++j) { + int blockSize = rand() % maxThreadsPerBlock; + test_cg_thread_block_type_via_public_api(blockSize); + } + + passed(); +}