Debug for remote benchmark

HermesCui · HermesCui · commit ec91df9a9bfd · 2025-07-29T19:00:21.000Z
diff --git a/gpu_driven/Makefile b/gpu_driven/Makefile
@@ -8,9 +8,10 @@ INC = -I./ -I../include -I$(HIP_HOME)/include -I${CONDA_LIB_HOME}/../include -L$
 LIBS = -lz -lelf -libverbs -L ${HIP_HOME}/lib -lamdhip64
 LIBS_SHARED = -lglog -lgflags -lgtest -lz -lelf -libverbs
 override CXXFLAGS += -O3 -g -std=c++17 -Wno-pointer-arith -Wno-interference-size -fPIC -D__HIP_PLATFORM_AMD__
+override HIPFLAGS += -O3 -g -std=c++17 -Wno-pointer-arith -Wno-interference-size -fPIC -D__HIP_PLATFORM_AMD__ --offload-arch=gfx942
 DEPS = *.h
-PLUGIN_SO = librccl-net-uccl.so
-NCCL_INC:= -I$(RCCL_HOME)/build/release/include -I$(RCCL_HOME)/src/include -I$(HIP_HOME)/include
+# PLUGIN_SO = librccl-net-uccl.so
+# NCCL_INC:= -I$(RCCL_HOME)/build/release/include -I$(RCCL_HOME)/src/include -I$(HIP_HOME)/include
 
 # HIP compiler
 HIPCC := $(HIP_HOME)/bin/hipcc
@@ -48,7 +49,7 @@ TARGET_LOCAL := benchmark_local
 TARGET_REMOTE := benchmark_remote
 
 .PHONY: build
-build: $(test_bin) $(lib_obj) $(PLUGIN_SO) librdma_hip.a $(TARGET_LOCAL) $(TARGET_REMOTE)
+build: $(test_bin) $(lib_obj) librdma_hip.a $(TARGET_LOCAL) $(TARGET_REMOTE)
 
 # Test binary compilation
 %_test: %_test.cc $(DEPS) $(lib_obj)
