fix(rdma): move real bytes on the staging and GPU-direct paths

The RDMA data plane reported success for transfers that moved nothing.
Four defects, each of which independently made a green run meaningless.

buildRdmaToken advertised mr->addr. ibv_reg_dmabuf_mr leaves that NULL,
and on the bounce path it is the staging buffer rather than the caller's
pointer, so the token described memory the peer could not use. BufferMap
now records the address to advertise at registration time.

Nothing copied between the caller's buffer and the staging buffer. When
the NIC cannot reach GPU memory directly, hipObject registers host
staging instead of the caller's pointer -- and then left it empty. Adds
hipObjBufSync: callers driving a transfer through hipObjGetRdmaToken call
it with HIPOBJ_SYNC_TO_HOST before a PUT and HIPOBJ_SYNC_TO_DEVICE after
a GET. It is a no-op for a directly registered buffer.

The minio-cpp bridge sent the caller's raw pointer in the token's
":<addr>:<len>" suffix, which rocm-ernic treats as an override of the
address in the token. On the staging path that address is not covered by
the MR, giving LOCAL_PROT_ERR and a zero-byte transfer that the bridge
reported as success. Send zeros -- "no override" -- and let the token
speak. rdmaPutWithRetry now stages to host before the request and
rdmaGetWithRetry stages back to device after a successful reply.

dmabuf support was predicted rather than attempted.
init_dmabuf_support_flag() grepped /boot/config-$(uname -r) for both
CONFIG_DMABUF_MOVE_NOTIFY=y and CONFIG_PCI_P2PDMA=y, and switched the
whole GPU-direct path off when either was missing. Measured on
7.2.4-070204-generic: PCI_P2PDMA is set, ibv_reg_dmabuf_mr resolves in
rdma-core 61.0, hsa_amd_portable_export_dmabuf succeeds -- and
DMABUF_MOVE_NOTIFY is simply absent from the config. One missing line
turned off a path that works. The gate is gone; the call is tried and its
result believed.

With that, a GPU-direct transfer that cannot be done directly must not be
quietly done some other way. HIPOBJ_REQUIRE_GPU_DIRECT=1 makes the client
refuse to fall back to host staging for a device pointer, and say so,
rather than return success for a transfer the caller did not ask for.
This is what turns the GPU Direct CI lane red: it was passing by faking.

Failures also say which failure. Every RDMA error used to surface as one
anonymous -1, which the HTTP layer rendered as "500 RDMA PUT failed" --
six distinct faults, from an unparseable token to a completion error,
indistinguishable from a server with no RDMA at all. pollCompletion now
reports the work-completion status, opcode and vendor_err, and each
failure path names itself.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>

Author: sbates130272

include/hipobj.h

