Fix ROCm 6.4 warnings (#47)

* Removed __AMDGCN_WAVEFRONT_SIZE

* Added unit test to validate WF_SIZE

src/containers/array_impl.hpp

@@ -26,6 +26,7 @@
 #include <cassert>
 
 #include "array.hpp"
+#include "../constants.hpp"
 
 namespace rocshmem {
 
@@ -148,8 +149,8 @@ __device__ void Array<TYPE>::zero_thread_dump() {
 template <typename TYPE>
 __device__ void Array<TYPE>::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) {
       }
 

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;

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;

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
 )
 
 ###############################################################################

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);
     }
   }
 }

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<void**>(&_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;
 };

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);
+  }
+}
+

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