From ac3c18bc7c9f1157f4cb86d495f8cfec9ffdb18e Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Tue, 2 Jul 2024 10:07:43 -0700 Subject: [PATCH] silence warnings in functional testsuite check the return code of hip functions in order to silence some warnings. [ROCm/rocshmem commit: 7e7bc4b0a9e20987038d73d8fecc092a26431025] --- .../functional_tests/amo_bitwise_tester.cpp | 4 +-- .../functional_tests/amo_extended_tester.cpp | 4 +-- .../functional_tests/amo_standard_tester.cpp | 4 +-- .../functional_tests/primitive_mr_tester.cpp | 2 +- .../functional_tests/random_access_tester.cpp | 18 ++++++------ .../functional_tests/shmem_ptr_tester.cpp | 4 +-- .../tests/functional_tests/swarm_tester.cpp | 2 +- .../team_ctx_infra_tester.cpp | 10 +++---- .../tests/functional_tests/test_driver.cpp | 4 +-- .../tests/functional_tests/tester.cpp | 28 +++++++++---------- .../tests/functional_tests/tester.hpp | 10 +++++++ 11 files changed, 50 insertions(+), 40 deletions(-) diff --git a/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp index 9e2e2337c2..625fd9cd7b 100644 --- a/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp @@ -40,7 +40,7 @@ __global__ void AMOBitwiseTest(int loop, int skip, uint64_t *timer, char *r_buf, *****************************************************************************/ template AMOBitwiseTester::AMOBitwiseTester(TesterArguments args) : Tester(args) { - hipMalloc((void **)&_ret_val, args.max_msg_size * args.num_wgs); + CHECK_HIP(hipMalloc((void **)&_ret_val, args.max_msg_size * args.num_wgs)); _r_buf = (char *)roc_shmem_malloc(args.max_msg_size); _s_buf = (T *)roc_shmem_malloc(args.max_msg_size * args.num_wgs); } @@ -48,7 +48,7 @@ AMOBitwiseTester::AMOBitwiseTester(TesterArguments args) : Tester(args) { template AMOBitwiseTester::~AMOBitwiseTester() { roc_shmem_free(_r_buf); - hipFree(_ret_val); + CHECK_HIP(hipFree(_ret_val)); } template diff --git a/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp index fd0e528de7..565e09f117 100644 --- a/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp @@ -40,7 +40,7 @@ __global__ void AMOExtendedTest(int loop, int skip, uint64_t *timer, *****************************************************************************/ template AMOExtendedTester::AMOExtendedTester(TesterArguments args) : Tester(args) { - hipMalloc((void **)&_ret_val, args.max_msg_size * args.num_wgs); + CHECK_HIP(hipMalloc((void **)&_ret_val, args.max_msg_size * args.num_wgs)); _r_buf = (char *)roc_shmem_malloc(args.max_msg_size); _s_buf = (T *)roc_shmem_malloc(args.max_msg_size * args.num_wgs); } @@ -48,7 +48,7 @@ AMOExtendedTester::AMOExtendedTester(TesterArguments args) : Tester(args) { template AMOExtendedTester::~AMOExtendedTester() { roc_shmem_free(_r_buf); - hipFree(_ret_val); + CHECK_HIP(hipFree(_ret_val)); } template diff --git a/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp index 0e4dc91fe5..17a69f9120 100644 --- a/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp @@ -40,7 +40,7 @@ __global__ void AMOStandardTest(int loop, int skip, uint64_t *timer, *****************************************************************************/ template AMOStandardTester::AMOStandardTester(TesterArguments args) : Tester(args) { - hipMalloc((void **)&_ret_val, args.max_msg_size * args.num_wgs); + CHECK_HIP(hipMalloc((void **)&_ret_val, args.max_msg_size * args.num_wgs)); _r_buf = (char *)roc_shmem_malloc(args.max_msg_size); _s_buf = (T *)roc_shmem_malloc(args.max_msg_size * args.num_wgs); } @@ -48,7 +48,7 @@ AMOStandardTester::AMOStandardTester(TesterArguments args) : Tester(args) { template AMOStandardTester::~AMOStandardTester() { roc_shmem_free(_r_buf); - hipFree(_ret_val); + CHECK_HIP(hipFree(_ret_val)); } template diff --git a/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp b/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp index 527cea9181..4152b5e10f 100644 --- a/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp @@ -87,7 +87,7 @@ void PrimitiveMRTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, hipLaunchKernelGGL(PrimitiveMRTest, gridSize, blockSize, shared_bytes, stream, loop, timer, s_buf, r_buf, size, _shmem_context); - hipDeviceSynchronize(); + CHECK_HIP(hipDeviceSynchronize()); num_msgs = (loop + args.skip) * gridSize.x; num_timed_msgs = loop * 64; diff --git a/projects/rocshmem/tests/functional_tests/random_access_tester.cpp b/projects/rocshmem/tests/functional_tests/random_access_tester.cpp index 298534e1ff..7f2891b268 100644 --- a/projects/rocshmem/tests/functional_tests/random_access_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/random_access_tester.cpp @@ -135,9 +135,9 @@ RandomAccessTester::RandomAccessTester(TesterArguments args) : Tester(args) { r_buf = (int *)roc_shmem_malloc(max_size * wg_size * space); h_buf = (int *)malloc(max_size * wg_size * space); h_dev_buf = (int *)malloc(max_size * wg_size * space); - hipMalloc((void **)&_threads_bins, sizeof(uint32_t) * _num_waves * _num_bins); - hipMalloc((void **)&_off_bins, sizeof(uint32_t) * _num_waves * _num_bins); - hipMalloc((void **)&_PE_bins, sizeof(uint32_t) * _num_waves * _num_bins); + CHECK_HIP(hipMalloc((void **)&_threads_bins, sizeof(uint32_t) * _num_waves * _num_bins)); + CHECK_HIP(hipMalloc((void **)&_off_bins, sizeof(uint32_t) * _num_waves * _num_bins)); + CHECK_HIP(hipMalloc((void **)&_PE_bins, sizeof(uint32_t) * _num_waves * _num_bins)); memset(_threads_bins, 0, sizeof(uint32_t) * _num_waves * _num_bins); memset(_off_bins, 0, sizeof(uint32_t) * _num_waves * _num_bins); memset(_PE_bins, 0, sizeof(uint32_t) * _num_waves * _num_bins); @@ -148,9 +148,9 @@ RandomAccessTester::~RandomAccessTester() { roc_shmem_free(r_buf); free(h_buf); free(h_dev_buf); - hipFree(_threads_bins); - hipFree(_off_bins); - hipFree(_PE_bins); + CHECK_HIP(hipFree(_threads_bins)); + CHECK_HIP(hipFree(_off_bins)); + CHECK_HIP(hipFree(_PE_bins)); } void RandomAccessTester::resetBuffers(uint64_t size) { @@ -209,9 +209,9 @@ void RandomAccessTester::verifyResults(uint64_t size) { } } - hipMemcpy(h_dev_buf, r_buf, space * args.wg_size * size, - hipMemcpyDeviceToHost); - hipDeviceSynchronize(); + CHECK_HIP(hipMemcpy(h_dev_buf, r_buf, space * args.wg_size * size, + hipMemcpyDeviceToHost)); + CHECK_HIP(hipDeviceSynchronize()); for (i = 0; i < (space * args.wg_size * size / sizeof(int)); i++) { if (h_dev_buf[i] != h_buf[i]) { printf("PE %d Got Data Validation: expecting %d got %d at %d \n", diff --git a/projects/rocshmem/tests/functional_tests/shmem_ptr_tester.cpp b/projects/rocshmem/tests/functional_tests/shmem_ptr_tester.cpp index 73535d2fd6..e715714ab9 100644 --- a/projects/rocshmem/tests/functional_tests/shmem_ptr_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/shmem_ptr_tester.cpp @@ -47,12 +47,12 @@ __global__ void ShmemPtrTest(char *r_buf, int *available) { * HOST TESTER CLASS METHODS *****************************************************************************/ ShmemPtrTester::ShmemPtrTester(TesterArguments args) : Tester(args) { - hipMalloc((void **)&_available, sizeof(int)); + CHECK_HIP(hipMalloc((void **)&_available, sizeof(int))); r_buf = (char *)roc_shmem_malloc(args.max_msg_size); } ShmemPtrTester::~ShmemPtrTester() { - hipFree(_available); + CHECK_HIP(hipFree(_available)); roc_shmem_free(r_buf); } diff --git a/projects/rocshmem/tests/functional_tests/swarm_tester.cpp b/projects/rocshmem/tests/functional_tests/swarm_tester.cpp index eefd3756d5..b6c9f3357a 100644 --- a/projects/rocshmem/tests/functional_tests/swarm_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/swarm_tester.cpp @@ -72,7 +72,7 @@ void GetSwarmTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, hipLaunchKernelGGL(GetSwarmTest, gridSize, blockSize, shared_bytes, stream, loop, args.skip, timer, s_buf, r_buf, size, - _shmem_context); + _shmem_context); num_msgs = (loop + args.skip) * gridSize.x * blockSize.x; num_timed_msgs = loop * gridSize.x * blockSize.x; diff --git a/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp b/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp index 6d39b9d4bc..6e0fe0c45f 100644 --- a/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp @@ -126,14 +126,14 @@ void TeamCtxInfraTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, /* Copy array of teams to device */ roc_shmem_team_t *teams_on_device; - hipMalloc(&teams_on_device, sizeof(roc_shmem_team_t) * NUM_TEAMS); - hipMemcpy(teams_on_device, team_world_dup, - sizeof(roc_shmem_team_t) * NUM_TEAMS, hipMemcpyHostToDevice); + CHECK_HIP(hipMalloc(&teams_on_device, sizeof(roc_shmem_team_t) * NUM_TEAMS)); + CHECK_HIP(hipMemcpy(teams_on_device, team_world_dup, + sizeof(roc_shmem_team_t) * NUM_TEAMS, hipMemcpyHostToDevice)); hipLaunchKernelGGL(TeamCtxInfraTest, gridSize, blockSize, shared_bytes, - stream, _shmem_context, teams_on_device); + stream, _shmem_context, teams_on_device); - hipFree(teams_on_device); + CHECK_HIP(hipFree(teams_on_device)); } void TeamCtxInfraTester::postLaunchKernel() { diff --git a/projects/rocshmem/tests/functional_tests/test_driver.cpp b/projects/rocshmem/tests/functional_tests/test_driver.cpp index 645e9d7f2a..7acef2dbc0 100644 --- a/projects/rocshmem/tests/functional_tests/test_driver.cpp +++ b/projects/rocshmem/tests/functional_tests/test_driver.cpp @@ -39,9 +39,9 @@ int main(int argc, char *argv[]) { */ int rank = roc_shmem_my_pe(); int ndevices, my_device = 0; - hipGetDeviceCount(&ndevices); + CHECK_HIP(hipGetDeviceCount(&ndevices)); my_device = rank % ndevices; - hipSetDevice(my_device); + CHECK_HIP(hipSetDevice(my_device)); /** * Must initialize rocshmem to access arguments needed by the tester. diff --git a/projects/rocshmem/tests/functional_tests/tester.cpp b/projects/rocshmem/tests/functional_tests/tester.cpp index 23a9aa01da..f35928eb53 100644 --- a/projects/rocshmem/tests/functional_tests/tester.cpp +++ b/projects/rocshmem/tests/functional_tests/tester.cpp @@ -56,17 +56,17 @@ Tester::Tester(TesterArguments args) : args(args) { _type = (TestType)args.algorithm; _shmem_context = args.shmem_context; - hipStreamCreate(&stream); - hipEventCreate(&start_event); - hipEventCreate(&stop_event); - hipMalloc((void**)&timer, sizeof(uint64_t) * args.num_wgs); + CHECK_HIP(hipStreamCreate(&stream)); + CHECK_HIP(hipEventCreate(&start_event)); + CHECK_HIP(hipEventCreate(&stop_event)); + CHECK_HIP(hipMalloc((void**)&timer, sizeof(uint64_t) * args.num_wgs)); } Tester::~Tester() { - hipFree(timer); - hipEventDestroy(stop_event); - hipEventDestroy(start_event); - hipStreamDestroy(stream); + CHECK_HIP(hipFree(timer)); + CHECK_HIP(hipEventDestroy(stop_event)); + CHECK_HIP(hipEventDestroy(start_event)); + CHECK_HIP(hipStreamDestroy(stream)); } std::vector Tester::create(TesterArguments args) { @@ -538,11 +538,11 @@ void Tester::execute() { const dim3 blockSize(args.wg_size, 1, 1); const dim3 gridSize(args.num_wgs, 1, 1); - hipEventRecord(start_event, stream); + CHECK_HIP(hipEventRecord(start_event, stream)); launchKernel(gridSize, blockSize, num_loops, size); - hipEventRecord(stop_event, stream); + CHECK_HIP(hipEventRecord(stop_event, stream)); hipError_t err = hipStreamSynchronize(stream); if (err != hipSuccess) { @@ -601,7 +601,7 @@ void Tester::print(uint64_t size) { double avg_msg_rate = num_timed_msgs / (timer_avg / 1e6); float total_kern_time_ms; - hipEventElapsedTime(&total_kern_time_ms, start_event, stop_event); + CHECK_HIP(hipEventElapsedTime(&total_kern_time_ms, start_event, stop_event)); float total_kern_time_s = total_kern_time_ms / 1000; double bandwidth_avg_gbs = num_msgs * size * bw_factor / total_kern_time_s / pow(2, 30); @@ -624,9 +624,9 @@ void Tester::print(uint64_t size) { void flush_hdp() { int hip_dev_id{}; unsigned int* hdp_flush_ptr_{nullptr}; - hipGetDevice(&hip_dev_id); - hipDeviceGetAttribute(reinterpret_cast(&hdp_flush_ptr_), - hipDeviceAttributeHdpMemFlushCntl, hip_dev_id); + CHECK_HIP(hipGetDevice(&hip_dev_id)); + CHECK_HIP(hipDeviceGetAttribute(reinterpret_cast(&hdp_flush_ptr_), + hipDeviceAttributeHdpMemFlushCntl, hip_dev_id)); __atomic_store_n(hdp_flush_ptr_, 0x1, __ATOMIC_SEQ_CST); } diff --git a/projects/rocshmem/tests/functional_tests/tester.hpp b/projects/rocshmem/tests/functional_tests/tester.hpp index e1082502e0..14a91435e5 100644 --- a/projects/rocshmem/tests/functional_tests/tester.hpp +++ b/projects/rocshmem/tests/functional_tests/tester.hpp @@ -143,4 +143,14 @@ class Tester { hipEvent_t stop_event; }; +#define CHECK_HIP(cmd) \ + { \ + hipError_t error = cmd; \ + if (error != hipSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), \ + error, __FILE__, __LINE__); \ + exit(EXIT_FAILURE); \ + } \ + } + #endif /* _TESTER_HPP */