@@ -64,27 +65,27 @@ build: $(test_bin) $(lib_obj) $(PLUGIN_SO) librdma_hip.a $(TARGET_LOCAL) $(TARGE
 
 # HIP compilation rule
 %.o: %.hip
-	$(HIPCC) $(CXXFLAGS) $(INCLUDES) -MMD -MP -c $< -o $@
+	$(HIPCC) $(HIPFLAGS) $(INCLUDES) -MMD -MP -c $< -o $@
 
-# Plugin shared library
-$(PLUGIN_SO): nccl_plugin.cc $(DEPS) $(lib_obj)
-	g++ $(NCCL_INC) -fPIC -shared -o $@ -Wl,-soname,$(PLUGIN_SO) nccl_plugin.cc $(lib_obj) $(INC) $(LIBS_SHARED) $(CXXFLAGS)
+# Plugin shared library - commented out as NCCL plugin is not part of this project
+# $(PLUGIN_SO): nccl_plugin.cc $(DEPS) $(lib_obj)
+# 	g++ $(NCCL_INC) -fPIC -shared -o $@ -Wl,-soname,$(PLUGIN_SO) nccl_plugin.cc $(lib_obj) $(INC) $(LIBS_SHARED) $(CXXFLAGS)
 
 # Static library
 librdma_hip.a: $(lib_obj)
 	ar rcs $@ $(lib_obj)
 
 # Benchmark targets - linking with HIP
 $(TARGET_LOCAL): $(OBJ_LOCAL)
-	$(HIPCC) $(CXXFLAGS) $(INCLUDES) $(OBJ_LOCAL) $(LIBS) -o $@
+	$(HIPCC) $(HIPFLAGS) $(INCLUDES) $(OBJ_LOCAL) $(LIBS) -o $@
 
 $(TARGET_REMOTE): $(OBJ_REMOTE)
-	$(HIPCC) $(CXXFLAGS) $(INCLUDES) $(OBJ_REMOTE) $(LIBS) -o $@
+	$(HIPCC) $(HIPFLAGS) $(INCLUDES) $(OBJ_REMOTE) $(LIBS) -o $@
 
 # Clean all generated files
 .PHONY: clean
 clean:
-	rm -f *.o $(test_bin) $(PLUGIN_SO) librdma_hip.a $(TARGET_LOCAL) $(TARGET_REMOTE) *.d src/*.d bench/*.d
+	rm -f *.o $(test_bin) librdma_hip.a $(TARGET_LOCAL) $(TARGET_REMOTE) *.d src/*.d bench/*.d
 
 # Automatically include dependency files if they exist
 -include $(OBJ_LOCAL:.o=.d) $(OBJ_REMOTE:.o=.d)
diff --git a/gpu_driven/bench/benchmark_local.hip b/gpu_driven/bench/benchmark_local.hip
@@ -15,13 +15,7 @@ int main(int argc, char** argv) {
     return 1;
   }
 
-  GdrSupportInitOnce();
-  if (!GdrSupportInitOnce()) {
-    printf(
-        "Error: GPUDirect RDMA module is not loaded. Please load "
-        "amdgpu or rocm_smi modules!\n");
-    exit(1);
-  }
+
 
   hipStream_t stream1;
   hipStreamCreate(&stream1);
diff --git a/gpu_driven/bench/benchmark_remote.hip b/gpu_driven/bench/benchmark_remote.hip
@@ -49,6 +49,25 @@ int main(int argc, char** argv) {
 
   printf("Allocated %zu bytes of GPU buffer at %p\n", total_size, gpu_buffer);
 
+  // Initialize per_GPU_device_buf array
+  for (int d = 0; d < NUM_GPUS; ++d) {
+    per_GPU_device_buf[d] = nullptr;
+  }
+
+  // Initialize per_GPU_device_buf for remote node
+  if (rank == 1) {
+    for (int d = 0; d < NUM_GPUS; ++d) {
+      hipSetDevice(d);
+      hipMalloc(&per_GPU_device_buf[d], total_size);
+      if (per_GPU_device_buf[d] == nullptr) {
+        fprintf(stderr, "Failed to allocate GPU buffer on GPU %d\n", d);
+        exit(1);
+      }
+      printf("Allocated %zu bytes for per_GPU_device_buf[%d] at %p\n", total_size, d, per_GPU_device_buf[d]);
+    }
+    hipSetDevice(0);
+  }
+
   // Initialize global RDMA resources
   RDMAConnectionInfo local_info;
   global_rdma_init(gpu_buffer, total_size, &local_info, rank);
diff --git a/gpu_driven/include/common.hpp b/gpu_driven/include/common.hpp
@@ -23,6 +23,15 @@
     }                                                               \
   } while (0)
 
+#define HIP_CHECK(cmd)                                                                             \
+    do {                                                                                           \
+        hipError_t err = cmd;                                                                      \
+        if (err != hipSuccess) {                                                                   \
+            fprintf(stderr, "HIP error (%s:%d): %s\n", __FILE__, __LINE__, hipGetErrorString(err)); \
+            exit(EXIT_FAILURE);                                                                    \
+        }                                                                                          \
+    } while (0)
+
 #define hipCheckErrors(msg)                                  \
   do {                                                        \
     hipError_t __err = hipGetLastError();                   \
diff --git a/gpu_driven/include/rdma.hpp b/gpu_driven/include/rdma.hpp
@@ -81,6 +81,7 @@ void handle_peer_copy(uint64_t wr_id, uint32_t imm, int src_dev, int dst_dev,
                       void* src_ptr, void* dst_ptr, size_t num_bytes);
 
 void per_thread_rdma_init(void* gpu_buf, size_t bytes, int rank, int block_idx);
+void fill_local_gid(RDMAConnectionInfo* local_info);
 
 extern void* per_GPU_device_buf[NUM_GPUS];
 
diff --git a/gpu_driven/include/util_simple.h b/gpu_driven/include/util_simple.h
@@ -0,0 +1,119 @@
+#pragma once
+
+#include <filesystem>
+#include <vector>
+#include <string>
+#include <algorithm>
+#include <regex>
+#include <fstream>
+#include <iostream>
+
+namespace uccl {
+
+namespace fs = std::filesystem;
+
+static bool is_bdf(std::string const& s) {
+  // Match full PCI BDF allowing hexadecimal digits
+  static const std::regex re(
+      R"([0-9a-fA-F]{4}:[0-9a-fA-F]{2}:[0-9a-fA-F]{2}\.[0-9a-fA-F])");
+  return std::regex_match(s, re);
+}
+
+static int cal_pcie_distance(fs::path const& devA, fs::path const& devB) {
+  auto devA_parent = devA.parent_path();
+  auto devB_parent = devB.parent_path();
+
+  auto build_chain = [](fs::path const& dev) {
+    std::vector<std::string> chain;
+    for (fs::path p = fs::canonical(dev);; p = p.parent_path()) {
+      std::string leaf = p.filename();
+      if (is_bdf(leaf)) chain.push_back(leaf);  // collect BDF components
+      if (p == p.root_path()) break;            // reached filesystem root
+    }
+    return chain; /* self → root */
+  };
+
+  auto chainA = build_chain(devA_parent);
+  auto chainB = build_chain(devB_parent);
+
+  // Walk back from root until paths diverge
+  size_t i = chainA.size();
+  size_t j = chainB.size();
+  while (i > 0 && j > 0 && chainA[i - 1] == chainB[j - 1]) {
+    --i;
+    --j;
+  }
+  // Distance = remaining unique hops in each chain
+  return static_cast<int>(i + j);
+}
+
+static std::vector<fs::path> get_gpu_cards() {
+  // Discover GPU BDF using /sys/class/drm/cardX/device symlinks
+  std::vector<fs::path> gpu_cards;
+  const fs::path drm_class{"/sys/class/drm"};
+  const std::regex card_re(R"(card(\d+))");
+
+  if (fs::exists(drm_class)) {
+    for (auto const& entry : fs::directory_iterator(drm_class)) {
+      const std::string name = entry.path().filename();
+      std::smatch m;
+      if (!std::regex_match(name, m, card_re)) continue;
+
+      fs::path dev_path = fs::canonical(entry.path() / "device");
+
+      // check vendor id
+      std::ifstream vf(dev_path / "vendor");
+      std::string vs;
+      if (!(vf >> vs)) continue;
+      uint32_t vendor = std::stoul(vs, nullptr, 0);  // handles "0x10de"
+
+      if (vendor != 0x10de && vendor != 0x1002) continue;  // NVIDIA or AMD
+
+      gpu_cards.push_back(dev_path);
+    }
+  }
+
+  const fs::path nvidia_gpus{"/proc/driver/nvidia/gpus"};
+  if (gpu_cards.empty() && fs::exists(nvidia_gpus)) {
+    for (auto const& entry : fs::directory_iterator(nvidia_gpus)) {
+      gpu_cards.push_back(entry.path());
+    }
+  }
+
+  std::sort(gpu_cards.begin(), gpu_cards.end(),
+            [](fs::path const& a, fs::path const& b) {
+              return a.filename() < b.filename();
+            });
+
+  return gpu_cards;
+}
+
+static std::vector<std::pair<std::string, fs::path>> get_rdma_nics() {
+  // Discover RDMA NICs under /sys/class/infiniband
+  std::vector<std::pair<std::string, fs::path>> ib_nics;
+  const fs::path ib_class{"/sys/class/infiniband"};
+  if (!fs::exists(ib_class)) {
+    std::cerr << "No /sys/class/infiniband directory found. Are RDMA drivers "
+                 "loaded?\n";
+    return ib_nics;
+  }
+
+  for (auto const& ib_entry : fs::directory_iterator(ib_class)) {
+    std::string ibdev = ib_entry.path().filename();
+    fs::path ib_device_path = fs::canonical(ib_entry.path() / "device");
+
+    // Collect interface names under RDMA device
+    fs::path netdir = ib_device_path / "net";
+    if (fs::exists(netdir) && fs::is_directory(netdir)) {
+      ib_nics.push_back(std::make_pair(ibdev, ib_device_path));
+    }
+  }
+  std::sort(ib_nics.begin(), ib_nics.end(),
+            [](std::pair<std::string, fs::path> const& a,
+               std::pair<std::string, fs::path> const& b) {
+              return a.first < b.first;
+            });
+  return ib_nics;
+}
+
+} // namespace uccl
diff --git a/gpu_driven/src/gpu_kernel.hip b/gpu_driven/src/gpu_kernel.hip
@@ -7,11 +7,14 @@
 
 // HIP equivalent of clock64() - using wall clock time
 __device__ __forceinline__ unsigned long long hip_clock64() {
-#if defined(__AMDGCN__)
-  // AMD specific clock implementation
+#if defined(__AMDGCN__) && defined(__GFX942__)
+  // AMD specific clock implementation for gfx942
   unsigned long long clock_val;
   asm volatile("s_memrealtime %0" : "=s"(clock_val));
   return clock_val;
+#elif defined(__AMDGCN__)
+  // More compatible AMD implementation
+  return __builtin_amdgcn_s_memrealtime();
 #else
   // Fallback implementation
   return static_cast<unsigned long long>(clock());
diff --git a/gpu_driven/src/peer_copy.hip b/gpu_driven/src/peer_copy.hip
@@ -108,11 +108,11 @@ hipError_t launch_peer_bulk_copy2(CopyTask const* host_tasks, int num_tasks,
     hipLaunchKernelGGL(peer_copy_kernel_vec, blocks, threads_per_block, 0, stream,
                        d_tasks, num_tasks);
   } else if (true) {
-    int tasks_per_block = num_tasks / NVLINK_SM_PER_PROCESS;
+    int tasks_per_block = (num_tasks + NVLINK_SM_PER_PROCESS - 1) / NVLINK_SM_PER_PROCESS;
     hipLaunchKernelGGL(peer_copy_kernel_vec_batched, blocks, threads_per_block, 0, stream,
         d_tasks, num_tasks, tasks_per_block);
   } else {
-    int tasks_per_block = num_tasks / NVLINK_SM_PER_PROCESS;
+    int tasks_per_block = (num_tasks + NVLINK_SM_PER_PROCESS - 1) / NVLINK_SM_PER_PROCESS;
     size_t shmem = threads_per_block * 2 /*PIPE_DEPTH*/ * sizeof(int4);
     hipLaunchKernelGGL(HIP_KERNEL_NAME(peer_copy_kernel_vec_pipelined<2, int4>),
         blocks, threads_per_block, shmem, stream, d_tasks, num_tasks, tasks_per_block);
diff --git a/gpu_driven/src/rdma.cpp b/gpu_driven/src/rdma.cpp
