add support for GPUs using wavefront size of 32 (#285)

* add gfx1100 support Add support for Radeon 7900 GPUs (RX and PRO), and 7800 PRO. I was contemplating to add gfx1101 and gfx1102 GPUs as well, but those are the lower end models that are more unlikely to be used for compute intensive jobs. In addition, I do not have access to them to test the support. * update WF_SIZe for different options Radeon systems use a WarpSize of 32, unlike current Instinct systems, which use a warp size of 64. For the device side, a gfx specific ifdef is sufficient. For the host side, we need to query the device properties. * adjust functional tests to wf_size of 32 * update unit tests to handle wf_size of 32 * address reviewer comments [ROCm/rocshmem commit: d0c2845031]
2025-10-22 16:04:58 -05:00
@@ -46,7 +46,7 @@ __device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) {
 #endif
 #if defined(__gfx908__)
 #endif
-#if defined(__gfx90a__)
+#if defined(__gfx90a__) || defined (__gfx1100__)
  asm volatile(
      "global_load_ubyte %0 %1 off glc slc \n"
      "s_waitcnt vmcnt(0)"
@@ -69,7 +69,7 @@ __device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_va
 #endif
 #if defined(__gfx908__)
 #endif
-#if defined(__gfx90a__)
+#if defined(__gfx90a__) || defined (__gfx1100__)
  asm volatile(
    "global_load_sbyte %0 %1 off glc slc\n "
    "s_waitcnt vmcnt(0)"
@@ -91,7 +91,7 @@ __device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assi
 #endif
 #if defined(__gfx908__)
 #endif
-#if defined(__gfx90a__)
+#if defined(__gfx90a__) || defined (__gfx1100__)
  asm volatile(
    "global_load_dwordx2 %0 %1 off glc slc\n "
    "s_waitcnt vmcnt(0)"
@@ -122,7 +122,7 @@ NOWARN(-Wdeprecated-volatile,
 #endif
 #if defined(__gfx908__)
 #endif
-#if defined(__gfx90a__)
+#if defined(__gfx90a__) || defined (__gfx1100__)
        asm volatile(
            "global_load_dword %0 %1 off glc slc \n"
            "s_waitcnt vmcnt(0)"
@@ -142,7 +142,7 @@ NOWARN(-Wdeprecated-volatile,
 #endif
 #if defined(__gfx908__)
 #endif
-#if defined(__gfx90a__)
+#if defined(__gfx90a__) || defined (__gfx1100__)
        asm volatile(
            "global_load_dwordx2 %0 %1 off glc slc \n"
            "s_waitcnt vmcnt(0)"
@@ -191,7 +191,7 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
 #endif
 #if defined(__gfx908__)
 #endif
-#if defined(__gfx90a__)
+#if defined(__gfx90a__) || defined (__gfx1100__)
      asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16));
 #endif
 #if defined(__gfx942__) || defined(__gfx950__)
@@ -205,7 +205,7 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
 #endif
 #if defined(__gfx908__)
 #endif
-#if defined(__gfx90a__)
+#if defined(__gfx90a__) || defined (__gfx1100__)
      asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32));
 #endif
 #if defined(__gfx942__) || defined(__gfx950__)
@@ -219,7 +219,7 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
 #endif
 #if defined(__gfx908__)
 #endif
-#if defined(__gfx90a__)
+#if defined(__gfx90a__) || defined (__gfx1100__)
      asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64));
 #endif
 #if defined(__gfx942__) || defined(__gfx950__)
@@ -57,7 +57,11 @@ inline const unsigned MAX_WG_SIZE{1024};
 *
 * @note Wavefront size on most systems is either 32 or 64.
 */
+#if defined(__gfx90a__) || defined(__gfx942__) || defined (__gfx950__)
 inline const int WF_SIZE{64};
+#else
+inline const int WF_SIZE{32};
+#endif

 }  // namespace rocshmem

@@ -28,6 +28,7 @@

 #include "backend_gda.hpp"
 #include "constants.hpp"
+#include "util.hpp"

 namespace rocshmem {

@@ -59,7 +60,10 @@ QueuePair::QueuePair(struct ibv_pd* pd, int gda_provider) {
    fetching_atomic_lkey = mr_fetching_atomic->lkey;
  }

-  for(int i{0}; i < FETCHING_ATOMIC_CNT; i+=WF_SIZE) {
+  int deviceId;
+  CHECK_HIP(hipGetDevice(&deviceId));
+  int wf_size = get_wf_size(deviceId);
+  for(int i{0}; i < FETCHING_ATOMIC_CNT; i+=wf_size) {
    fetching_atomic_freelist->push_back(fetching_atomic + i);
  }

@@ -56,14 +56,11 @@ ROBackend::ROBackend(MPI_Comm comm)
  profiler_proxy_ = ProfilerProxyT(envvar::max_num_contexts);

  int device_id;
-  hipDeviceProp_t device_props;
-
  CHECK_HIP(hipGetDevice(&device_id));
-  CHECK_HIP(hipGetDeviceProperties(&device_props, device_id));

-  max_wg_size_ = device_props.maxThreadsPerBlock;
+  max_wg_size_ = get_threads_per_block(device_id);

-  wf_size_ = device_props.warpSize;
+  wf_size_ = get_wf_size(device_id);

  setup_default_ctx_buffers();

@@ -44,6 +44,21 @@ typedef struct device_agent {
 std::vector<device_agent_t> gpu_agents;
 std::vector<device_agent_t> cpu_agents;

+std::vector<device_prop_t> device_properties;
+
+static void device_properties_init(void) {
+  int numDevices;
+  CHECK_HIP(hipGetDeviceCount(&numDevices));
+
+  device_prop_t prop;
+  hipDeviceProp_t hipprop;
+  for (int i=0; i<numDevices; i++) {
+    CHECK_HIP(hipGetDeviceProperties(&hipprop, i));
+    prop.warpSize = hipprop.warpSize;
+    prop.maxThreadsPerBlock = hipprop.maxThreadsPerBlock;
+    device_properties.push_back(prop);
+  }
+}
 hsa_status_t rocm_hsa_amd_memory_pool_callback(
    hsa_amd_memory_pool_t memory_pool, void* data) {
  hsa_amd_memory_pool_global_flag_t pool_flag{};
@@ -108,6 +123,8 @@ int rocm_init() {
    return 1;
  }

+  device_properties_init();
+
  return 0;
 }

@@ -30,6 +30,8 @@
 #include <hsa/hsa_ext_amd.h>

 #include <cstdio>
+#include <cassert>
+#include <vector>

 #include "rocshmem/rocshmem_config.h"  // NOLINT(build/include_subdir)
 #include "constants.hpp"
@@ -146,6 +148,24 @@ do {                                                                        \

 extern const int gpu_clock_freq_mhz;

+
+typedef struct device_prop {
+  int warpSize;
+  int maxThreadsPerBlock;
+} device_prop_t;
+
+extern std::vector<device_prop_t> device_properties;
+
+static int get_threads_per_block(int device_id) {
+  assert(device_properties.size() > device_id);
+  return device_properties[device_id].maxThreadsPerBlock;
+}
+
+static int get_wf_size(int device_id) {
+  assert(device_properties.size() > device_id);
+  return device_properties[device_id].warpSize;
+}
+
 /* Device-side internal functions */
 __device__ __forceinline__ uint32_t lowerID() {
  return __ffsll(__ballot(1)) - 1;