From 3b65453acf1b477c3abb67c441d9eb69fe88bf2e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E4=BA=91=E6=8C=9A?= Date: Sun, 2 Aug 2026 03:03:14 -0500 Subject: [PATCH 1/2] perf: pipeline GPU-initiated PUT completions --- P0_PERF_JOURNAL.md | 92 +++++++++++++ .../gpunetio_verbs_put_bw_kernel.cu | 125 +++++++++++++++++- .../gpunetio_verbs_put_bw_main.cpp | 57 +++++++- .../gpunetio_verbs_put_bw_sample.cpp | 24 +++- examples/verbs_common.h | 3 + 5 files changed, 284 insertions(+), 17 deletions(-) create mode 100644 P0_PERF_JOURNAL.md diff --git a/P0_PERF_JOURNAL.md b/P0_PERF_JOURNAL.md new file mode 100644 index 0000000..0c566a6 --- /dev/null +++ b/P0_PERF_JOURNAL.md @@ -0,0 +1,92 @@ +# 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. + +## 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. diff --git a/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_kernel.cu b/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_kernel.cu index 380a1e3..d171156 100644 --- a/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_kernel.cu +++ b/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( + 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( + qp, ticket) != 0) { +#if ENABLE_DEBUG == 1 + printf("Error CQE!\n"); +#endif + } +} + +template +__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 +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<<>>( + 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<<>>( - 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 + <<>>( + 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<<>>( qp, num_iters, data_size, src_buf, src_buf_mkey, dst_buf, dst_buf_mkey); + } result = cudaGetLastError(); if (cudaSuccess != result) { diff --git a/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_main.cpp b/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_main.cpp index d2e90f5..f21bddf 100644 --- a/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_main.cpp +++ b/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 (Client only)\n" @@ -103,11 +107,52 @@ int main(int argc, char **argv) { << " -l \n" << " -p \n" + << " -t \n" + << " -w \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) { diff --git a/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_sample.cpp b/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_sample.cpp index 133bfeb..cef43ac 100644 --- a/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_sample.cpp +++ b/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_sample.cpp @@ -298,6 +298,8 @@ doca_error_t verbs_client(struct verbs_config *cfg) { pthread_t thread_id; struct cpu_proxy_args args; struct doca_gpu_dev_verbs_qp *qp_gpu; + uint32_t cuda_blocks; + uint32_t cuda_threads_per_block; int ret; resources.conn_socket = -1; @@ -366,10 +368,18 @@ doca_error_t verbs_client(struct verbs_config *cfg) { goto destroy_events; } + cuda_blocks = (resources.scope == DOCA_GPUNETIO_VERBS_EXEC_SCOPE_THREAD && + resources.cuda_threads >= VERBS_CUDA_BLOCK && + resources.cuda_threads % VERBS_CUDA_BLOCK == 0) + ? VERBS_CUDA_BLOCK + : 1; + cuda_threads_per_block = resources.cuda_threads / cuda_blocks; + DOCA_LOG(LOG_INFO, - "Launching gpunetio_verbs_put_bw kernel with %d CUDA Blocks, %d CUDA threads each, %d " - "total number of iterations, %d iterations per cuda thread, %d nic handler, %s scope", - VERBS_CUDA_BLOCK, resources.cuda_threads / VERBS_CUDA_BLOCK, resources.num_iters, + "Launching gpunetio_verbs_put_bw kernel with %d CUDA Blocks, %d CUDA threads each, " + "window depth %d, %d total number of iterations, %d iterations per cuda thread, %d " + "nic handler, %s scope", + cuda_blocks, cuda_threads_per_block, cfg->put_window_depth, resources.num_iters, resources.num_iters / resources.cuda_threads, resources.nic_handler, (resources.scope == DOCA_GPUNETIO_VERBS_EXEC_SCOPE_THREAD) ? "THREAD" : "WARP"); @@ -399,8 +409,8 @@ doca_error_t verbs_client(struct verbs_config *cfg) { for (int idx = 0; idx < NUM_MSG_SIZE; idx++) { /* Warmup per size*/ status = gpunetio_verbs_put_bw( - cstream, qp_gpu, resources.num_iters, VERBS_CUDA_BLOCK, - resources.cuda_threads / VERBS_CUDA_BLOCK, message_size[idx], resources.data_buf[idx], + cstream, qp_gpu, resources.num_iters, cuda_blocks, cuda_threads_per_block, + cfg->put_window_depth, message_size[idx], resources.data_buf[idx], htobe32(resources.data_mr[idx]->lkey), (uint8_t *)(resources.remote_data_buf[idx]), htobe32(resources.remote_data_mkey[idx]), resources.scope); if (status != DOCA_SUCCESS) { @@ -418,8 +428,8 @@ doca_error_t verbs_client(struct verbs_config *cfg) { } status = gpunetio_verbs_put_bw( - cstream, qp_gpu, resources.num_iters, VERBS_CUDA_BLOCK, - resources.cuda_threads / VERBS_CUDA_BLOCK, message_size[idx], resources.data_buf[idx], + cstream, qp_gpu, resources.num_iters, cuda_blocks, cuda_threads_per_block, + cfg->put_window_depth, message_size[idx], resources.data_buf[idx], htobe32(resources.data_mr[idx]->lkey), (uint8_t *)(resources.remote_data_buf[idx]), htobe32(resources.remote_data_mkey[idx]), resources.scope); if (status != DOCA_SUCCESS) { diff --git a/examples/verbs_common.h b/examples/verbs_common.h index a9b829d..c550e9d 100644 --- a/examples/verbs_common.h +++ b/examples/verbs_common.h @@ -89,6 +89,7 @@ struct verbs_config { std::string server_ip_addr; /* DOCA device name */ uint32_t num_iters; /* total number orations per cuda kernel */ uint32_t cuda_threads; /* cuda threads per cuda block */ + uint32_t put_window_depth; /* outstanding PUTs per CUDA thread */ enum doca_gpu_dev_verbs_nic_handler nic_handler; uint8_t exec_scope; enum doca_gpu_verbs_send_dbr_mode_ext @@ -266,6 +267,7 @@ extern "C" { * @num_iters [in]: Total number of iterations * @cuda_blocks [in]: Number of CUDA blocks to launch the kernel * @cuda_threads [in]: Number of CUDA threads to launch the kernel + * @window_depth [in]: Maximum outstanding PUTs per CUDA thread * @data_size [in]: Data buffer size (number of bytes) * @src_buf [in]: Source GPU data buffer address * @src_buf_mkey [in]: Source GPU data buffer memory key @@ -322,6 +324,7 @@ doca_error_t gpunetio_verbs_write_lat( */ 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); From 5d428f43eaf6ccfdb6db99ef43e9d4a0e0bc49b5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E4=BA=91=E6=8C=9A?= Date: Sun, 2 Aug 2026 07:03:29 -0500 Subject: [PATCH 2/2] fix: validate PUT peer workload configuration --- P0_PERF_JOURNAL.md | 16 +++ .../gpunetio_verbs_put_bw_sample.cpp | 119 +++++++++++++++--- 2 files changed, 121 insertions(+), 14 deletions(-) diff --git a/P0_PERF_JOURNAL.md b/P0_PERF_JOURNAL.md index 0c566a6..70f2a70 100644 --- a/P0_PERF_JOURNAL.md +++ b/P0_PERF_JOURNAL.md @@ -80,6 +80,22 @@ - 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. diff --git a/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_sample.cpp b/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_sample.cpp index cef43ac..459bdd2 100644 --- a/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_sample.cpp +++ b/examples/gpunetio_verbs_put_bw/gpunetio_verbs_put_bw_sample.cpp @@ -40,6 +40,90 @@ int message_size[NUM_MSG_SIZE] = {1, 64, 128, 256, 512, 1024, 204 8192, 16384, 32768, 65536, 131072, 262144, 524288, 1048576}; volatile bool server_force_quit = false; +#define PUT_BW_PROTOCOL_MAGIC 0x47505550U +#define PUT_BW_PROTOCOL_VERSION 1U + +struct put_bw_peer_config { + uint32_t magic; + uint32_t version; + uint32_t cuda_threads; + uint32_t num_qps; + uint32_t message_size; +}; + +static bool socket_send_all(int socket_fd, const void *buffer, size_t length) { + const uint8_t *cursor = (const uint8_t *)buffer; + + while (length != 0) { + ssize_t sent = send(socket_fd, cursor, length, MSG_NOSIGNAL); + if (sent < 0 && errno == EINTR) + continue; + if (sent <= 0) + return false; + cursor += sent; + length -= (size_t)sent; + } + + return true; +} + +static bool socket_recv_all(int socket_fd, void *buffer, size_t length) { + uint8_t *cursor = (uint8_t *)buffer; + + while (length != 0) { + ssize_t received = recv(socket_fd, cursor, length, 0); + if (received < 0 && errno == EINTR) + continue; + if (received <= 0) + return false; + cursor += received; + length -= (size_t)received; + } + + return true; +} + +static doca_error_t exchange_peer_config(struct verbs_resources *resources) { + const struct put_bw_peer_config local = { + .magic = htonl(PUT_BW_PROTOCOL_MAGIC), + .version = htonl(PUT_BW_PROTOCOL_VERSION), + .cuda_threads = htonl(resources->cuda_threads), + .num_qps = htonl(1), + .message_size = htonl(0), + }; + struct put_bw_peer_config remote = {0}; + + if (!socket_send_all(resources->conn_socket, &local, sizeof(local)) || + !socket_recv_all(resources->conn_socket, &remote, sizeof(remote))) { + DOCA_LOG(LOG_ERR, "Failed to exchange PUT benchmark configuration"); + return DOCA_ERROR_CONNECTION_ABORTED; + } + + remote.magic = ntohl(remote.magic); + remote.version = ntohl(remote.version); + remote.cuda_threads = ntohl(remote.cuda_threads); + remote.num_qps = ntohl(remote.num_qps); + remote.message_size = ntohl(remote.message_size); + + if (remote.magic != PUT_BW_PROTOCOL_MAGIC || remote.version != PUT_BW_PROTOCOL_VERSION) { + DOCA_LOG(LOG_ERR, "Unsupported PUT benchmark peer protocol (magic 0x%x, version %u)", + remote.magic, remote.version); + return DOCA_ERROR_NOT_SUPPORTED; + } + + if (remote.cuda_threads != resources->cuda_threads || remote.num_qps != 1 || + remote.message_size != 0) { + DOCA_LOG(LOG_ERR, + "PUT peer configuration mismatch: threads local=%u remote=%u, QPs local=1 " + "remote=%u, message size local=0 remote=%u", + resources->cuda_threads, remote.cuda_threads, remote.num_qps, + remote.message_size); + return DOCA_ERROR_INVALID_VALUE; + } + + return DOCA_SUCCESS; +} + /* * Server validates data from client at the end of the test */ @@ -155,17 +239,21 @@ static doca_error_t create_local_memory_object(struct verbs_resources *resources } static doca_error_t exchange_params_with_remote_peer(struct verbs_resources *resources) { + doca_error_t status = exchange_peer_config(resources); + if (status != DOCA_SUCCESS) + return status; + if (resources->cfg->is_server) { // Server sends local info for (int idx = 0; idx < NUM_MSG_SIZE; idx++) { uint64_t local_addr = (uint64_t)resources->data_buf[idx]; - if (send(resources->conn_socket, &local_addr, sizeof(uint64_t), 0) < 0) { + if (!socket_send_all(resources->conn_socket, &local_addr, sizeof(local_addr))) { DOCA_LOG(LOG_ERR, "Failed to send local buffer address"); return DOCA_ERROR_CONNECTION_ABORTED; } - if (send(resources->conn_socket, &resources->data_mr[idx]->rkey, sizeof(uint32_t), 0) < - 0) { + if (!socket_send_all(resources->conn_socket, &resources->data_mr[idx]->rkey, + sizeof(resources->data_mr[idx]->rkey))) { DOCA_LOG(LOG_ERR, "Failed to send local MKEY"); return DOCA_ERROR_CONNECTION_ABORTED; } @@ -173,47 +261,50 @@ static doca_error_t exchange_params_with_remote_peer(struct verbs_resources *res } else { // Client waits for server info for (int idx = 0; idx < NUM_MSG_SIZE; idx++) { - if (recv(resources->conn_socket, &resources->remote_data_buf[idx], sizeof(uint64_t), - 0) < 0) { + if (!socket_recv_all(resources->conn_socket, &resources->remote_data_buf[idx], + sizeof(resources->remote_data_buf[idx]))) { DOCA_LOG(LOG_ERR, "Failed to receive remote buffer address "); return DOCA_ERROR_CONNECTION_ABORTED; } - if (recv(resources->conn_socket, &resources->remote_data_mkey[idx], sizeof(uint32_t), - 0) < 0) { + if (!socket_recv_all(resources->conn_socket, &resources->remote_data_mkey[idx], + sizeof(resources->remote_data_mkey[idx]))) { DOCA_LOG(LOG_ERR, "Failed to receive remote MKEY, err = %d", errno); return DOCA_ERROR_CONNECTION_ABORTED; } } } - if (send(resources->conn_socket, &resources->local_qp_number, sizeof(uint32_t), 0) < 0) { + if (!socket_send_all(resources->conn_socket, &resources->local_qp_number, + sizeof(resources->local_qp_number))) { DOCA_LOG(LOG_ERR, "Failed to send local QP number"); return DOCA_ERROR_CONNECTION_ABORTED; } - if (recv(resources->conn_socket, &resources->remote_qp_number, sizeof(uint32_t), 0) < 0) { + if (!socket_recv_all(resources->conn_socket, &resources->remote_qp_number, + sizeof(resources->remote_qp_number))) { DOCA_LOG(LOG_ERR, "Failed to receive remote QP number, err = %d", errno); return DOCA_ERROR_CONNECTION_ABORTED; } - if (send(resources->conn_socket, &resources->gid.raw, sizeof(resources->gid.raw), 0) < 0) { + if (!socket_send_all(resources->conn_socket, &resources->gid.raw, + sizeof(resources->gid.raw))) { DOCA_LOG(LOG_ERR, "Failed to send local GID address"); return DOCA_ERROR_CONNECTION_ABORTED; } - if (recv(resources->conn_socket, &resources->remote_gid.raw, sizeof(resources->gid.raw), 0) < - 0) { + if (!socket_recv_all(resources->conn_socket, &resources->remote_gid.raw, + sizeof(resources->remote_gid.raw))) { DOCA_LOG(LOG_ERR, "Failed to receive remote GID address, err = %d", errno); return DOCA_ERROR_CONNECTION_ABORTED; } - if (send(resources->conn_socket, &resources->lid, sizeof(uint32_t), 0) < 0) { + if (!socket_send_all(resources->conn_socket, &resources->lid, sizeof(resources->lid))) { DOCA_LOG(LOG_ERR, "Failed to send local GID address"); return DOCA_ERROR_CONNECTION_ABORTED; } - if (recv(resources->conn_socket, &resources->dlid, sizeof(uint32_t), 0) < 0) { + if (!socket_recv_all(resources->conn_socket, &resources->dlid, sizeof(resources->dlid))) { DOCA_LOG(LOG_ERR, "Failed to receive remote GID address, err = %d", errno); return DOCA_ERROR_CONNECTION_ABORTED; }