AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Use C++ iota function to reset buffers and use its values for verification

Przeglądaj źródła 
* Update functional test script to include new tests
This commit is contained in:
avinashkethineedi
2024-10-14 23:49:40 +00:00
rodzic e981f61693
commit 18a1bdd0ac
8 zmienionych plików z 161 dodań i 56 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
scripts/build_configs/ro_net
+1 -2
Wyświetl plik
@@ -17,9 +17,8 @@ cmake \
-DDEBUG=OFF \
-DPROFILE=OFF \
-DUSE_GPU_IB=OFF \
-DUSE_RO=ON \
-DUSE_DC=OFF \
-DUSE_IPC=OFF \
-DUSE_IPC=ON \
-DUSE_THREADS=ON \
-DUSE_WF_COAL=OFF \
-DUSE_COHERENT_HEAP=OFF \
scripts/functional_tests/driver.sh
+64 -16
Wyświetl plik
@@ -56,52 +56,52 @@ case $2 in
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 1 -s 1048576 -a 3 > $3/putnbi_n2_w1_z1_1MB.log
check putnbi_n2_w1_z1_1MB
echo "wg_get_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 28 > $3/wg_get_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 28 > $3/wg_get_n2_w1_z64_1MB.log
check wg_get_n2_w1_z1_1MB
echo "wg_getnbi_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 29 > $3/wg_getnbi_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 29 > $3/wg_getnbi_n2_w1_z64_1MB.log
check wg_getnbi_n2_w1_z1_1MB
echo "wg_put_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 30 > $3/wg_put_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 30 > $3/wg_put_n2_w1_z64_1MB.log
check wg_put_n2_w1_z1_1MB
echo "wg_putnbi_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 31 > $3/wg_putnbi_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 31 > $3/wg_putnbi_n2_w1_z64_1MB.log
check wg_putnbi_n2_w1_z1_1MB
echo "wg_get_tiled_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 64 -s 1048576 -a 28 > $3/wg_get_tiled_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 64 -s 1048576 -a 28 > $3/wg_get_tiled_n2_w2_z64_1MB.log
check wg_get_tiled_n2_w1_z1_1MB
echo "wg_getnbi_tiled_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 64 -s 1048576 -a 29 > $3/wg_getnbi_tiled_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 64 -s 1048576 -a 29 > $3/wg_getnbi_tiled_n2_w2_z64_1MB.log
check wg_getnbi_tiled_n2_w1_z1_1MB
echo "wg_put_tiled_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 64 -s 1048576 -a 30 > $3/wg_put_tiled_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 64 -s 1048576 -a 30 > $3/wg_put_tiled_n2_w2_z64_1MB.log
check wg_put_tiled_n2_w1_z1_1MB
echo "wg_putnbi_tiled_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 64 -s 1048576 -a 31 > $3/wg_putnbi_tiled_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 64 -s 1048576 -a 31 > $3/wg_putnbi_tiled_n2_w2_z64_1MB.log
check wg_putnbi_tiled_n2_w1_z1_1MB
echo "wave_get_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 32 > $3/wave_get_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 32 > $3/wave_get_n2_w1_z64_1MB.log
check wave_get_n2_w1_z1_1MB
echo "wave_getnbi_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 33 > $3/wave_getnbi_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 33 > $3/wave_getnbi_n2_w1_z64_1MB.log
check wave_getnbi_n2_w1_z1_1MB
echo "wave_put_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 34 > $3/wave_put_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 34 > $3/wave_put_n2_w1_z64_1MB.log
check wave_put_n2_w1_z1_1MB
echo "wave_putnbi_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 35 > $3/wave_putnbi_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 1048576 -a 35 > $3/wave_putnbi_n2_w1_z64_1MB.log
check wave_putnbi_n2_w1_z1_1MB
echo "wave_get_tiled_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 128 -s 1048576 -a 32 > $3/wave_get_tiled_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 128 -s 1048576 -a 32 > $3/wave_get_tiled_n2_w2_z128_1MB.log
check wave_get_tiled_n2_w1_z1_1MB
echo "wave_getnbi_tiled_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 128 -s 1048576 -a 33 > $3/wave_getnbi_tiled_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 128 -s 1048576 -a 33 > $3/wave_getnbi_tiled_n2_w2_z128_1MB.log
check wave_getnbi_tiled_n2_w1_z1_1MB
echo "wave_put_tiled_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 128 -s 1048576 -a 34 > $3/wave_put_tiled_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 128 -s 1048576 -a 34 > $3/wave_put_tiled_n2_w2_z128_1MB.log
check wave_put_tiled_n2_w1_z1_1MB
echo "wave_putnbi_tiled_n2_w1_z1_1MB"
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 128 -s 1048576 -a 35 > $3/wave_putnbi_tiled_n2_w1_z1_1MB.log
ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 128 -s 1048576 -a 35 > $3/wave_putnbi_tiled_n2_w2_z128_1MB.log
check wave_putnbi_tiled_n2_w1_z1_1MB
echo "amofadd_n2_w1_z1"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 1 -a 6 > $3/amofadd_n2_w1_z1.log
@@ -145,6 +145,54 @@ case $2 in
echo "putnbi_n2_w16_z128_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=16 mpirun -np 2 $1 -w 16 -z 128 -s 8 -a 3 > $3/putnbi_n2_w16_z128_8B.log
check putnbi_n2_w16_z128_8B
echo "wg_get_n2_w1_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 8 -a 28 > $3/wg_get_n2_w1_z64_8B.log
check wg_get_n2_w1_z64_8B
echo "wg_getnbi_n2_w1_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 8 -a 29 > $3/wg_getnbi_n2_w1_z64_8B.log
check wg_getnbi_n2_w1_z64_8B
echo "wg_put_n2_w1_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 8 -a 30 > $3/wg_put_n2_w1_z64_8B.log
check wg_put_n2_w1_z64_8B
echo "wg_putnbi_n2_w1_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 8 -a 31 > $3/wg_putnbi_n2_w1_z64_8B.log
check wg_putnbi_n2_w1_z64_8B
echo "wg_get_tiled_n2_w16_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=16 mpirun -np 2 $1 -w 16 -z 64 -s 8 -a 28 > $3/wg_get_tiled_n2_w16_z64_8B.log
check wg_get_tiled_n2_w16_z64_8B
echo "wg_getnbi_tiled_n2_w16_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=16 mpirun -np 2 $1 -w 16 -z 64 -s 8 -a 29 > $3/wg_getnbi_tiled_n2_w16_z64_8B.log
check wg_getnbi_tiled_n2_w16_z64_8B
echo "wg_put_tiled_n2_w16_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=16 mpirun -np 2 $1 -w 16 -z 64 -s 8 -a 30 > $3/wg_put_tiled_n2_w16_z64_8B.log
check wg_put_tiled_n2_w16_z64_8B
echo "wg_putnbi_tiled_n2_w16_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=16 mpirun -np 2 $1 -w 16 -z 64 -s 8 -a 31 > $3/wg_putnbi_tiled_n2_w16_z64_8B.log
check wg_putnbi_tiled_n2_w16_z64_8B
echo "wave_get_n2_w1_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 8 -a 32 > $3/wave_get_n2_w1_z64_8B.log
check wave_get_n2_w1_z64_8B
echo "wave_getnbi_n2_w1_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 8 -a 33 > $3/wave_getnbi_n2_w1_z64_8B.log
check wave_getnbi_n2_w1_z64_8B
echo "wave_put_n2_w1_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 8 -a 34 > $3/wave_put_n2_w1_z64_8B.log
check wave_put_n2_w1_z64_8B
echo "wave_putnbi_n2_w1_z64_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 64 -s 8 -a 35 > $3/wave_putnbi_n2_w1_z64_8B.log
check wave_putnbi_n2_w1_z64_8B
echo "wave_get_tiled_n2_w16_z128_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=16 mpirun -np 2 $1 -w 16 -z 128 -s 8 -a 32 > $3/wave_get_tiled_n2_w16_z128_8B.log
check wave_get_tiled_n2_w16_z128_8B
echo "wave_getnbi_tiled_n2_w16_z128_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=16 mpirun -np 2 $1 -w 16 -z 128 -s 8 -a 33 > $3/wave_getnbi_tiled_n2_w16_z128_8B.log
check wave_getnbi_tiled_n2_w16_z128_8B
echo "wave_put_tiled_n2_w16_z128_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=16 mpirun -np 2 $1 -w 16 -z 128 -s 8 -a 34 > $3/wave_put_tiled_n2_w16_z128_8B.log
check wave_put_tiled_n2_w16_z128_8B
echo "wave_putnbi_tiled_n2_w16_z128_8B"
ROC_SHMEM_MAX_NUM_CONTEXTS=16 mpirun -np 2 $1 -w 16 -z 128 -s 8 -a 35 > $3/wave_putnbi_tiled_n2_w16_z128_8B.log
check wave_putnbi_tiled_n2_w16_z128_8B
echo "amofadd_n2_w8_z1"
ROC_SHMEM_MAX_NUM_CONTEXTS=8 mpirun -np 2 $1 -w 8 -z 1 -a 6 > $3/amofadd_n2_w8_z1.log
check amofadd_n2_w8_z1
tests/functional_tests/extended_primitives.cpp
+18 -10
Wyświetl plik
@@ -24,10 +24,12 @@
#include <roc_shmem/roc_shmem.hpp>
#include <numeric>
using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNELS
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer,
char *s_buf, char *r_buf, int size,
@@ -37,6 +39,11 @@ __global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer,
roc_shmem_wg_init();
roc_shmem_wg_ctx_create(ctx_type, &ctx);
/**
* Calculate start index for each work group for tiled version
* If the number of work groups is greater than 1, this kernel performs a
* tiled functional test
*/
uint64_t start;
uint64_t idx = size * get_flat_grid_id();
s_buf += idx;
@@ -78,8 +85,8 @@ __global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer,
*****************************************************************************/
ExtendedPrimitiveTester::ExtendedPrimitiveTester(TesterArguments args)
: Tester(args) {
s_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.num_wgs);
r_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.num_wgs);
s_buf = static_cast<int*>(roc_shmem_malloc(args.max_msg_size * args.num_wgs));
r_buf = static_cast<int*>(roc_shmem_malloc(args.max_msg_size * args.num_wgs));
}
ExtendedPrimitiveTester::~ExtendedPrimitiveTester() {
@@ -88,8 +95,9 @@ ExtendedPrimitiveTester::~ExtendedPrimitiveTester() {
}
void ExtendedPrimitiveTester::resetBuffers(uint64_t size) {
memset(s_buf, '0', size * args.num_wgs);
memset(r_buf, '1', size * args.num_wgs);
num_elems = (size * args.num_wgs) / sizeof(int);
std::iota(s_buf, s_buf + num_elems, 0);
memset(r_buf, 0, size * args.num_wgs);
}
void ExtendedPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
@@ -97,8 +105,8 @@ void ExtendedPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
size_t shared_bytes = 0;
hipLaunchKernelGGL(ExtendedPrimitiveTest, gridSize, blockSize, shared_bytes,
stream, loop, args.skip, timer, s_buf, r_buf, size, _type,
_shmem_context);
stream, loop, args.skip, timer, (char*)s_buf,
(char*)r_buf, size, _type, _shmem_context);
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop * gridSize.x;
@@ -110,10 +118,10 @@ void ExtendedPrimitiveTester::verifyResults(uint64_t size) {
: 1;
if (args.myid == check_id) {
for (int i = 0; i < size * args.num_wgs; i++) {
if (r_buf[i] != '0') {
for (int i = 0; i < num_elems; i++) {
if (r_buf[i] != i) {
fprintf(stderr, "Data validation error at idx %d\n", i);
fprintf(stderr, "Got %c, Expected %c \n", r_buf[i], '0');
fprintf(stderr, "Got %d, Expected %d \n", r_buf[i], i);
exit(-1);
}
}
tests/functional_tests/extended_primitives.hpp
+3 -2
Wyświetl plik
@@ -42,8 +42,9 @@ class ExtendedPrimitiveTester : public Tester {
virtual void verifyResults(uint64_t size) override;
char *s_buf = nullptr;
char *r_buf = nullptr;
int *s_buf = nullptr;
int *r_buf = nullptr;
int num_elems = 0;
};
#endif
tests/functional_tests/tester.cpp
+37 -8
Wyświetl plik
@@ -59,7 +59,7 @@ Tester::Tester(TesterArguments args) : args(args) {
_shmem_context = args.shmem_context;
CHECK_HIP(hipGetDevice(&device_id));
CHECK_HIP(hipGetDeviceProperties(&deviceProps, device_id));
num_warps = args.wg_size / deviceProps.warpSize;
num_warps = (args.wg_size - 1) / deviceProps.warpSize + 1;
CHECK_HIP(hipStreamCreate(&stream));
CHECK_HIP(hipEventCreate(&start_event));
CHECK_HIP(hipEventCreate(&stop_event));
@@ -76,6 +76,11 @@ Tester::~Tester() {
std::vector<Tester*> Tester::create(TesterArguments args) {
int rank = args.myid;
std::vector<Tester*> testers;
hipDeviceProp_t deviceProps;
int device_id, numWarps;
CHECK_HIP(hipGetDevice(&device_id));
CHECK_HIP(hipGetDeviceProperties(&deviceProps, device_id));
numWarps = (args.wg_size - 1) / deviceProps.warpSize + 1;
if (rank == 0) std::cout << "*** Creating Test: ";
@@ -510,7 +515,7 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
return testers;
case WAVEGetTestType:
if (rank == 0) {
if (args.num_wgs > 1 || args.wg_size / 64 > 1)
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Blocking WAVE level Gets***" << std::endl;
else std::cout << "Blocking WAVE level Gets***" << std::endl;
}
@@ -518,7 +523,7 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
return testers;
case WAVEGetNBITestType:
if (rank == 0) {
if (args.num_wgs > 1 || args.wg_size / 64 > 1)
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Non-Blocking WAVE level Gets***" << std::endl;
else std::cout << "Non-Blocking WAVE level Gets***" << std::endl;
}
@@ -526,7 +531,7 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
return testers;
case WAVEPutTestType:
if (rank == 0) {
if (args.num_wgs > 1 || args.wg_size / 64 > 1)
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Blocking WAVE level Puts***" << std::endl;
else std::cout << "Blocking WAVE level Puts***" << std::endl;
}
@@ -534,7 +539,7 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
return testers;
case WAVEPutNBITestType:
if (rank == 0) {
if (args.num_wgs > 1 || args.wg_size / 64 > 1)
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Non-Blocking WAVE level Puts***" << std::endl;
else std::cout << "Non-Blocking WAVE level Puts***" << std::endl;
}
@@ -612,10 +617,32 @@ void Tester::execute() {
// data validation
verifyResults(size);
/**
* Adjust size for *_wg and *_wave functions
*/
uint64_t size_ = size;
TestType type = (TestType)args.algorithm;
switch (type) {
case WAVEGetTestType:
case WAVEGetNBITestType:
case WAVEPutTestType:
case WAVEPutNBITestType:
size_ *= (args.num_wgs * num_warps);
break;
case WGGetTestType:
case WGGetNBITestType:
case WGPutTestType:
case WGPutNBITestType:
size_ *= args.num_wgs;
break;
default:
break;
}
barrier();
if (_type != TeamCtxInfraTestType) {
print(size);
print(size_);
}
}
}
@@ -707,8 +734,10 @@ uint64_t Tester::gpuCyclesToMicroseconds(uint64_t cycles) {
uint64_t Tester::timerAvgInMicroseconds() {
uint64_t sum = 0;
//TODO: Modify the calcuation for the Tiled version of puts and gets at
// wavefront level (bpotter/avinash)
/**
* TODO: (bpotter/avinash) Modify the calcuation for the Tiled version of
* puts and gets at wavefront level
*/
for (int i = 0; i < args.num_wgs; i++) {
sum += gpuCyclesToMicroseconds(timer[i]);
}
tests/functional_tests/tester_arguments.cpp
+10
Wyświetl plik
@@ -103,6 +103,16 @@ TesterArguments::TesterArguments(int argc, char *argv[]) {
case PutNBIMRTestType:
min_msg_size = max_msg_size;
break;
case WAVEGetTestType:
case WAVEGetNBITestType:
case WAVEPutTestType:
case WAVEPutNBITestType:
case WGGetTestType:
case WGGetNBITestType:
case WGPutTestType:
case WGPutNBITestType:
min_msg_size = 4;
break;
default:
break;
}
tests/functional_tests/wave_level_primitives.cpp
+25 -16
Wyświetl plik
@@ -24,10 +24,12 @@
#include <roc_shmem/roc_shmem.hpp>
#include <numeric>
using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNELS
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer,
char *s_buf, char *r_buf, int size,
@@ -37,10 +39,15 @@ __global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer,
roc_shmem_wg_init();
roc_shmem_wg_ctx_create(ctx_type, &ctx);
/**
* Calculate start index for each wavefront for tiled version
* If the number of wavefronts is greater than 1, this kernel performs a
* tiled functional test
*/
uint64_t start;
int wf_id = get_flat_block_id() / wf_size;
int offset = size * get_flat_grid_id() * (get_flat_block_size() / wf_size);
int idx = wf_id * size + offset;
int wg_offset = size * get_flat_grid_id() * (get_flat_block_size() / wf_size);
int idx = wf_id * size + wg_offset;
s_buf += idx;
r_buf += idx;
@@ -80,10 +87,10 @@ __global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer,
*****************************************************************************/
WaveLevelPrimitiveTester::WaveLevelPrimitiveTester(TesterArguments args)
: Tester(args) {
s_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.num_wgs
* num_warps);
r_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.num_wgs
* num_warps);
s_buf = static_cast<int*>(
roc_shmem_malloc(args.max_msg_size * args.num_wgs * num_warps));
r_buf = static_cast<int*>(
roc_shmem_malloc(args.max_msg_size * args.num_wgs * num_warps));
}
WaveLevelPrimitiveTester::~WaveLevelPrimitiveTester() {
@@ -92,8 +99,9 @@ WaveLevelPrimitiveTester::~WaveLevelPrimitiveTester() {
}
void WaveLevelPrimitiveTester::resetBuffers(uint64_t size) {
memset(s_buf, '0', size * args.num_wgs * num_warps);
memset(r_buf, '1', size * args.num_wgs * num_warps);
num_elems = (size * args.num_wgs * num_warps) / sizeof(int);
std::iota(s_buf, s_buf + num_elems, 0);
memset(r_buf, 0, size * args.num_wgs * num_warps);
}
void WaveLevelPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
@@ -101,11 +109,12 @@ void WaveLevelPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
size_t shared_bytes = 0;
hipLaunchKernelGGL(WaveLevelPrimitiveTest, gridSize, blockSize, shared_bytes,
stream, loop, args.skip, timer, s_buf, r_buf, size, _type,
_shmem_context, deviceProps.warpSize);
stream, loop, args.skip, timer, (char*)s_buf,
(char*)r_buf, size, _type, _shmem_context,
deviceProps.warpSize);
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop * gridSize.x;
num_msgs = (loop + args.skip) * gridSize.x * num_warps;
num_timed_msgs = loop * gridSize.x * num_warps;
}
void WaveLevelPrimitiveTester::verifyResults(uint64_t size) {
@@ -114,10 +123,10 @@ void WaveLevelPrimitiveTester::verifyResults(uint64_t size) {
: 1;
if (args.myid == check_id) {
for (int i = 0; i < size * args.num_wgs * num_warps; i++) {
if (r_buf[i] != '0') {
for (int i = 0; i < num_elems; i++) {
if (r_buf[i] != i) {
fprintf(stderr, "Data validation error at idx %d\n", i);
fprintf(stderr, "Got %c, Expected %c \n", r_buf[i], '0');
fprintf(stderr, "Got %d, Expected %d \n", r_buf[i], i);
exit(-1);
}
}
tests/functional_tests/wave_level_primitives.hpp
+3 -2
Wyświetl plik
@@ -42,8 +42,9 @@ class WaveLevelPrimitiveTester : public Tester {
virtual void verifyResults(uint64_t size) override;
char *s_buf = nullptr;
char *r_buf = nullptr;
int *s_buf = nullptr;
int *r_buf = nullptr;
int num_elems = 0;
};
#endif