@@ -403,6 +403,29 @@ HIPOBJ_API hipObjError_t hipObjPutV2(const char* bucket, const char* key,
 
 #endif /* HIPOBJECT_V2_API */
 
+/*! @brief hipObjBufSync direction: device -> staging buffer @ingroup io */
+#define HIPOBJ_SYNC_TO_HOST 0
+/*! @brief hipObjBufSync direction: staging buffer -> device @ingroup io */
+#define HIPOBJ_SYNC_TO_DEVICE 1
+
+/*!
+ * @brief Stage a registered buffer between device and host memory
+ *
+ * When a device pointer cannot be registered with the NIC directly (no
+ * dmabuf support, or a fabric with no peer-to-peer path to the GPU),
+ * hipObjBufRegister() registers a host staging buffer instead and the
+ * peer's RDMA reads and writes land there. Callers driving a transfer
+ * through hipObjGetRdmaToken() must therefore call this with
+ * HIPOBJ_SYNC_TO_HOST before a PUT and HIPOBJ_SYNC_TO_DEVICE after a
+ * GET. It is a no-op, reporting success, for a directly registered
+ * buffer; hipObjGet()/hipObjPut() and the V2 entry points do it
+ * themselves.
+ *
+ * @ingroup io
+ */
+HIPOBJ_API hipObjError_t hipObjBufSync(void* devPtr, size_t size, off_t offset,
+                                       int direction);
+
 /*!
  * @brief Mint a hex-encoded RC RDMA token for a registered buffer
  *

integrations/minio-cpp/examples/getput-rdma.cpp

@@ -5,19 +5,62 @@
 
 /* AMD port of minio-cpp GetPutRDMA: PUT + GET over hipObject RDMA */
 
+#include <cstdint>
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
+#include <ctime>
 #include <fstream>
 #include <iostream>
 #include <string>
+#include <vector>
 
 #include <hip/hip_runtime.h>
 
 #include <miniocpp/client.h>
+#include <unistd.h>
 
 #include "hipobj_minio/client.h"
 
+namespace {
+
+/* A position-dependent pattern, not a constant fill.
+ *
+ * This used to memset the buffer to 'A' and check afterwards that every byte
+ * was still 'A', which cannot distinguish a correct transfer from a
+ * byte-swapped, misaligned, short, duplicated or scattered one -- every one
+ * of those returns a buffer full of 'A'. A cheap LCG keyed on the offset
+ * makes each byte depend on where it is, so any of those failures shows up
+ * as a mismatch at a specific offset, which is also the first thing you want
+ * to know when it does. */
+uint8_t PatternByte(size_t offset, uint32_t seed) {
+  uint64_t x = (static_cast<uint64_t>(offset) + 1) * 6364136223846793005ULL +
+               seed;
+  x ^= x >> 33;
+  x *= 0xff51afd7ed558ccdULL;
+  x ^= x >> 33;
+  return static_cast<uint8_t>(x & 0xff);
+}
+
+void FillPattern(char* buf, size_t size, uint32_t seed) {
+  for (size_t i = 0; i < size; ++i) {
+    buf[i] = static_cast<char>(PatternByte(i, seed));
+  }
+}
+
+/* Returns the offset of the first mismatch, or size if the blob read back
+ * matches the blob written byte for byte. */
+size_t FirstMismatch(const char* buf, size_t size, uint32_t seed) {
+  for (size_t i = 0; i < size; ++i) {
+    if (static_cast<uint8_t>(buf[i]) != PatternByte(i, seed)) {
+      return i;
+    }
+  }
+  return size;
+}
+
+} // namespace
+
 int main(int argc, char* argv[]) {
   if (argc < 4) {
     std::cerr << "usage: " << argv[0]
@@ -38,6 +81,15 @@ int main(int argc, char* argv[]) {
     gpu_enabled = std::string(argv[5]) == "gpu";
   }
 
+  /* Seeded per run, and printed, so a stale object left in the store by an
+   * earlier run cannot satisfy the readback. */
+  const uint32_t seed = static_cast<uint32_t>(::getpid()) ^
+                        static_cast<uint32_t>(std::time(nullptr));
+  std::cout << "Pattern seed " << seed << "\n";
+
+  std::vector<char> expected(bufsize);
+  FillPattern(expected.data(), bufsize, seed);
+
   minio::s3::BaseUrl base_url(host, false, "us-east-1");
   minio::creds::StaticProvider provider(access_key, secret_key);
   hipobj::minio::Client client(base_url, &provider);
@@ -51,9 +103,9 @@ int main(int argc, char* argv[]) {
       std::cerr << "hipMalloc failed: " << err << std::endl;
       return 1;
     }
-    err = hipMemset(dev_ptr, 'A', bufsize);
+    err = hipMemcpy(dev_ptr, expected.data(), bufsize, hipMemcpyHostToDevice);
     if (err != hipSuccess) {
-      std::cerr << "hipMemset failed: " << err << std::endl;
+      std::cerr << "hipMemcpy H2D failed: " << err << std::endl;
       (void)hipFree(dev_ptr);
       return 1;
     }
@@ -67,7 +119,7 @@ int main(int argc, char* argv[]) {
       std::cerr << "posix_memalign failed\n";
       return 1;
     }
-    std::memset(bufptr, 'A', bufsize);
+    std::memcpy(bufptr, expected.data(), bufsize);
     std::cout << "Host buffer " << bufsize << " bytes\n";
   }
 
@@ -89,11 +141,14 @@ int main(int argc, char* argv[]) {
   }
   std::cout << "PUT ok etag=" << presp.etag << std::endl;
 
+  /* Clobber the buffer before the GET. Without this the readback could be
+   * satisfied by whatever the PUT left behind, and a GET that transferred
+   * nothing at all would still verify. */
   if (gpu_enabled) {
-    (void)hipMemset(dev_ptr, 'U', bufsize);
+    (void)hipMemset(dev_ptr, 0x55, bufsize);
     (void)hipDeviceSynchronize();
   } else {
-    std::memset(bufptr, 'U', bufsize);
+    std::memset(bufptr, 0x55, bufsize);
   }
 
   minio::s3::GetObjectArgs gargs;
@@ -153,15 +208,25 @@ int main(int argc, char* argv[]) {
   out.close();
   std::cout << "Wrote output.bin (" << bufsize << " bytes)\n";
 
-  bool ok = true;
-  for (size_t i = 0; i < bufsize; ++i) {
-    if (hostptr[i] != 'A') {
-      ok = false;
-      break;
-    }
+  size_t bad = FirstMismatch(hostptr, bufsize, seed);
+  bool ok = bad == bufsize;
+  if (ok) {
+    std::cout << "Data integrity check passed (" << bufsize << " bytes, seed "
+              << seed << ")\n";
+  } else {
+    std::cout << "Data integrity check FAILED at offset " << bad
+              << ": expected 0x" << std::hex
+              << static_cast<int>(PatternByte(bad, seed)) << " got 0x"
+              << static_cast<int>(static_cast<uint8_t>(hostptr[bad]))
+              << std::dec << "\n";
   }
-  std::cout << (ok ? "Data integrity check passed\n"
-                   : "Data integrity check FAILED\n");
+
+  /* Whether the transfers actually used RDMA. The client falls back to
+   * ordinary HTTP silently when no NIC is available, so without this a lane
+   * that is meant to be exercising the RDMA data path passes over TCP. */
+  std::cout << hipobj::minio::TransferStatsLine(
+                 hipobj::minio::TransferStatsSnapshot())
+            << "\n";
 
   free(hostptr);
   if (gpu_enabled) {

integrations/minio-cpp/include/hipobj_minio/client.h

@@ -5,26 +5,67 @@
 
 #pragma once
 
+#include <cstddef>
+#include <cstdint>
+#include <string>
+
 #include <miniocpp/args.h>
 #include <miniocpp/client.h>
 #include <miniocpp/credentials.h>
 #include <miniocpp/response.h>
 
 namespace hipobj::minio {
 
+/* How the last transfers actually went out.
+ *
+ * Client::PutObject and Client::GetObject fall back to ordinary HTTP in two
+ * places -- when hipObjInit cannot find a NIC, and when an RDMA transfer
+ * returns an error -- and the caller sees a successful response either way.
+ * That is the right behaviour for a client and the wrong behaviour for a
+ * test: an integration lane that is supposed to be proving the RDMA data
+ * path will pass just as happily over TCP. These counters are how a caller
+ * tells the two apart after the fact.
+ *
+ * Process-wide and monotonic. Snapshot with TransferStatsSnapshot() before
+ * the transfers you care about and subtract. */
+struct TransferStats {
+  uint64_t rdmaPuts = 0;
+  uint64_t rdmaGets = 0;
+  uint64_t httpPuts = 0;
+  uint64_t httpGets = 0;
+  uint64_t rdmaBytes = 0;
+  uint64_t httpBytes = 0;
+};
+
+/* A snapshot of the process-wide counters. */
+TransferStats TransferStatsSnapshot();
+
+/* Reset the counters to zero. */
+void TransferStatsReset();
+
+/* One line, stable and greppable, for a CI step to assert on:
+ *
+ *   hipobj-stats: rdma_put=1 rdma_get=1 http_put=0 http_get=0
+ *                 rdma_bytes=131072 http_bytes=0
+ *
+ * A lane that requires RDMA asserts http_put=0 and http_get=0; a lane that
+ * cannot have it (no verbs device, e.g. a plain container) asserts the
+ * reverse, so the expectation is written down either way. */
+std::string TransferStatsLine(const TransferStats& stats);
+
 class Client {
 public:
-  Client(minio::s3::BaseUrl base_url, minio::creds::Provider* provider);
+  Client(::minio::s3::BaseUrl base_url, ::minio::creds::Provider* provider);
 
-  minio::s3::PutObjectResponse PutObject(minio::s3::PutObjectArgs args);
-  minio::s3::GetObjectResponse GetObject(minio::s3::GetObjectArgs args);
+  ::minio::s3::PutObjectResponse PutObject(::minio::s3::PutObjectArgs args);
+  ::minio::s3::GetObjectResponse GetObject(::minio::s3::GetObjectArgs args);
 
   bool RdmaAvailable() const;
 
 private:
-  minio::s3::BaseUrl base_url_;
-  minio::creds::Provider* provider_ = nullptr;
-  minio::s3::Client s3_client_;
+  ::minio::s3::BaseUrl base_url_;
+  ::minio::creds::Provider* provider_ = nullptr;
+  ::minio::s3::Client s3_client_;
 };
 
 } // namespace hipobj::minio

integrations/minio-cpp/include/hipobj_minio/context.h

@@ -14,10 +14,10 @@
 namespace hipobj::minio {
 
 struct S3RdmaContext {
-  minio::creds::Provider* provider = nullptr;
+  ::minio::creds::Provider* provider = nullptr;
   std::string bucket;
   std::string object;
-  minio::s3::BaseUrl url;
+  ::minio::s3::BaseUrl url;
   std::string region;
   std::string uploadId;
   unsigned int partNumber = 0;