diff --git a/tests/functional_tests/amo_bitwise_tester.cpp b/tests/functional_tests/amo_bitwise_tester.cpp index 318386a756..446074b45a 100644 --- a/tests/functional_tests/amo_bitwise_tester.cpp +++ b/tests/functional_tests/amo_bitwise_tester.cpp @@ -55,7 +55,7 @@ AMOBitwiseTester::~AMOBitwiseTester() { } template -void AMOBitwiseTester::resetBuffers(uint64_t size) { +void AMOBitwiseTester::resetBuffers(size_t size) { memset(_r_buf, 0, args.max_msg_size); memset(_ret_val, 0, args.max_msg_size * args.num_wgs); memset(_s_buf, 0, args.max_msg_size * args.num_wgs); @@ -63,7 +63,7 @@ void AMOBitwiseTester::resetBuffers(uint64_t size) { template void AMOBitwiseTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(AMOBitwiseTest, gridsize, blocksize, shared_bytes, stream, @@ -82,7 +82,7 @@ void AMOBitwiseTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, #endif template -void AMOBitwiseTester::verifyResults(uint64_t size) { +void AMOBitwiseTester::verifyResults(size_t size) { T ret; if(DISABLE_IPC_TEST) { printf("AMO binary ops not implemented for IPC: values were not verified\n"); diff --git a/tests/functional_tests/amo_bitwise_tester.hpp b/tests/functional_tests/amo_bitwise_tester.hpp index f9818e4d03..ed627d60ce 100644 --- a/tests/functional_tests/amo_bitwise_tester.hpp +++ b/tests/functional_tests/amo_bitwise_tester.hpp @@ -37,12 +37,12 @@ class AMOBitwiseTester : public Tester { virtual ~AMOBitwiseTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; dim3 _gridSize{}; char *_r_buf; diff --git a/tests/functional_tests/amo_extended_tester.cpp b/tests/functional_tests/amo_extended_tester.cpp index 4fab12fcbe..15652048b4 100644 --- a/tests/functional_tests/amo_extended_tester.cpp +++ b/tests/functional_tests/amo_extended_tester.cpp @@ -55,7 +55,7 @@ AMOExtendedTester::~AMOExtendedTester() { } template -void AMOExtendedTester::resetBuffers(uint64_t size) { +void AMOExtendedTester::resetBuffers(size_t size) { memset(_r_buf, 0, args.max_msg_size); memset(_ret_val, 0, args.max_msg_size * args.num_wgs); memset(_s_buf, 0, args.max_msg_size * args.num_wgs); @@ -63,7 +63,7 @@ void AMOExtendedTester::resetBuffers(uint64_t size) { template void AMOExtendedTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(AMOExtendedTest, gridsize, blocksize, shared_bytes, stream, @@ -82,7 +82,7 @@ void AMOExtendedTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, #endif template -void AMOExtendedTester::verifyResults(uint64_t size) { +void AMOExtendedTester::verifyResults(size_t size) { T ret; if (args.myid == 0) { T expected_val = 0; diff --git a/tests/functional_tests/amo_extended_tester.hpp b/tests/functional_tests/amo_extended_tester.hpp index 5b593abbc0..3e76610307 100644 --- a/tests/functional_tests/amo_extended_tester.hpp +++ b/tests/functional_tests/amo_extended_tester.hpp @@ -37,12 +37,12 @@ class AMOExtendedTester : public Tester { virtual ~AMOExtendedTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; dim3 _gridSize{}; char *_r_buf; diff --git a/tests/functional_tests/amo_standard_tester.cpp b/tests/functional_tests/amo_standard_tester.cpp index 2f8c221337..028cb4447e 100644 --- a/tests/functional_tests/amo_standard_tester.cpp +++ b/tests/functional_tests/amo_standard_tester.cpp @@ -55,7 +55,7 @@ AMOStandardTester::~AMOStandardTester() { } template -void AMOStandardTester::resetBuffers(uint64_t size) { +void AMOStandardTester::resetBuffers(size_t size) { memset(_r_buf, 0, args.max_msg_size); memset(_ret_val, 0, args.max_msg_size * args.num_wgs); memset(_s_buf, 0, args.max_msg_size * args.num_wgs); @@ -63,7 +63,7 @@ void AMOStandardTester::resetBuffers(uint64_t size) { template void AMOStandardTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(AMOStandardTest, gridsize, blocksize, shared_bytes, stream, @@ -76,7 +76,7 @@ void AMOStandardTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, } template -void AMOStandardTester::verifyResults(uint64_t size) { +void AMOStandardTester::verifyResults(size_t size) { T ret; if (args.myid == 0) { T expected_val = 0; diff --git a/tests/functional_tests/amo_standard_tester.hpp b/tests/functional_tests/amo_standard_tester.hpp index b1dfe8316e..cec63882fc 100644 --- a/tests/functional_tests/amo_standard_tester.hpp +++ b/tests/functional_tests/amo_standard_tester.hpp @@ -37,12 +37,12 @@ class AMOStandardTester : public Tester { virtual ~AMOStandardTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; dim3 _gridSize{}; char *_r_buf; diff --git a/tests/functional_tests/barrier_all_tester.cpp b/tests/functional_tests/barrier_all_tester.cpp index ac112f4430..260d73319a 100644 --- a/tests/functional_tests/barrier_all_tester.cpp +++ b/tests/functional_tests/barrier_all_tester.cpp @@ -95,7 +95,7 @@ BarrierAllTester::BarrierAllTester(TesterArguments args) : Tester(args) {} BarrierAllTester::~BarrierAllTester() {} void BarrierAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(BarrierAllTest, gridSize, blockSize, shared_bytes, stream, @@ -105,6 +105,6 @@ void BarrierAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, num_timed_msgs = loop; } -void BarrierAllTester::resetBuffers(uint64_t size) {} +void BarrierAllTester::resetBuffers(size_t size) {} -void BarrierAllTester::verifyResults(uint64_t size) {} +void BarrierAllTester::verifyResults(size_t size) {} diff --git a/tests/functional_tests/barrier_all_tester.hpp b/tests/functional_tests/barrier_all_tester.hpp index bd515686f4..84e3bc1624 100644 --- a/tests/functional_tests/barrier_all_tester.hpp +++ b/tests/functional_tests/barrier_all_tester.hpp @@ -41,12 +41,12 @@ class BarrierAllTester : public Tester { virtual ~BarrierAllTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; }; #endif diff --git a/tests/functional_tests/default_ctx_primitive_tester.cpp b/tests/functional_tests/default_ctx_primitive_tester.cpp index 7e1424bfdd..e64ae844ca 100644 --- a/tests/functional_tests/default_ctx_primitive_tester.cpp +++ b/tests/functional_tests/default_ctx_primitive_tester.cpp @@ -34,7 +34,7 @@ __global__ void DefaultCTXPrimitiveTest(int loop, int skip, long long int *start_time, long long int *end_time, char *source, - char *dest, int size, TestType type, + char *dest, size_t size, TestType type, ShmemContextType ctx_type, int wf_size) { int wg_id = get_flat_grid_id(); int t_id = get_flat_block_id(); @@ -51,7 +51,7 @@ /** * Calculate start index for each thread within the grid */ - uint64_t offset = size * get_flat_id(); + size_t offset = size * get_flat_id(); source += offset; dest += offset; @@ -155,13 +155,13 @@ rocshmem_free(dest); } - void DefaultCTXPrimitiveTester::resetBuffers(uint64_t size) { + void DefaultCTXPrimitiveTester::resetBuffers(size_t size) { size_t buff_size = size * args.wg_size * args.num_wgs; memset(dest, '1', buff_size); } void DefaultCTXPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { + int loop, size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(DefaultCTXPrimitiveTest, gridSize, blockSize, @@ -173,7 +173,7 @@ num_timed_msgs = loop * gridSize.x * blockSize.x; } - void DefaultCTXPrimitiveTester::verifyResults(uint64_t size) { + void DefaultCTXPrimitiveTester::verifyResults(size_t size) { int check_id = (_type == DefaultCTXGetTestType || _type == DefaultCTXGetNBITestType || _type == DefaultCTXGTestType) @@ -182,7 +182,7 @@ if (args.myid == check_id) { size_t buff_size = size * args.wg_size * args.num_wgs; - for (uint64_t i = 0; i < buff_size; i++) { + for (size_t i = 0; i < buff_size; i++) { if (dest[i] != source[i]) { std::cerr << "Data validation error at idx " << i << std::endl; std::cerr << " Got " << dest[i] << ", Expected " diff --git a/tests/functional_tests/default_ctx_primitive_tester.hpp b/tests/functional_tests/default_ctx_primitive_tester.hpp index ae4de9b1c4..4b42d621c0 100644 --- a/tests/functional_tests/default_ctx_primitive_tester.hpp +++ b/tests/functional_tests/default_ctx_primitive_tester.hpp @@ -36,12 +36,12 @@ class DefaultCTXPrimitiveTester : public Tester { virtual ~DefaultCTXPrimitiveTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; char *source = nullptr; char *dest = nullptr; diff --git a/tests/functional_tests/empty_tester.cpp b/tests/functional_tests/empty_tester.cpp index 76376c578c..f244d27712 100644 --- a/tests/functional_tests/empty_tester.cpp +++ b/tests/functional_tests/empty_tester.cpp @@ -49,10 +49,10 @@ EmptyTester::EmptyTester(TesterArguments args) : Tester(args) {} EmptyTester::~EmptyTester() {} -void EmptyTester::resetBuffers(uint64_t size) {} +void EmptyTester::resetBuffers(size_t size) {} void EmptyTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(EmptyTest, gridSize, blockSize, shared_bytes, stream, @@ -60,4 +60,4 @@ void EmptyTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, _shmem_context); } -void EmptyTester::verifyResults(uint64_t size) {} +void EmptyTester::verifyResults(size_t size) {} diff --git a/tests/functional_tests/empty_tester.hpp b/tests/functional_tests/empty_tester.hpp index 894963dec8..1e7cdd93a3 100644 --- a/tests/functional_tests/empty_tester.hpp +++ b/tests/functional_tests/empty_tester.hpp @@ -36,12 +36,12 @@ class EmptyTester : public Tester { virtual ~EmptyTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; }; #endif // ROCSHMEM_CLIENTS_FUNCTIONAL_TESTS_EMPTY_TESTER_HPP diff --git a/tests/functional_tests/ping_all_tester.cpp b/tests/functional_tests/ping_all_tester.cpp index a89205c6dd..0bdcc0b840 100644 --- a/tests/functional_tests/ping_all_tester.cpp +++ b/tests/functional_tests/ping_all_tester.cpp @@ -77,13 +77,13 @@ PingAllTester::PingAllTester(TesterArguments args) : Tester(args) { PingAllTester::~PingAllTester() { rocshmem_free(r_buf); } -void PingAllTester::resetBuffers(uint64_t size) { +void PingAllTester::resetBuffers(size_t size) { int num_pes {rocshmem_n_pes()}; memset(r_buf, 0, sizeof(int) * args.num_wgs * num_pes); } void PingAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(PingAllTest, gridSize, blockSize, shared_bytes, stream, @@ -94,4 +94,4 @@ void PingAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, num_timed_msgs = loop; } -void PingAllTester::verifyResults(uint64_t size) {} +void PingAllTester::verifyResults(size_t size) {} diff --git a/tests/functional_tests/ping_all_tester.hpp b/tests/functional_tests/ping_all_tester.hpp index 64d2494517..8edf0ccd32 100644 --- a/tests/functional_tests/ping_all_tester.hpp +++ b/tests/functional_tests/ping_all_tester.hpp @@ -42,12 +42,12 @@ class PingAllTester : public Tester { virtual ~PingAllTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; int *r_buf; }; diff --git a/tests/functional_tests/ping_pong_tester.cpp b/tests/functional_tests/ping_pong_tester.cpp index 0fb1fc12b8..c3e9ead963 100644 --- a/tests/functional_tests/ping_pong_tester.cpp +++ b/tests/functional_tests/ping_pong_tester.cpp @@ -77,12 +77,12 @@ PingPongTester::PingPongTester(TesterArguments args) : Tester(args) { PingPongTester::~PingPongTester() { rocshmem_free(r_buf); } -void PingPongTester::resetBuffers(uint64_t size) { +void PingPongTester::resetBuffers(size_t size) { memset(r_buf, 0, sizeof(int) * args.num_wgs); } void PingPongTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(PingPongTest, gridSize, blockSize, shared_bytes, stream, @@ -93,4 +93,4 @@ void PingPongTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, num_timed_msgs = loop; } -void PingPongTester::verifyResults(uint64_t size) {} +void PingPongTester::verifyResults(size_t size) {} diff --git a/tests/functional_tests/ping_pong_tester.hpp b/tests/functional_tests/ping_pong_tester.hpp index f00cc0ec81..03ce056370 100644 --- a/tests/functional_tests/ping_pong_tester.hpp +++ b/tests/functional_tests/ping_pong_tester.hpp @@ -42,12 +42,12 @@ class PingPongTester : public Tester { virtual ~PingPongTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; int *r_buf; }; diff --git a/tests/functional_tests/primitive_mr_tester.cpp b/tests/functional_tests/primitive_mr_tester.cpp index dadc0a70ce..22483f73b7 100644 --- a/tests/functional_tests/primitive_mr_tester.cpp +++ b/tests/functional_tests/primitive_mr_tester.cpp @@ -33,7 +33,7 @@ using namespace rocshmem; *****************************************************************************/ __global__ void PrimitiveMRTest(int loop, long long int *start_time, long long int *end_time, char *s_buf, - char *r_buf, int size, + char *r_buf, size_t size, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); @@ -79,7 +79,7 @@ void PrimitiveMRTester::resetBuffers(size_t size) { } void PrimitiveMRTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; /* Warmup */ @@ -98,7 +98,7 @@ void PrimitiveMRTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, num_timed_msgs = loop * 64; } -void PrimitiveMRTester::verifyResults(uint64_t size) { +void PrimitiveMRTester::verifyResults(size_t size) { int check_id = (_type == GetTestType || _type == GetNBITestType || _type == GTestType) ? 0 diff --git a/tests/functional_tests/primitive_mr_tester.hpp b/tests/functional_tests/primitive_mr_tester.hpp index d341bc1ece..7e66b3d3f9 100644 --- a/tests/functional_tests/primitive_mr_tester.hpp +++ b/tests/functional_tests/primitive_mr_tester.hpp @@ -39,9 +39,9 @@ class PrimitiveMRTester : public Tester { virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; char *s_buf = nullptr; char *r_buf = nullptr; diff --git a/tests/functional_tests/primitive_tester.cpp b/tests/functional_tests/primitive_tester.cpp index adecd00dfd..bdba1a95b5 100644 --- a/tests/functional_tests/primitive_tester.cpp +++ b/tests/functional_tests/primitive_tester.cpp @@ -33,7 +33,7 @@ using namespace rocshmem; *****************************************************************************/ __global__ void PrimitiveTest(int loop, int skip, long long int *start_time, long long int *end_time, char *source, - char *dest, int size, TestType type, + char *dest, size_t size, TestType type, ShmemContextType ctx_type, int wf_size) { __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); @@ -52,7 +52,7 @@ __global__ void PrimitiveTest(int loop, int skip, long long int *start_time, /** * Calculate start index for each thread within the grid */ - uint64_t offset = size * get_flat_id(); + size_t offset = size * get_flat_id(); source += offset; dest += offset; @@ -156,13 +156,13 @@ PrimitiveTester::~PrimitiveTester() { rocshmem_free(dest); } -void PrimitiveTester::resetBuffers(uint64_t size) { +void PrimitiveTester::resetBuffers(size_t size) { size_t buff_size = size * args.wg_size * args.num_wgs; memset(dest, '1', buff_size); } void PrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(PrimitiveTest, gridSize, blockSize, shared_bytes, stream, @@ -173,7 +173,7 @@ void PrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, num_timed_msgs = loop * gridSize.x * blockSize.x; } -void PrimitiveTester::verifyResults(uint64_t size) { +void PrimitiveTester::verifyResults(size_t size) { int check_id = (_type == GetTestType || _type == GetNBITestType || _type == GTestType) ? 0 diff --git a/tests/functional_tests/primitive_tester.hpp b/tests/functional_tests/primitive_tester.hpp index f7d7b0a7af..a76b0abb16 100644 --- a/tests/functional_tests/primitive_tester.hpp +++ b/tests/functional_tests/primitive_tester.hpp @@ -36,12 +36,12 @@ class PrimitiveTester : public Tester { virtual ~PrimitiveTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; char *source = nullptr; char *dest = nullptr; diff --git a/tests/functional_tests/random_access_tester.cpp b/tests/functional_tests/random_access_tester.cpp index 0467c30b95..b195259e76 100644 --- a/tests/functional_tests/random_access_tester.cpp +++ b/tests/functional_tests/random_access_tester.cpp @@ -35,7 +35,8 @@ using namespace rocshmem; __device__ bool thread_passing(int num_bins, uint32_t *bin_threads, uint32_t *off_bins, uint32_t *PE_bins, - int *offset, int *PE, int coal_coef, int size) { + size_t *offset, int *PE, int coal_coef, + size_t size) { bool pass = false; int wave_id = ((hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x) / 64); // get_global_wave_id(); @@ -55,7 +56,7 @@ __device__ bool thread_passing(int num_bins, uint32_t *bin_threads, __global__ void RandomAccessTest(int loop, int skip, long long int *start_time, long long int *end_time, int *s_buf, - int *r_buf, int size, OpType type, + int *r_buf, size_t size, OpType type, int coal_coef, int num_bins, int num_waves, uint32_t *threads_bins, uint32_t *off_bins, uint32_t *PE_bins, ShmemContextType ctx_type) { @@ -65,7 +66,7 @@ __global__ void RandomAccessTest(int loop, int skip, long long int *start_time, rocshmem_wg_ctx_create(ctx_type, &ctx); int pe = rocshmem_ctx_my_pe(ctx); - int offset; + size_t offset; int PE; if (thread_passing(num_bins, threads_bins, off_bins, PE_bins, &offset, &PE, @@ -166,7 +167,7 @@ RandomAccessTester::~RandomAccessTester() { CHECK_HIP(hipFree(_PE_bins)); } -void RandomAccessTester::resetBuffers(uint64_t size) { +void RandomAccessTester::resetBuffers(size_t size) { for (size_t i = 0; i < args.max_msg_size / sizeof(int) * args.wg_size * space; i++) { s_buf[i] = 1; @@ -176,7 +177,7 @@ void RandomAccessTester::resetBuffers(uint64_t size) { } void RandomAccessTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; int _thread_access = args.thread_access; @@ -200,7 +201,7 @@ void RandomAccessTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, num_timed_msgs = loop * _num_waves * _thread_access; } -void RandomAccessTester::verifyResults(uint64_t size) { +void RandomAccessTester::verifyResults(size_t size) { uint64_t offset; for (int k = 0; k < _num_waves; k++) { for (int i = 0; i < _num_bins; i++) { diff --git a/tests/functional_tests/random_access_tester.hpp b/tests/functional_tests/random_access_tester.hpp index 1d011eef3d..d52c7bec27 100644 --- a/tests/functional_tests/random_access_tester.hpp +++ b/tests/functional_tests/random_access_tester.hpp @@ -46,12 +46,12 @@ class RandomAccessTester : public Tester { virtual ~RandomAccessTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; int *r_buf; int *s_buf; diff --git a/tests/functional_tests/shmem_ptr_tester.cpp b/tests/functional_tests/shmem_ptr_tester.cpp index 4dde64c094..019493829b 100644 --- a/tests/functional_tests/shmem_ptr_tester.cpp +++ b/tests/functional_tests/shmem_ptr_tester.cpp @@ -59,13 +59,13 @@ ShmemPtrTester::~ShmemPtrTester() { rocshmem_free(r_buf); } -void ShmemPtrTester::resetBuffers(uint64_t size) { +void ShmemPtrTester::resetBuffers(size_t size) { memset(r_buf, '0', args.max_msg_size); memset(_available, 0, sizeof(int)); } void ShmemPtrTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(ShmemPtrTest, gridSize, blockSize, shared_bytes, stream, @@ -75,7 +75,7 @@ void ShmemPtrTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, num_timed_msgs = 0; } -void ShmemPtrTester::verifyResults(uint64_t size) { +void ShmemPtrTester::verifyResults(size_t size) { if (args.myid == 0) { if (*_available == 0) { fprintf(stderr, "SHMEM_PTR NOT AVAILBLE \n"); diff --git a/tests/functional_tests/shmem_ptr_tester.hpp b/tests/functional_tests/shmem_ptr_tester.hpp index 8a1f44c228..89fa573814 100644 --- a/tests/functional_tests/shmem_ptr_tester.hpp +++ b/tests/functional_tests/shmem_ptr_tester.hpp @@ -36,12 +36,12 @@ class ShmemPtrTester : public Tester { virtual ~ShmemPtrTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; char *r_buf = nullptr; int *_available = nullptr; diff --git a/tests/functional_tests/signaling_operations_tester.cpp b/tests/functional_tests/signaling_operations_tester.cpp index 1028f54f6b..ff37945cf0 100644 --- a/tests/functional_tests/signaling_operations_tester.cpp +++ b/tests/functional_tests/signaling_operations_tester.cpp @@ -33,8 +33,9 @@ using namespace rocshmem; *****************************************************************************/ __global__ void PutmemSignalTest(int loop, int skip, long long int *start_time, long long int *end_time, char *s_buf, - char *r_buf, int size, uint64_t *sig_addr, - TestType type, ShmemContextType ctx_type, int sig_op) { + char *r_buf, size_t size, uint64_t *sig_addr, + TestType type, ShmemContextType ctx_type, + int sig_op) { __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); rocshmem_wg_init(); @@ -50,22 +51,28 @@ __global__ void PutmemSignalTest(int loop, int skip, long long int *start_time, switch (type) { case PutSignalTestType: - rocshmem_ctx_putmem_signal(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + rocshmem_ctx_putmem_signal(ctx, r_buf, s_buf, size, sig_addr, + signal, sig_op, 1); break; case WGPutSignalTestType: - rocshmem_ctx_putmem_signal_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + rocshmem_ctx_putmem_signal_wg(ctx, r_buf, s_buf, size, sig_addr, + signal, sig_op, 1); break; case WAVEPutSignalTestType: - rocshmem_ctx_putmem_signal_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + rocshmem_ctx_putmem_signal_wave(ctx, r_buf, s_buf, size, sig_addr, + signal, sig_op, 1); break; case PutSignalNBITestType: - rocshmem_ctx_putmem_signal_nbi(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + rocshmem_ctx_putmem_signal_nbi(ctx, r_buf, s_buf, size, sig_addr, + signal, sig_op, 1); break; case WGPutSignalNBITestType: - rocshmem_ctx_putmem_signal_nbi_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + rocshmem_ctx_putmem_signal_nbi_wg(ctx, r_buf, s_buf, size, sig_addr, + signal, sig_op, 1); break; case WAVEPutSignalNBITestType: - rocshmem_ctx_putmem_signal_nbi_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + rocshmem_ctx_putmem_signal_nbi_wave(ctx, r_buf, s_buf, size, sig_addr, + signal, sig_op, 1); break; default: break; @@ -124,15 +131,17 @@ __global__ void SignalFetchTest(int loop, int skip, long long int *start_time, /****************************************************************************** * HOST TESTER CLASS METHODS *****************************************************************************/ -SignalingOperationsTester::SignalingOperationsTester(TesterArguments args) : Tester(args) { +SignalingOperationsTester::SignalingOperationsTester(TesterArguments args) + : Tester(args) { s_buf = (char *)rocshmem_malloc(args.max_msg_size * args.wg_size); r_buf = (char *)rocshmem_malloc(args.max_msg_size * args.wg_size); sig_addr = (uint64_t *)rocshmem_malloc(sizeof(uint64_t)); CHECK_HIP(hipMallocManaged(&fetched_value, sizeof(uint64_t), hipMemAttachHost)); } -SignalingOperationsTester::SignalingOperationsTester(TesterArguments args, int signal_op) - : SignalingOperationsTester(args) { +SignalingOperationsTester::SignalingOperationsTester(TesterArguments args, + int signal_op) + : SignalingOperationsTester(args) { sig_op = signal_op; } @@ -143,7 +152,7 @@ SignalingOperationsTester::~SignalingOperationsTester() { CHECK_HIP(hipFree(fetched_value)); } -void SignalingOperationsTester::resetBuffers(uint64_t size) { +void SignalingOperationsTester::resetBuffers(size_t size) { memset(s_buf, '0', args.max_msg_size * args.wg_size); memset(r_buf, '1', args.max_msg_size * args.wg_size); *fetched_value = -1; @@ -151,7 +160,7 @@ void SignalingOperationsTester::resetBuffers(uint64_t size) { } void SignalingOperationsTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; @@ -170,7 +179,7 @@ void SignalingOperationsTester::launchKernel(dim3 gridSize, dim3 blockSize, int num_timed_msgs = loop; } -void SignalingOperationsTester::verifyResults(uint64_t size) { +void SignalingOperationsTester::verifyResults(size_t size) { if (_type == SignalFetchTestType || _type == WAVESignalFetchTestType || _type == WGSignalFetchTestType) { diff --git a/tests/functional_tests/signaling_operations_tester.hpp b/tests/functional_tests/signaling_operations_tester.hpp index 6412c45c7d..d2772773da 100644 --- a/tests/functional_tests/signaling_operations_tester.hpp +++ b/tests/functional_tests/signaling_operations_tester.hpp @@ -37,12 +37,12 @@ class SignalingOperationsTester : public Tester { virtual ~SignalingOperationsTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; int sig_op; char *s_buf = nullptr; diff --git a/tests/functional_tests/sync_all_tester.cpp b/tests/functional_tests/sync_all_tester.cpp index 902787b313..df26626041 100644 --- a/tests/functional_tests/sync_all_tester.cpp +++ b/tests/functional_tests/sync_all_tester.cpp @@ -95,7 +95,7 @@ SyncAllTester::SyncAllTester(TesterArguments args) : Tester(args) {} SyncAllTester::~SyncAllTester() {} void SyncAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(SyncAllTest, gridSize, blockSize, shared_bytes, stream, @@ -105,6 +105,6 @@ void SyncAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, num_timed_msgs = loop; } -void SyncAllTester::resetBuffers(uint64_t size) {} +void SyncAllTester::resetBuffers(size_t size) {} -void SyncAllTester::verifyResults(uint64_t size) {} +void SyncAllTester::verifyResults(size_t size) {} diff --git a/tests/functional_tests/sync_all_tester.hpp b/tests/functional_tests/sync_all_tester.hpp index d42ca4cd43..dd1cec53f2 100644 --- a/tests/functional_tests/sync_all_tester.hpp +++ b/tests/functional_tests/sync_all_tester.hpp @@ -41,12 +41,12 @@ class SyncAllTester : public Tester { virtual ~SyncAllTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; }; #endif diff --git a/tests/functional_tests/sync_tester.cpp b/tests/functional_tests/sync_tester.cpp index ca021784e5..9bd9bd6f0a 100644 --- a/tests/functional_tests/sync_tester.cpp +++ b/tests/functional_tests/sync_tester.cpp @@ -90,7 +90,7 @@ SyncTester::~SyncTester() { CHECK_HIP(hipFree(team_sync_world_dup)); } -void SyncTester::resetBuffers(uint64_t size) {} +void SyncTester::resetBuffers(size_t size) {} void SyncTester::preLaunchKernel() { int n_pes = rocshmem_team_n_pes(ROCSHMEM_TEAM_WORLD); @@ -107,7 +107,7 @@ void SyncTester::preLaunchKernel() { } void SyncTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; int n_pes = rocshmem_team_n_pes(ROCSHMEM_TEAM_WORLD); @@ -126,4 +126,4 @@ void SyncTester::postLaunchKernel() { } } -void SyncTester::verifyResults(uint64_t size) {} +void SyncTester::verifyResults(size_t size) {} diff --git a/tests/functional_tests/sync_tester.hpp b/tests/functional_tests/sync_tester.hpp index 853d2769b7..6da8e1fc95 100644 --- a/tests/functional_tests/sync_tester.hpp +++ b/tests/functional_tests/sync_tester.hpp @@ -40,16 +40,16 @@ class SyncTester : public Tester { virtual ~SyncTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void preLaunchKernel() override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; virtual void postLaunchKernel() override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; private: /** diff --git a/tests/functional_tests/team_alltoall_tester.cpp b/tests/functional_tests/team_alltoall_tester.cpp index 6c89d672af..6c11859a81 100644 --- a/tests/functional_tests/team_alltoall_tester.cpp +++ b/tests/functional_tests/team_alltoall_tester.cpp @@ -153,7 +153,7 @@ void TeamAlltoallTester::preLaunchKernel() { template void TeamAlltoallTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { + int loop, size_t size) { size_t shared_bytes = 0; int num_elems = size / sizeof(T1); @@ -175,7 +175,7 @@ void TeamAlltoallTester::postLaunchKernel() { } template -void TeamAlltoallTester::resetBuffers(uint64_t size) { +void TeamAlltoallTester::resetBuffers(size_t size) { int num_elems = size / sizeof(T1); int buff_size = num_elems * sizeof(T1) * args.num_wgs * n_pes; @@ -204,7 +204,7 @@ void TeamAlltoallTester::resetBuffers(uint64_t size) { } template -void TeamAlltoallTester::verifyResults(uint64_t size) { +void TeamAlltoallTester::verifyResults(size_t size) { int num_elems = size / sizeof(T1); int idx = 0; diff --git a/tests/functional_tests/team_alltoall_tester.hpp b/tests/functional_tests/team_alltoall_tester.hpp index c30f41deea..98368c7bea 100644 --- a/tests/functional_tests/team_alltoall_tester.hpp +++ b/tests/functional_tests/team_alltoall_tester.hpp @@ -42,16 +42,16 @@ class TeamAlltoallTester : public Tester { virtual ~TeamAlltoallTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void preLaunchKernel() override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; virtual void postLaunchKernel() override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; T1 *source_buf = nullptr; T1 *dest_buf = nullptr; diff --git a/tests/functional_tests/team_barrier_tester.cpp b/tests/functional_tests/team_barrier_tester.cpp index 7a1d16e58f..bd0430537d 100644 --- a/tests/functional_tests/team_barrier_tester.cpp +++ b/tests/functional_tests/team_barrier_tester.cpp @@ -106,7 +106,7 @@ void TeamBarrierTester::preLaunchKernel() { } void TeamBarrierTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { + int loop, size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(TeamBarrierTest, gridSize, blockSize, shared_bytes, @@ -124,6 +124,6 @@ void TeamBarrierTester::postLaunchKernel() { } } -void TeamBarrierTester::resetBuffers(uint64_t size) {} +void TeamBarrierTester::resetBuffers(size_t size) {} -void TeamBarrierTester::verifyResults(uint64_t size) {} +void TeamBarrierTester::verifyResults(size_t size) {} diff --git a/tests/functional_tests/team_barrier_tester.hpp b/tests/functional_tests/team_barrier_tester.hpp index 3fe95ee4d1..ea8052ddfb 100644 --- a/tests/functional_tests/team_barrier_tester.hpp +++ b/tests/functional_tests/team_barrier_tester.hpp @@ -41,16 +41,16 @@ class TeamBarrierTester : public Tester { virtual ~TeamBarrierTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void preLaunchKernel() override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; virtual void postLaunchKernel() override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; private: int my_pe = 0; diff --git a/tests/functional_tests/team_broadcast_tester.cpp b/tests/functional_tests/team_broadcast_tester.cpp index a155afb74d..393916a814 100644 --- a/tests/functional_tests/team_broadcast_tester.cpp +++ b/tests/functional_tests/team_broadcast_tester.cpp @@ -153,7 +153,7 @@ void TeamBroadcastTester::preLaunchKernel() { template void TeamBroadcastTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { + int loop, size_t size) { size_t shared_bytes = 0; int num_elems = size / sizeof(T1); @@ -175,7 +175,7 @@ void TeamBroadcastTester::postLaunchKernel() { } template -void TeamBroadcastTester::resetBuffers(uint64_t size) { +void TeamBroadcastTester::resetBuffers(size_t size) { int num_elems = size / sizeof(T1); int buff_size = num_elems * sizeof(T1) * args.num_wgs; @@ -203,7 +203,7 @@ void TeamBroadcastTester::resetBuffers(uint64_t size) { } template -void TeamBroadcastTester::verifyResults(uint64_t size) { +void TeamBroadcastTester::verifyResults(size_t size) { int num_elems = size / sizeof(T1); int idx = 0; diff --git a/tests/functional_tests/team_broadcast_tester.hpp b/tests/functional_tests/team_broadcast_tester.hpp index c4224891b2..ce607a5c46 100644 --- a/tests/functional_tests/team_broadcast_tester.hpp +++ b/tests/functional_tests/team_broadcast_tester.hpp @@ -42,16 +42,16 @@ class TeamBroadcastTester : public Tester { virtual ~TeamBroadcastTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void preLaunchKernel() override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; virtual void postLaunchKernel() override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; T1 *source_buf; T1 *dest_buf; diff --git a/tests/functional_tests/team_ctx_infra_tester.cpp b/tests/functional_tests/team_ctx_infra_tester.cpp index d1847607bb..c1af21a4c4 100644 --- a/tests/functional_tests/team_ctx_infra_tester.cpp +++ b/tests/functional_tests/team_ctx_infra_tester.cpp @@ -113,7 +113,7 @@ TeamCtxInfraTester::TeamCtxInfraTester(TesterArguments args) : Tester(args) {} TeamCtxInfraTester::~TeamCtxInfraTester() {} -void TeamCtxInfraTester::resetBuffers(uint64_t size) {} +void TeamCtxInfraTester::resetBuffers(size_t size) {} void TeamCtxInfraTester::preLaunchKernel() { int n_pes = rocshmem_team_n_pes(ROCSHMEM_TEAM_WORLD); @@ -149,7 +149,7 @@ void TeamCtxInfraTester::preLaunchKernel() { } void TeamCtxInfraTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { + size_t size) { size_t shared_bytes = 0; /* Copy array of teams to device */ @@ -170,4 +170,4 @@ void TeamCtxInfraTester::postLaunchKernel() { } } -void TeamCtxInfraTester::verifyResults(uint64_t size) {} +void TeamCtxInfraTester::verifyResults(size_t size) {} diff --git a/tests/functional_tests/team_ctx_infra_tester.hpp b/tests/functional_tests/team_ctx_infra_tester.hpp index b1e3f5760e..2b108ad980 100644 --- a/tests/functional_tests/team_ctx_infra_tester.hpp +++ b/tests/functional_tests/team_ctx_infra_tester.hpp @@ -36,16 +36,16 @@ class TeamCtxInfraTester : public Tester { virtual ~TeamCtxInfraTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void preLaunchKernel() override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; virtual void postLaunchKernel() override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; char *s_buf = nullptr; char *r_buf = nullptr; diff --git a/tests/functional_tests/team_ctx_primitive_tester.cpp b/tests/functional_tests/team_ctx_primitive_tester.cpp index e5caeb9135..8e5e10ea97 100644 --- a/tests/functional_tests/team_ctx_primitive_tester.cpp +++ b/tests/functional_tests/team_ctx_primitive_tester.cpp @@ -35,7 +35,7 @@ rocshmem_team_t team_primitive_world_dup; *****************************************************************************/ __global__ void TeamCtxPrimitiveTest(int loop, int skip, long long int *start_time, long long int *end_time, char *source, - char *dest, int size, TestType type, + char *dest, size_t size, TestType type, ShmemContextType ctx_type, int wf_size, rocshmem_team_t team) { __shared__ rocshmem_ctx_t ctx; @@ -56,7 +56,7 @@ __global__ void TeamCtxPrimitiveTest(int loop, int skip, long long int *start_ti /** * Calculate start index for each thread within the grid */ - uint64_t offset = size * get_flat_id(); + size_t offset = size * get_flat_id(); source += offset; dest += offset; @@ -148,7 +148,7 @@ TeamCtxPrimitiveTester::~TeamCtxPrimitiveTester() { rocshmem_free(dest); } -void TeamCtxPrimitiveTester::resetBuffers(uint64_t size) { +void TeamCtxPrimitiveTester::resetBuffers(size_t size) { size_t buff_size = size * args.wg_size * args.num_wgs; memset(dest, '1', buff_size); } @@ -162,7 +162,7 @@ void TeamCtxPrimitiveTester::preLaunchKernel() { } void TeamCtxPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { + int loop, size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(TeamCtxPrimitiveTest, gridSize, blockSize, shared_bytes, @@ -178,7 +178,7 @@ void TeamCtxPrimitiveTester::postLaunchKernel() { rocshmem_team_destroy(team_primitive_world_dup); } -void TeamCtxPrimitiveTester::verifyResults(uint64_t size) { +void TeamCtxPrimitiveTester::verifyResults(size_t size) { int check_id = (_type == TeamCtxGetTestType || _type == TeamCtxGetNBITestType) ? 0 : 1; diff --git a/tests/functional_tests/team_ctx_primitive_tester.hpp b/tests/functional_tests/team_ctx_primitive_tester.hpp index 8d1a9681f6..1e30aee7c8 100644 --- a/tests/functional_tests/team_ctx_primitive_tester.hpp +++ b/tests/functional_tests/team_ctx_primitive_tester.hpp @@ -36,16 +36,16 @@ class TeamCtxPrimitiveTester : public Tester { virtual ~TeamCtxPrimitiveTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void preLaunchKernel() override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; virtual void postLaunchKernel() override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; char *source = nullptr; char *dest = nullptr; diff --git a/tests/functional_tests/team_fcollect_tester.cpp b/tests/functional_tests/team_fcollect_tester.cpp index 353d8f3142..244556e25d 100644 --- a/tests/functional_tests/team_fcollect_tester.cpp +++ b/tests/functional_tests/team_fcollect_tester.cpp @@ -168,7 +168,7 @@ void TeamFcollectTester::preLaunchKernel() { template void TeamFcollectTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { + int loop, size_t size) { size_t shared_bytes = 0; int num_elems = size / sizeof(T1); @@ -193,7 +193,7 @@ void TeamFcollectTester::postLaunchKernel() { } template -void TeamFcollectTester::resetBuffers(uint64_t size) { +void TeamFcollectTester::resetBuffers(size_t size) { int num_elems = (size / sizeof(T1)); int buff_size = num_elems * sizeof(T1) * args.num_wgs * n_pes; @@ -201,7 +201,7 @@ void TeamFcollectTester::resetBuffers(uint64_t size) { } template -void TeamFcollectTester::verifyResults(uint64_t size) { +void TeamFcollectTester::verifyResults(size_t size) { int num_elems = size / sizeof(T1); int idx = 0; diff --git a/tests/functional_tests/team_fcollect_tester.hpp b/tests/functional_tests/team_fcollect_tester.hpp index 2d51d9b24f..83b5107d9d 100644 --- a/tests/functional_tests/team_fcollect_tester.hpp +++ b/tests/functional_tests/team_fcollect_tester.hpp @@ -42,16 +42,16 @@ class TeamFcollectTester : public Tester { virtual ~TeamFcollectTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void preLaunchKernel() override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; virtual void postLaunchKernel() override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; T1 *source_buf; T1 *dest_buf; diff --git a/tests/functional_tests/team_reduction_tester.cpp b/tests/functional_tests/team_reduction_tester.cpp index cd88b9f803..a826678fd2 100644 --- a/tests/functional_tests/team_reduction_tester.cpp +++ b/tests/functional_tests/team_reduction_tester.cpp @@ -77,7 +77,7 @@ rocshmem_team_t team_reduce_world_dup; template __global__ void TeamReductionTest(int loop, int skip, long long int *start_time, long long int *end_time, T1 *s_buf, T1 *r_buf, - int size, TestType type, + size_t size, TestType type, ShmemContextType ctx_type, rocshmem_team_t team) { __shared__ rocshmem_ctx_t ctx; @@ -136,7 +136,7 @@ void TeamReductionTester::preLaunchKernel() { template void TeamReductionTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { + int loop, size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(HIP_KERNEL_NAME(TeamReductionTest), gridSize, @@ -154,14 +154,14 @@ void TeamReductionTester::postLaunchKernel() { } template -void TeamReductionTester::resetBuffers(uint64_t size) { +void TeamReductionTester::resetBuffers(size_t size) { for (uint64_t i = 0; i < args.max_msg_size; i++) { init_buf(s_buf[i], r_buf[i]); } } template -void TeamReductionTester::verifyResults(uint64_t size) { +void TeamReductionTester::verifyResults(size_t size) { int n_pes = rocshmem_n_pes(); for (uint64_t i = 0; i < size; i++) { auto r = verify_buf(r_buf[i], (T1)n_pes); diff --git a/tests/functional_tests/team_reduction_tester.hpp b/tests/functional_tests/team_reduction_tester.hpp index 6cbb1ea414..8466854b10 100644 --- a/tests/functional_tests/team_reduction_tester.hpp +++ b/tests/functional_tests/team_reduction_tester.hpp @@ -42,16 +42,16 @@ class TeamReductionTester : public Tester { virtual ~TeamReductionTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void preLaunchKernel() override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; virtual void postLaunchKernel() override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; T1 *s_buf; T1 *r_buf; diff --git a/tests/functional_tests/tester.cpp b/tests/functional_tests/tester.cpp index 2661cce5a0..8a8ee222a9 100644 --- a/tests/functional_tests/tester.cpp +++ b/tests/functional_tests/tester.cpp @@ -478,7 +478,7 @@ void Tester::execute() { * Some tests loop through data sizes in powers of 2 and report the * results for those ranges. */ - for (uint64_t size = args.min_msg_size; size <= args.max_msg_size; + for (size_t size = args.min_msg_size; size <= args.max_msg_size; size <<= 1) { resetBuffers(size); diff --git a/tests/functional_tests/tester_arguments.hpp b/tests/functional_tests/tester_arguments.hpp index dc66663f73..61478fb1e6 100644 --- a/tests/functional_tests/tester_arguments.hpp +++ b/tests/functional_tests/tester_arguments.hpp @@ -54,8 +54,8 @@ class TesterArguments { unsigned num_wgs = 1; unsigned num_threads = 1; unsigned algorithm = 0; - uint64_t min_msg_size = 1; - uint64_t max_msg_size = 1 << 20; + size_t min_msg_size = 1; + size_t max_msg_size = 1 << 20; unsigned wg_size = 64; unsigned thread_access = 64; unsigned coal_coef = 64; @@ -74,7 +74,7 @@ class TesterArguments { int loop = 10; int skip = 10; int loop_large = 10; - uint64_t large_message_size = 32768; + size_t large_message_size = 32768; }; #endif diff --git a/tests/functional_tests/wavefront_primitives.cpp b/tests/functional_tests/wavefront_primitives.cpp index 4a208ffc40..be18c1e822 100644 --- a/tests/functional_tests/wavefront_primitives.cpp +++ b/tests/functional_tests/wavefront_primitives.cpp @@ -36,7 +36,7 @@ using namespace rocshmem; __global__ void WaveFrontPrimitiveTest(int loop, int skip, long long int *start_time, long long int *end_time, char *source, - char *dest, int size, TestType type, + char *dest, size_t size, TestType type, ShmemContextType ctx_type, int wf_size) { __shared__ rocshmem_ctx_t ctx; @@ -49,7 +49,7 @@ __global__ void WaveFrontPrimitiveTest(int loop, int skip, int wf_id = get_flat_block_id() / wf_size; int wg_offset = wg_id * ((get_flat_block_size() - 1 ) / wf_size + 1); int idx = wf_id + wg_offset; - int offset = size * idx; + size_t offset = size * idx; source += offset; dest += offset; @@ -120,13 +120,13 @@ WaveFrontPrimitiveTester::~WaveFrontPrimitiveTester() { rocshmem_free(dest); } -void WaveFrontPrimitiveTester::resetBuffers(uint64_t size) { +void WaveFrontPrimitiveTester::resetBuffers(size_t size) { size_t buff_size = size * args.num_wgs * num_warps; memset(dest, '1', buff_size); } void WaveFrontPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { + int loop, size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(WaveFrontPrimitiveTest, gridSize, blockSize, shared_bytes, @@ -138,7 +138,7 @@ void WaveFrontPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, num_timed_msgs = loop * gridSize.x * num_warps; } -void WaveFrontPrimitiveTester::verifyResults(uint64_t size) { +void WaveFrontPrimitiveTester::verifyResults(size_t size) { int check_id = (_type == WAVEGetTestType || _type == WAVEGetNBITestType) ? 0 : 1; diff --git a/tests/functional_tests/wavefront_primitives.hpp b/tests/functional_tests/wavefront_primitives.hpp index 2c19beb0e2..04f88d669d 100644 --- a/tests/functional_tests/wavefront_primitives.hpp +++ b/tests/functional_tests/wavefront_primitives.hpp @@ -36,12 +36,12 @@ class WaveFrontPrimitiveTester : public Tester { virtual ~WaveFrontPrimitiveTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; char *source = nullptr; char *dest = nullptr; diff --git a/tests/functional_tests/workgroup_primitives.cpp b/tests/functional_tests/workgroup_primitives.cpp index 27fa8c961f..f4e14b6d4d 100644 --- a/tests/functional_tests/workgroup_primitives.cpp +++ b/tests/functional_tests/workgroup_primitives.cpp @@ -36,7 +36,7 @@ using namespace rocshmem; __global__ void WorkGroupPrimitiveTest(int loop, int skip, long long int *start_time, long long int *end_time, char *source, - char *dest, int size, TestType type, + char *dest, size_t size, TestType type, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; int wg_id = get_flat_grid_id(); @@ -44,7 +44,7 @@ __global__ void WorkGroupPrimitiveTest(int loop, int skip, rocshmem_wg_ctx_create(ctx_type, &ctx); // Calculate start index for each work group - uint64_t offset = size * wg_id; + size_t offset = size * wg_id; source += offset; dest += offset; @@ -116,13 +116,13 @@ WorkGroupPrimitiveTester::~WorkGroupPrimitiveTester() { rocshmem_free(dest); } -void WorkGroupPrimitiveTester::resetBuffers(uint64_t size) { +void WorkGroupPrimitiveTester::resetBuffers(size_t size) { size_t buff_size = size * args.num_wgs; memset(dest, '1', buff_size); } void WorkGroupPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { + int loop, size_t size) { size_t shared_bytes = 0; hipLaunchKernelGGL(WorkGroupPrimitiveTest, gridSize, blockSize, shared_bytes, @@ -133,7 +133,7 @@ void WorkGroupPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, num_timed_msgs = loop * gridSize.x; } -void WorkGroupPrimitiveTester::verifyResults(uint64_t size) { +void WorkGroupPrimitiveTester::verifyResults(size_t size) { int check_id = (_type == WGGetTestType || _type == WGGetNBITestType) ? 0 : 1; diff --git a/tests/functional_tests/workgroup_primitives.hpp b/tests/functional_tests/workgroup_primitives.hpp index f293dabcf2..abcbccdf54 100644 --- a/tests/functional_tests/workgroup_primitives.hpp +++ b/tests/functional_tests/workgroup_primitives.hpp @@ -36,12 +36,12 @@ class WorkGroupPrimitiveTester : public Tester { virtual ~WorkGroupPrimitiveTester(); protected: - virtual void resetBuffers(uint64_t size) override; + virtual void resetBuffers(size_t size) override; virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) override; + size_t size) override; - virtual void verifyResults(uint64_t size) override; + virtual void verifyResults(size_t size) override; char *source = nullptr; char *dest = nullptr;