AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

Support public apis for cooperative group types.

Sfoglia il codice sorgente 
Change-Id: I01346afde485e82c34b7868b9241b34c542d0cf9
This commit is contained in:
mshivama
2020-08-09 13:30:08 +05:30
parent 5fb155dbf5
commit bfb54cc5e9
10 ha cambiato i file con 1670 aggiunte e 98 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
hipamd/include/hip/hcc_detail/hip_cooperative_groups.h
+30 -1
Vedi File
@@ -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
hipamd/tests/src/cg/hipCGGridGroupType.cpp
+157
Vedi File
@@ -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();
}
hipamd/tests/src/cg/hipCGGridGroupTypeViaBaseType.cpp
+157
Vedi File
@@ -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();
}
hipamd/tests/src/cg/hipCGGridGroupTypeViaPublicApi.cpp
+157
Vedi File
@@ -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();
}
hipamd/tests/src/cg/hipCGMultiGridGroupType.cpp
+264
Vedi File
@@ -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();
}
hipamd/tests/src/cg/hipCGMultiGridGroupTypeViaBaseType.cpp
+238
Vedi File
@@ -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();
}
hipamd/tests/src/cg/hipCGMultiGridGroupTypeViaPublicApi.cpp
+238
Vedi File
@@ -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();
}
hipamd/tests/src/cg/hipCGThreadBlockType.cpp
+119 -97
Vedi File
@@ -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();
}
hipamd/tests/src/cg/hipCGThreadBlockTypeViaBaseType.cpp
+155
Vedi File
@@ -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();
}
hipamd/tests/src/cg/hipCGThreadBlockTypeViaPublicApi.cpp
+155
Vedi File
@@ -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();
}