gpunetio_verbs_put_bw: pipeline PUT completions with a fixed in-flight window

P0_PERF_JOURNAL.md

@@ -0,0 +1,108 @@
+# GPUNetIO P0 Performance Journal
+
+## Scope
+
+- Goal: identify and optimize one P0 steady-state performance bottleneck.
+- Target: 4x NVIDIA GB200 host; use one local GPU/NIC path for controlled runs first.
+- H20 comparison branch: `tyh` at `aa7f5d7af3910f7343120902c79cc54865a50446`.
+- Implementation branch: `p0/put-window`, based on official main
+  `df883ff6b5b52f793a7a7eea434485e38891c12f` (live-checked 2026-08-01).
+- Local coroutine baseline is based on upstream `f9525001f694c472ba2d901df20325c9aafde8b4`.
+
+## Prior-context hypotheses (not evidence)
+
+- The H20 coroutine proof of concept improves throughput by keeping multiple PUTs outstanding before polling completion.
+- Its scheduler is much heavier than the required fixed FIFO window: global task queues, global task contexts, and per-call CUDA allocations/copies.
+- It has known correctness/accounting hazards, so all claims require fresh GB200 validation.
+
+## Host inventory
+
+- CUDA toolkit: 13.1 (`nvcc` 13.1.80).
+- GPU: 4x NVIDIA GB200, each 189471 MiB, idle at initial inspection.
+- GPU interconnect: NV18 between every GPU pair.
+- NICs visible to NVML: `mlx5_0`, `mlx5_1`, `mlx5_4`, `mlx5_5`, `mlx5_bond_0`, `mlx5_bond_1`.
+- GPU0/GPU1 are NUMA-local (`NODE`) to `mlx5_0`, `mlx5_1`, and `mlx5_bond_0`; GPU2/GPU3 are local to `mlx5_4`, `mlx5_5`, and `mlx5_bond_1`.
+
+## Experiment log
+
+### E0 - Build readiness
+
+- Status: pass.
+- Command: `make -j16 CUDA_ARCH=100`.
+- Result: full library and all official examples build with CUDA 13.1 for `sm_100`.
+
+### E1 - Reproduce current sync and coroutine baselines
+
+- Status: pass on the H20-derived `tyh` branch, two repetitions per cell on GB200.
+- Configuration: GPU0 `08:01:00.0`, `mlx5_0`, RoCEv2 GID 3, thread scope,
+  one CUDA thread, 4096 operations, coroutine depth 1/2/4/8.
+- Selected median Gbps (`sync`, `coro2`, `coro4`, `coro8`):
+  - 4 KiB: 3.796, 5.560, 4.225, 4.311.
+  - 64 KiB: 52.311, 92.884, 68.570, 73.235.
+  - 1 MiB: 290.979, 378.329, 391.060, 389.704.
+- All completed server runs passed full-buffer validation.
+- Conclusion: deferring completion is useful, but the global task scheduler is not a good
+  production mechanism and depth beyond four does not help large-message bandwidth.
+
+### E2 - Fixed-window kernel
+
+- Status: pass on official v4.
+- Implementation: compile-time FIFO depths 2/4/8; PUT tickets are scalarized into registers;
+  poll the oldest ticket, immediately refill that slot, then drain exactly at the end.
+- Depth sweep: one CUDA thread, 8192 operations, two repetitions per depth. Depth four reached
+  392.884 Gbps at 1 MiB; depths 8/16/32 in the exploratory version were all within 0.02 Gbps,
+  so the production interface is capped at 8.
+- Compiler evidence for depth four on `sm_100`: 60 registers, zero stack, zero local memory,
+  zero shared memory, zero barriers.
+
+### E3 - Final controlled comparison
+
+- Status: pass.
+- Configuration: official v4, GPU0 `08:01:00.0`, `mlx5_0`, GID 3, direct GPU doorbell,
+  thread scope, 8192 operations, three repetitions per final cell.
+- Representative commands (start server first, then client; `127.0.0.1` is OOB control only):
+  - Server: `./examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw -g 08:01:00.0 -d mlx5_0 -l 3 -t 1 -w 4 -i 8192`
+  - Client: `./examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw -g 08:01:00.0 -d mlx5_0 -l 3 -c 127.0.0.1 -e 0 -t 1 -w 4 -i 8192`
+
+| Message | 1 thread, depth 1 | 1 thread, depth 4 | Gain |
+|---:|---:|---:|---:|
+| 4 KiB | 3.658 Gbps | 5.603 Gbps | +53.2% |
+| 64 KiB | 50.632 Gbps | 89.657 Gbps | +77.1% |
+| 1 MiB | 287.670 Gbps | 392.902 Gbps | +36.6% |
+
+- 1 MiB depth-four range: 392.892-392.915 Gbps across three final runs. An earlier five-run
+  repetition measured 392.867-392.907 Gbps (392.896 Gbps median).
+- Official default path (512 threads, depth 1, 2048 operations): 389.574 Gbps median at 1 MiB.
+  The one-thread depth-four path reaches 100.85% of that throughput while launching one partial
+  warp instead of 16 full warps.
+- WARP-scope compatibility check (32 threads, depth 1) passed validation and reached
+  388.801 Gbps at 1 MiB.
+- Every final server log reported `Validation successfull! Data received correctly from client`.
+- Invalid depths, zero threads, and invalid WARP thread counts fail before resource creation.
+
+### E4 - Independent review hardening (2026-08-02)
+
+- Status: pass.
+- Upstream check: NVIDIA `main` remains `df883ff`; the checkpoint has no rebase drift.
+- Confirmed issue: P0 added the `-t` control, but the peers previously allocated and addressed
+  buffers from independently selected thread counts. A mismatched client could therefore address
+  beyond the server registration.
+- Fix: exchange a versioned PUT workload descriptor before RDMA metadata, reject mismatched
+  producer geometry, and use exact-length socket transfers so EOF and short I/O are failures.
+  The descriptor reserves stable fields for later QP-count and message-size selection.
+- Build: `make -j16 CUDA_ARCH=100` passed for the library and all examples.
+- Matched regression: one thread, depth four, 8192 operations passed; representative throughput
+  was 5.604 Gbps at 4 KiB, 89.656 Gbps at 64 KiB, and 392.916 Gbps at 1 MiB.
+- Negative test: server `-t 1` versus client `-t 2` made both peers exit nonzero with the explicit
+  configuration-mismatch diagnostic before QP connection or kernel launch.
+
+## Decisions
+
+- Optimize one steady-state data-plane bottleneck, not broad control-plane cleanup.
+- Keep source claims separate from measured results.
+- Do not count allocation/setup removal as network throughput unless it is inside the timed interval.
+- Keep upstream depth-one launch on the original kernel so the default hot path and its block
+  barrier are unchanged.
+- Select depth four as the GB200 optimum; keep 2 and 8 available for other message sizes/hosts.
+- Remaining boundary: small/medium messages plateau on per-WQE submission/doorbell cost. This
+  patch fixes completion-depth starvation; it does not claim to solve submission batching.

examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_kernel.cu

@@ -108,10 +108,106 @@ __global__ void put_bw(struct doca_gpu_dev_verbs_qp *qp, uint32_t num_iters, uin
     }
 }
 
+__device__ static __forceinline__ void post_put_bw(
+    struct doca_gpu_dev_verbs_qp *qp, uint32_t data_size, uint8_t *src_buf,
+    uint32_t src_buf_mkey, uint8_t *dst_buf, uint32_t dst_buf_mkey, uint32_t tidx,
+    doca_gpu_dev_verbs_ticket_t *ticket) {
+    doca_gpu_dev_verbs_put<DOCA_GPUNETIO_VERBS_RESOURCE_SHARING_MODE_GPU,
+                           DOCA_GPUNETIO_VERBS_NIC_HANDLER_AUTO,
+                           DOCA_GPUNETIO_VERBS_EXEC_SCOPE_THREAD>(
+        qp,
+        doca_gpu_dev_verbs_addr{.addr = (uint64_t)(dst_buf + (data_size * tidx)),
+                                .key = (uint32_t)dst_buf_mkey},
+        doca_gpu_dev_verbs_addr{.addr = (uint64_t)(src_buf + (data_size * tidx)),
+                                .key = (uint32_t)src_buf_mkey},
+        data_size, ticket);
+}
+
+__device__ static __forceinline__ void poll_put_bw(struct doca_gpu_dev_verbs_qp *qp,
+                                                    doca_gpu_dev_verbs_ticket_t ticket) {
+    if (doca_gpu_dev_verbs_poll_cq_at<DOCA_GPUNETIO_VERBS_RESOURCE_SHARING_MODE_GPU>(
+            qp, ticket) != 0) {
+#if ENABLE_DEBUG == 1
+        printf("Error CQE!\n");
+#endif
+    }
+}
+
+template <uint32_t window_depth>
+__global__ void put_bw_window(struct doca_gpu_dev_verbs_qp *qp, uint32_t num_iters,
+                              uint32_t data_size, uint8_t *src_buf, uint32_t src_buf_mkey,
+                              uint8_t *dst_buf, uint32_t dst_buf_mkey) {
+    doca_gpu_dev_verbs_ticket_t tickets[window_depth];
+    const uint32_t tidx = threadIdx.x + (blockIdx.x * blockDim.x);
+    const uint32_t stride = blockDim.x * gridDim.x;
+    uint32_t next_idx = tidx;
+    uint32_t initial_depth = 0;
+
+#pragma unroll
+    for (uint32_t slot = 0; slot < window_depth; ++slot) {
+        if (next_idx >= num_iters)
+            break;
+
+        post_put_bw(qp, data_size, src_buf, src_buf_mkey, dst_buf, dst_buf_mkey, tidx,
+                    &tickets[slot]);
+        next_idx += stride;
+        ++initial_depth;
+    }
+
+    if (initial_depth != window_depth) {
+#pragma unroll
+        for (uint32_t slot = 0; slot < window_depth; ++slot) {
+            if (slot < initial_depth)
+                poll_put_bw(qp, tickets[slot]);
+        }
+        return;
+    }
+
+    while ((uint64_t)next_idx + ((uint64_t)window_depth - 1) * stride < num_iters) {
+#pragma unroll
+        for (uint32_t slot = 0; slot < window_depth; ++slot) {
+            poll_put_bw(qp, tickets[slot]);
+            post_put_bw(qp, data_size, src_buf, src_buf_mkey, dst_buf, dst_buf_mkey, tidx,
+                        &tickets[slot]);
+            next_idx += stride;
+        }
+    }
+
+    uint32_t tail_depth = 0;
+#pragma unroll
+    for (uint32_t slot = 0; slot < window_depth; ++slot) {
+        poll_put_bw(qp, tickets[slot]);
+
+        if (next_idx < num_iters) {
+            post_put_bw(qp, data_size, src_buf, src_buf_mkey, dst_buf, dst_buf_mkey, tidx,
+                        &tickets[slot]);
+            next_idx += stride;
+            ++tail_depth;
+        }
+    }
+
+#pragma unroll
+    for (uint32_t slot = 0; slot < window_depth; ++slot) {
+        if (slot < tail_depth)
+            poll_put_bw(qp, tickets[slot]);
+    }
+}
+
+template <uint32_t window_depth>
+static void launch_put_bw_window(cudaStream_t stream, struct doca_gpu_dev_verbs_qp *qp,
+                                 uint32_t num_iters, uint32_t cuda_blocks,
+                                 uint32_t cuda_threads, uint32_t data_size, uint8_t *src_buf,
+                                 uint32_t src_buf_mkey, uint8_t *dst_buf,
+                                 uint32_t dst_buf_mkey) {
+    put_bw_window<window_depth><<<cuda_blocks, cuda_threads, 0, stream>>>(
+        qp, num_iters, data_size, src_buf, src_buf_mkey, dst_buf, dst_buf_mkey);
+}
+
 extern "C" {
 
 doca_error_t gpunetio_verbs_put_bw(cudaStream_t stream, struct doca_gpu_dev_verbs_qp *qp,
                                    uint32_t num_iters, uint32_t cuda_blocks, uint32_t cuda_threads,
+                                   uint32_t window_depth,
                                    uint32_t data_size, uint8_t *src_buf, uint32_t src_buf_mkey,
                                    uint8_t *dst_buf, uint32_t dst_buf_mkey,
                                    enum doca_gpu_dev_verbs_exec_scope scope) {
@@ -125,12 +221,33 @@ doca_error_t gpunetio_verbs_put_bw(cudaStream_t stream, struct doca_gpu_dev_verb
         return DOCA_ERROR_BAD_STATE;
     }
 
-    if (scope == DOCA_GPUNETIO_VERBS_EXEC_SCOPE_THREAD)
-        put_bw<DOCA_GPUNETIO_VERBS_EXEC_SCOPE_THREAD><<<cuda_blocks, cuda_threads, 0, stream>>>(
-            qp, num_iters, data_size, src_buf, src_buf_mkey, dst_buf, dst_buf_mkey);
-    else if (scope == DOCA_GPUNETIO_VERBS_EXEC_SCOPE_WARP)
+    if (scope == DOCA_GPUNETIO_VERBS_EXEC_SCOPE_THREAD) {
+        switch (window_depth) {
+            case 1:
+                put_bw<DOCA_GPUNETIO_VERBS_EXEC_SCOPE_THREAD>
+                    <<<cuda_blocks, cuda_threads, 0, stream>>>(
+                        qp, num_iters, data_size, src_buf, src_buf_mkey, dst_buf, dst_buf_mkey);
+                break;
+            case 2:
+                launch_put_bw_window<2>(stream, qp, num_iters, cuda_blocks, cuda_threads, data_size,
+                                        src_buf, src_buf_mkey, dst_buf, dst_buf_mkey);
+                break;
+            case 4:
+                launch_put_bw_window<4>(stream, qp, num_iters, cuda_blocks, cuda_threads, data_size,
+                                        src_buf, src_buf_mkey, dst_buf, dst_buf_mkey);
+                break;
+            case 8:
+                launch_put_bw_window<8>(stream, qp, num_iters, cuda_blocks, cuda_threads, data_size,
+                                        src_buf, src_buf_mkey, dst_buf, dst_buf_mkey);
+                break;
+            default:
+                DOCA_LOG(LOG_ERR, "Unsupported PUT window depth %u", window_depth);
+                return DOCA_ERROR_INVALID_VALUE;
+        }
+    } else if (scope == DOCA_GPUNETIO_VERBS_EXEC_SCOPE_WARP) {
         put_bw<DOCA_GPUNETIO_VERBS_EXEC_SCOPE_WARP><<<cuda_blocks, cuda_threads, 0, stream>>>(
             qp, num_iters, data_size, src_buf, src_buf_mkey, dst_buf, dst_buf_mkey);
+    }
 
     result = cudaGetLastError();
     if (cudaSuccess != result) {

examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_main.cpp

@@ -44,10 +44,11 @@ int main(int argc, char **argv) {
     verbs_cfg.gid_index = DEFAULT_GID_INDEX;
     verbs_cfg.num_iters = NUM_ITERS;
     verbs_cfg.cuda_threads = CUDA_THREADS_BW;
+    verbs_cfg.put_window_depth = 1;
     verbs_cfg.nic_handler = DOCA_GPUNETIO_VERBS_NIC_HANDLER_AUTO;
     verbs_cfg.exec_scope = DOCA_GPUNETIO_VERBS_EXEC_SCOPE_THREAD;
 
-    while ((option = getopt(argc, argv, "c:d:e:g:i:l:p:")) != -1) {
+    while ((option = getopt(argc, argv, "c:d:e:g:i:l:p:t:w:")) != -1) {
         switch (option) {
             case 'c': {
                 verbs_cfg.server_ip_addr = optarg;
@@ -74,11 +75,6 @@ int main(int argc, char **argv) {
             }
             case 'i': {
                 verbs_cfg.num_iters = std::atoi(optarg);
-                if ((verbs_cfg.num_iters % verbs_cfg.cuda_threads) != 0) {
-                    DOCA_LOG(LOG_ERR, "Iterations must be a multiple of CUDA threads number (%d)",
-                             verbs_cfg.cuda_threads);
-                    return 1;
-                }
                 break;
             }
             case 'l': {
@@ -93,6 +89,14 @@ int main(int argc, char **argv) {
                 }
                 break;
             }
+            case 't': {
+                verbs_cfg.cuda_threads = std::atoi(optarg);
+                break;
+            }
+            case 'w': {
+                verbs_cfg.put_window_depth = std::atoi(optarg);
+                break;
+            }
             default:
                 std::cerr << "Usage: " << argv[0] << "\n"
                           << " -c <server_ip> (Client only)\n"
@@ -103,11 +107,52 @@ int main(int argc, char **argv) {
                           << " -l <GID Index (default: 0)>\n"
                           << " -p <NIC handler. 0: AUTO 1: CPU PROXY 2: GPU SM_DB 6: GPU BF "
                              "(default: 0)>\n"
+                          << " -t <total CUDA threads (default: 512)>\n"
+                          << " -w <outstanding PUTs per thread: 1, 2, 4, or 8 (default: 1)>\n"
                           << std::endl;
                 return 1;
         }
     }
 
+    if (verbs_cfg.cuda_threads == 0 || verbs_cfg.num_iters == 0 ||
+        verbs_cfg.num_iters % verbs_cfg.cuda_threads != 0) {
+        DOCA_LOG(LOG_ERR,
+                 "Iterations and CUDA thread count must be non-zero, and iterations must be a "
+                 "multiple of the thread count (%d)",
+                 verbs_cfg.cuda_threads);
+        return 1;
+    }
+    if (verbs_cfg.put_window_depth != 1 && verbs_cfg.put_window_depth != 2 &&
+        verbs_cfg.put_window_depth != 4 && verbs_cfg.put_window_depth != 8) {
+        DOCA_LOG(LOG_ERR, "PUT window must be one of: 1, 2, 4, 8");
+        return 1;
+    }
+    if ((uint64_t)verbs_cfg.cuda_threads * verbs_cfg.put_window_depth >
+        VERBS_TEST_QUEUE_SIZE) {
+        DOCA_LOG(LOG_ERR, "CUDA threads x PUT window must not exceed SQ size (%d)",
+                 VERBS_TEST_QUEUE_SIZE);
+        return 1;
+    }
+    if (verbs_cfg.exec_scope == DOCA_GPUNETIO_VERBS_EXEC_SCOPE_WARP &&
+        verbs_cfg.put_window_depth != 1) {
+        DOCA_LOG(LOG_ERR, "PUT windows greater than one currently require THREAD scope");
+        return 1;
+    }
+    if (verbs_cfg.exec_scope == DOCA_GPUNETIO_VERBS_EXEC_SCOPE_WARP &&
+        (verbs_cfg.cuda_threads % DOCA_GPUNETIO_VERBS_WARP_SIZE != 0 ||
+         verbs_cfg.cuda_threads > 1024)) {
+        DOCA_LOG(LOG_ERR, "WARP scope requires a multiple of 32 threads in one block");
+        return 1;
+    }
+    if (verbs_cfg.exec_scope == DOCA_GPUNETIO_VERBS_EXEC_SCOPE_THREAD &&
+        (verbs_cfg.cuda_threads > 2048 ||
+         (verbs_cfg.cuda_threads > 1024 && verbs_cfg.cuda_threads % 2 != 0))) {
+        DOCA_LOG(LOG_ERR,
+                 "THREAD scope supports one block up to 1024 threads or two equal blocks up to "
+                 "2048 total threads");
+        return 1;
+    }
+
     if (verbs_cfg.is_server) {
         status = verbs_server(&verbs_cfg);
         if (status != DOCA_SUCCESS) {