From 487e5b7d0f6170a1ee96efae3d857c85827816d8 Mon Sep 17 00:00:00 2001 From: Yiltan Date: Mon, 24 Feb 2025 13:34:13 -0500 Subject: [PATCH] Fix ROCm 6.4 warnings (#47) * Removed __AMDGCN_WAVEFRONT_SIZE * Added unit test to validate WF_SIZE --- src/containers/array_impl.hpp | 5 +-- src/containers/share_strategy.cpp | 8 ++--- src/reverse_offload/context_ro_device.cpp | 6 ++-- tests/unit_tests/CMakeLists.txt | 1 + tests/unit_tests/bitwise_gtest.cpp | 12 +++---- tests/unit_tests/bitwise_gtest.hpp | 34 ++++++------------- tests/unit_tests/wavefront_size_gtest.cpp | 41 +++++++++++++++++++++++ tests/unit_tests/wavefront_size_gtest.hpp | 34 +++++++++++++++++++ 8 files changed, 100 insertions(+), 41 deletions(-) create mode 100644 tests/unit_tests/wavefront_size_gtest.cpp create mode 100644 tests/unit_tests/wavefront_size_gtest.hpp diff --git a/src/containers/array_impl.hpp b/src/containers/array_impl.hpp index c0a750d100..73a660e6c3 100644 --- a/src/containers/array_impl.hpp +++ b/src/containers/array_impl.hpp @@ -26,6 +26,7 @@ #include #include "array.hpp" +#include "../constants.hpp" namespace rocshmem { @@ -148,8 +149,8 @@ __device__ void Array::zero_thread_dump() { template __device__ void Array::any_thread_dump() { Identity id{}; - for (int i = 0; i < __AMDGCN_WAVEFRONT_SIZE; i++) { - if ((id.local_thread_id() % __AMDGCN_WAVEFRONT_SIZE) == i) { + for (int i = 0; i < WF_SIZE; i++) { + if ((id.local_thread_id() % WF_SIZE) == i) { while (atomicCAS(GLOBAL_DEVICE_PRINT_LOCK, 0, 1) == 1) { } diff --git a/src/containers/share_strategy.cpp b/src/containers/share_strategy.cpp index d557600189..85d07ebd16 100644 --- a/src/containers/share_strategy.cpp +++ b/src/containers/share_strategy.cpp @@ -21,14 +21,10 @@ *****************************************************************************/ #include "share_strategy.hpp" +#include "../constants.hpp" namespace rocshmem { -/* - * @brief The warp (wave-front) size. - */ -static constexpr uint64_t WARP_SIZE = __AMDGCN_WAVEFRONT_SIZE; - __device__ uint64_t Block::lane_id() { /* * amd_detail/device_functions.h::__lane_id() @@ -202,7 +198,7 @@ __device__ uint64_t Block::active_logical_lane_id() { } __device__ uint64_t Block::broadcast_up(uint64_t fetch_value) { - for (unsigned i = 0; i < WARP_SIZE; i++) { + for (unsigned i = 0; i < WF_SIZE; i++) { uint64_t temp = __shfl_up(fetch_value, i); if (temp) { fetch_value = temp; diff --git a/src/reverse_offload/context_ro_device.cpp b/src/reverse_offload/context_ro_device.cpp index 8b82a97578..e9e42c4afe 100644 --- a/src/reverse_offload/context_ro_device.cpp +++ b/src/reverse_offload/context_ro_device.cpp @@ -470,9 +470,9 @@ __device__ uint64_t active_logical_lane_id() { } __device__ uint64_t broadcast_lds(bool lowest_active, uint64_t value) { - constexpr size_t SIZE = 1024 / __AMDGCN_WAVEFRONT_SIZE; + constexpr size_t SIZE = 1024 / WF_SIZE; __shared__ uint64_t value_per_warp[SIZE]; - auto wavefront_number {get_flat_block_id() / __AMDGCN_WAVEFRONT_SIZE}; + auto wavefront_number {get_flat_block_id() / WF_SIZE}; if (lowest_active) { value_per_warp[wavefront_number] = value; __threadfence_block(); @@ -481,7 +481,7 @@ __device__ uint64_t broadcast_lds(bool lowest_active, uint64_t value) { } __device__ uint64_t broadcast_shfl_up(uint64_t value) { - for (unsigned i{0}; i < __AMDGCN_WAVEFRONT_SIZE; i++) { + for (unsigned i{0}; i < WF_SIZE; i++) { uint64_t temp{__shfl_up(value, i)}; if (temp) { value = temp; diff --git a/tests/unit_tests/CMakeLists.txt b/tests/unit_tests/CMakeLists.txt index 67114785bc..aab5d1a5e5 100644 --- a/tests/unit_tests/CMakeLists.txt +++ b/tests/unit_tests/CMakeLists.txt @@ -94,6 +94,7 @@ target_sources( ipc_impl_simple_coarse_gtest.cpp ipc_impl_simple_fine_gtest.cpp ipc_impl_tiled_fine_gtest.cpp + wavefront_size_gtest.cpp ) ############################################################################### diff --git a/tests/unit_tests/bitwise_gtest.cpp b/tests/unit_tests/bitwise_gtest.cpp index 73723098ba..df7f8bd806 100644 --- a/tests/unit_tests/bitwise_gtest.cpp +++ b/tests/unit_tests/bitwise_gtest.cpp @@ -88,7 +88,7 @@ __global__ void activate_lane_helper_kernel( * Index into warp matrix to save return value to be read by host. */ if (is_an_active_lane) { - size_t warp_index = hipThreadIdx_x / device_methods->warp_size(); + size_t warp_index = hipThreadIdx_x / WF_SIZE; size_t block_index = hipBlockIdx_x; auto *elem = warp_matrix->access(warp_index, block_index); *elem = hipThreadIdx_x; @@ -145,7 +145,7 @@ TEST_F(BitwiseTestFixture, verify_host_warp_matrix_init_1024_8) { TEST_F(BitwiseTestFixture, verify_warp_size_64) { setup_fixture({1, 1, 1}, {1, 1, 1}); - ASSERT_EQ(_warp_size, 64); + ASSERT_EQ(WF_SIZE, 64); } /***************************************************************************** @@ -1400,7 +1400,7 @@ TEST_F(BitwiseTestFixture, fetch_incr_kernel_4_1) { for (size_t i = 0; i < _warp_matrix->rows(); i++) { for (size_t j = 0; j < _warp_matrix->columns(); j++) { auto *elem = _warp_matrix->access(i, j); - ASSERT_EQ(*elem % _warp_size, 0); + ASSERT_EQ(*elem % WF_SIZE, 0); } } } @@ -1424,7 +1424,7 @@ TEST_F(BitwiseTestFixture, fetch_incr_kernel_64_1) { for (size_t i = 0; i < _warp_matrix->rows(); i++) { for (size_t j = 0; j < _warp_matrix->columns(); j++) { auto *elem = _warp_matrix->access(i, j); - ASSERT_EQ(*elem % _warp_size, 0); + ASSERT_EQ(*elem % WF_SIZE, 0); } } } @@ -1496,7 +1496,7 @@ TEST_F(BitwiseTestFixture, fetch_incr_kernel_1024_1024) { for (size_t i = 0; i < _warp_matrix->rows(); i++) { for (size_t j = 0; j < _warp_matrix->columns(); j++) { auto *elem = _warp_matrix->access(i, j); - ASSERT_EQ(*elem % _warp_size, 0); + ASSERT_EQ(*elem % WF_SIZE, 0); } } } @@ -1520,7 +1520,7 @@ TEST_F(BitwiseTestFixture, fetch_incr_logical_1_kernel_1024_1024) { for (size_t i = 0; i < _warp_matrix->rows(); i++) { for (size_t j = 0; j < _warp_matrix->columns(); j++) { auto *elem = _warp_matrix->access(i, j); - ASSERT_EQ(*elem % _warp_size, 1); + ASSERT_EQ(*elem % WF_SIZE, 1); } } } diff --git a/tests/unit_tests/bitwise_gtest.hpp b/tests/unit_tests/bitwise_gtest.hpp index 0b040daabf..e6770a7c7f 100644 --- a/tests/unit_tests/bitwise_gtest.hpp +++ b/tests/unit_tests/bitwise_gtest.hpp @@ -58,7 +58,7 @@ class BitwiseDeviceMethods Block block {}; if (activate_lane_helper(lanes_bitfield)) { auto low_lane = block.lowest_active_lane(); - size_t warp_index = hipThreadIdx_x / _warp_size; + size_t warp_index = hipThreadIdx_x / WF_SIZE; size_t block_index = hipBlockIdx_x; auto *elem = warp_matrix->access(warp_index, block_index); *elem = low_lane; @@ -73,7 +73,7 @@ class BitwiseDeviceMethods Block block {}; if (activate_lane_helper(lanes_bitfield)) { if (block.is_lowest_active_lane()) { - size_t warp_index = hipThreadIdx_x / _warp_size; + size_t warp_index = hipThreadIdx_x / WF_SIZE; size_t block_index = hipBlockIdx_x; auto *elem = warp_matrix->access(warp_index, block_index); *elem = block.lane_id(); @@ -89,7 +89,7 @@ class BitwiseDeviceMethods Block block {}; if (activate_lane_helper(lanes_bitfield)) { if (block.active_logical_lane_id() == 2) { - size_t warp_index = hipThreadIdx_x / _warp_size; + size_t warp_index = hipThreadIdx_x / WF_SIZE; size_t block_index = hipBlockIdx_x; auto *elem = warp_matrix->access(warp_index, block_index); *elem = block.lane_id(); @@ -105,7 +105,7 @@ class BitwiseDeviceMethods Block block {}; if (activate_lane_helper(lanes_bitfield)) { auto lane_id = block.lane_id(); - size_t warp_index = hipThreadIdx_x / _warp_size; + size_t warp_index = hipThreadIdx_x / WF_SIZE; size_t block_index = hipBlockIdx_x; auto *elem = warp_matrix->access(warp_index, block_index); *elem = lane_id; @@ -120,7 +120,7 @@ class BitwiseDeviceMethods Block block {}; if (activate_lane_helper(lanes_bitfield)) { auto number_active_lanes = block.number_active_lanes(); - size_t warp_index = hipThreadIdx_x / _warp_size; + size_t warp_index = hipThreadIdx_x / WF_SIZE; size_t block_index = hipBlockIdx_x; auto *elem = warp_matrix->access(warp_index, block_index); *elem = number_active_lanes; @@ -139,7 +139,7 @@ class BitwiseDeviceMethods value = 42; } value = block.broadcast_up(value); - size_t warp_index = hipThreadIdx_x / _warp_size; + size_t warp_index = hipThreadIdx_x / WF_SIZE; size_t block_index = hipBlockIdx_x; auto *elem = warp_matrix->access(warp_index, block_index); *elem = value; @@ -155,7 +155,7 @@ class BitwiseDeviceMethods if (activate_lane_helper(lanes_bitfield)) { auto orig = block.fetch_incr(_fetch_value); if (block.is_lowest_active_lane()) { - size_t warp_index = hipThreadIdx_x / _warp_size; + size_t warp_index = hipThreadIdx_x / WF_SIZE; size_t block_index = hipBlockIdx_x; auto *elem = warp_matrix->access(warp_index, block_index); *elem = orig; @@ -172,7 +172,7 @@ class BitwiseDeviceMethods if (activate_lane_helper(lanes_bitfield)) { auto orig = block.fetch_incr(_fetch_value); if (block.active_logical_lane_id() == 1) { - size_t warp_index = hipThreadIdx_x / _warp_size; + size_t warp_index = hipThreadIdx_x / WF_SIZE; size_t block_index = hipBlockIdx_x; auto *elem = warp_matrix->access(warp_index, block_index); *elem = orig; @@ -201,7 +201,7 @@ class BitwiseDeviceMethods * warp_bit_id := 66 % 64 * warp_bit_id := 2 */ - uint64_t warp_bit_id = hipThreadIdx_x % _warp_size; + uint64_t warp_bit_id = hipThreadIdx_x % WF_SIZE; /* * Example (continued): @@ -224,20 +224,7 @@ class BitwiseDeviceMethods return is_an_active_lane; } - __device__ - uint64_t - warp_size() const - { - return _warp_size; - } - long long unsigned *_fetch_value = nullptr; - - private: - /************************************************************************* - ********************** Implementation Variables ************************* - *************************************************************************/ - static constexpr uint64_t _warp_size = __AMDGCN_WAVEFRONT_SIZE; }; /***************************************************************************** @@ -287,7 +274,7 @@ class BitwiseTestFixture : public ::testing::Test _hip_allocator.allocate(reinterpret_cast(&_warp_matrix), sizeof(WarpMatrix)); - size_t warps_per_block = ceil(float(_hip_block_dim.x) / _warp_size); + size_t warps_per_block = ceil(float(_hip_block_dim.x) / WF_SIZE); const ObjectStrategy *default_object_strategy = DefaultObjectStrategy::instance()->get(); @@ -354,7 +341,6 @@ class BitwiseTestFixture : public ::testing::Test dim3 _hip_block_dim {}; dim3 _hip_grid_dim {}; HIPAllocator _hip_allocator {}; - static constexpr uint64_t _warp_size = __AMDGCN_WAVEFRONT_SIZE; WarpMatrix *_warp_matrix = nullptr; BitwiseDeviceMethods *_device_methods = nullptr; }; diff --git a/tests/unit_tests/wavefront_size_gtest.cpp b/tests/unit_tests/wavefront_size_gtest.cpp new file mode 100644 index 0000000000..114ecd2662 --- /dev/null +++ b/tests/unit_tests/wavefront_size_gtest.cpp @@ -0,0 +1,41 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#include "wavefront_size_gtest.hpp" + +#include "../src/util.hpp" + +using namespace rocshmem; + +TEST_F(WavefrontSizeTestFixture, constant_matches_runtime) { + int device_count = 0; + hipDeviceProp_t prop; + + CHECK_HIP(hipGetDeviceCount(&device_count)); + ASSERT_GT(device_count, 0); + + for (int i = 0; i < device_count; i++) { + CHECK_HIP(hipGetDeviceProperties(&prop, i)); + ASSERT_EQ(WF_SIZE, prop.warpSize); + } +} + diff --git a/tests/unit_tests/wavefront_size_gtest.hpp b/tests/unit_tests/wavefront_size_gtest.hpp new file mode 100644 index 0000000000..0ea683b14b --- /dev/null +++ b/tests/unit_tests/wavefront_size_gtest.hpp @@ -0,0 +1,34 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef ROCSHMEM_WAVEFRONT_SIZE_GTEST_HPP +#define ROCSHMEM_WAVEFRONT_SIZE_GTEST_HPP + +#include "gtest/gtest.h" + +namespace rocshmem { + +class WavefrontSizeTestFixture : public ::testing::Test { }; + +} // namespace rocshmem + +#endif // ROCSHMEM_WAVEFRONT_SIZE_GTEST_HPP