AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

Merge 'amd-master-next' into 'amd-npi-next'

Procházet zdrojové kódy 
Change-Id: I914aa748c92849f9f7ab796129a4617753d728db
Tento commit je obsažen v:
Jenkins
2020-08-25 11:10:53 -04:00
rodič 1d83fef5da 15aff89ba1
revize f9ebdacdfa
15 změnil soubory, kde provedl 1698 přidání a 121 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
bin/hipcc
+1 -1
Zobrazit soubor
@@ -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) {
bin/hipconfig
+2 -2
Zobrazit soubor
@@ -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;
@@ -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";
include/hip/hcc_detail/hip_cooperative_groups.h
+30 -1
Zobrazit soubor
@@ -230,7 +230,8 @@ __CG_QUALIFIER__ uint32_t thread_group::thread_rank() const {
return (static_cast<const thread_block*>(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<const thread_block*>(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<const thread_block*>(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 <class CGTy>
__CG_QUALIFIER__ uint32_t group_size(CGTy const &g) {
return g.size();
}
template <class CGTy>
__CG_QUALIFIER__ uint32_t thread_rank(CGTy const &g) {
return g.thread_rank();
}
template <class CGTy>
__CG_QUALIFIER__ bool is_valid(CGTy const &g) {
return g.is_valid();
}
template <class CGTy>
__CG_QUALIFIER__ void sync(CGTy const &g) {
g.sync();
}
} // namespace cooperative_groups
#endif // __cplusplus
rocclr/hip_code_object.cpp
+1 -1
Zobrazit soubor
@@ -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 */
rocclr/hip_global.cpp
+23 -18
Zobrazit soubor
@@ -14,20 +14,20 @@ DeviceVar::DeviceVar(std::string name, hipModule_t hmod) : shadowVptr(nullptr),
amd::Program* program = as_amd(reinterpret_cast<cl_program>(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<size_t>(deviceId) < modules_->size());
guarantee(modules_ != nullptr && "Module not initialized");
guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero");
guarantee((static_cast<size_t>(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<size_t>(deviceId) < modules_->size());
guarantee((modules_ != nullptr) && "Module not initialized");
guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero");
guarantee((static_cast<size_t>(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<size_t>(deviceId) < g_devices.size());
guarantee(dVar_.size() == g_devices.size());
guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero");
guarantee((static_cast<size_t>(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<size_t>(deviceId) < g_devices.size());
guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero");
guarantee((static_cast<size_t>(deviceId) < g_devices.size())
&& "Invalid DeviceId, greater than no of code objects");
hipModule_t module = (*modules_)[deviceId].first;
FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second;
rocclr/hip_platform.cpp
+1 -1
Zobrazit soubor
@@ -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");
}
}
}
tests/src/cg/hipCGGridGroupType.cpp
+157
Zobrazit soubor
@@ -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 <cmath>
#include <cstdlib>
#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();
}
tests/src/cg/hipCGGridGroupTypeViaBaseType.cpp
+157
Zobrazit soubor
@@ -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 <cmath>
#include <cstdlib>
#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();
}
tests/src/cg/hipCGGridGroupTypeViaPublicApi.cpp
+157
Zobrazit soubor
@@ -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 <cmath>
#include <cstdlib>
#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();
}
tests/src/cg/hipCGMultiGridGroupType.cpp
+264
Zobrazit soubor
@@ -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 <cmath>
#include <cstdlib>
#include <climits>
#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<void*>(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();
}
tests/src/cg/hipCGMultiGridGroupTypeViaBaseType.cpp
+238
Zobrazit soubor
@@ -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 <cmath>
#include <cstdlib>
#include <climits>
#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<void*>(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();
}
tests/src/cg/hipCGMultiGridGroupTypeViaPublicApi.cpp
+238
Zobrazit soubor
@@ -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 <cmath>
#include <cstdlib>
#include <climits>
#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<void*>(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();
}
tests/src/cg/hipCGThreadBlockType.cpp
+119 -97
Zobrazit soubor
@@ -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 <cmath>
#include <cstdlib>
#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();
}
tests/src/cg/hipCGThreadBlockTypeViaBaseType.cpp
+155
Zobrazit soubor
@@ -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 <cmath>
#include <cstdlib>
#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();
}
tests/src/cg/hipCGThreadBlockTypeViaPublicApi.cpp
+155
Zobrazit soubor
@@ -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 <cmath>
#include <cstdlib>
#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();
}