diff --git a/scripts/build_configs/ro_net b/scripts/build_configs/ro_net index 6abc528ca6..95ea2950a8 100755 --- a/scripts/build_configs/ro_net +++ b/scripts/build_configs/ro_net @@ -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 \ diff --git a/scripts/functional_tests/driver.sh b/scripts/functional_tests/driver.sh index f93045fae8..e0e0d45fa6 100755 --- a/scripts/functional_tests/driver.sh +++ b/scripts/functional_tests/driver.sh @@ -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 diff --git a/tests/functional_tests/extended_primitives.cpp b/tests/functional_tests/extended_primitives.cpp index cc73d9c502..d6c6042faa 100644 --- a/tests/functional_tests/extended_primitives.cpp +++ b/tests/functional_tests/extended_primitives.cpp @@ -24,10 +24,12 @@ #include +#include + 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(roc_shmem_malloc(args.max_msg_size * args.num_wgs)); + r_buf = static_cast(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); } } diff --git a/tests/functional_tests/extended_primitives.hpp b/tests/functional_tests/extended_primitives.hpp index 33fdac2f8f..76225f07b0 100644 --- a/tests/functional_tests/extended_primitives.hpp +++ b/tests/functional_tests/extended_primitives.hpp @@ -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 diff --git a/tests/functional_tests/tester.cpp b/tests/functional_tests/tester.cpp index cbd1ba89a5..e802c0b289 100644 --- a/tests/functional_tests/tester.cpp +++ b/tests/functional_tests/tester.cpp @@ -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::create(TesterArguments args) { int rank = args.myid; std::vector 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::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::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::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::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]); } diff --git a/tests/functional_tests/tester_arguments.cpp b/tests/functional_tests/tester_arguments.cpp index 88e4b12cd6..1167fafaeb 100644 --- a/tests/functional_tests/tester_arguments.cpp +++ b/tests/functional_tests/tester_arguments.cpp @@ -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; } diff --git a/tests/functional_tests/wave_level_primitives.cpp b/tests/functional_tests/wave_level_primitives.cpp index bb7c299ef9..56b263ebf4 100644 --- a/tests/functional_tests/wave_level_primitives.cpp +++ b/tests/functional_tests/wave_level_primitives.cpp @@ -24,10 +24,12 @@ #include +#include + 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( + roc_shmem_malloc(args.max_msg_size * args.num_wgs * num_warps)); + r_buf = static_cast( + 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); } } diff --git a/tests/functional_tests/wave_level_primitives.hpp b/tests/functional_tests/wave_level_primitives.hpp index 3b8a98541d..af65ac0a30 100644 --- a/tests/functional_tests/wave_level_primitives.hpp +++ b/tests/functional_tests/wave_level_primitives.hpp @@ -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