From cac0598c671ba9505909c902d6027c1ae89244eb Mon Sep 17 00:00:00 2001 From: Ramya Gurunathan Date: Wed, 8 Jul 2026 10:05:53 -0400 Subject: [PATCH 01/18] #15 - Add fixed-size GEMM comparison to DGX Spark perf report Pin the GEMM dimension (n=1024, --workload-gemm-n) so the RoCE-vs-raw comparison isolates transport from problem size. The prior batch sweep confounded the two (n scaled with batch), which produced the misleading "gap is pipelining depth" reading. The window-matched fixed-n data shows the gap is real and receive-path/threading bound, not the GPU; adds the GEMMs/s + SM% table, flags the batch sweep as superseded, and documents the reproduce commands. Assisted-by: Claude Opus 4.8 Signed-off-by: Ramya Gurunathan --- docs/benchmarks/performance-dgx-spark.md | 86 ++++++++++++++++++++---- 1 file changed, 72 insertions(+), 14 deletions(-) diff --git a/docs/benchmarks/performance-dgx-spark.md b/docs/benchmarks/performance-dgx-spark.md index b0f5014..14f9cec 100644 --- a/docs/benchmarks/performance-dgx-spark.md +++ b/docs/benchmarks/performance-dgx-spark.md @@ -317,21 +317,60 @@ compute), so compute overlaps with receive on both and every cell is drop-free. fixed-batch rows above are the 8 MB native operating point, batch-matched to the raw path's ~8.19 MB reorder window (1024 packets × 8000 B). -The remaining RoCE-vs-raw gap on the heavy GEMM is **pipelining depth, not transport**: -at the 8 MB default, one RoCE message is a single GEMM with no neighbor to overlap, while -a raw burst packs ~10 GEMMs back-to-back. Sub-dividing the message into a smaller compute -batch packs several GEMMs per message and recovers most of the throughput — at a **4 MB -batch (2 GEMMs/message), 3-rep**: - -| Workload | Throughput | Drops | GPU SM% | vs 8 MB default | -| -------- | ---------: | ----- | ------: | --------------- | -| GEMM (FP32) | 76.3 ±0.2 Gb/s | 0 | 77.7% | 2.2× (was 35.0) | -| GEMM (FP16 tensor) | 92.3 ±2.0 Gb/s | 0 | 38.7% | within ~5% of raw (was 85.8) | - -The full curve is in the [batch-size sweep](#workload-batch-size-sweep) below. +### Fixed-size GEMM comparison + +The tables above size the GEMM from the working set, so sweeping the batch also changes +the matrix dimension (FLOPs ∝ n³) and conflates problem size with transport. To compare +transports cleanly we **pin the GEMM dimension** with `--workload-gemm-n 1024` +(env `GEMM_N` in `run_spark_bench.sh`): every call is an identical **1024×1024×1024 +matmul = 2.15 GFLOP**, and both transports consume **exactly 4 MB per GEMM** (one GEMM +per 4 MB window), so throughput *and* GEMMs/s are directly comparable. Fixed n=1024, +3 reps, 30 s each; RoCE at its 8 MB message (2 GEMMs/message), raw at a 4 MB reorder +window (4096 B × 1024 packets). + +| Transport | Workload | Throughput | GEMMs/s | GPU SM% | +| --------- | -------- | ---------: | ------: | ------: | +| RoCE (RC) | GEMM FP32 | 48.1 Gb/s | 1434 | 78% | +| Raw (DPDK) | GEMM FP32 | 103.8 Gb/s | 3095 | 40% | +| RoCE (RC) | GEMM FP16 (tensor) | 88.3 Gb/s | 2631 | 56% | +| Raw (DPDK) | GEMM FP16 (tensor) | 105.6 Gb/s | 3148 | 15% | + +Three things fall out, and they revise the "pipelining depth" reading of the earlier +batch sweep: + +- **The RoCE↔raw gap is real, and it is not the GPU.** At the *identical* fixed workload + raw sustains **2.2× the GEMM rate at FP32** (and ~1.2× at FP16) while using *less* than + half the SM. Raw has GPU headroom (15–40% SM) and is wire-limited (~104–106 Gb/s for + both precisions); RoCE is the side leaving the GPU idle-but-stalled. +- **GEMMs/s is the transport-fair metric, not Gb/s** — because an engine that consumes more + bytes per GEMM carries more throughput at the same compute rate. (Matching the window to + 4 MB on both sides is what makes the two columns comparable.) Achieved rates are + ~3.1 TFLOP/s (FP32) / ~5.7 TFLOP/s (FP16), only ~10% of peak, so the small matrix is + **launch/sync-latency-bound**, not FLOP-bound. +- **GPU SM% is a duty-cycle metric, anti-correlated with throughput here.** RoCE is + "stall-busy" (78% SM at 48 Gb/s); raw is efficient (40% SM at 104 Gb/s). High SM% means + the GPU work is *stretched thin* against a slower receive path, not that more useful work + is done — so SM% should not be read as compute efficiency. + +The bottleneck is the **single-threaded RoCE receive+compute loop**: one thread runs +`poll completion → gather → issue GEMM → stream-sync → free → repeat`, and the periodic +`cudaStreamSynchronize` (required before a received buffer can be safely reused) blocks +that thread — so while the GPU drains, no receives are posted and the message rate falls. +The heavier the GEMM, the longer each stall, which is why FP32 (48 Gb/s) suffers more than +FP16 (88 Gb/s) while raw stays wire-limited for both. Raw proves one thread *can* feed the +GPU at ~3100 GEMM/s, so the remedy is to decouple receive from compute (a dedicated +GPU-worker thread) or thin the stalls (`--workload-sync-interval`); a sync-interval sweep +quantifying the latter is in progress. ### Workload batch-size sweep +!!! warning "Superseded — conflates problem size with pipelining" + This sweep varies `--workload-batch-bytes`, which also sets the GEMM dimension + (`n = √(batch/4)`), so each column changes the FLOP count *and* the pipelining depth + at once. The [fixed-size GEMM comparison](#fixed-size-gemm-comparison) above pins n and + shows the RoCE↔raw gap is a receive-path / threading effect, not pipelining depth. The + curve is kept here for transparency. + The workload's compute working set is **decoupled from the I/O unit** via `--workload-batch-bytes` (env `WORKLOAD_BATCH` in `run_spark_bench.sh`): it sets the bytes fed to one compute call — the GEMM dimension is `n = √(batch/4)` — independent of @@ -347,8 +386,10 @@ the GPU stays fed regardless of the per-window size; the workload is never the b one 8 MB message yields a single GEMM with nothing to overlap (underpipelined, 85 Gb/s), sub-dividing it to 2–4 MB packs 2–4 GEMMs per message and recovers **~92 Gb/s — on par with raw**, and tiny 512 KB slices collapse to 37 Gb/s on per-call launch/sync overhead -(16 GEMMs/message). So the RoCE↔raw gap at the default batch is purely pipelining depth: -give RoCE a batch that packs a few GEMMs per message and the two paths converge. +(16 GEMMs/message). The apparent RoCE recovery here is partly an artifact of the growing +matrix (bigger batch → bigger n → a more efficient, more GPU-bound GEMM on both paths); the +[fixed-size GEMM comparison](#fixed-size-gemm-comparison) above pins n to separate that from +the transport effect and shows the gap is the RoCE receive path, not pipelining depth. | Batch | RoCE Gb/s | RoCE GPU SM% | Raw Gb/s | Raw GPU SM% | | ----: | --------: | -----------: | -------: | ----------: | @@ -448,6 +489,23 @@ done # raw: netns down, ETH_DST_ADDR exported; same loop with `dpdk` ``` +**Fixed-size GEMM comparison** pins the matrix dimension with `GEMM_N` +(`--workload-gemm-n`) so the FLOP count is constant across transports; both sides run one +GEMM per 4 MB window. Lands in the CSV `post_process_gemm_n` column: + +```bash +# RoCE (netns up): 8 MB message, 4 MB chunk -> 2 fixed 1024^3 GEMMs/message +for WL in gemm gemm_fp16; do + WORKLOAD=$WL GEMM_N=1024 WORKLOAD_BATCH=4194304 REPEATS=3 \ + PAYLOADS_OVERRIDE="8388608" ./examples/run_spark_bench.sh rdma sweep +done +# Raw (netns down, ETH_DST_ADDR exported): 4096 B x 1024 packets = 4 MB reorder window +for WL in gemm gemm_fp16; do + WORKLOAD=$WL GEMM_N=1024 WORKLOAD_BATCH=4194304 REPEATS=3 \ + PAYLOADS_OVERRIDE="4096" BATCHES_OVERRIDE="1024" ./examples/run_spark_bench.sh dpdk sweep +done +``` + Each run writes `bench-results/--/runs.csv`. See [Socket and RDMA Benchmarking](socket_benchmarking.md) and [Raw Ethernet Benchmarking](raw_benchmarking.md) for the namespace setup and From 90c83fb6cbf5dc6c0fce250131767b98528edd48 Mon Sep 17 00:00:00 2001 From: Ramya Gurunathan Date: Wed, 8 Jul 2026 10:12:48 -0400 Subject: [PATCH 02/18] #15 - Add --workload-gemm-n and --workload-sync-interval bench flags Pin the square GEMM dimension (--workload-gemm-n) so the FLOP count per call is fixed while the I/O unit is swept, isolating transport from problem size; init logs the chosen GEMM/FFT shape and FLOP count, and rejects a pinned n whose A operand exceeds the received working set. Add --workload-sync-interval to vary GPU stream drain depth for characterizing the single-threaded receive+compute ceiling. Thread both through all bench mains and expose them via run_spark_bench.sh (GEMM_N, SYNC_INTERVAL envs; post_process_gemm_n, post_process_sync CSV columns). Assisted-by: Claude Opus 4.8 Signed-off-by: Ramya Gurunathan --- examples/bench_workload.cu | 64 +++++++++++++++++++++++++++++--- examples/bench_workload.h | 20 +++++++++- examples/raw_bench_common.cpp | 4 +- examples/raw_bench_common.h | 4 +- examples/raw_gpudirect_bench.cpp | 5 ++- examples/raw_hds_bench.cpp | 7 +++- examples/rdma_bench.cpp | 13 +++++-- examples/run_spark_bench.sh | 59 +++++++++++++++++++++++++++-- examples/socket_bench.cpp | 15 +++++--- 9 files changed, 165 insertions(+), 26 deletions(-) diff --git a/examples/bench_workload.cu b/examples/bench_workload.cu index 4d602da..eb4dbb5 100644 --- a/examples/bench_workload.cu +++ b/examples/bench_workload.cu @@ -85,6 +85,30 @@ size_t parse_workload_batch_bytes(int argc, char** argv) { return 0; } +int parse_workload_gemm_n(int argc, char** argv) { + for (int i = 2; i + 1 < argc; i += 2) { + if (std::string(argv[i]) == "--workload-gemm-n") { + const long long v = std::atoll(argv[i + 1]); + if (v > 0) { + return static_cast(v); + } + } + } + return 0; +} + +int parse_workload_sync_interval(int argc, char** argv) { + for (int i = 2; i + 1 < argc; i += 2) { + if (std::string(argv[i]) == "--workload-sync-interval") { + const long long v = std::atoll(argv[i + 1]); + if (v > 0) { + return static_cast(v); + } + } + } + return 2; // default sync depth +} + const char* workload_name(BenchWorkload workload) { switch (workload) { case BenchWorkload::Fft: @@ -131,7 +155,8 @@ void GpuWorkload::destroy() { ok_ = false; } -bool GpuWorkload::init(BenchWorkload kind, size_t batch_bytes, int sync_interval) { +bool GpuWorkload::init(BenchWorkload kind, size_t batch_bytes, int sync_interval, + int gemm_n_override) { kind_ = kind; sync_interval_ = sync_interval > 0 ? sync_interval : 1; run_count_ = 0; @@ -176,16 +201,36 @@ bool GpuWorkload::init(BenchWorkload kind, size_t batch_bytes, int sync_interval destroy(); return false; } + // Explicit shape so every published FFT number carries its compute size. + // ~5·N·log2(N) flops per length-N C2C transform, times `batch` transforms. + const double flops = 5.0 * kFftLen * std::log2(static_cast(kFftLen)) * batch; + std::cerr << "GpuWorkload: fft shape = " << batch << " x C2C length-" << kFftLen + << " (working set " << (bytes >> 10) << " KiB, " << (flops / 1e6) << " MFLOP/call)\n"; } else { // Gemm or GemmFp16 - // Square matmul whose matrices match the working set. Size n from FP32 in - // both cases so gemm and gemm_fp16 use the SAME dimension -> identical FLOP - // count, isolating the precision / tensor-core effect. The A operand is the - // caller's input buffer (n*n*elem_size <= bytes); B and C are owned scratch. - int n = static_cast(std::sqrt(static_cast(bytes) / sizeof(float))); + // Square matmul. By default size n from the working set (FP32 in both cases so + // gemm and gemm_fp16 use the SAME dimension -> identical FLOP count, isolating + // the precision / tensor-core effect). A caller-supplied gemm_n_override pins n + // directly so the FLOP count per call stays FIXED while the I/O unit is swept -- + // this isolates pipelining depth from problem size in the RoCE-vs-raw study. + // The A operand is the caller's input buffer (n*n*elem_size must be <= bytes, + // enforced below); B and C are owned scratch. + int n = gemm_n_override > 0 + ? gemm_n_override + : static_cast(std::sqrt(static_cast(bytes) / sizeof(float))); n = std::max(64, (n / 8) * 8); // multiple of 8, sane floor gemm_n_ = n; const size_t elems = static_cast(n) * n; const size_t elem_size = kind_ == BenchWorkload::GemmFp16 ? sizeof(__half) : sizeof(float); + // The A operand is read from the caller's received-data buffer, which holds + // `bytes`. A pinned n must fit -- otherwise cuBLAS reads past the buffer. Reject + // rather than silently OOB-read; the caller should raise --workload-batch-bytes. + if (elems * elem_size > bytes) { + std::cerr << "GpuWorkload: pinned gemm n=" << n << " needs " << (elems * elem_size >> 10) + << " KiB for the A operand but the working set is only " << (bytes >> 10) + << " KiB; raise --workload-batch-bytes to >= n*n*elem_size\n"; + destroy(); + return false; + } if (cudaMalloc(&gemm_b_, elems * elem_size) != cudaSuccess || cudaMalloc(&gemm_c_, elems * elem_size) != cudaSuccess || cudaMemset(gemm_b_, 0, elems * elem_size) != cudaSuccess) { @@ -205,6 +250,13 @@ bool GpuWorkload::init(BenchWorkload kind, size_t batch_bytes, int sync_interval destroy(); return false; } + // Explicit shape so every published GEMM number carries its compute size. A + // square n×n×n matmul is 2·n³ flops; note whether n was pinned or derived. + const double flops = 2.0 * static_cast(n) * n * n; + std::cerr << "GpuWorkload: " << workload_name(kind_) << " shape = " << n << "x" << n << "x" << n + << (gemm_n_override > 0 ? " (pinned)" : " (derived)") << ", " + << (elem_size == sizeof(__half) ? "fp16" : "fp32") << " A/B, " << (flops / 1e9) + << " GFLOP/call, matrix " << (elems * elem_size >> 10) << " KiB\n"; } // Warm up and validate the chosen op once against a transient zeroed input diff --git a/examples/bench_workload.h b/examples/bench_workload.h index 63ccbea..4fa6052 100644 --- a/examples/bench_workload.h +++ b/examples/bench_workload.h @@ -38,6 +38,20 @@ BenchWorkload parse_workload(int argc, char** argv); // then falls back to its backend-default batch). Mirrors parse_workload's stride. size_t parse_workload_batch_bytes(int argc, char** argv); +// Parse "--workload-gemm-n N" from argv: pin the square GEMM dimension directly, +// independent of the compute batch bytes, so the FLOP count per call is FIXED as +// the I/O unit (message / burst window) is swept. Isolates pipelining depth from +// problem size in the RoCE-vs-raw comparison. Returns 0 if unset (the GEMM +// dimension then derives from the working-set size). Mirrors parse_workload's stride. +int parse_workload_gemm_n(int argc, char** argv); + +// Parse "--workload-sync-interval N" from argv: drain the GPU stream every N compute +// calls (bounds outstanding GPU work). Larger N = deeper async queue, fewer CPU +// stalls waiting on the GPU; N=1 is fully synchronous. Used to characterize how much +// the single-threaded receive+compute loop is limited by sync stalls. Returns 2 (the +// default) if unset. Mirrors parse_workload's stride. +int parse_workload_sync_interval(int argc, char** argv); + // Lower-case name ("none"/"fft"/"gemm"); used for the run_spark_bench.sh // post_process CSV column and log lines. const char* workload_name(BenchWorkload workload); @@ -69,9 +83,13 @@ class GpuWorkload { // batch_bytes the caller will pass to run() (0 => an internal default). The op // reads at most batch_bytes from that buffer. kind == None leaves the object an // inert no-op. sync_interval bounds outstanding GPU work (sync every N runs). + // gemm_n_override (>0) pins the square GEMM dimension directly, holding the FLOP + // count per call fixed regardless of batch_bytes (0 => derive n from batch_bytes). + // Logs the chosen problem shape (GEMM n / FFT length+batch) and FLOP count so + // every published benchmark number carries its explicit compute size. // Returns false on CUDA / library error; the caller may warn and continue with // the workload disabled (enabled() will report false). - bool init(BenchWorkload kind, size_t batch_bytes, int sync_interval = 2); + bool init(BenchWorkload kind, size_t batch_bytes, int sync_interval = 2, int gemm_n_override = 0); // Enqueue one representative FFT/SGEMM on the internal stream, reading `input` // (a device pointer to >= batch_bytes valid bytes). No-op unless enabled(). diff --git a/examples/raw_bench_common.cpp b/examples/raw_bench_common.cpp index 1d514a1..ef686a1 100644 --- a/examples/raw_bench_common.cpp +++ b/examples/raw_bench_common.cpp @@ -549,7 +549,7 @@ void print_queue_stats(const char *direction, const std::string &interface_name, } void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, BenchWorkload workload, - const ReorderGeometry& geom) { + const ReorderGeometry& geom, int workload_gemm_n, int workload_sync_interval) { if (!set_current_thread_affinity(cfg.cpu_core, "bench_rx")) { stop.store(true); return; @@ -564,7 +564,7 @@ void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, Bench GpuWorkload gpu_workload; ReorderPipeline pipeline; if (workload != BenchWorkload::None) { - if (!gpu_workload.init(workload, batch_bytes) || + if (!gpu_workload.init(workload, batch_bytes, workload_sync_interval, workload_gemm_n) || !pipeline.init(ReorderMode::SeqReorder, geom.packets_per_batch, out_payload_len, geom.payload_byte_offset, geom.seq_bit_offset, geom.seq_bit_width, /*staging_needed=*/false, gpu_workload.stream())) { diff --git a/examples/raw_bench_common.h b/examples/raw_bench_common.h index c79a605..496a2c7 100644 --- a/examples/raw_bench_common.h +++ b/examples/raw_bench_common.h @@ -160,7 +160,7 @@ struct ReorderGeometry { // is freed only after the stream drains (so the reorder has read the buffers). // workload == None leaves the bare-loopback count-only path untouched. void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, - BenchWorkload workload = BenchWorkload::None, - const ReorderGeometry& geom = {}); + BenchWorkload workload = BenchWorkload::None, const ReorderGeometry& geom = {}, + int workload_gemm_n = 0, int workload_sync_interval = 2); } // namespace daqiri::bench diff --git a/examples/raw_gpudirect_bench.cpp b/examples/raw_gpudirect_bench.cpp index 9aaf574..0181454 100644 --- a/examples/raw_gpudirect_bench.cpp +++ b/examples/raw_gpudirect_bench.cpp @@ -173,6 +173,8 @@ int main(int argc, char **argv) { const double target_gbps = daqiri::bench::parse_target_gbps(argc, argv); const auto workload = daqiri::bench::parse_workload(argc, argv); const size_t workload_batch_bytes = daqiri::bench::parse_workload_batch_bytes(argc, argv); + const int workload_gemm_n = daqiri::bench::parse_workload_gemm_n(argc, argv); + const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); std::vector rx_configs; @@ -223,7 +225,8 @@ int main(int argc, char **argv) { } rx_threads.reserve(rx_configs.size()); for (const auto &cfg : rx_configs) { - rx_threads.emplace_back(daqiri::bench::rx_count_worker, cfg, std::ref(stop), workload, geom); + rx_threads.emplace_back(daqiri::bench::rx_count_worker, cfg, std::ref(stop), workload, geom, + workload_gemm_n, workload_sync_interval); } tx_threads.reserve(tx_configs.size()); for (const auto &cfg : tx_configs) { diff --git a/examples/raw_hds_bench.cpp b/examples/raw_hds_bench.cpp index 4a1abca..743bc13 100644 --- a/examples/raw_hds_bench.cpp +++ b/examples/raw_hds_bench.cpp @@ -148,6 +148,8 @@ int main(int argc, char **argv) { const int run_seconds = daqiri::bench::parse_run_seconds(argc, argv); const auto workload = daqiri::bench::parse_workload(argc, argv); const size_t workload_batch_bytes = daqiri::bench::parse_workload_batch_bytes(argc, argv); + const int workload_gemm_n = daqiri::bench::parse_workload_gemm_n(argc, argv); + const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); if (daqiri::daqiri_init(argv[1]) != daqiri::Status::SUCCESS) { std::cerr << "daqiri_init failed\n"; @@ -186,8 +188,9 @@ int main(int argc, char **argv) { : 1024; geom.packets_per_batch = std::min(ppb, tx.batch_size); } - rx_thread = std::thread(daqiri::bench::rx_count_worker, daqiri::bench::parse_rx(root), - std::ref(stop), workload, geom); + rx_thread = + std::thread(daqiri::bench::rx_count_worker, daqiri::bench::parse_rx(root), std::ref(stop), + workload, geom, workload_gemm_n, workload_sync_interval); } if (has_tx) { tx_thread = diff --git a/examples/rdma_bench.cpp b/examples/rdma_bench.cpp index 9cec488..00d7266 100644 --- a/examples/rdma_bench.cpp +++ b/examples/rdma_bench.cpp @@ -81,7 +81,8 @@ RdmaBenchConfig parse_rdma_cfg(const YAML::Node& node) { void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pacer, std::atomic& stop, RdmaWorkerStats& stats, - daqiri::bench::BenchWorkload workload, size_t workload_batch_bytes) { + daqiri::bench::BenchWorkload workload, size_t workload_batch_bytes, + int workload_gemm_n, int workload_sync_interval) { const char *thread_name = cfg.server ? "rdma_bench_server" : "rdma_bench_client"; if (!daqiri::bench::set_current_thread_affinity(cfg.cpu_core, thread_name)) { @@ -107,7 +108,7 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa daqiri::bench::GpuWorkload gpu_workload; daqiri::bench::ReorderPipeline pipeline; if (workload != daqiri::bench::BenchWorkload::None) { - if (!gpu_workload.init(workload, chunk) || + if (!gpu_workload.init(workload, chunk, workload_sync_interval, workload_gemm_n) || !pipeline.init(daqiri::bench::ReorderMode::GatherOnly, /*packets_per_batch=*/1, msg, /*payload_byte_offset=*/0, /*seq_bit_offset=*/0, @@ -323,6 +324,8 @@ int main(int argc, char** argv) { } const auto workload = daqiri::bench::parse_workload(argc, argv); const size_t workload_batch_bytes = daqiri::bench::parse_workload_batch_bytes(argc, argv); + const int workload_gemm_n = daqiri::bench::parse_workload_gemm_n(argc, argv); + const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); if (daqiri::daqiri_init(argv[1]) != daqiri::Status::SUCCESS) { @@ -359,11 +362,13 @@ int main(int argc, char** argv) { if (run_server) { server_thread = std::thread(rdma_worker, server_cfg, std::ref(server_pacer), std::ref(stop), - std::ref(server_stats), workload, workload_batch_bytes); + std::ref(server_stats), workload, workload_batch_bytes, + workload_gemm_n, workload_sync_interval); } if (run_client) { client_thread = std::thread(rdma_worker, client_cfg, std::ref(client_pacer), std::ref(stop), - std::ref(client_stats), workload, workload_batch_bytes); + std::ref(client_stats), workload, workload_batch_bytes, + workload_gemm_n, workload_sync_interval); } if (!server_thread.joinable() && !client_thread.joinable()) { diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index 4e97e7c..c0c0606 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -29,6 +29,18 @@ # tensor-core matmul). Honoured by all backends (dpdk, rdma, # socket-udp, socket-tcp); recorded in the CSV post_process # column. +# WORKLOAD_BATCH — GPU compute working-set bytes per call, decoupled from the +# I/O unit (empty = backend default). Recorded in +# post_process_batch. +# GEMM_N — pin the square GEMM dimension n directly (--workload-gemm-n), +# holding FLOPs/call fixed (2·n³) while the I/O unit is swept +# via PAYLOADS_OVERRIDE. The fixed-n message-size sweep that +# isolates pipelining depth from problem size. Recorded in +# post_process_gemm_n. +# SYNC_INTERVAL — drain the GPU stream every N compute calls +# (--workload-sync-interval; default 2). Sweep it (1 2 4 8 16 32) +# to see how much of the receive+compute ceiling is single-thread +# GPU sync-stall. Recorded in post_process_sync. # # Optional (dpdk only): DPDK_{TX,RX}_PCI / DPDK_{TX,RX}_NETDEV override the p0/p1 # ports used for the per-cell *_phy wire-transit check (defaults p0 0000:01:00.0 / @@ -68,9 +80,11 @@ CSV="$OUT_DIR/runs.csv" # `pairs` = number of concurrent client/server process pairs (socket backends sweep # this; dpdk/rdma are always 1). `gbps` is aggregate App TX, `rx_gbps` aggregate App RX # (summed across pairs); App-level loss is (gbps - rx_gbps) / gbps. -# post_process_batch (last column) = WORKLOAD_BATCH bytes, or "default" when unset. +# post_process_batch = WORKLOAD_BATCH bytes, or "default" when unset. +# post_process_gemm_n = GEMM_N pinned dimension, or "derived" when unset. +# post_process_sync (last column) = SYNC_INTERVAL, or "default" (2) when unset. # Appended at the end so existing column positions are unchanged. -echo "lang,backend,post_process,payload,batch,pairs,target_gbps,rep,seconds,packets,bytes,pps,gbps,rx_gbps,drops,drops_kind,cpu_master_pct,cpu_tx_pct,cpu_rx_pct,gpu_sm_pct,gpu_mem_pct,post_process_batch" > "$CSV" +echo "lang,backend,post_process,payload,batch,pairs,target_gbps,rep,seconds,packets,bytes,pps,gbps,rx_gbps,drops,drops_kind,cpu_master_pct,cpu_tx_pct,cpu_rx_pct,gpu_sm_pct,gpu_mem_pct,post_process_batch,post_process_gemm_n,post_process_sync" > "$CSV" # Capture slow-moving environment state once per result set. "$SCRIPT_DIR/bench_capture_environment.sh" "$OUT_DIR" @@ -98,6 +112,25 @@ WORKLOAD_BATCH="${WORKLOAD_BATCH:-}" if [[ -n "$WORKLOAD_BATCH" && ! "$WORKLOAD_BATCH" =~ ^[0-9]+$ ]]; then echo "Invalid WORKLOAD_BATCH '$WORKLOAD_BATCH' (expected a positive byte count)" >&2; exit 1 fi +# GEMM_N: pin the square GEMM dimension n directly (--workload-gemm-n), holding the +# FLOP count per call FIXED (2·n³) while the I/O unit (RoCE message via +# PAYLOADS_OVERRIDE, raw reorder window) is swept. This isolates pipelining depth +# from problem size -- the fixed-n message-size sweep that proves the RoCE↔raw gap +# is pipelining, not transport. Empty = derive n from the working set. Recorded in +# the CSV post_process_gemm_n column. +GEMM_N="${GEMM_N:-}" +if [[ -n "$GEMM_N" && ! "$GEMM_N" =~ ^[0-9]+$ ]]; then + echo "Invalid GEMM_N '$GEMM_N' (expected a positive integer)" >&2; exit 1 +fi +# SYNC_INTERVAL: drain the GPU stream every N compute calls (--workload-sync-interval). +# Larger N lets more GEMMs queue asynchronously before the single receive+compute +# thread blocks on the GPU, so sweeping it (e.g. 1 2 4 8 16 32) shows how much of the +# ceiling is CPU sync-stall vs the receive path. Empty = default (2). Recorded in the +# CSV post_process_sync column. +SYNC_INTERVAL="${SYNC_INTERVAL:-}" +if [[ -n "$SYNC_INTERVAL" && ! "$SYNC_INTERVAL" =~ ^[0-9]+$ ]]; then + echo "Invalid SYNC_INTERVAL '$SYNC_INTERVAL' (expected a positive integer)" >&2; exit 1 +fi DRIVER_LOG="$OUT_DIR/last_run.stderr" FAILURES=0 @@ -187,6 +220,20 @@ case "$BACKEND" in *) echo "Unknown backend: $BACKEND" >&2; exit 1 ;; esac +# Optional space-separated overrides for the sweep ladders, so a one-off experiment +# can re-target the payload/batch matrix without editing the per-backend defaults +# above. Used by the fixed-n workload sweep: pin the GEMM dimension with GEMM_N +# (constant 2·n³ FLOPs/call) and sweep PAYLOADS_OVERRIDE = the RoCE message size, so +# num_chunks = message/batch varies pipelining depth at a FIXED matrix size. This is +# the "send an 80 MB message, launch N same-size GEMMs" experiment that isolates +# pipelining depth from problem size -- proving the RoCE↔raw gap is pipelining, not +# transport. Hold WORKLOAD_BATCH at the per-chunk size (= the GEMM working set). +# e.g. WORKLOAD=gemm_fp16 GEMM_N=1024 WORKLOAD_BATCH=4194304 \ +# PAYLOADS_OVERRIDE="4194304 8388608 16777216 33554432 83886080" \ +# ./run_spark_bench.sh rdma sweep +[[ -n "${PAYLOADS_OVERRIDE:-}" ]] && read -r -a PAYLOADS_SWEEP <<< "$PAYLOADS_OVERRIDE" +[[ -n "${BATCHES_OVERRIDE:-}" ]] && read -r -a BATCHES_SWEEP <<< "$BATCHES_OVERRIDE" + # All backends (dpdk, rdma, socket-udp, socket-tcp) now run the workload on real # received data, so the CSV post_process column records the requested workload # for every backend. @@ -400,6 +447,8 @@ run_cell() { if [[ "$WORKLOAD_EFF" != "none" ]]; then bench_extra+=(--workload "$WORKLOAD_EFF") [[ -n "$WORKLOAD_BATCH" ]] && bench_extra+=(--workload-batch-bytes "$WORKLOAD_BATCH") + [[ -n "$GEMM_N" ]] && bench_extra+=(--workload-gemm-n "$GEMM_N") + [[ -n "$SYNC_INTERVAL" ]] && bench_extra+=(--workload-sync-interval "$SYNC_INTERVAL") fi # Shutdown ordering: the server must keep receiving until the client has fully # stopped sending, otherwise the client's last in-flight messages have no peer @@ -574,7 +623,11 @@ run_cell() { local pp_batch="${WORKLOAD_BATCH:-default}" [[ -z "$pp_batch" ]] && pp_batch="default" - echo "$lang,$BACKEND,$WORKLOAD_EFF,$payload,$batch,$pairs,$target_gbps,$rep,$secs,$pkts,$bytes,$pps,$gbps,$rx_gbps,$drops,$drops_kind,$cpu_master_pct,$cpu_tx_pct,$cpu_rx_pct,$gpu_sm,$gpu_mem,$pp_batch" \ + local pp_gemm_n="${GEMM_N:-derived}" + [[ -z "$pp_gemm_n" ]] && pp_gemm_n="derived" + local pp_sync="${SYNC_INTERVAL:-default}" + [[ -z "$pp_sync" ]] && pp_sync="default" + echo "$lang,$BACKEND,$WORKLOAD_EFF,$payload,$batch,$pairs,$target_gbps,$rep,$secs,$pkts,$bytes,$pps,$gbps,$rx_gbps,$drops,$drops_kind,$cpu_master_pct,$cpu_tx_pct,$cpu_rx_pct,$gpu_sm,$gpu_mem,$pp_batch,$pp_gemm_n,$pp_sync" \ | tee -a "$CSV" } diff --git a/examples/socket_bench.cpp b/examples/socket_bench.cpp index 05d1339..dfc64ad 100644 --- a/examples/socket_bench.cpp +++ b/examples/socket_bench.cpp @@ -102,8 +102,8 @@ bool socket_transport_is_tcp(const YAML::Node& root) { void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pacer, std::atomic& stop, SocketWorkerStats& stats, - daqiri::bench::BenchWorkload workload, bool is_tcp, - size_t workload_batch_bytes) { + daqiri::bench::BenchWorkload workload, bool is_tcp, size_t workload_batch_bytes, + int workload_gemm_n, int workload_sync_interval) { const char *thread_name = cfg.server ? "socket_bench_server" : "socket_bench_client"; if (!daqiri::bench::set_current_thread_affinity(cfg.cpu_core, thread_name)) { @@ -126,7 +126,8 @@ void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer daqiri::bench::GpuWorkload gpu_workload; daqiri::bench::ReorderPipeline pipeline; if (workload != daqiri::bench::BenchWorkload::None) { - if (!gpu_workload.init(workload, static_cast(packets_per_batch) * msg) || + if (!gpu_workload.init(workload, static_cast(packets_per_batch) * msg, + workload_sync_interval, workload_gemm_n) || !pipeline.init(is_tcp ? daqiri::bench::ReorderMode::GatherOnly : daqiri::bench::ReorderMode::SeqReorder, packets_per_batch, msg, /*payload_byte_offset=*/0, @@ -259,6 +260,8 @@ int main(int argc, char** argv) { } const auto workload = daqiri::bench::parse_workload(argc, argv); const size_t workload_batch_bytes = daqiri::bench::parse_workload_batch_bytes(argc, argv); + const int workload_gemm_n = daqiri::bench::parse_workload_gemm_n(argc, argv); + const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); const bool is_tcp = socket_transport_is_tcp(root); @@ -296,11 +299,13 @@ int main(int argc, char** argv) { if (run_server) { server_thread = std::thread(socket_worker, server_cfg, std::ref(server_pacer), std::ref(stop), - std::ref(server_stats), workload, is_tcp, workload_batch_bytes); + std::ref(server_stats), workload, is_tcp, workload_batch_bytes, + workload_gemm_n, workload_sync_interval); } if (run_client) { client_thread = std::thread(socket_worker, client_cfg, std::ref(client_pacer), std::ref(stop), - std::ref(client_stats), workload, is_tcp, workload_batch_bytes); + std::ref(client_stats), workload, is_tcp, workload_batch_bytes, + workload_gemm_n, workload_sync_interval); } if (!server_thread.joinable() && !client_thread.joinable()) { From e86a7530b9c5603c97725afb4edbd61c9df4fc3c Mon Sep 17 00:00:00 2001 From: Ramya Gurunathan Date: Wed, 8 Jul 2026 10:35:53 -0400 Subject: [PATCH 03/18] #15 - Record sync-interval result in fixed-GEMM perf section The --workload-sync-interval sweep (2->32, both precisions) shows no measurable change in RoCE throughput, and fully-synchronous matches deep-async. Sync cadence is not the ceiling; the RoCE path already drains once per held-receive batch, masking the knob, so the fix is structural (dedicated GPU-worker thread) rather than sync tuning. Updates the fixed-size GEMM section accordingly. Assisted-by: Claude Opus 4.8 Signed-off-by: Ramya Gurunathan --- docs/benchmarks/performance-dgx-spark.md | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/docs/benchmarks/performance-dgx-spark.md b/docs/benchmarks/performance-dgx-spark.md index 14f9cec..e23df43 100644 --- a/docs/benchmarks/performance-dgx-spark.md +++ b/docs/benchmarks/performance-dgx-spark.md @@ -358,9 +358,15 @@ The bottleneck is the **single-threaded RoCE receive+compute loop**: one thread that thread — so while the GPU drains, no receives are posted and the message rate falls. The heavier the GEMM, the longer each stall, which is why FP32 (48 Gb/s) suffers more than FP16 (88 Gb/s) while raw stays wire-limited for both. Raw proves one thread *can* feed the -GPU at ~3100 GEMM/s, so the remedy is to decouple receive from compute (a dedicated -GPU-worker thread) or thin the stalls (`--workload-sync-interval`); a sync-interval sweep -quantifying the latter is in progress. +GPU at ~3100 GEMM/s, so the ceiling is the RoCE receive path, not the GPU. + +Tuning the sync cadence does **not** recover it: sweeping `--workload-sync-interval` from 2 +to 32 leaves throughput flat (FP32 ~48 Gb/s / ~1450 GEMMs/s, FP16 ~88 Gb/s / ~2650 GEMMs/s), +and fully-synchronous (interval 1) matches deep-async — the path already drains the stream +once per held-receive batch, which masks the knob, and the real limit is the single-threaded +single-QP receive itself. The remedy is therefore **structural**: decouple receive from +compute onto a dedicated GPU-worker thread (and/or fan out across multiple QPs/threads), +not sync tuning. ### Workload batch-size sweep From 23617a6caffa8d381f4088eb37fb050c35fdfde5 Mon Sep 17 00:00:00 2001 From: Ramya Gurunathan Date: Wed, 8 Jul 2026 14:46:50 -0400 Subject: [PATCH 04/18] #15 - Pin GEMM dimension and drop workload batch-bytes/payload-override flags The receive-path GPU workload sized its GEMM from a --workload-batch-bytes knob (deriving n = sqrt(batch/4)) and sub-divided each RoCE message into chunk-sized slices, which conflated problem size with pipelining depth and made the fixed-n RoCE-vs-raw comparison confusing to run (juggling WORKLOAD_BATCH + PAYLOADS_OVERRIDE). Pin the square GEMM side length directly instead: - Rename --workload-gemm-n -> --workload-gemm-dim (default 1024); the compute working set is exactly n*n*elem_size, read from the front of each received I/O unit. GpuWorkload rejects a unit too small to hold the operand. - Delete --workload-batch-bytes / parse_workload_batch_bytes and the derive-n-from-bytes branch. Each backend keeps its default reorder window (raw 1024 packets, socket 8 MB, RoCE whole message); RoCE now runs one fixed-dimension GEMM per received message. - run_spark_bench.sh: GEMM_N -> GEMM_DIM (default 1024); drop WORKLOAD_BATCH and PAYLOADS_OVERRIDE. When a workload is active the payload sweep is restricted to the headline size (the small cells cannot hold the operand). CSV: post_process_gemm_n -> post_process_gemm_dim; drop post_process_batch. - Docs: update AGENTS.md bench-flag summary and the DGX Spark perf report methodology/reproduce; remove the superseded batch-size-sweep section and mark the fixed-GEMM table pending refresh. Example-only change; no library rebuild required. Signed-off-by: Ramya Gurunathan Assisted-by: Claude Opus 4.8 --- AGENTS.md | 2 +- docs/benchmarks/performance-dgx-spark.md | 89 ++++++------------------ examples/bench_workload.cu | 60 ++++++---------- examples/bench_workload.h | 46 ++++++------ examples/raw_bench_common.cpp | 5 +- examples/raw_bench_common.h | 2 +- examples/raw_gpudirect_bench.cpp | 20 +++--- examples/raw_hds_bench.cpp | 16 ++--- examples/rdma_bench.cpp | 41 +++++------ examples/run_spark_bench.sh | 75 +++++++------------- examples/socket_bench.cpp | 28 ++++---- 11 files changed, 142 insertions(+), 242 deletions(-) diff --git a/AGENTS.md b/AGENTS.md index d9d639a..3eea151 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -44,7 +44,7 @@ There is no unit test suite. Verification is done via the benchmark executables | `daqiri_bench_socket` | `socket_bench.cpp` | `daqiri_bench_socket_{udp,tcp}_tx_rx.yaml`, `daqiri_bench_socket_{udp,tcp}_tx_rx_spark_netns.yaml` (combined-role netns bases) | | `daqiri_pool_ring_bench` | `pool_ring_bench.cpp` | none — microbenchmark comparing `daqiri::Ring`/`daqiri::ObjectPool` vs DPDK `rte_ring`/`rte_mempool` (SPSC/MPMC, single/bulk, thread sweep). The `rte_*` comparison arm compiles only in a DPDK-enabled build; takes no YAML/CLI args | -The four `raw_*` benches share `raw_bench_common.{cpp,h}` and accept `--seconds N`. `daqiri_bench_rdma` and `daqiri_bench_socket` also take `--mode {tx,rx,both}`. `daqiri_bench_raw_gpudirect`, `daqiri_bench_raw_hds`, `daqiri_bench_rdma`, and `daqiri_bench_socket` additionally accept `--workload none|fft|gemm|gemm_fp16` — a reusable representative GPU workload (`examples/bench_workload.{h,cu}`, cuFFT/cuBLAS) run once per received reorder window on the **actual received payload**. Each backend first assembles the burst's payloads into one contiguous GPU buffer via `examples/bench_pipeline.{h,cu}` (`ReorderPipeline`): a sequence-number reorder kernel for the out-of-order transports (DPDK raw, UDP) and an arrival-order gather for the in-order ones (RoCE RC, TCP); sockets stage host→device first since their payloads land in pageable host memory. The reorder/gather kernels (`packet_reorder_copy_payload_by_sequence`, `packet_gather_copy_payload`) live in `src/kernels.cu`. `gemm` is FP32 `cublasSgemm`; `gemm_fp16` is the same-size mixed-precision FP16/tensor-core `cublasGemmEx` (inference-style); the contiguous buffer supplies the FFT input / GEMM A operand. `--workload-batch-bytes N` decouples the compute working set from the I/O unit (RoCE sub-chunks each message; raw sizes its reorder window), enabling a batch-size sweep. Used by `run_spark_bench.sh`'s `WORKLOAD` / `WORKLOAD_BATCH` env (all backends) to fill the CSV `post_process` / `post_process_batch` columns (issue #15). +The four `raw_*` benches share `raw_bench_common.{cpp,h}` and accept `--seconds N`. `daqiri_bench_rdma` and `daqiri_bench_socket` also take `--mode {tx,rx,both}`. `daqiri_bench_raw_gpudirect`, `daqiri_bench_raw_hds`, `daqiri_bench_rdma`, and `daqiri_bench_socket` additionally accept `--workload none|fft|gemm|gemm_fp16` — a reusable representative GPU workload (`examples/bench_workload.{h,cu}`, cuFFT/cuBLAS) run once per received reorder window on the **actual received payload**. Each backend first assembles the burst's payloads into one contiguous GPU buffer via `examples/bench_pipeline.{h,cu}` (`ReorderPipeline`): a sequence-number reorder kernel for the out-of-order transports (DPDK raw, UDP) and an arrival-order gather for the in-order ones (RoCE RC, TCP); sockets stage host→device first since their payloads land in pageable host memory. The reorder/gather kernels (`packet_reorder_copy_payload_by_sequence`, `packet_gather_copy_payload`) live in `src/kernels.cu`. `gemm` is FP32 `cublasSgemm`; `gemm_fp16` is the same-size mixed-precision FP16/tensor-core `cublasGemmEx` (inference-style); the contiguous buffer supplies the FFT input / GEMM A operand. `--workload-gemm-dim N` pins the square GEMM side length (default 1024), so the FLOP count per call (2·n³) is FIXED and the compute working set is exactly n·n·elem_size, read from the front of each received I/O unit (the unit must be at least that large). Used by `run_spark_bench.sh`'s `WORKLOAD` / `GEMM_DIM` env (all backends) to fill the CSV `post_process` / `post_process_gemm_dim` columns (issue #15). ```bash ./build/examples/daqiri_bench_raw_gpudirect ./build/examples/daqiri_bench_raw_tx_rx.yaml --seconds 10 diff --git a/docs/benchmarks/performance-dgx-spark.md b/docs/benchmarks/performance-dgx-spark.md index e23df43..9a1320b 100644 --- a/docs/benchmarks/performance-dgx-spark.md +++ b/docs/benchmarks/performance-dgx-spark.md @@ -319,14 +319,18 @@ path's ~8.19 MB reorder window (1024 packets × 8000 B). ### Fixed-size GEMM comparison -The tables above size the GEMM from the working set, so sweeping the batch also changes -the matrix dimension (FLOPs ∝ n³) and conflates problem size with transport. To compare -transports cleanly we **pin the GEMM dimension** with `--workload-gemm-n 1024` -(env `GEMM_N` in `run_spark_bench.sh`): every call is an identical **1024×1024×1024 -matmul = 2.15 GFLOP**, and both transports consume **exactly 4 MB per GEMM** (one GEMM -per 4 MB window), so throughput *and* GEMMs/s are directly comparable. Fixed n=1024, -3 reps, 30 s each; RoCE at its 8 MB message (2 GEMMs/message), raw at a 4 MB reorder -window (4096 B × 1024 packets). +To compare transports cleanly we **pin the GEMM dimension** with `--workload-gemm-dim 1024` +(env `GEMM_DIM` in `run_spark_bench.sh`): every call is an identical **1024×1024×1024 +matmul = 2.15 GFLOP** reading the first **4 MB** (n²·4 B, FP32) of each received I/O unit, +so throughput *and* GEMMs/s are directly comparable across transports. One fixed-n GEMM +per received unit — RoCE per 8 MB message, raw per reorder window. Fixed n=1024, 3 reps, +30 s each. + +!!! note "Numbers pending refresh" + The table below was measured under the earlier per-message chunking (2 GEMMs per 8 MB + message) and the blocking per-batch `cudaStreamSynchronize`. The receive path now runs + **one GEMM per message** and recycles recv buffers with non-blocking CUDA events; these + cells will be re-measured under that model. | Transport | Workload | Throughput | GEMMs/s | GPU SM% | | --------- | -------- | ---------: | ------: | ------: | @@ -368,46 +372,6 @@ single-QP receive itself. The remedy is therefore **structural**: decouple recei compute onto a dedicated GPU-worker thread (and/or fan out across multiple QPs/threads), not sync tuning. -### Workload batch-size sweep - -!!! warning "Superseded — conflates problem size with pipelining" - This sweep varies `--workload-batch-bytes`, which also sets the GEMM dimension - (`n = √(batch/4)`), so each column changes the FLOP count *and* the pipelining depth - at once. The [fixed-size GEMM comparison](#fixed-size-gemm-comparison) above pins n and - shows the RoCE↔raw gap is a receive-path / threading effect, not pipelining depth. The - curve is kept here for transparency. - -The workload's compute working set is **decoupled from the I/O unit** via -`--workload-batch-bytes` (env `WORKLOAD_BATCH` in `run_spark_bench.sh`): it sets the -bytes fed to one compute call — the GEMM dimension is `n = √(batch/4)` — independent of -the 8 MB RoCE message or 8 KB raw frame. RoCE sub-divides each message into batch-sized -slices (one compute per slice); raw sizes its reorder window to the batch. Sweep run on -`gemm_fp16` (cuBLAS takes the dimension per call); the tables above are the fixed-batch -(8 MB) operating points. - -The two paths respond very differently, which is the whole point. **Raw is flat at line -rate (~98 Gb/s) across every batch** — a raw burst always packs ~10 reorder windows, so -the GPU stays fed regardless of the per-window size; the workload is never the bottleneck. -**RoCE is strongly batch-sensitive** and **non-monotonic with a sweet spot at 2–4 MB**: -one 8 MB message yields a single GEMM with nothing to overlap (underpipelined, 85 Gb/s), -sub-dividing it to 2–4 MB packs 2–4 GEMMs per message and recovers **~92 Gb/s — on par -with raw**, and tiny 512 KB slices collapse to 37 Gb/s on per-call launch/sync overhead -(16 GEMMs/message). The apparent RoCE recovery here is partly an artifact of the growing -matrix (bigger batch → bigger n → a more efficient, more GPU-bound GEMM on both paths); the -[fixed-size GEMM comparison](#fixed-size-gemm-comparison) above pins n to separate that from -the transport effect and shows the gap is the RoCE receive path, not pipelining depth. - -| Batch | RoCE Gb/s | RoCE GPU SM% | Raw Gb/s | Raw GPU SM% | -| ----: | --------: | -----------: | -------: | ----------: | -| 512 KB | 37.0 | 76.9% | 98.0 | 24.8% | -| 1 MB | 76.9 | 75.9% | 98.3 | 16.8% | -| 2 MB | 90.8 | 54.0% | 98.2 | 13.6% | -| 4 MB | 92.5 | 40.4% | 97.8 | 15.1% | -| 8 MB | 85.3 | 52.2% | 96.7 | 16.0% | - -(`gemm_fp16`, single rep per point. Raw cells use the nearest whole-packet window to the -batch size; RoCE caps the batch at the 8 MB message.) - ## Reproduce Run inside the project container (privileged, GPUs passed through, hugepages @@ -484,31 +448,20 @@ WORKLOAD=gemm ./examples/run_spark_bench.sh rdma smoke The chosen workload lands in the CSV `post_process` column; compare `gbps` / `gpu_sm_pct` against the matching `WORKLOAD=none` baseline. -**Workload batch-size sweep** decouples the compute working set from the I/O unit -by also exporting `WORKLOAD_BATCH` (bytes); it lands in the CSV `post_process_batch` -column: - -```bash -for B in 262144 524288 1048576 2097152 4194304 8388608; do - WORKLOAD=gemm_fp16 WORKLOAD_BATCH=$B ./examples/run_spark_bench.sh rdma smoke # netns up -done -# raw: netns down, ETH_DST_ADDR exported; same loop with `dpdk` -``` - -**Fixed-size GEMM comparison** pins the matrix dimension with `GEMM_N` -(`--workload-gemm-n`) so the FLOP count is constant across transports; both sides run one -GEMM per 4 MB window. Lands in the CSV `post_process_gemm_n` column: +**Fixed-size GEMM comparison** pins the matrix dimension with `GEMM_DIM` +(`--workload-gemm-dim`, default 1024) so the FLOP count is constant across transports; +each side runs one fixed 1024³ GEMM per received I/O unit, reading its first 4 MB. Lands +in the CSV `post_process_gemm_dim` column (the sweep is restricted to the headline +message size automatically when a workload is active): ```bash -# RoCE (netns up): 8 MB message, 4 MB chunk -> 2 fixed 1024^3 GEMMs/message +# RoCE (netns up) for WL in gemm gemm_fp16; do - WORKLOAD=$WL GEMM_N=1024 WORKLOAD_BATCH=4194304 REPEATS=3 \ - PAYLOADS_OVERRIDE="8388608" ./examples/run_spark_bench.sh rdma sweep + WORKLOAD=$WL GEMM_DIM=1024 REPEATS=3 ./examples/run_spark_bench.sh rdma sweep done -# Raw (netns down, ETH_DST_ADDR exported): 4096 B x 1024 packets = 4 MB reorder window +# Raw (netns down, ETH_DST_ADDR exported) for WL in gemm gemm_fp16; do - WORKLOAD=$WL GEMM_N=1024 WORKLOAD_BATCH=4194304 REPEATS=3 \ - PAYLOADS_OVERRIDE="4096" BATCHES_OVERRIDE="1024" ./examples/run_spark_bench.sh dpdk sweep + WORKLOAD=$WL GEMM_DIM=1024 REPEATS=3 ./examples/run_spark_bench.sh dpdk sweep done ``` diff --git a/examples/bench_workload.cu b/examples/bench_workload.cu index eb4dbb5..d89019b 100644 --- a/examples/bench_workload.cu +++ b/examples/bench_workload.cu @@ -73,28 +73,16 @@ BenchWorkload parse_workload(int argc, char** argv) { return workload; } -size_t parse_workload_batch_bytes(int argc, char** argv) { +int parse_workload_gemm_dim(int argc, char** argv) { for (int i = 2; i + 1 < argc; i += 2) { - if (std::string(argv[i]) == "--workload-batch-bytes") { - const long long v = std::atoll(argv[i + 1]); - if (v > 0) { - return static_cast(v); - } - } - } - return 0; -} - -int parse_workload_gemm_n(int argc, char** argv) { - for (int i = 2; i + 1 < argc; i += 2) { - if (std::string(argv[i]) == "--workload-gemm-n") { + if (std::string(argv[i]) == "--workload-gemm-dim") { const long long v = std::atoll(argv[i + 1]); if (v > 0) { return static_cast(v); } } } - return 0; + return 1024; // default square GEMM side length } int parse_workload_sync_interval(int argc, char** argv) { @@ -155,8 +143,8 @@ void GpuWorkload::destroy() { ok_ = false; } -bool GpuWorkload::init(BenchWorkload kind, size_t batch_bytes, int sync_interval, - int gemm_n_override) { +bool GpuWorkload::init(BenchWorkload kind, size_t input_bytes, int sync_interval, + int gemm_dim) { kind_ = kind; sync_interval_ = sync_interval > 0 ? sync_interval : 1; run_count_ = 0; @@ -165,7 +153,7 @@ bool GpuWorkload::init(BenchWorkload kind, size_t batch_bytes, int sync_interval return true; // inert no-op object; not an error } - const size_t bytes = batch_bytes > 0 ? batch_bytes : kDefaultBytes; + const size_t bytes = input_bytes > 0 ? input_bytes : kDefaultBytes; cudaStream_t stream = nullptr; if (cudaStreamCreate(&stream) != cudaSuccess) { @@ -207,27 +195,26 @@ bool GpuWorkload::init(BenchWorkload kind, size_t batch_bytes, int sync_interval std::cerr << "GpuWorkload: fft shape = " << batch << " x C2C length-" << kFftLen << " (working set " << (bytes >> 10) << " KiB, " << (flops / 1e6) << " MFLOP/call)\n"; } else { // Gemm or GemmFp16 - // Square matmul. By default size n from the working set (FP32 in both cases so - // gemm and gemm_fp16 use the SAME dimension -> identical FLOP count, isolating - // the precision / tensor-core effect). A caller-supplied gemm_n_override pins n - // directly so the FLOP count per call stays FIXED while the I/O unit is swept -- - // this isolates pipelining depth from problem size in the RoCE-vs-raw study. - // The A operand is the caller's input buffer (n*n*elem_size must be <= bytes, - // enforced below); B and C are owned scratch. - int n = gemm_n_override > 0 - ? gemm_n_override - : static_cast(std::sqrt(static_cast(bytes) / sizeof(float))); + // Square matmul with the side length n pinned directly (FP32 and FP16 use the + // SAME dimension -> identical FLOP count, isolating the precision / tensor-core + // effect). n is held FIXED while the I/O unit is swept, so 2*n^3 FLOPs/call + // stays constant -- this isolates pipelining depth from problem size in the + // RoCE-vs-raw study. The A operand is the caller's input buffer (n*n*elem_size + // must be <= input_bytes, enforced below); B and C are owned scratch. + int n = gemm_dim > 0 ? gemm_dim : 1024; n = std::max(64, (n / 8) * 8); // multiple of 8, sane floor gemm_n_ = n; const size_t elems = static_cast(n) * n; const size_t elem_size = kind_ == BenchWorkload::GemmFp16 ? sizeof(__half) : sizeof(float); // The A operand is read from the caller's received-data buffer, which holds - // `bytes`. A pinned n must fit -- otherwise cuBLAS reads past the buffer. Reject - // rather than silently OOB-read; the caller should raise --workload-batch-bytes. + // `bytes`. The pinned n must fit -- otherwise cuBLAS reads past the buffer. + // Reject rather than silently OOB-read; the caller should use a larger I/O unit + // (message / reorder window) for this GEMM dimension. if (elems * elem_size > bytes) { - std::cerr << "GpuWorkload: pinned gemm n=" << n << " needs " << (elems * elem_size >> 10) - << " KiB for the A operand but the working set is only " << (bytes >> 10) - << " KiB; raise --workload-batch-bytes to >= n*n*elem_size\n"; + std::cerr << "GpuWorkload: gemm n=" << n << " needs " << (elems * elem_size >> 10) + << " KiB for the A operand but the input buffer is only " << (bytes >> 10) + << " KiB; use a larger I/O unit (>= n*n*elem_size) or a smaller " + "--workload-gemm-dim\n"; destroy(); return false; } @@ -251,12 +238,11 @@ bool GpuWorkload::init(BenchWorkload kind, size_t batch_bytes, int sync_interval return false; } // Explicit shape so every published GEMM number carries its compute size. A - // square n×n×n matmul is 2·n³ flops; note whether n was pinned or derived. + // square n×n×n matmul is 2·n³ flops; n is pinned via --workload-gemm-dim. const double flops = 2.0 * static_cast(n) * n * n; std::cerr << "GpuWorkload: " << workload_name(kind_) << " shape = " << n << "x" << n << "x" << n - << (gemm_n_override > 0 ? " (pinned)" : " (derived)") << ", " - << (elem_size == sizeof(__half) ? "fp16" : "fp32") << " A/B, " << (flops / 1e9) - << " GFLOP/call, matrix " << (elems * elem_size >> 10) << " KiB\n"; + << " (pinned), " << (elem_size == sizeof(__half) ? "fp16" : "fp32") << " A/B, " + << (flops / 1e9) << " GFLOP/call, matrix " << (elems * elem_size >> 10) << " KiB\n"; } // Warm up and validate the chosen op once against a transient zeroed input diff --git a/examples/bench_workload.h b/examples/bench_workload.h index 4fa6052..7d5163c 100644 --- a/examples/bench_workload.h +++ b/examples/bench_workload.h @@ -32,18 +32,13 @@ enum class BenchWorkload { None, Fft, Gemm, GemmFp16 }; // the flag/value stride used by parse_run_seconds / parse_target_gbps. BenchWorkload parse_workload(int argc, char** argv); -// Parse "--workload-batch-bytes N" from argv: the working-set size (bytes) fed to -// one compute call, decoupled from the I/O unit (RoCE message / raw frame). The -// GEMM matrix dimension and FFT batch scale from it. Returns 0 if unset (the bench -// then falls back to its backend-default batch). Mirrors parse_workload's stride. -size_t parse_workload_batch_bytes(int argc, char** argv); - -// Parse "--workload-gemm-n N" from argv: pin the square GEMM dimension directly, -// independent of the compute batch bytes, so the FLOP count per call is FIXED as -// the I/O unit (message / burst window) is swept. Isolates pipelining depth from -// problem size in the RoCE-vs-raw comparison. Returns 0 if unset (the GEMM -// dimension then derives from the working-set size). Mirrors parse_workload's stride. -int parse_workload_gemm_n(int argc, char** argv); +// Parse "--workload-gemm-dim N" from argv: the square GEMM side length n, held +// FIXED so the FLOP count per call (2*n^3) is constant as the I/O unit (message / +// burst window) is swept -- this isolates pipelining depth from problem size in +// the RoCE-vs-raw comparison. The compute working set is n*n*elem_size, derived +// entirely from n; the assembled input buffer must be at least that large. +// Returns 1024 (the default) if unset. Mirrors parse_workload's stride. +int parse_workload_gemm_dim(int argc, char** argv); // Parse "--workload-sync-interval N" from argv: drain the GPU stream every N compute // calls (bounds outstanding GPU work). Larger N = deeper async queue, fewer CPU @@ -79,20 +74,23 @@ class GpuWorkload { GpuWorkload(const GpuWorkload&) = delete; GpuWorkload& operator=(const GpuWorkload&) = delete; - // Build the plan/handle and size the problem to the contiguous input buffer of - // batch_bytes the caller will pass to run() (0 => an internal default). The op - // reads at most batch_bytes from that buffer. kind == None leaves the object an - // inert no-op. sync_interval bounds outstanding GPU work (sync every N runs). - // gemm_n_override (>0) pins the square GEMM dimension directly, holding the FLOP - // count per call fixed regardless of batch_bytes (0 => derive n from batch_bytes). - // Logs the chosen problem shape (GEMM n / FFT length+batch) and FLOP count so - // every published benchmark number carries its explicit compute size. - // Returns false on CUDA / library error; the caller may warn and continue with - // the workload disabled (enabled() will report false). - bool init(BenchWorkload kind, size_t batch_bytes, int sync_interval = 2, int gemm_n_override = 0); + // Build the plan/handle for `kind`. `input_bytes` is the size of the contiguous + // buffer the caller will pass to run() (its reorder window / message size, 0 => + // an internal default); it bounds the op's read and is validated against the + // pinned GEMM working set. `gemm_dim` is the square GEMM side length n, held + // fixed so 2*n^3 FLOPs/call stays constant across the I/O sweep (>0 required for + // GEMM; <=0 falls back to the 1024 default). kind == None leaves the object an + // inert no-op. sync_interval bounds outstanding GPU work (sync every N runs) for + // the run()+sync() callers (DPDK/socket); the RoCE path bounds work with events + // instead and ignores it. Logs the chosen problem shape (GEMM n / FFT + // length+batch) and FLOP count so every published benchmark number carries its + // explicit compute size. Returns false on CUDA / library error; the caller may + // warn and continue with the workload disabled (enabled() will report false). + bool init(BenchWorkload kind, size_t input_bytes, int sync_interval = 2, int gemm_dim = 1024); // Enqueue one representative FFT/SGEMM on the internal stream, reading `input` - // (a device pointer to >= batch_bytes valid bytes). No-op unless enabled(). + // (a device pointer to >= the GEMM working set of valid bytes). No-op unless + // enabled(). void run(const void* input); // Every sync_interval runs, block until the stream drains so the GPU stays on diff --git a/examples/raw_bench_common.cpp b/examples/raw_bench_common.cpp index ef686a1..3d266e5 100644 --- a/examples/raw_bench_common.cpp +++ b/examples/raw_bench_common.cpp @@ -549,7 +549,8 @@ void print_queue_stats(const char *direction, const std::string &interface_name, } void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, BenchWorkload workload, - const ReorderGeometry& geom, int workload_gemm_n, int workload_sync_interval) { + const ReorderGeometry& geom, int workload_gemm_dim, + int workload_sync_interval) { if (!set_current_thread_affinity(cfg.cpu_core, "bench_rx")) { stop.store(true); return; @@ -564,7 +565,7 @@ void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, Bench GpuWorkload gpu_workload; ReorderPipeline pipeline; if (workload != BenchWorkload::None) { - if (!gpu_workload.init(workload, batch_bytes, workload_sync_interval, workload_gemm_n) || + if (!gpu_workload.init(workload, batch_bytes, workload_sync_interval, workload_gemm_dim) || !pipeline.init(ReorderMode::SeqReorder, geom.packets_per_batch, out_payload_len, geom.payload_byte_offset, geom.seq_bit_offset, geom.seq_bit_width, /*staging_needed=*/false, gpu_workload.stream())) { diff --git a/examples/raw_bench_common.h b/examples/raw_bench_common.h index 496a2c7..bc006dd 100644 --- a/examples/raw_bench_common.h +++ b/examples/raw_bench_common.h @@ -161,6 +161,6 @@ struct ReorderGeometry { // workload == None leaves the bare-loopback count-only path untouched. void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, BenchWorkload workload = BenchWorkload::None, const ReorderGeometry& geom = {}, - int workload_gemm_n = 0, int workload_sync_interval = 2); + int workload_gemm_dim = 1024, int workload_sync_interval = 2); } // namespace daqiri::bench diff --git a/examples/raw_gpudirect_bench.cpp b/examples/raw_gpudirect_bench.cpp index 0181454..06d26b4 100644 --- a/examples/raw_gpudirect_bench.cpp +++ b/examples/raw_gpudirect_bench.cpp @@ -163,7 +163,7 @@ int main(int argc, char **argv) { if (argc < 2) { std::cerr << "Usage: " << argv[0] << " [--seconds N] [--target-gbps G] " - "[--workload none|fft|gemm|gemm_fp16] [--workload-batch-bytes N]\n"; + "[--workload none|fft|gemm|gemm_fp16] [--workload-gemm-dim N]\n"; return 1; } @@ -172,8 +172,7 @@ int main(int argc, char **argv) { const int run_seconds = daqiri::bench::parse_run_seconds(argc, argv); const double target_gbps = daqiri::bench::parse_target_gbps(argc, argv); const auto workload = daqiri::bench::parse_workload(argc, argv); - const size_t workload_batch_bytes = daqiri::bench::parse_workload_batch_bytes(argc, argv); - const int workload_gemm_n = daqiri::bench::parse_workload_gemm_n(argc, argv); + const int workload_gemm_dim = daqiri::bench::parse_workload_gemm_dim(argc, argv); const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); @@ -214,19 +213,16 @@ int main(int argc, char **argv) { geom.seq_bit_offset = static_cast(tx.header_size * 8); geom.seq_bit_width = 32; geom.out_payload_len = tx.payload_size; - // Reorder window = the workload batch. --workload-batch-bytes sets it - // (rounded to whole packets); default ~8 MB (1024 packets at the native - // shape). Capped to one burst's packet count. - const uint32_t ppb = - workload_batch_bytes > 0 - ? std::max(1, static_cast(workload_batch_bytes / tx.payload_size)) - : 1024; - geom.packets_per_batch = std::min(ppb, tx.batch_size); + // Reorder window: ~8 MB (1024 packets at the native shape), matching the RoCE + // working set. Capped to one burst's packet count. The compute reads the first + // n*n*elem_size bytes of this window (GEMM dimension pinned via + // --workload-gemm-dim), so the window must be at least that large. + geom.packets_per_batch = std::min(1024, tx.batch_size); } rx_threads.reserve(rx_configs.size()); for (const auto &cfg : rx_configs) { rx_threads.emplace_back(daqiri::bench::rx_count_worker, cfg, std::ref(stop), workload, geom, - workload_gemm_n, workload_sync_interval); + workload_gemm_dim, workload_sync_interval); } tx_threads.reserve(tx_configs.size()); for (const auto &cfg : tx_configs) { diff --git a/examples/raw_hds_bench.cpp b/examples/raw_hds_bench.cpp index 743bc13..8d936ee 100644 --- a/examples/raw_hds_bench.cpp +++ b/examples/raw_hds_bench.cpp @@ -141,14 +141,13 @@ int main(int argc, char **argv) { if (argc < 2) { std::cerr << "Usage: " << argv[0] << " [--seconds N] " - "[--workload none|fft|gemm|gemm_fp16] [--workload-batch-bytes N]\n"; + "[--workload none|fft|gemm|gemm_fp16] [--workload-gemm-dim N]\n"; return 1; } const int run_seconds = daqiri::bench::parse_run_seconds(argc, argv); const auto workload = daqiri::bench::parse_workload(argc, argv); - const size_t workload_batch_bytes = daqiri::bench::parse_workload_batch_bytes(argc, argv); - const int workload_gemm_n = daqiri::bench::parse_workload_gemm_n(argc, argv); + const int workload_gemm_dim = daqiri::bench::parse_workload_gemm_dim(argc, argv); const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); if (daqiri::daqiri_init(argv[1]) != daqiri::Status::SUCCESS) { @@ -182,15 +181,14 @@ int main(int argc, char **argv) { geom.seq_bit_offset = 0; geom.seq_bit_width = 32; geom.out_payload_len = tx.payload_size; - const uint32_t ppb = - workload_batch_bytes > 0 - ? std::max(1, static_cast(workload_batch_bytes / tx.payload_size)) - : 1024; - geom.packets_per_batch = std::min(ppb, tx.batch_size); + // Reorder window: ~8 MB (1024 packets at the native shape), capped to one + // burst. The compute reads the first n*n*elem_size bytes (GEMM dimension + // pinned via --workload-gemm-dim). + geom.packets_per_batch = std::min(1024, tx.batch_size); } rx_thread = std::thread(daqiri::bench::rx_count_worker, daqiri::bench::parse_rx(root), std::ref(stop), - workload, geom, workload_gemm_n, workload_sync_interval); + workload, geom, workload_gemm_dim, workload_sync_interval); } if (has_tx) { tx_thread = diff --git a/examples/rdma_bench.cpp b/examples/rdma_bench.cpp index 00d7266..4c0e127 100644 --- a/examples/rdma_bench.cpp +++ b/examples/rdma_bench.cpp @@ -81,8 +81,8 @@ RdmaBenchConfig parse_rdma_cfg(const YAML::Node& node) { void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pacer, std::atomic& stop, RdmaWorkerStats& stats, - daqiri::bench::BenchWorkload workload, size_t workload_batch_bytes, - int workload_gemm_n, int workload_sync_interval) { + daqiri::bench::BenchWorkload workload, int workload_gemm_dim, + int workload_sync_interval) { const char *thread_name = cfg.server ? "rdma_bench_server" : "rdma_bench_client"; if (!daqiri::bench::set_current_thread_affinity(cfg.cpu_core, thread_name)) { @@ -96,19 +96,14 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa // GPU-accessible: the gather-only pipeline is a zero-copy pass-through that // hands the recv pointer straight to the compute. // - // The compute working set is decoupled from the 8 MB message: --workload-batch- - // bytes sub-divides each message into chunk-sized slices, one compute per slice - // (the per-message sync then amortizes over num_chunks computes). 0 or >= the - // message size means one compute per message. + // One fixed-dimension compute per received message, reading the first + // n*n*elem_size bytes of the message (GpuWorkload::init rejects a message too + // small to hold the pinned GEMM operand). const uint32_t msg = static_cast(cfg.message_size); - const uint32_t chunk = (workload_batch_bytes > 0 && workload_batch_bytes < msg) - ? static_cast(workload_batch_bytes) - : msg; - const uint32_t num_chunks = chunk > 0 ? msg / chunk : 1; daqiri::bench::GpuWorkload gpu_workload; daqiri::bench::ReorderPipeline pipeline; if (workload != daqiri::bench::BenchWorkload::None) { - if (!gpu_workload.init(workload, chunk, workload_sync_interval, workload_gemm_n) || + if (!gpu_workload.init(workload, msg, workload_sync_interval, workload_gemm_dim) || !pipeline.init(daqiri::bench::ReorderMode::GatherOnly, /*packets_per_batch=*/1, msg, /*payload_byte_offset=*/0, /*seq_bit_offset=*/0, @@ -259,13 +254,10 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa if (run_workload) { pipeline.reset_batch(); pipeline.add_device_packet(daqiri::get_segment_packet_ptr(completion, 0, 0)); - const auto* base = static_cast(pipeline.finish_batch()); - // One compute per chunk-sized slice of the message (num_chunks == 1 when - // the batch is the whole message). Enqueue only; the GPU work overlaps - // with subsequent receives and is drained per batch (below). - for (uint32_t k = 0; k < num_chunks; ++k) { - gpu_workload.run(base + static_cast(k) * chunk); - } + // One fixed-dimension compute per message on its real (zero-copy) data. + // Enqueue only; the GPU work overlaps with subsequent receives and is + // drained per batch (below). + gpu_workload.run(pipeline.finish_batch()); // Hold the completion (and thus its recv buffer) until the batch drains, // so the pass-through compute reads valid data without a per-message sync. held_recv.push_back(completion); @@ -306,7 +298,7 @@ int main(int argc, char** argv) { std::cerr << "Usage: " << argv[0] << " [--seconds N] [--mode server|client|both] " "[--target-gbps G] [--workload none|fft|gemm|gemm_fp16] " - "[--workload-batch-bytes N]\n"; + "[--workload-gemm-dim N]\n"; return 1; } @@ -323,8 +315,7 @@ int main(int argc, char** argv) { } } const auto workload = daqiri::bench::parse_workload(argc, argv); - const size_t workload_batch_bytes = daqiri::bench::parse_workload_batch_bytes(argc, argv); - const int workload_gemm_n = daqiri::bench::parse_workload_gemm_n(argc, argv); + const int workload_gemm_dim = daqiri::bench::parse_workload_gemm_dim(argc, argv); const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); @@ -362,13 +353,13 @@ int main(int argc, char** argv) { if (run_server) { server_thread = std::thread(rdma_worker, server_cfg, std::ref(server_pacer), std::ref(stop), - std::ref(server_stats), workload, workload_batch_bytes, - workload_gemm_n, workload_sync_interval); + std::ref(server_stats), workload, workload_gemm_dim, + workload_sync_interval); } if (run_client) { client_thread = std::thread(rdma_worker, client_cfg, std::ref(client_pacer), std::ref(stop), - std::ref(client_stats), workload, workload_batch_bytes, - workload_gemm_n, workload_sync_interval); + std::ref(client_stats), workload, workload_gemm_dim, + workload_sync_interval); } if (!server_thread.joinable() && !client_thread.joinable()) { diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index c0c0606..ad56c66 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -29,14 +29,10 @@ # tensor-core matmul). Honoured by all backends (dpdk, rdma, # socket-udp, socket-tcp); recorded in the CSV post_process # column. -# WORKLOAD_BATCH — GPU compute working-set bytes per call, decoupled from the -# I/O unit (empty = backend default). Recorded in -# post_process_batch. -# GEMM_N — pin the square GEMM dimension n directly (--workload-gemm-n), -# holding FLOPs/call fixed (2·n³) while the I/O unit is swept -# via PAYLOADS_OVERRIDE. The fixed-n message-size sweep that -# isolates pipelining depth from problem size. Recorded in -# post_process_gemm_n. +# GEMM_DIM — the square GEMM side length n (--workload-gemm-dim; default +# 1024), held fixed so FLOPs/call (2·n³) is constant. The +# compute working set is n·n·elem_size, read from the front of +# each received I/O unit. Recorded in post_process_gemm_dim. # SYNC_INTERVAL — drain the GPU stream every N compute calls # (--workload-sync-interval; default 2). Sweep it (1 2 4 8 16 32) # to see how much of the receive+compute ceiling is single-thread @@ -80,11 +76,9 @@ CSV="$OUT_DIR/runs.csv" # `pairs` = number of concurrent client/server process pairs (socket backends sweep # this; dpdk/rdma are always 1). `gbps` is aggregate App TX, `rx_gbps` aggregate App RX # (summed across pairs); App-level loss is (gbps - rx_gbps) / gbps. -# post_process_batch = WORKLOAD_BATCH bytes, or "default" when unset. -# post_process_gemm_n = GEMM_N pinned dimension, or "derived" when unset. +# post_process_gemm_dim = GEMM_DIM pinned dimension (default 1024). # post_process_sync (last column) = SYNC_INTERVAL, or "default" (2) when unset. -# Appended at the end so existing column positions are unchanged. -echo "lang,backend,post_process,payload,batch,pairs,target_gbps,rep,seconds,packets,bytes,pps,gbps,rx_gbps,drops,drops_kind,cpu_master_pct,cpu_tx_pct,cpu_rx_pct,gpu_sm_pct,gpu_mem_pct,post_process_batch,post_process_gemm_n,post_process_sync" > "$CSV" +echo "lang,backend,post_process,payload,batch,pairs,target_gbps,rep,seconds,packets,bytes,pps,gbps,rx_gbps,drops,drops_kind,cpu_master_pct,cpu_tx_pct,cpu_rx_pct,gpu_sm_pct,gpu_mem_pct,post_process_gemm_dim,post_process_sync" > "$CSV" # Capture slow-moving environment state once per result set. "$SCRIPT_DIR/bench_capture_environment.sh" "$OUT_DIR" @@ -104,23 +98,14 @@ case "$WORKLOAD" in none|fft|gemm|gemm_fp16) ;; *) echo "Invalid WORKLOAD '$WORKLOAD' (expected none|fft|gemm|gemm_fp16)" >&2; exit 1 ;; esac -# WORKLOAD_BATCH (bytes): the GPU compute working set per call, decoupled from the -# I/O unit (RoCE sub-chunks each message; raw sets the reorder window). Empty = -# backend default. Sweep it (e.g. 262144 524288 1048576 2097152 4194304 8388608) -# to trace the compute-intensity vs throughput curve. Recorded in post_process_batch. -WORKLOAD_BATCH="${WORKLOAD_BATCH:-}" -if [[ -n "$WORKLOAD_BATCH" && ! "$WORKLOAD_BATCH" =~ ^[0-9]+$ ]]; then - echo "Invalid WORKLOAD_BATCH '$WORKLOAD_BATCH' (expected a positive byte count)" >&2; exit 1 -fi -# GEMM_N: pin the square GEMM dimension n directly (--workload-gemm-n), holding the -# FLOP count per call FIXED (2·n³) while the I/O unit (RoCE message via -# PAYLOADS_OVERRIDE, raw reorder window) is swept. This isolates pipelining depth -# from problem size -- the fixed-n message-size sweep that proves the RoCE↔raw gap -# is pipelining, not transport. Empty = derive n from the working set. Recorded in -# the CSV post_process_gemm_n column. -GEMM_N="${GEMM_N:-}" -if [[ -n "$GEMM_N" && ! "$GEMM_N" =~ ^[0-9]+$ ]]; then - echo "Invalid GEMM_N '$GEMM_N' (expected a positive integer)" >&2; exit 1 +# GEMM_DIM: the square GEMM side length n (--workload-gemm-dim), held FIXED so the +# FLOP count per call (2·n³) is constant. The compute working set is n·n·elem_size, +# read from the front of each received I/O unit (RoCE message / raw reorder window), +# so that unit must be at least that large. Default 1024. Recorded in the CSV +# post_process_gemm_dim column. +GEMM_DIM="${GEMM_DIM:-1024}" +if [[ ! "$GEMM_DIM" =~ ^[0-9]+$ ]]; then + echo "Invalid GEMM_DIM '$GEMM_DIM' (expected a positive integer)" >&2; exit 1 fi # SYNC_INTERVAL: drain the GPU stream every N compute calls (--workload-sync-interval). # Larger N lets more GEMMs queue asynchronously before the single receive+compute @@ -220,18 +205,16 @@ case "$BACKEND" in *) echo "Unknown backend: $BACKEND" >&2; exit 1 ;; esac -# Optional space-separated overrides for the sweep ladders, so a one-off experiment -# can re-target the payload/batch matrix without editing the per-backend defaults -# above. Used by the fixed-n workload sweep: pin the GEMM dimension with GEMM_N -# (constant 2·n³ FLOPs/call) and sweep PAYLOADS_OVERRIDE = the RoCE message size, so -# num_chunks = message/batch varies pipelining depth at a FIXED matrix size. This is -# the "send an 80 MB message, launch N same-size GEMMs" experiment that isolates -# pipelining depth from problem size -- proving the RoCE↔raw gap is pipelining, not -# transport. Hold WORKLOAD_BATCH at the per-chunk size (= the GEMM working set). -# e.g. WORKLOAD=gemm_fp16 GEMM_N=1024 WORKLOAD_BATCH=4194304 \ -# PAYLOADS_OVERRIDE="4194304 8388608 16777216 33554432 83886080" \ -# ./run_spark_bench.sh rdma sweep -[[ -n "${PAYLOADS_OVERRIDE:-}" ]] && read -r -a PAYLOADS_SWEEP <<< "$PAYLOADS_OVERRIDE" +# When a GPU workload is active, its pinned GEMM operand (n·n·elem_size, from +# GEMM_DIM) must fit inside each received I/O unit -- the small entries in the +# default payload sweep can't hold it. Restrict the sweep to the headline +# (native-shape) size so the fixed-n comparison is a single clean point per +# backend instead of silently disabling the workload on the small cells. +# e.g. WORKLOAD=gemm_fp16 GEMM_DIM=1024 REPEATS=3 ./run_spark_bench.sh rdma sweep +if [[ "$WORKLOAD" != "none" ]]; then + PAYLOADS_SWEEP=("${PAYLOADS_HEADLINE[@]}") +fi +# Optional space-separated override for the batch ladder (one-off experiments). [[ -n "${BATCHES_OVERRIDE:-}" ]] && read -r -a BATCHES_SWEEP <<< "$BATCHES_OVERRIDE" # All backends (dpdk, rdma, socket-udp, socket-tcp) now run the workload on real @@ -446,8 +429,7 @@ run_cell() { # Every backend honours --workload (runs it on real received data); none = skip. if [[ "$WORKLOAD_EFF" != "none" ]]; then bench_extra+=(--workload "$WORKLOAD_EFF") - [[ -n "$WORKLOAD_BATCH" ]] && bench_extra+=(--workload-batch-bytes "$WORKLOAD_BATCH") - [[ -n "$GEMM_N" ]] && bench_extra+=(--workload-gemm-n "$GEMM_N") + bench_extra+=(--workload-gemm-dim "$GEMM_DIM") [[ -n "$SYNC_INTERVAL" ]] && bench_extra+=(--workload-sync-interval "$SYNC_INTERVAL") fi # Shutdown ordering: the server must keep receiving until the client has fully @@ -621,13 +603,10 @@ run_cell() { gpu_mem="$(awk '/^ *[0-9]/ { count++; sum += $6 } END { if (count) printf "%.1f", sum/count; else print 0 }' \ "$cell_dir/nvidia_smi_dmon.txt" 2>/dev/null || echo 0)" - local pp_batch="${WORKLOAD_BATCH:-default}" - [[ -z "$pp_batch" ]] && pp_batch="default" - local pp_gemm_n="${GEMM_N:-derived}" - [[ -z "$pp_gemm_n" ]] && pp_gemm_n="derived" + local pp_gemm_dim="$GEMM_DIM" local pp_sync="${SYNC_INTERVAL:-default}" [[ -z "$pp_sync" ]] && pp_sync="default" - echo "$lang,$BACKEND,$WORKLOAD_EFF,$payload,$batch,$pairs,$target_gbps,$rep,$secs,$pkts,$bytes,$pps,$gbps,$rx_gbps,$drops,$drops_kind,$cpu_master_pct,$cpu_tx_pct,$cpu_rx_pct,$gpu_sm,$gpu_mem,$pp_batch,$pp_gemm_n,$pp_sync" \ + echo "$lang,$BACKEND,$WORKLOAD_EFF,$payload,$batch,$pairs,$target_gbps,$rep,$secs,$pkts,$bytes,$pps,$gbps,$rx_gbps,$drops,$drops_kind,$cpu_master_pct,$cpu_tx_pct,$cpu_rx_pct,$gpu_sm,$gpu_mem,$pp_gemm_dim,$pp_sync" \ | tee -a "$CSV" } diff --git a/examples/socket_bench.cpp b/examples/socket_bench.cpp index dfc64ad..40048de 100644 --- a/examples/socket_bench.cpp +++ b/examples/socket_bench.cpp @@ -102,8 +102,8 @@ bool socket_transport_is_tcp(const YAML::Node& root) { void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pacer, std::atomic& stop, SocketWorkerStats& stats, - daqiri::bench::BenchWorkload workload, bool is_tcp, size_t workload_batch_bytes, - int workload_gemm_n, int workload_sync_interval) { + daqiri::bench::BenchWorkload workload, bool is_tcp, int workload_gemm_dim, + int workload_sync_interval) { const char *thread_name = cfg.server ? "socket_bench_server" : "socket_bench_client"; if (!daqiri::bench::set_current_thread_affinity(cfg.cpu_core, thread_name)) { @@ -116,18 +116,17 @@ void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer // copy on the workload stream -- the honest cost of the socket path). UDP can // reorder by sequence number; TCP is in-order, so it gathers per chunk. const uint32_t msg = static_cast(cfg.message_size); - // --workload-batch-bytes sets the UDP ordered-buffer / working-set size - // (default ~8 MB, matching the RoCE working set); the UDP reorder window is that - // many datagrams. TCP gathers one chunk at a time (one compute per recv). - const uint32_t target_bytes = - workload_batch_bytes > 0 ? static_cast(workload_batch_bytes) : 8u * 1024u * 1024u; + // The UDP reorder window is a fixed ~8 MB of datagrams (matching the RoCE + // working set); TCP gathers one chunk at a time (one compute per recv). The GEMM + // reads the first n*n*elem_size bytes of that window (dimension pinned below). + const uint32_t target_bytes = 8u * 1024u * 1024u; const uint32_t packets_per_batch = is_tcp ? 1u : std::max(1u, std::min(8192u, target_bytes / std::max(1u, msg))); daqiri::bench::GpuWorkload gpu_workload; daqiri::bench::ReorderPipeline pipeline; if (workload != daqiri::bench::BenchWorkload::None) { if (!gpu_workload.init(workload, static_cast(packets_per_batch) * msg, - workload_sync_interval, workload_gemm_n) || + workload_sync_interval, workload_gemm_dim) || !pipeline.init(is_tcp ? daqiri::bench::ReorderMode::GatherOnly : daqiri::bench::ReorderMode::SeqReorder, packets_per_batch, msg, /*payload_byte_offset=*/0, @@ -242,7 +241,7 @@ int main(int argc, char** argv) { std::cerr << "Usage: " << argv[0] << " [--seconds N] [--mode server|client|both] " "[--target-gbps G] [--workload none|fft|gemm|gemm_fp16] " - "[--workload-batch-bytes N]\n"; + "[--workload-gemm-dim N]\n"; return 1; } @@ -259,8 +258,7 @@ int main(int argc, char** argv) { } } const auto workload = daqiri::bench::parse_workload(argc, argv); - const size_t workload_batch_bytes = daqiri::bench::parse_workload_batch_bytes(argc, argv); - const int workload_gemm_n = daqiri::bench::parse_workload_gemm_n(argc, argv); + const int workload_gemm_dim = daqiri::bench::parse_workload_gemm_dim(argc, argv); const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); @@ -299,13 +297,13 @@ int main(int argc, char** argv) { if (run_server) { server_thread = std::thread(socket_worker, server_cfg, std::ref(server_pacer), std::ref(stop), - std::ref(server_stats), workload, is_tcp, workload_batch_bytes, - workload_gemm_n, workload_sync_interval); + std::ref(server_stats), workload, is_tcp, workload_gemm_dim, + workload_sync_interval); } if (run_client) { client_thread = std::thread(socket_worker, client_cfg, std::ref(client_pacer), std::ref(stop), - std::ref(client_stats), workload, is_tcp, workload_batch_bytes, - workload_gemm_n, workload_sync_interval); + std::ref(client_stats), workload, is_tcp, workload_gemm_dim, + workload_sync_interval); } if (!server_thread.joinable() && !client_thread.joinable()) { From bd36f79ac3f4486273c2b8abd6e48f08b02d7724 Mon Sep 17 00:00:00 2001 From: Ramya Gurunathan Date: Wed, 8 Jul 2026 14:51:12 -0400 Subject: [PATCH 05/18] #15 - Recycle RoCE recv buffers with CUDA events instead of blocking drain The RoCE workload bench read each recv buffer zero-copy and, to know when the buffer was safe to free/repost, drained the whole GPU stream (cudaStreamSynchronize) once per held-receive batch -- and run() also drained every sync_interval GEMMs. On the single receive thread that drain parks the thread on the GPU, so it stops reaping completions and reposting receives while the GEMM runs. That serialization (not the GPU, buffers, or transport) is why the fixed-n RoCE point ran ~2x slower than raw at the identical workload. Replace the blocking drain with non-blocking, event-based recycling: - GpuWorkload gains a lazily-created pool of cudaEvents (cudaEventDisableTiming) and record_event / event_done / wait_event / release_event, plus run_async (enqueue with no maybe_sync). Events complete in stream order, so an event stamped after a message's GEMM signals that the read of that recv buffer is done. The pool costs nothing for the run()+sync() callers that never record. - rdma_bench holds each completion with its event in a FIFO, frees buffers as soon as event_done() reports the GEMM finished (reclaim_completed), and keeps reaping completions meanwhile. Outstanding GPU work is bounded by max_inflight (backpressure: block on the oldest event only at the cap), which replaces the periodic sync_interval drain as the measurement-honesty bound. Shutdown / connection loss drains all events fully before teardown. DPDK/socket are unchanged (they copy into the pipeline's shared output buffer, so overlapping GEMMs would need that buffer ring-buffered first); a comment in raw_bench_common notes the prerequisite. --workload-sync-interval still governs the DPDK/socket run()+sync() path; the RoCE path bounds work with events and ignores it. Example-only change; no library rebuild required. Signed-off-by: Ramya Gurunathan Assisted-by: Claude Opus 4.8 --- examples/bench_workload.cu | 69 ++++++++++++++++++++++++++++ examples/bench_workload.h | 28 ++++++++++++ examples/raw_bench_common.cpp | 5 +++ examples/rdma_bench.cpp | 85 +++++++++++++++++++++++------------ 4 files changed, 159 insertions(+), 28 deletions(-) diff --git a/examples/bench_workload.cu b/examples/bench_workload.cu index d89019b..bd4b100 100644 --- a/examples/bench_workload.cu +++ b/examples/bench_workload.cu @@ -37,6 +37,9 @@ namespace { constexpr int kFftLen = 1024; // Default working-set size when the caller passes bytes_per_burst == 0. constexpr size_t kDefaultBytes = 1u << 16; // 64 KiB +// Fixed CUDA-event pool for record_event(); comfortably above the RoCE recv-path +// in-flight cap (min(rx_depth/2, 8)). Created lazily on first record_event(). +constexpr int kEventPoolSize = 32; cufftHandle as_fft_plan(int p) { return static_cast(p); @@ -47,6 +50,9 @@ cudaStream_t as_stream(void* s) { cublasHandle_t as_cublas(void* h) { return static_cast(h); } +cudaEvent_t as_event(void* e) { + return static_cast(e); +} } // namespace @@ -136,6 +142,11 @@ void GpuWorkload::destroy() { cudaFree(gemm_c_); gemm_c_ = nullptr; } + for (void* ev : event_pool_) { + cudaEventDestroy(as_event(ev)); + } + event_pool_.clear(); + free_events_.clear(); if (stream_ != nullptr) { cudaStreamDestroy(as_stream(stream_)); stream_ = nullptr; @@ -298,6 +309,64 @@ void GpuWorkload::run(const void* input) { maybe_sync(); } +void GpuWorkload::run_async(const void* input) { + if (!enabled() || input == nullptr) { + return; + } + issue_op(input); + ++run_count_; + // No maybe_sync: the caller bounds outstanding work with events. +} + +void* GpuWorkload::record_event() { + if (!enabled()) { + return nullptr; + } + // Lazily create the pool on first use so the run()+sync() callers (DPDK/socket) + // that never record events pay nothing. + if (event_pool_.empty()) { + event_pool_.reserve(kEventPoolSize); + free_events_.reserve(kEventPoolSize); + for (int i = 0; i < kEventPoolSize; ++i) { + cudaEvent_t ev = nullptr; + if (cudaEventCreateWithFlags(&ev, cudaEventDisableTiming) != cudaSuccess) { + break; // partial pool is fine; record_event just returns null when empty + } + event_pool_.push_back(ev); + free_events_.push_back(ev); + } + } + if (free_events_.empty()) { + return nullptr; // caller must drain (event_done / wait_event) and retry + } + void* ev = free_events_.back(); + free_events_.pop_back(); + if (cudaEventRecord(as_event(ev), as_stream(stream_)) != cudaSuccess) { + free_events_.push_back(ev); + return nullptr; + } + return ev; +} + +bool GpuWorkload::event_done(void* ev) const { + if (ev == nullptr) { + return true; + } + return cudaEventQuery(as_event(ev)) == cudaSuccess; +} + +void GpuWorkload::wait_event(void* ev) { + if (ev != nullptr) { + cudaEventSynchronize(as_event(ev)); + } +} + +void GpuWorkload::release_event(void* ev) { + if (ev != nullptr) { + free_events_.push_back(ev); + } +} + void GpuWorkload::maybe_sync() { if (!enabled()) { return; diff --git a/examples/bench_workload.h b/examples/bench_workload.h index 7d5163c..36d8f60 100644 --- a/examples/bench_workload.h +++ b/examples/bench_workload.h @@ -18,6 +18,7 @@ #pragma once #include +#include namespace daqiri::bench { @@ -93,6 +94,11 @@ class GpuWorkload { // enabled(). void run(const void* input); + // Like run(), but enqueue only -- never blocks (no maybe_sync). For callers that + // bound outstanding GPU work with events (record_event) instead of the periodic + // maybe_sync drain, so the receive thread is never parked on the GPU. + void run_async(const void* input); + // Every sync_interval runs, block until the stream drains so the GPU stays on // the critical path without unbounded queueing. void maybe_sync(); @@ -100,6 +106,22 @@ class GpuWorkload { // Drain any remaining queued work (call once on shutdown). void sync(); + // Event-based buffer recycling for callers that read a still-owned input buffer + // zero-copy (the RoCE recv path): record an event on the stream after the ops + // for one input have been enqueued; the event completes only once those ops + // (and thus the reads of that input) finish. Poll it non-blocking with + // event_done() to free/repost the input, or block on it (wait_event) for + // backpressure / shutdown. Handles come from a fixed internal pool; release_event + // returns one for reuse. All are no-ops / return true / null when disabled. + // + // Enabled callers must keep at most pool-capacity events outstanding (wait+ + // release the oldest before exceeding it); record_event returns null if the pool + // is momentarily exhausted so the caller can drain and retry. + void* record_event(); + bool event_done(void* ev) const; + void wait_event(void* ev); + void release_event(void* ev); + bool enabled() const { return kind_ != BenchWorkload::None && ok_; } @@ -137,6 +159,12 @@ class GpuWorkload { void* gemm_b_ = nullptr; // B operand void* gemm_c_ = nullptr; // C output int gemm_n_ = 0; + + // Fixed pool of CUDA events (cudaEvent_t, cast in the .cu) for record_event(). + // Created lazily on first use; freed in destroy(). free_events_ holds the + // currently-available handles. + std::vector event_pool_; // all owned events (for teardown) + std::vector free_events_; // subset available to hand out }; } // namespace daqiri::bench diff --git a/examples/raw_bench_common.cpp b/examples/raw_bench_common.cpp index 3d266e5..28db189 100644 --- a/examples/raw_bench_common.cpp +++ b/examples/raw_bench_common.cpp @@ -622,6 +622,11 @@ void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, Bench // is freed, so process complete batches within this burst (the remainder // < packets_per_batch is dropped), then drain the stream so the reorder // kernel has finished reading the buffers before they are recycled. + // Unlike the RoCE path (which reads the recv buffer zero-copy and recycles + // it with CUDA events), the reorder here copies into the pipeline's single + // shared output buffer, so overlapping GEMMs across bursts would need that + // buffer double/ring-buffered first -- hence the per-burst blocking drain. + // The raw path already sustains line rate, so events are not wired in here. pipeline.reset_batch(); for (int i = 0; i < num_pkts; ++i) { pipeline.add_device_packet( diff --git a/examples/rdma_bench.cpp b/examples/rdma_bench.cpp index 4c0e127..e5a8dee 100644 --- a/examples/rdma_bench.cpp +++ b/examples/rdma_bench.cpp @@ -22,6 +22,7 @@ #include #include #include +#include #include #include #include @@ -113,23 +114,40 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa } const bool run_workload = gpu_workload.enabled() && pipeline.enabled(); - // To overlap receive with compute (apples-to-apples with the raw path, which - // holds a whole burst and drains the GPU once per burst), hold a small batch of - // RECEIVE completions instead of draining + freeing after each one. Holding a - // completion keeps its recv buffer alive, so the pass-through compute can read it - // while later messages keep arriving into other pool buffers; we drain the stream - // once per batch, then free them all. Bounded well under rx_depth so receives are - // not starved. - std::vector held_recv; - const size_t recv_hold_batch = + // To overlap receive with compute (apples-to-apples with the raw path), hold each + // RECEIVE completion until the GPU has finished reading its recv buffer, then + // free/repost it -- WITHOUT blocking the single receive thread. The compute reads + // the recv buffer zero-copy, so the buffer must stay alive until its GEMM + // completes. We stamp a CUDA event on the stream after each message's GEMM and + // free the buffer only once event_done() reports the read is finished, so the + // thread keeps reaping completions and reposting meanwhile. Outstanding GPU work + // is bounded by max_inflight (backpressure, not a periodic hard drain), which is + // what keeps the throughput measurement honest. Bounded well under rx_depth (and + // the event pool) so receives are not starved. + struct HeldRecv { + daqiri::BurstParams* buf; + void* event; + }; + std::deque held_recv; + const size_t max_inflight = run_workload ? std::max(1, std::min(cfg.rx_depth / 2, 8)) : 0; - auto flush_held_recv = [&]() { - if (held_recv.empty()) { - return; + // Free every held buffer whose GEMM has completed (non-blocking); keeps the recv + // pool drained without parking the thread on the GPU. + auto reclaim_completed = [&]() { + while (!held_recv.empty() && gpu_workload.event_done(held_recv.front().event)) { + daqiri::free_tx_burst(held_recv.front().buf); + gpu_workload.release_event(held_recv.front().event); + held_recv.pop_front(); } - gpu_workload.sync(); - for (auto* held : held_recv) { - daqiri::free_tx_burst(held); + }; + // Block until every held buffer's GEMM is done, then free them all (shutdown / + // connection loss). Events complete in stream order, so waiting front-to-back + // drains the whole queue. + auto drain_held_recv = [&]() { + for (auto& held : held_recv) { + gpu_workload.wait_event(held.event); + daqiri::free_tx_burst(held.buf); + gpu_workload.release_event(held.event); } held_recv.clear(); }; @@ -255,16 +273,27 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa pipeline.reset_batch(); pipeline.add_device_packet(daqiri::get_segment_packet_ptr(completion, 0, 0)); // One fixed-dimension compute per message on its real (zero-copy) data. - // Enqueue only; the GPU work overlaps with subsequent receives and is - // drained per batch (below). - gpu_workload.run(pipeline.finish_batch()); - // Hold the completion (and thus its recv buffer) until the batch drains, - // so the pass-through compute reads valid data without a per-message sync. - held_recv.push_back(completion); - if (held_recv.size() >= recv_hold_batch) { - flush_held_recv(); + // Enqueue only (never blocks); the GPU work overlaps with subsequent + // receives and is reclaimed once its event completes. + gpu_workload.run_async(pipeline.finish_batch()); + void* ev = gpu_workload.record_event(); + if (ev == nullptr) { + // Event pool momentarily exhausted (rare: pool >> max_inflight). Drain + // and free now so we never read/repost a buffer the GPU is still using. + gpu_workload.sync(); + daqiri::free_tx_burst(completion); + } else { + // Hold the completion (its recv buffer) until the event says the GEMM + // has finished reading it. + held_recv.push_back({completion, ev}); + } + // Backpressure: bound outstanding GPU work so the CPU cannot race ahead of + // the timer. Only blocks when at the cap, not on every message. + if (held_recv.size() >= max_inflight) { + gpu_workload.wait_event(held_recv.front().event); } - continue; // do not free yet; freed by flush_held_recv() + reclaim_completed(); + continue; // do not free yet; freed by reclaim_completed()/drain_held_recv() } } daqiri::free_tx_burst(completion); @@ -281,13 +310,13 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa } if (!got_completion && !posted_work) { - // Idle: return any held recv buffers so they are not pinned while traffic - // pauses, then back off. - flush_held_recv(); + // Idle: traffic has paused, so drain the held recv buffers (blocking is fine + // here -- there are no receives to starve) and back off. + drain_held_recv(); std::this_thread::sleep_for(std::chrono::microseconds(100)); } } - flush_held_recv(); + drain_held_recv(); gpu_workload.sync(); } From 5683fbe10bff73d18fabf7bc29c7cefff95728ec Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Thu, 9 Jul 2026 17:25:08 -0400 Subject: [PATCH 06/18] #15 - Fix RoCE workload receive-path GPU blocking; add one-way + profiling knobs The RoCE GPU-workload receive thread parked on the GPU every message -- a blocking drain of held recv buffers on the idle path, plus an eager cap wait -- which starved the GPU and throttled the sender. Reclaim held recv buffers non-blocking instead: drop the idle-path drain, reclaim before the last-resort cap block, and only block when the event pool would otherwise starve. nsys confirmed the per-message cudaEventSynchronize -- not the GPU, memory, or transport -- was the cap. Add the bench knobs used to root-cause and re-measure this: - run_spark_bench.sh: UNIDIR (one-way RoCE, matching the DPDK and socket benches which are already one-way), NSYS/NSYS_BIN (wrap the rdma server in nsys), RUN_SECONDS override, MAX_INFLIGHT, and RDMA_BUDGET_GIB/RX_NB/TX_NB for the buffer-depth experiment. - --workload-max-inflight bench flag and a larger CUDA event pool. The residual RoCE Signed-off-by: rgurunathan --- examples/bench_workload.cu | 20 +++++++++- examples/bench_workload.h | 15 ++++++++ examples/rdma_bench.cpp | 54 +++++++++++++++++++------- examples/run_spark_bench.sh | 76 ++++++++++++++++++++++++++++++++++--- 4 files changed, 143 insertions(+), 22 deletions(-) diff --git a/examples/bench_workload.cu b/examples/bench_workload.cu index bd4b100..2f055d1 100644 --- a/examples/bench_workload.cu +++ b/examples/bench_workload.cu @@ -38,8 +38,12 @@ constexpr int kFftLen = 1024; // Default working-set size when the caller passes bytes_per_burst == 0. constexpr size_t kDefaultBytes = 1u << 16; // 64 KiB // Fixed CUDA-event pool for record_event(); comfortably above the RoCE recv-path -// in-flight cap (min(rx_depth/2, 8)). Created lazily on first record_event(). -constexpr int kEventPoolSize = 32; +// in-flight cap (default min(rx_depth/2, 32), overridable up to kMaxWorkloadInflight +// via --workload-max-inflight). Created lazily on first record_event(). +constexpr int kEventPoolSize = 64; +static_assert(kEventPoolSize > daqiri::bench::kMaxWorkloadInflight, + "event pool must stay above the max in-flight cap so record_event() " + "never starves exactly at the cap"); cufftHandle as_fft_plan(int p) { return static_cast(p); @@ -103,6 +107,18 @@ int parse_workload_sync_interval(int argc, char** argv) { return 2; // default sync depth } +int parse_workload_max_inflight(int argc, char** argv) { + for (int i = 2; i + 1 < argc; i += 2) { + if (std::string(argv[i]) == "--workload-max-inflight") { + const long long v = std::atoll(argv[i + 1]); + if (v > 0) { + return static_cast(v); + } + } + } + return 0; // unset -> caller applies its computed default +} + const char* workload_name(BenchWorkload workload) { switch (workload) { case BenchWorkload::Fft: diff --git a/examples/bench_workload.h b/examples/bench_workload.h index 36d8f60..d4f8ed0 100644 --- a/examples/bench_workload.h +++ b/examples/bench_workload.h @@ -48,6 +48,21 @@ int parse_workload_gemm_dim(int argc, char** argv); // default) if unset. Mirrors parse_workload's stride. int parse_workload_sync_interval(int argc, char** argv); +// Upper bound for the RoCE recv-path in-flight cap (--workload-max-inflight, +// below). Kept one below the internal CUDA-event pool (kEventPoolSize in +// bench_workload.cu) so record_event() never starves exactly at the cap and +// falls back to a per-message sync. +constexpr int kMaxWorkloadInflight = 63; + +// Parse "--workload-max-inflight N" from argv: for the event-recycling RoCE recv +// path, the max number of recv buffers that may have in-flight GPU work before +// the receive thread blocks on the oldest event (backpressure). Larger N absorbs +// more receive/compute jitter before stalling reposts, bounded by the event pool +// (kMaxWorkloadInflight) and the posted-receive window (rx_depth). Returns 0 when +// unset, so the caller applies its computed default (min(rx_depth/2, 32)). Mirrors +// parse_workload's stride. Ignored by the run()+sync() DPDK/socket path. +int parse_workload_max_inflight(int argc, char** argv); + // Lower-case name ("none"/"fft"/"gemm"); used for the run_spark_bench.sh // post_process CSV column and log lines. const char* workload_name(BenchWorkload workload); diff --git a/examples/rdma_bench.cpp b/examples/rdma_bench.cpp index e5a8dee..fe4a19b 100644 --- a/examples/rdma_bench.cpp +++ b/examples/rdma_bench.cpp @@ -83,7 +83,7 @@ RdmaBenchConfig parse_rdma_cfg(const YAML::Node& node) { void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pacer, std::atomic& stop, RdmaWorkerStats& stats, daqiri::bench::BenchWorkload workload, int workload_gemm_dim, - int workload_sync_interval) { + int workload_sync_interval, int workload_max_inflight) { const char *thread_name = cfg.server ? "rdma_bench_server" : "rdma_bench_client"; if (!daqiri::bench::set_current_thread_affinity(cfg.cpu_core, thread_name)) { @@ -129,8 +129,17 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa void* event; }; std::deque held_recv; - const size_t max_inflight = - run_workload ? std::max(1, std::min(cfg.rx_depth / 2, 8)) : 0; + // Default cap min(rx_depth/2, 32); --workload-max-inflight overrides it. Clamp to + // the event pool (kMaxWorkloadInflight -- else record_event() starves at the cap + // and falls back to a per-message sync) and to rx_depth (else we would block + // reposting more receives than are posted). + size_t inflight_cap = workload_max_inflight > 0 + ? static_cast(workload_max_inflight) + : std::min(cfg.rx_depth / 2, 32); + inflight_cap = + std::min(inflight_cap, static_cast(daqiri::bench::kMaxWorkloadInflight)); + inflight_cap = std::min(inflight_cap, static_cast(cfg.rx_depth)); + const size_t max_inflight = run_workload ? std::max(1, inflight_cap) : 0; // Free every held buffer whose GEMM has completed (non-blocking); keeps the recv // pool drained without parking the thread on the GPU. auto reclaim_completed = [&]() { @@ -278,8 +287,9 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa gpu_workload.run_async(pipeline.finish_batch()); void* ev = gpu_workload.record_event(); if (ev == nullptr) { - // Event pool momentarily exhausted (rare: pool >> max_inflight). Drain - // and free now so we never read/repost a buffer the GPU is still using. + // Event pool momentarily exhausted (defensive: max_inflight is clamped + // below the pool, so steady state never starves). Drain and free now so + // we never read/repost a buffer the GPU is still using. gpu_workload.sync(); daqiri::free_tx_burst(completion); } else { @@ -287,12 +297,21 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa // has finished reading it. held_recv.push_back({completion, ev}); } - // Backpressure: bound outstanding GPU work so the CPU cannot race ahead of - // the timer. Only blocks when at the cap, not on every message. + // Reclaim completed GEMMs FIRST (non-blocking) so finished work frees its + // held slot before we ever consider blocking. In steady state the GPU keeps + // up (GEMM ~337us << message interval), so this drops held below the cap + // every iteration and the block below never fires -- the thread keeps + // launching GEMMs and reposting receives, so the GPU stays fed and the wire + // stays full. (Parking here instead is what starved the GPU and throttled + // the sender: it stalled launches/reposts every message.) + reclaim_completed(); + // Last-resort backpressure: only park if, even after reclaiming, held is + // still at the cap -- i.e. the GPU is genuinely behind (rare; bounds CPU + // run-ahead for measurement honesty), not the per-message stall it was. if (held_recv.size() >= max_inflight) { gpu_workload.wait_event(held_recv.front().event); + reclaim_completed(); // free the one we just drained } - reclaim_completed(); continue; // do not free yet; freed by reclaim_completed()/drain_held_recv() } } @@ -310,10 +329,16 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa } if (!got_completion && !posted_work) { - // Idle: traffic has paused, so drain the held recv buffers (blocking is fine - // here -- there are no receives to starve) and back off. - drain_held_recv(); - std::this_thread::sleep_for(std::chrono::microseconds(100)); + // Out of completions for the moment. Do NOT touch the GPU here: neither block + // (drain_held_recv parks the thread on each in-flight GEMM) NOR hot-spin + // reclaim_completed() -- polling cudaEventQuery millions of times stole ~40% of + // the thread from wire servicing, throttling the sender just as badly. The main + // receive path already reclaims finished GEMMs on every message, so here we just + // keep polling get_rx_burst (cheap, no CUDA) for the next completion, and back + // off only when there is genuinely nothing in flight. + if (held_recv.empty()) { + std::this_thread::sleep_for(std::chrono::microseconds(100)); + } } } drain_held_recv(); @@ -346,6 +371,7 @@ int main(int argc, char** argv) { const auto workload = daqiri::bench::parse_workload(argc, argv); const int workload_gemm_dim = daqiri::bench::parse_workload_gemm_dim(argc, argv); const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); + const int workload_max_inflight = daqiri::bench::parse_workload_max_inflight(argc, argv); const auto root = YAML::LoadFile(argv[1]); if (daqiri::daqiri_init(argv[1]) != daqiri::Status::SUCCESS) { @@ -383,12 +409,12 @@ int main(int argc, char** argv) { if (run_server) { server_thread = std::thread(rdma_worker, server_cfg, std::ref(server_pacer), std::ref(stop), std::ref(server_stats), workload, workload_gemm_dim, - workload_sync_interval); + workload_sync_interval, workload_max_inflight); } if (run_client) { client_thread = std::thread(rdma_worker, client_cfg, std::ref(client_pacer), std::ref(stop), std::ref(client_stats), workload, workload_gemm_dim, - workload_sync_interval); + workload_sync_interval, workload_max_inflight); } if (!server_thread.joinable() && !client_thread.joinable()) { diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index ad56c66..0041243 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -83,7 +83,10 @@ echo "lang,backend,post_process,payload,batch,pairs,target_gbps,rep,seconds,pack # Capture slow-moving environment state once per result set. "$SCRIPT_DIR/bench_capture_environment.sh" "$OUT_DIR" -RUN_SECONDS=30 +RUN_SECONDS="${RUN_SECONDS:-30}" +if [[ ! "$RUN_SECONDS" =~ ^[0-9]+$ || "$RUN_SECONDS" -lt 1 ]]; then + echo "Invalid RUN_SECONDS '$RUN_SECONDS' (expected a positive integer)" >&2; exit 1 +fi # Repeats per cell for error bars. Each rep is a full independent run with its own # capture dir (-r) and CSV row; the perf-doc tables report mean +/- std # across reps. Default 1; set REPEATS=3 for the published re-run. @@ -116,6 +119,39 @@ SYNC_INTERVAL="${SYNC_INTERVAL:-}" if [[ -n "$SYNC_INTERVAL" && ! "$SYNC_INTERVAL" =~ ^[0-9]+$ ]]; then echo "Invalid SYNC_INTERVAL '$SYNC_INTERVAL' (expected a positive integer)" >&2; exit 1 fi +# MAX_INFLIGHT: for the RoCE event-recycling recv path, the max recv buffers with +# in-flight GPU work before the receive thread blocks on the oldest event +# (--workload-max-inflight). Larger absorbs more jitter before stalling reposts; +# clamped in the bench to the event pool (63) and rx_depth. Empty = bench default +# (min(rx_depth/2, 32)). Ignored by the run()+sync() DPDK/socket path. +MAX_INFLIGHT="${MAX_INFLIGHT:-}" +if [[ -n "$MAX_INFLIGHT" && ! "$MAX_INFLIGHT" =~ ^[0-9]+$ ]]; then + echo "Invalid MAX_INFLIGHT '$MAX_INFLIGHT' (expected a positive integer)" >&2; exit 1 +fi +# NSYS: when set (NSYS=1), wrap the RoCE *server* (the receive + GPU-workload +# process) in `nsys profile` to capture the CUDA/GPU timeline and thread states, +# so we can see whether the GPU stream idles between GEMMs (receive-thread cadence +# gaps) or the SM is genuinely busy. Only the server is traced -- the client is a +# pure traffic generator. Per-cell report lands as /roce_server.nsys-rep. +# Use `rdma smoke` + a short RUN_SECONDS so the report stays small. +NSYS="${NSYS:-}" +NSYS_BIN="${NSYS_BIN:-}" +if [[ -n "$NSYS" ]]; then + if [[ -z "$NSYS_BIN" ]]; then + # Resolve nsys: PATH first, then the usual CUDA toolkit location. + if command -v nsys >/dev/null 2>&1; then + NSYS_BIN="$(command -v nsys)" + elif [[ -x /usr/local/cuda/bin/nsys ]]; then + NSYS_BIN="/usr/local/cuda/bin/nsys" + fi + fi + if [[ -z "$NSYS_BIN" || ! -x "$NSYS_BIN" ]]; then + echo "NSYS set but no nsys binary found (looked on PATH and /usr/local/cuda/bin)." >&2 + echo "Set NSYS_BIN=/path/to/nsys explicitly." >&2 + exit 1 + fi + echo "nsys profiling enabled: $NSYS_BIN" >&2 +fi DRIVER_LOG="$OUT_DIR/last_run.stderr" FAILURES=0 @@ -330,10 +366,14 @@ generate_yaml() { # (rx 512 / tx 128) but are capped so num_bufs * buf_size stays within a memory # budget per region -- small messages get the full deep window (where RNR NACKs # bite), large messages stay memory-bounded (where window depth is irrelevant). - local budget=1073741824 # 1 GiB pinned per memory region + # Overridable for the buffer-depth experiment: RDMA_BUDGET_GIB raises the + # per-region pinned-memory cap, RDMA_RX_NB / RDMA_TX_NB raise the window + # ceilings (num_bufs + {rx,tx}_depth). Defaults reproduce the shipped sizing + # (1 GiB, rx 512 / tx 128). + local budget=$(( ${RDMA_BUDGET_GIB:-1} * 1073741824 )) # pinned per memory region local cap=$(( budget / payload )); (( cap < 1 )) && cap=1 - local rx_nb=512; (( rx_nb > cap )) && rx_nb=$cap - local tx_nb=128; (( tx_nb > cap )) && tx_nb=$cap + local rx_nb=${RDMA_RX_NB:-512}; (( rx_nb > cap )) && rx_nb=$cap + local tx_nb=${RDMA_TX_NB:-128}; (( tx_nb > cap )) && tx_nb=$cap # Split the combined base per role, then apply the per-message-size window # rewrite to each. Server -> $out, client -> ${out%.yaml}_client.yaml. local role dst @@ -342,8 +382,13 @@ generate_yaml() { # name-anchored num_bufs rewrite: RX regions -> rx_nb, TX regions -> tx_nb. # Depths are clamped to their region's num_bufs so the window never exceeds # the buffers backing it. + # UNIDIR=1 makes the flow one-way to match the DPDK bench (which only + # receives+computes on the RX port): server receive-only (runs the GEMM), + # client send-only (no GEMM). Removes the reverse-direction DMA and the + # second GEMM, so RoCE and raw do the same GPU + memory work. python3 "$NETNS_GEN" "$BASE_YAML" --role "$role" | \ - awk -v p="$payload" -v bs="$payload" -v rxnb="$rx_nb" -v txnb="$tx_nb" ' + awk -v p="$payload" -v bs="$payload" -v rxnb="$rx_nb" -v txnb="$tx_nb" \ + -v uni="${UNIDIR:-}" -v role="$role" ' /^[[:space:]]*- name:/ { region = $0 } /^[[:space:]]*num_bufs:/ { if (region ~ /RX/) { sub(/num_bufs:.*/, "num_bufs: " rxnb) } @@ -354,6 +399,8 @@ generate_yaml() { /^[[:space:]]*message_size:/ { sub(/message_size:.*/, "message_size: " p); print; next } /^[[:space:]]*rx_depth:/ { sub(/rx_depth:.*/, "rx_depth: " rxnb); print; next } /^[[:space:]]*tx_depth:/ { sub(/tx_depth:.*/, "tx_depth: " txnb); print; next } + /^[[:space:]]*send:/ { if (uni != "" && role == "server") { sub(/send:.*/, "send: false") } print; next } + /^[[:space:]]*receive:/ { if (uni != "" && role == "client") { sub(/receive:.*/, "receive: false") } print; next } { print } ' > "$dst" done @@ -431,6 +478,7 @@ run_cell() { bench_extra+=(--workload "$WORKLOAD_EFF") bench_extra+=(--workload-gemm-dim "$GEMM_DIM") [[ -n "$SYNC_INTERVAL" ]] && bench_extra+=(--workload-sync-interval "$SYNC_INTERVAL") + [[ -n "$MAX_INFLIGHT" ]] && bench_extra+=(--workload-max-inflight "$MAX_INFLIGHT") fi # Shutdown ordering: the server must keep receiving until the client has fully # stopped sending, otherwise the client's last in-flight messages have no peer @@ -487,7 +535,14 @@ run_cell() { # addresses over the wire rather than short-cutting the kernel's local table. local yaml="$cell_dir/config.yaml" generate_yaml "$yaml" "$payload" "$batch" - ip netns exec dq_wire_server "$BENCH_BIN" "$yaml" \ + # Optionally wrap the server (receive + GPU-workload) in nsys. Placed AFTER the + # netns exec so nsys traces the bench, not `ip netns exec`. + local -a nsys_pre=() + if [[ -n "$NSYS" ]]; then + nsys_pre=("$NSYS_BIN" profile --trace=cuda,osrt --sample=cpu --cpuctxsw=process-tree + --force-overwrite=true --output "$cell_dir/roce_server") + fi + ip netns exec dq_wire_server "${nsys_pre[@]}" "$BENCH_BIN" "$yaml" \ --seconds "$server_seconds" "${bench_extra[@]}" --mode server \ > "$cell_dir/server_stdout.txt" 2> "$cell_dir/server_stderr.txt" & local server_pid=$! @@ -498,6 +553,15 @@ run_cell() { wait "$server_pid" 2>/dev/null || true cat "$cell_dir/server_stdout.txt" >> "$stdout" cat "$cell_dir/server_stderr.txt" >> "$stderr" + if [[ -n "$NSYS" && -f "$cell_dir/roce_server.nsys-rep" ]]; then + echo "=== nsys GPU kernel summary ($cell) ===" >&2 + "$NSYS_BIN" stats --report cuda_gpu_kern_sum --format table \ + "$cell_dir/roce_server.nsys-rep" >&2 || true + echo "=== nsys CUDA API summary ($cell) ===" >&2 + "$NSYS_BIN" stats --report cuda_api_sum --format table \ + "$cell_dir/roce_server.nsys-rep" >&2 || true + echo "nsys report: $cell_dir/roce_server.nsys-rep" >&2 + fi # RDMA prints "Client complete: ... send_completions=N send_bytes=N seconds=S". pkts="$(extract_field 'Client complete' send_completions "$stdout")" bytes="$(extract_field 'Client complete' send_bytes "$stdout")" From b394196f74a585eb4247a7fd580928b013fd45c7 Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Thu, 9 Jul 2026 17:25:21 -0400 Subject: [PATCH 07/18] #15 - Rewrite DGX Spark GPU-workload section with one-way, same-session results Replace the stale, pre-fix bidirectional numbers with a single same-session table comparing DPDK and one-way RoCE (none / FFT / GEMM FP32); both hold ~line rate under the GPU workload. Delete the obsolete "single-threaded RoCE stall" narrative and the GEMMs/s + duty-cycle-SM analysis -- the gap was the bidirectional benchmark, now removed. Drop the FP16 GEMM result, clarify both workloads are FP32, and remove the misleading "native shape" wording (8 KB is a chosen operating point, not the throughput peak). Fix the Reproduce commands to set UNIDIR=1 for the RoCE runs. Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- docs/benchmarks/performance-dgx-spark.md | 156 +++++++---------------- 1 file changed, 47 insertions(+), 109 deletions(-) diff --git a/docs/benchmarks/performance-dgx-spark.md b/docs/benchmarks/performance-dgx-spark.md index 9a1320b..af67e08 100644 --- a/docs/benchmarks/performance-dgx-spark.md +++ b/docs/benchmarks/performance-dgx-spark.md @@ -27,27 +27,27 @@ loopback). The exact commands are collected under [Reproduce](#reproduce) below. ## Results Summary (C++ loopback) -Each transport at its best-case **native operation size**. Raw/RoCE are +Each transport at its best-case **operation size**. Raw/RoCE are single-stream; socket TCP/UDP scale with the number of client/server pairs, so the four-pair aggregate is shown. | Stream / Protocol | Best case | Throughput | Drops | Notes | | ----------------- | --------- | ---------: | ----- | ----- | -| Raw Ethernet / GPUDirect | 4 KB packet | **105.5 ±0.9 Gb/s** | 0 | 98.5 Gb/s single-queue at the 8 KB native shape | +| Raw Ethernet / GPUDirect | 4 KB packet | **105.5 ±0.9 Gb/s** | 0 | 98.5 Gb/s single-queue at 8 KB | | Socket / RoCE (SEND) | 8 MB message | **102.2 ±0.3 Gb/s** | 0 | Single QP, batch 1 | | Socket / TCP | 8 KB × 4 pairs | **97.2 ±2.8 Gb/s** | ~0 | Flow-controlled (App TX = App RX) | | Socket / UDP | 8 KB × 4 pairs | **29.8 ±0.2 Gb/s** | ~51% loss | Receiver goodput; unpaced sender | -Each transport is best read at its own native operation size (see the per-transport +Each transport is best read at its own best-case operation size (see the per-transport tables below); a single cross-transport unit of work isn't meaningful here, since RoCE at 8 KB is op-rate-bound well below its large-message peak and TCP has no operation boundary. ## Raw Ethernet / GPUDirect (DPDK) -Physical port-to-port loopback, GPU-resident payloads. Native 8 KB packets run -at **98.5 Gb/s** drop-free across all batch sizes; the throughput peak is -**105.5 Gb/s** at a 4 KB payload. Packet handling is CPU-bound (see the CPU +Physical port-to-port loopback, GPU-resident payloads. Throughput peaks at +**105.5 Gb/s** at a 4 KB payload; 8 KB packets hold **98.5 Gb/s** drop-free +across all batch sizes. Packet handling is CPU-bound (see the CPU utilization table below). Throughput is flat across batch size and stable run-to-run (3 reps per cell, ≤1% spread). @@ -98,7 +98,7 @@ not a compute engine. ### Multi-queue core scaling -Each packet-handling core spins in poll-mode. At the native 8 KB shape a single +Each packet-handling core spins in poll-mode. At 8 KB a single TX core caps throughput near ~98 Gb/s while a single RX core already drains the line, so the second TX core is the lever: it scales `(1,1)` to `(2,1)` from 97.5 to 108.8 Gb/s, while a second RX core adds little. The matrix sweeps @@ -129,7 +129,7 @@ At small payloads the path is packet-rate-bound, so **RX cores** are the lever: at 64 B a second RX core lifts throughput from 20.3 to 26.8 Gb/s (~20 M → ~26 M pps) while a second TX core does nothing. At large payloads it inverts to the byte/line-rate-bound regime where **TX cores** are the lever: at 8 KB the second -TX core takes 97.5 to 108.8 Gb/s (the native-shape result above) while a second +TX core takes 97.5 to 108.8 Gb/s (the 8 KB result above) while a second RX core adds nothing. The curves cross around 1–4 KB, where the link saturates and all four cells converge near ~104–109 Gb/s. Every cell is drop-free. Generated by `examples/run_spark_mq_bench.sh` (30 s per point) and @@ -242,13 +242,13 @@ the GPU also crunches the incoming data. The benchmarks accept `--workload none|fft|gemm|gemm_fp16`, exposed by `run_spark_bench.sh` as the `WORKLOAD` env var (recorded in the CSV `post_process` column); more workload kinds can be added to the same reusable component over time. The workload runs on the -**actual received packet data** — every backend first assembles the burst's +received packet data — every backend first assembles the burst's payloads into one contiguous GPU buffer (a sequence-number **reorder** on the out-of-order transports, an arrival-order **gather** on the in-order ones) and the compute consumes that buffer. -**What the two workloads compute** (the reusable component -`examples/bench_workload.{h,cu}`): +**What the two workloads compute** — both in **FP32** (single precision) — from the +reusable component `examples/bench_workload.{h,cu}`: - **FFT** — a batched 1-D **complex-to-complex forward FFT** via cuFFT (`cufftExecC2C`). The reordered buffer is treated as an array of single-precision @@ -257,12 +257,11 @@ compute consumes that buffer. signal-processing receiver — channelization or spectral analysis that FFTs every frame as it arrives. - **GEMM** — a dense **matrix multiply** `C = A·B` via cuBLAS on square *n×n* - matrices, with the reordered buffer supplying the *A* operand and *n* chosen so - the matrices match the working-set size. Two precisions: **`gemm`** is FP32 - (`cublasSgemm`); **`gemm_fp16`** is the *same-size* mixed-precision matmul (FP16 - inputs, FP32 accumulate) on the **tensor cores** (`cublasGemmEx`, - `CUBLAS_GEMM_DEFAULT_TENSOR_OP`) — the core op of GPU inference. Running both at - one matrix size isolates the precision / tensor-core effect. This models a + matrices, with the reordered buffer supplying the *A* operand. The side length is + **pinned at n=1024** (`--workload-gemm-dim`, env `GEMM_DIM`), so every call is an + identical **2.15 GFLOP** matmul reading the first **4 MB** (n²·4 B, FP32) of each + received unit — the compute is fixed regardless of message size, which is what + makes it comparable across transports. The matmul is FP32 (`cublasSgemm`). This models a receiver feeding incoming data into a dense linear-algebra or neural-network stage (beamforming, correlation, an inference layer). @@ -276,12 +275,10 @@ chosen to be representative for each transport: | UDP sockets | host RX buffers | **host→device stage**, then **seq reorder** | | TCP sockets | host RX buffers | **host→device stage**, then **gather** (in-order stream) | -Each compute runs **once per reorder window** on a dedicated CUDA stream (shared -with the reorder/gather kernel so the two serialize without an extra sync), with up -to two kernels in flight (sync depth 2) to overlap GPU work with ingest. The -reorder window is sized so the contiguous buffer is ~8 MB on every backend, giving -a comparable GPU working set across transports. GPU SM% is from `nvidia-smi dmon` -across the run. Throughput is mean ± std over 3 reps, 30 s each. +Each compute runs **once per reorder window** on a dedicated CUDA stream, shared +with the reorder/gather kernel so the two serialize without an extra sync and +compute overlaps ingest. The reorder window is sized so the contiguous buffer is +~8 MB on every backend, giving a comparable GPU working set across transports. !!! note "Where the data lives, per backend" On the integrated GB10 the GPU shares memory with the CPU, so the raw and RoCE @@ -292,85 +289,26 @@ across the run. Throughput is mean ± std over 3 reps, 30 s each. copy on the measured path that the raw/RoCE paths avoid. Lost packets (raw/UDP) leave their reorder slots zero-filled; the FLOP/copy volume is unchanged. -Raw / GPUDirect, 8 KB native shape (batch 10240), GPU-resident payloads, seq reorder: - -| Workload | Throughput | Drops | GPU SM% | Notes | -| -------- | ---------: | ----- | ------: | ----- | -| none (baseline) | 98.4 ±0.2 Gb/s | 0 | ~0 | Bare loopback (no GPU compute) | -| FFT | 95.5 ±1.7 Gb/s | 0 | 16.4% | Light compute; line rate held within ~3% | -| GEMM (FP32) | 94.9 ±0.0 Gb/s | 0 | 55.9% | FP32 cores; GPU-bound end, still **drop-free** | -| GEMM (FP16 tensor) | 97.7 ±0.2 Gb/s | 0 | 16.8% | Same matrix, tensor cores; near line rate at a third of the SM | - -RoCE, 8 MB native message (single QP, batch 1), gather pass-through: - -| Workload | Throughput | Drops | GPU SM% | Notes | -| -------- | ---------: | ----- | ------: | ----- | -| none (baseline) | 101.7 ±0.2 Gb/s | 0 | ~0 | Pass-through, no compute | -| FFT | 93.6 ±1.6 Gb/s | 0 | 36.0% | Light compute; ~8% off line rate | -| GEMM (FP32) | 35.0 ±0.3 Gb/s | 0 | 79.0% | GPU-bound at this batch (one GEMM/message); see sweep | -| GEMM (FP16 tensor) | 85.8 ±0.2 Gb/s | 0 | 53.6% | Same matrix, tensor cores | - -Both paths drive the GPU the same way (each holds a batch of received data and drains -the GPU stream **once per batch** — the raw path a burst of ~10 reorder windows, the -RoCE path a small batch of messages whose recv buffers stay live during the pass-through -compute), so compute overlaps with receive on both and every cell is drop-free. The -fixed-batch rows above are the 8 MB native operating point, batch-matched to the raw -path's ~8.19 MB reorder window (1024 packets × 8000 B). - -### Fixed-size GEMM comparison - -To compare transports cleanly we **pin the GEMM dimension** with `--workload-gemm-dim 1024` -(env `GEMM_DIM` in `run_spark_bench.sh`): every call is an identical **1024×1024×1024 -matmul = 2.15 GFLOP** reading the first **4 MB** (n²·4 B, FP32) of each received I/O unit, -so throughput *and* GEMMs/s are directly comparable across transports. One fixed-n GEMM -per received unit — RoCE per 8 MB message, raw per reorder window. Fixed n=1024, 3 reps, -30 s each. - -!!! note "Numbers pending refresh" - The table below was measured under the earlier per-message chunking (2 GEMMs per 8 MB - message) and the blocking per-batch `cudaStreamSynchronize`. The receive path now runs - **one GEMM per message** and recycles recv buffers with non-blocking CUDA events; these - cells will be re-measured under that model. - -| Transport | Workload | Throughput | GEMMs/s | GPU SM% | -| --------- | -------- | ---------: | ------: | ------: | -| RoCE (RC) | GEMM FP32 | 48.1 Gb/s | 1434 | 78% | -| Raw (DPDK) | GEMM FP32 | 103.8 Gb/s | 3095 | 40% | -| RoCE (RC) | GEMM FP16 (tensor) | 88.3 Gb/s | 2631 | 56% | -| Raw (DPDK) | GEMM FP16 (tensor) | 105.6 Gb/s | 3148 | 15% | - -Three things fall out, and they revise the "pipelining depth" reading of the earlier -batch sweep: - -- **The RoCE↔raw gap is real, and it is not the GPU.** At the *identical* fixed workload - raw sustains **2.2× the GEMM rate at FP32** (and ~1.2× at FP16) while using *less* than - half the SM. Raw has GPU headroom (15–40% SM) and is wire-limited (~104–106 Gb/s for - both precisions); RoCE is the side leaving the GPU idle-but-stalled. -- **GEMMs/s is the transport-fair metric, not Gb/s** — because an engine that consumes more - bytes per GEMM carries more throughput at the same compute rate. (Matching the window to - 4 MB on both sides is what makes the two columns comparable.) Achieved rates are - ~3.1 TFLOP/s (FP32) / ~5.7 TFLOP/s (FP16), only ~10% of peak, so the small matrix is - **launch/sync-latency-bound**, not FLOP-bound. -- **GPU SM% is a duty-cycle metric, anti-correlated with throughput here.** RoCE is - "stall-busy" (78% SM at 48 Gb/s); raw is efficient (40% SM at 104 Gb/s). High SM% means - the GPU work is *stretched thin* against a slower receive path, not that more useful work - is done — so SM% should not be read as compute efficiency. - -The bottleneck is the **single-threaded RoCE receive+compute loop**: one thread runs -`poll completion → gather → issue GEMM → stream-sync → free → repeat`, and the periodic -`cudaStreamSynchronize` (required before a received buffer can be safely reused) blocks -that thread — so while the GPU drains, no receives are posted and the message rate falls. -The heavier the GEMM, the longer each stall, which is why FP32 (48 Gb/s) suffers more than -FP16 (88 Gb/s) while raw stays wire-limited for both. Raw proves one thread *can* feed the -GPU at ~3100 GEMM/s, so the ceiling is the RoCE receive path, not the GPU. - -Tuning the sync cadence does **not** recover it: sweeping `--workload-sync-interval` from 2 -to 32 leaves throughput flat (FP32 ~48 Gb/s / ~1450 GEMMs/s, FP16 ~88 Gb/s / ~2650 GEMMs/s), -and fully-synchronous (interval 1) matches deep-async — the path already drains the stream -once per held-receive batch, which masks the knob, and the real limit is the single-threaded -single-QP receive itself. The remedy is therefore **structural**: decouple receive from -compute onto a dedicated GPU-worker thread (and/or fan out across multiple QPs/threads), -not sync tuning. +Fixed **n=1024**, one GEMM (or a length-1024 batched FFT) per received unit. DPDK runs +at an **8 KB payload** (~8 MB reorder window, 1024 packets × 8000 B), matched to RoCE's +**8 MB message** so the GPU working set and per-unit compute are the same on both. RoCE +runs **unidirectional** (`UNIDIR=1`) — the server receives and computes, the client only +sends — matching DPDK's one-way TX→RX flow. 3 reps, 30 s each, GPU SM% from +`nvidia-smi dmon`; 0 drops on every cell. + +| Workload | DPDK (Raw / GPUDirect) | RoCE (RC) | +| -------- | ---------------------: | --------: | +| none (baseline) | 98.4 ±0.2 | 108.0 ±0.1 | +| FFT | 94.4 ±0.3 | 107.2 ±0.2 | +| GEMM (FP32) | 95.4 ±0.3 | 103.4 ±0.1 | + +Throughput in Gb/s. The two `none` baselines differ (98.4 vs 108.0) because each transport +runs at its own operation size — DPDK's 8 KB packets carry more per-packet overhead than +RoCE's large messages — not because of the workload. + +**GPU compute barely dents line rate on either transport:** both the GEMM and FFT leave +GPU headroom (SM well under 100%), so throughput stays wire-limited, not compute-limited. +The reorder/gather step assembles each unit's payload into one contiguous GPU buffer. ## Reproduce @@ -440,9 +378,9 @@ modes and netns setup as above (dpdk in the default namespace, rdma in the # Raw / GPUDirect (netns down, ETH_DST_ADDR exported as above) WORKLOAD=fft ./examples/run_spark_bench.sh dpdk smoke WORKLOAD=gemm ./examples/run_spark_bench.sh dpdk smoke -# Socket / RoCE (netns up) -WORKLOAD=fft ./examples/run_spark_bench.sh rdma smoke -WORKLOAD=gemm ./examples/run_spark_bench.sh rdma smoke +# Socket / RoCE (netns up; UNIDIR=1 for the one-way, apples-to-apples comparison) +UNIDIR=1 WORKLOAD=fft ./examples/run_spark_bench.sh rdma smoke +UNIDIR=1 WORKLOAD=gemm ./examples/run_spark_bench.sh rdma smoke ``` The chosen workload lands in the CSV `post_process` column; compare `gbps` / @@ -455,12 +393,12 @@ in the CSV `post_process_gemm_dim` column (the sweep is restricted to the headli message size automatically when a workload is active): ```bash -# RoCE (netns up) -for WL in gemm gemm_fp16; do - WORKLOAD=$WL GEMM_DIM=1024 REPEATS=3 ./examples/run_spark_bench.sh rdma sweep +# RoCE (netns up; UNIDIR=1 for the one-way comparison) +for WL in none fft gemm; do + UNIDIR=1 WORKLOAD=$WL GEMM_DIM=1024 REPEATS=3 ./examples/run_spark_bench.sh rdma sweep done # Raw (netns down, ETH_DST_ADDR exported) -for WL in gemm gemm_fp16; do +for WL in none fft gemm; do WORKLOAD=$WL GEMM_DIM=1024 REPEATS=3 ./examples/run_spark_bench.sh dpdk sweep done ``` From 8beaca75e3e37711a7935776cc167669498e4cbb Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Fri, 10 Jul 2026 09:05:31 -0400 Subject: [PATCH 08/18] #15 - Make the RoCE bench one-way by default, matching DPDK and sockets The RoCE netns bench was the only transport running bidirectionally (both roles send + receive), which understated its one-way rate and made it inconsistent with the DPDK bench (inherently TX->RX one-way) and the socket bench (explicitly one-way). Flip the combined base config so the server is receive-only (runs the GEMM) and the client is send-only, mirroring the socket netns base; gen_spark_netns_config.py splits each role's section verbatim, so the split configs inherit the one-way flags. Remove the now-redundant UNIDIR env knob and its awk send:/receive: rewrite from run_spark_bench.sh -- one-way is the default. Drop UNIDIR=1 from the perf-doc Reproduce commands and reword the methodology note accordingly. The RoCE result numbers are re-measured in a follow-up commit. Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- docs/benchmarks/performance-dgx-spark.md | 20 ++++++++++--------- .../daqiri_bench_rdma_tx_rx_spark_netns.yaml | 9 +++++++-- examples/run_spark_bench.sh | 13 ++++-------- 3 files changed, 22 insertions(+), 20 deletions(-) diff --git a/docs/benchmarks/performance-dgx-spark.md b/docs/benchmarks/performance-dgx-spark.md index af67e08..78c8255 100644 --- a/docs/benchmarks/performance-dgx-spark.md +++ b/docs/benchmarks/performance-dgx-spark.md @@ -9,6 +9,10 @@ Measured C++-loopback throughput for each stream/protocol on a single DGX Spark (GB10), driven over a physical cabled loopback on one ConnectX-7. Numbers are from a Release build via `examples/run_spark_bench.sh` (30 s per cell). +All backends are measured **one-way** (unidirectional) by default: one side sends, +the other receives and computes. For a bidirectional test, set `send: true` on the +receiving role (and `receive: true` on the sending role) in the bench config. + For the loopback setup these numbers depend on and the per-transport benchmarking procedure, see [Socket and RDMA Benchmarking](socket_benchmarking.md) (the `dq_wire_*` network-namespace wire loopback used by RoCE and sockets) and @@ -291,10 +295,8 @@ compute overlaps ingest. The reorder window is sized so the contiguous buffer is Fixed **n=1024**, one GEMM (or a length-1024 batched FFT) per received unit. DPDK runs at an **8 KB payload** (~8 MB reorder window, 1024 packets × 8000 B), matched to RoCE's -**8 MB message** so the GPU working set and per-unit compute are the same on both. RoCE -runs **unidirectional** (`UNIDIR=1`) — the server receives and computes, the client only -sends — matching DPDK's one-way TX→RX flow. 3 reps, 30 s each, GPU SM% from -`nvidia-smi dmon`; 0 drops on every cell. +**8 MB message** so the GPU working set and per-unit compute are the same on both. +3 reps, 30 s each, GPU SM% from `nvidia-smi dmon`; 0 drops on every cell. | Workload | DPDK (Raw / GPUDirect) | RoCE (RC) | | -------- | ---------------------: | --------: | @@ -378,9 +380,9 @@ modes and netns setup as above (dpdk in the default namespace, rdma in the # Raw / GPUDirect (netns down, ETH_DST_ADDR exported as above) WORKLOAD=fft ./examples/run_spark_bench.sh dpdk smoke WORKLOAD=gemm ./examples/run_spark_bench.sh dpdk smoke -# Socket / RoCE (netns up; UNIDIR=1 for the one-way, apples-to-apples comparison) -UNIDIR=1 WORKLOAD=fft ./examples/run_spark_bench.sh rdma smoke -UNIDIR=1 WORKLOAD=gemm ./examples/run_spark_bench.sh rdma smoke +# Socket / RoCE (netns up) +WORKLOAD=fft ./examples/run_spark_bench.sh rdma smoke +WORKLOAD=gemm ./examples/run_spark_bench.sh rdma smoke ``` The chosen workload lands in the CSV `post_process` column; compare `gbps` / @@ -393,9 +395,9 @@ in the CSV `post_process_gemm_dim` column (the sweep is restricted to the headli message size automatically when a workload is active): ```bash -# RoCE (netns up; UNIDIR=1 for the one-way comparison) +# RoCE (netns up) for WL in none fft gemm; do - UNIDIR=1 WORKLOAD=$WL GEMM_DIM=1024 REPEATS=3 ./examples/run_spark_bench.sh rdma sweep + WORKLOAD=$WL GEMM_DIM=1024 REPEATS=3 ./examples/run_spark_bench.sh rdma sweep done # Raw (netns down, ETH_DST_ADDR exported) for WL in none fft gemm; do diff --git a/examples/daqiri_bench_rdma_tx_rx_spark_netns.yaml b/examples/daqiri_bench_rdma_tx_rx_spark_netns.yaml index 510309b..0a11087 100644 --- a/examples/daqiri_bench_rdma_tx_rx_spark_netns.yaml +++ b/examples/daqiri_bench_rdma_tx_rx_spark_netns.yaml @@ -114,7 +114,10 @@ rdma_bench_server: # peer's posted receives and trigger RNR NACKs. Keep <= the MR num_bufs above. rx_depth: 128 tx_depth: 128 - send: true + # One-way (unidirectional) by default, matching the DPDK and socket benches: + # the server only receives, the client only sends. Set send: true here (and + # receive: true on the client) for a bidirectional test. + send: false receive: true server: true @@ -131,6 +134,8 @@ rdma_bench_client: client_address: 10.250.0.1 server_address: 10.250.0.2 server_port: 4096 - receive: true + # One-way (see server note above): the client only sends. Set receive: true for + # a bidirectional test. + receive: false send: true server: false diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index 0041243..a1603db 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -381,14 +381,11 @@ generate_yaml() { if [[ "$role" == server ]]; then dst="$out"; else dst="${out%.yaml}_client.yaml"; fi # name-anchored num_bufs rewrite: RX regions -> rx_nb, TX regions -> tx_nb. # Depths are clamped to their region's num_bufs so the window never exceeds - # the buffers backing it. - # UNIDIR=1 makes the flow one-way to match the DPDK bench (which only - # receives+computes on the RX port): server receive-only (runs the GEMM), - # client send-only (no GEMM). Removes the reverse-direction DMA and the - # second GEMM, so RoCE and raw do the same GPU + memory work. + # the buffers backing it. One-way (server receive-only + GEMM, client + # send-only) is baked into the base config, matching the DPDK and socket + # benches -- see the send:/receive: notes in the base YAML. python3 "$NETNS_GEN" "$BASE_YAML" --role "$role" | \ - awk -v p="$payload" -v bs="$payload" -v rxnb="$rx_nb" -v txnb="$tx_nb" \ - -v uni="${UNIDIR:-}" -v role="$role" ' + awk -v p="$payload" -v bs="$payload" -v rxnb="$rx_nb" -v txnb="$tx_nb" ' /^[[:space:]]*- name:/ { region = $0 } /^[[:space:]]*num_bufs:/ { if (region ~ /RX/) { sub(/num_bufs:.*/, "num_bufs: " rxnb) } @@ -399,8 +396,6 @@ generate_yaml() { /^[[:space:]]*message_size:/ { sub(/message_size:.*/, "message_size: " p); print; next } /^[[:space:]]*rx_depth:/ { sub(/rx_depth:.*/, "rx_depth: " rxnb); print; next } /^[[:space:]]*tx_depth:/ { sub(/tx_depth:.*/, "tx_depth: " txnb); print; next } - /^[[:space:]]*send:/ { if (uni != "" && role == "server") { sub(/send:.*/, "send: false") } print; next } - /^[[:space:]]*receive:/ { if (uni != "" && role == "client") { sub(/receive:.*/, "receive: false") } print; next } { print } ' > "$dst" done From da4c64016fbc6546cea5159a96f1d48a3983d13c Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Fri, 26 Jun 2026 15:49:36 -0400 Subject: [PATCH 09/18] #15 - Auto-detect cabled ports in Spark netns wire-loopback setup The setup script hardcoded CLIENT_IF/SERVER_IF (enp1s0f0np0 / enP2p1s0f1np1), which no longer exist after PCIe re-enumeration on the DGX Spark. With dead names, detect_macs aborts, the dq_wire_{client,server} namespaces are never created, and the netns-based socket/RoCE benches (launched via `ip netns exec`) silently have nowhere to run. Make the cabled ports auto-detected and robust to enumeration drift: - CLIENT_IF/SERVER_IF/CLIENT_RDMA/SERVER_RDMA default to blank, env-overridable. - autodetect_ports() picks the two carrier-up RoCE-backed netdevs in init_net and derives the RDMA devices from sysfs (rdma_for_netdev); errors with guidance if not exactly two are found. - down() now moves every non-lo netdev in each namespace back to init_net, so teardown no longer depends on a hardcoded interface name. - resolve_ports() lets verify()/monitor() read the in-namespace netdev names after up() has moved the ports, falling back to auto-detect otherwise. Co-Authored-By: Claude Opus 4.8 Signed-off-by: rgurunathan --- scripts/setup_spark_wire_loopback_netns.sh | 82 +++++++++++++++++++--- 1 file changed, 73 insertions(+), 9 deletions(-) diff --git a/scripts/setup_spark_wire_loopback_netns.sh b/scripts/setup_spark_wire_loopback_netns.sh index 67f9ac0..822f078 100755 --- a/scripts/setup_spark_wire_loopback_netns.sh +++ b/scripts/setup_spark_wire_loopback_netns.sh @@ -55,13 +55,22 @@ CLIENT_NS=dq_wire_client SERVER_NS=dq_wire_server -CLIENT_IF=enp1s0f0np0 # p0 via PCI segment 0000:01:00.0 (carrier up) -SERVER_IF=enP2p1s0f1np1 # p1 via PCI segment 0002:01:00.1 (carrier up) -- the - # cabled peer of CLIENT_IF (p0<->p1 QSFP loopback). NOT - # enP2p1s0f0np0, which is p0 again (same physical port). - -CLIENT_RDMA=rocep1s0f0 # rdma dev for enp1s0f0np0 -SERVER_RDMA=roceP2p1s0f1 # rdma dev for enP2p1s0f1np1 +# Leave CLIENT_IF / SERVER_IF blank to AUTO-DETECT the cabled data-plane ports +# (recommended): the two RoCE-backed netdevs that report carrier up -- exactly +# the pair the QSFP loopback lights. PCIe enumeration / interface names drift +# across reboots and kernel upgrades on this box, so hardcoding them is fragile +# (a stale name aborts setup, the netns never gets created, and the netns-based +# benches silently have nowhere to run). Override from the environment only if +# auto-detect picks the wrong pair, e.g.: +# CLIENT_IF=enp1s0f1np1 SERVER_IF=enP2p1s0f0np0 sudo -E ./setup_spark_wire_loopback_netns.sh up +CLIENT_IF="${CLIENT_IF:-}" +SERVER_IF="${SERVER_IF:-}" + +# RDMA devices backing the netdevs above. Blank = derive from the netdev via +# sysfs (/sys/class/net//device/infiniband). Only needed for RDMA exclusive +# netns mode; the socket/UDP/TCP benches do not use these. +CLIENT_RDMA="${CLIENT_RDMA:-}" +SERVER_RDMA="${SERVER_RDMA:-}" # Leave these blank to auto-detect from the interfaces (recommended). Only set # them if you need to override the peer MAC for some reason. @@ -83,6 +92,56 @@ require_root() { fi } +rdma_for_netdev() { # netdev -> backing RDMA device name (sysfs); empty if none + local ifc="$1" d + for d in /sys/class/net/"$ifc"/device/infiniband/*; do + [[ -e "$d" ]] || continue + basename "$d" + return 0 + done + return 1 +} + +autodetect_ports() { + # Pick the cabled data-plane ports automatically when not overridden: the + # RoCE-backed netdevs in init_net that report carrier up -- exactly the pair + # the QSFP loopback lights. Ports must be in init_net here (up() calls down() + # first, returning any moved netdevs), same precondition as detect_macs. + if [[ -z "$CLIENT_IF" || -z "$SERVER_IF" ]]; then + local found=() ifc + for ifc in $(ls /sys/class/net 2>/dev/null | sort); do + [[ "$ifc" == lo ]] && continue + rdma_for_netdev "$ifc" >/dev/null 2>&1 || continue + [[ "$(cat "/sys/class/net/$ifc/carrier" 2>/dev/null || echo 0)" == 1 ]] || continue + found+=("$ifc") + done + if [[ "${#found[@]}" -ne 2 ]]; then + echo "ERROR: auto-detect expected exactly 2 carrier-up RoCE ports, found ${#found[@]}: ${found[*]:-none}." >&2 + echo " Cable the loopback, or set CLIENT_IF / SERVER_IF explicitly, e.g.:" >&2 + echo " CLIENT_IF=enp1s0f1np1 SERVER_IF=enP2p1s0f0np0 sudo -E $0 up" >&2 + echo " List candidates with: ip -br link show | grep -E 'enp|enP'" >&2 + exit 1 + fi + [[ -z "$CLIENT_IF" ]] && CLIENT_IF="${found[0]}" + [[ -z "$SERVER_IF" ]] && SERVER_IF="${found[1]}" + fi + [[ -z "$CLIENT_RDMA" ]] && CLIENT_RDMA="$(rdma_for_netdev "$CLIENT_IF" || true)" + [[ -z "$SERVER_RDMA" ]] && SERVER_RDMA="$(rdma_for_netdev "$SERVER_IF" || true)" + echo "Using ports: CLIENT_IF=$CLIENT_IF (rdma=${CLIENT_RDMA:-none}) SERVER_IF=$SERVER_IF (rdma=${SERVER_RDMA:-none})" +} + +resolve_ports() { + # For verify()/monitor(), which run AFTER up() with the ports already moved + # into the namespaces: read each ns's non-lo netdev name. If the namespaces + # are absent, fall back to init_net auto-detect. Honours explicit overrides. + if ip netns list 2>/dev/null | grep -qw "$CLIENT_NS"; then + [[ -z "$CLIENT_IF" ]] && CLIENT_IF="$(ip netns exec "$CLIENT_NS" ls /sys/class/net 2>/dev/null | grep -vx lo | head -1)" + [[ -z "$SERVER_IF" ]] && SERVER_IF="$(ip netns exec "$SERVER_NS" ls /sys/class/net 2>/dev/null | grep -vx lo | head -1)" + else + autodetect_ports + fi +} + detect_macs() { # Interfaces must be in the current (init) namespace at this point -- up() # calls down() first, which returns any moved netdevs to init_net. @@ -108,8 +167,10 @@ down() { while read -r dev _; do [[ -n "$dev" ]] && ip netns exec "$NS" rdma dev set "$dev" netns 1 2>/dev/null || true done < <(ip netns exec "$NS" rdma dev show 2>/dev/null | awk -F': ' '/link/ {print $2}' | awk '{print $1}') - # Move netdevs back to init_net. - for IF in "$CLIENT_IF" "$SERVER_IF"; do + # Move every (non-lo) netdev back to init_net. Name-independent so + # teardown can't break on a stale/auto-detected interface name. + for IF in $(ip netns exec "$NS" ls /sys/class/net 2>/dev/null); do + [[ "$IF" == lo ]] && continue ip netns exec "$NS" ip link set "$IF" netns 1 2>/dev/null || true done ip netns delete "$NS" 2>/dev/null || true @@ -123,6 +184,7 @@ up() { echo "== Tearing down any previous state ==" down + autodetect_ports detect_macs echo "== Clearing shared-namespace IPs on the cabled ports ==" @@ -184,6 +246,7 @@ up() { } verify() { + resolve_ports echo "== RDMA device per namespace (each should show exactly its own dev) ==" echo "-- $CLIENT_NS --"; ip netns exec "$CLIENT_NS" rdma dev show || true echo "-- $SERVER_NS --"; ip netns exec "$SERVER_NS" rdma dev show || true @@ -230,6 +293,7 @@ monitor() { # enP2p1s0f0np0), not the rdma device (rocep1s0f0). # Auto-detects whether the ports live in the netns (this script's setup) or # the default namespace (the policy-routed setup_spark_rdma_loopback.sh setup). + resolve_ports local cns sns if ip netns list 2>/dev/null | grep -qw "$CLIENT_NS"; then cns="$CLIENT_NS"; sns="$SERVER_NS" From cd244e195cc66b543cc9df29b1b3549b4f5df38d Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Fri, 10 Jul 2026 09:45:21 -0400 Subject: [PATCH 10/18] #15 - Group Spark netns auto-detect by phys_port_name, not carrier count autodetect_ports() required exactly 2 carrier-up RoCE netdevs, but Spark's single CX-7 exposes each of its 2 physical ports over both PCIe segments (socket-direct), so a cabled loopback lights carrier on up to 4 PFs at once. The count check then aborts ("found 3/4"), and it could also pick two PCIe reps of the SAME physical port -- the on-chip eswitch shortcut instead of an over-the-wire path. Group carrier-up RoCE netdevs by phys_port_name (the hardware port label, stable across PCIe re-enumeration and interface renames), keep one rep per physical port (lowest BDF, matching the doc's canonical enp1s0f{0,1}np{0,1} names), and require exactly 2 distinct ports. CLIENT_IF/SERVER_IF therefore always straddle p0 and p1 (over-the-wire). Falls back to per-netdev grouping when phys_port_name is unavailable so the 2-port guard still holds. Refresh the header comments and the manual-override example accordingly. See docs/tutorials/system_configuration.md "Port topology: 4 PFs, 2 ports". Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- scripts/setup_spark_wire_loopback_netns.sh | 67 +++++++++++++++------- 1 file changed, 46 insertions(+), 21 deletions(-) diff --git a/scripts/setup_spark_wire_loopback_netns.sh b/scripts/setup_spark_wire_loopback_netns.sh index 822f078..5473da0 100755 --- a/scripts/setup_spark_wire_loopback_netns.sh +++ b/scripts/setup_spark_wire_loopback_netns.sh @@ -41,10 +41,11 @@ # ibdev2netdev # (mlnx) same mapping, friendlier # ethtool | grep -i 'link detected' # the two CABLED ports show "yes" # -# CLIENT_IF / SERVER_IF must be the TWO PORTS YOUR QSFP CABLE PHYSICALLY -# CONNECTS. Your colleague's box cabled f0np0<->f1np1; the repo YAMLs assume the -# two f0np0 ports across PCI segments. Confirm with `ethtool ... link detected` -# -- exactly the two cabled ports report carrier up. +# CLIENT_IF / SERVER_IF must be on the TWO DIFFERENT PHYSICAL PORTS (p0, p1) of +# the CX-7 -- the on-chip loopback the QSFP cable closes. Confirm each netdev's +# physical port with `cat /sys/class/net//phys_port_name` (p0 / p1); a cabled +# loopback lights carrier on all four PFs, so carrier alone does not distinguish +# them. See docs/tutorials/system_configuration.md "Port topology: 4 PFs, 2 ports". # # CLIENT_RDMA / SERVER_RDMA are the RDMA devices bound to those netdevs # (from `rdma link show` / `ibdev2netdev`). @@ -56,13 +57,14 @@ CLIENT_NS=dq_wire_client SERVER_NS=dq_wire_server # Leave CLIENT_IF / SERVER_IF blank to AUTO-DETECT the cabled data-plane ports -# (recommended): the two RoCE-backed netdevs that report carrier up -- exactly -# the pair the QSFP loopback lights. PCIe enumeration / interface names drift +# (recommended): autodetect_ports() collects the carrier-up RoCE netdevs, groups +# them by phys_port_name, and keeps one per physical port -- so the pair always +# straddles both ports (over-the-wire). PCIe enumeration / interface names drift # across reboots and kernel upgrades on this box, so hardcoding them is fragile # (a stale name aborts setup, the netns never gets created, and the netns-based -# benches silently have nowhere to run). Override from the environment only if -# auto-detect picks the wrong pair, e.g.: -# CLIENT_IF=enp1s0f1np1 SERVER_IF=enP2p1s0f0np0 sudo -E ./setup_spark_wire_loopback_netns.sh up +# benches silently have nowhere to run); phys_port_name does not drift. Override +# from the environment only if auto-detect picks the wrong pair, e.g.: +# CLIENT_IF=enp1s0f1np1 SERVER_IF=enp1s0f0np0 sudo -E ./setup_spark_wire_loopback_netns.sh up CLIENT_IF="${CLIENT_IF:-}" SERVER_IF="${SERVER_IF:-}" @@ -103,27 +105,50 @@ rdma_for_netdev() { # netdev -> backing RDMA device name (sysfs); empty if none } autodetect_ports() { - # Pick the cabled data-plane ports automatically when not overridden: the - # RoCE-backed netdevs in init_net that report carrier up -- exactly the pair - # the QSFP loopback lights. Ports must be in init_net here (up() calls down() - # first, returning any moved netdevs), same precondition as detect_macs. + # Pick the cabled data-plane ports automatically when not overridden. Spark's + # single ConnectX-7 exposes each of its TWO physical ports (p0, p1) over BOTH + # PCIe segments (socket-direct), so it presents up to 4 PFs and a cabled + # loopback lights carrier on ALL of them at once -- not just two. Counting + # carrier-up netdevs therefore over-counts the physical ports. Instead group by + # phys_port_name (the hardware port label, stable across PCIe re-enumeration and + # interface renames) and keep ONE netdev per physical port, so CLIENT_IF and + # SERVER_IF always land on DIFFERENT physical ports (over-the-wire, never the + # on-chip same-port eswitch shortcut). See docs/tutorials/system_configuration.md + # "Port topology: 4 PFs, 2 ports". Ports must be in init_net here (up() calls + # down() first, returning any moved netdevs), same precondition as detect_macs. if [[ -z "$CLIENT_IF" || -z "$SERVER_IF" ]]; then - local found=() ifc + local ifc port pci + declare -A rep_for_port=() pci_for_port=() for ifc in $(ls /sys/class/net 2>/dev/null | sort); do [[ "$ifc" == lo ]] && continue rdma_for_netdev "$ifc" >/dev/null 2>&1 || continue [[ "$(cat "/sys/class/net/$ifc/carrier" 2>/dev/null || echo 0)" == 1 ]] || continue - found+=("$ifc") + port="$(cat "/sys/class/net/$ifc/phys_port_name" 2>/dev/null || true)" + # Without phys_port_name we cannot fold the two PCIe reps of a port + # together; fall back to the netdev name as its own group so the 2-port + # check below still guards (it just won't dedupe). + [[ -z "$port" ]] && port="$ifc" + pci="$(basename "$(readlink -f "/sys/class/net/$ifc/device" 2>/dev/null || echo "$ifc")")" + # Prefer the lowest-BDF rep per port (the 0000: segment, matching the doc's + # canonical enp1s0f{0,1}np{0,1} names) for a stable, least-surprise pick. + if [[ -z "${rep_for_port[$port]:-}" || "$pci" < "${pci_for_port[$port]}" ]]; then + rep_for_port["$port"]="$ifc" + pci_for_port["$port"]="$pci" + fi done - if [[ "${#found[@]}" -ne 2 ]]; then - echo "ERROR: auto-detect expected exactly 2 carrier-up RoCE ports, found ${#found[@]}: ${found[*]:-none}." >&2 + local ports + mapfile -t ports < <(printf '%s\n' "${!rep_for_port[@]}" | LC_ALL=C sort) + if [[ "${#ports[@]}" -ne 2 ]]; then + echo "ERROR: auto-detect expected exactly 2 carrier-up physical ports, found ${#ports[@]}: ${ports[*]:-none}." >&2 + echo " (netdevs are grouped by phys_port_name; Spark exposes each port over 2 PCIe segments.)" >&2 echo " Cable the loopback, or set CLIENT_IF / SERVER_IF explicitly, e.g.:" >&2 - echo " CLIENT_IF=enp1s0f1np1 SERVER_IF=enP2p1s0f0np0 sudo -E $0 up" >&2 - echo " List candidates with: ip -br link show | grep -E 'enp|enP'" >&2 + echo " CLIENT_IF=enp1s0f1np1 SERVER_IF=enp1s0f0np0 sudo -E $0 up" >&2 + echo " List candidates and their ports with:" >&2 + echo " for d in /sys/class/net/*/device; do n=\${d%/device}; n=\${n##*/}; printf '%s port=%s\\n' \"\$n\" \"\$(cat /sys/class/net/\$n/phys_port_name 2>/dev/null)\"; done" >&2 exit 1 fi - [[ -z "$CLIENT_IF" ]] && CLIENT_IF="${found[0]}" - [[ -z "$SERVER_IF" ]] && SERVER_IF="${found[1]}" + [[ -z "$CLIENT_IF" ]] && CLIENT_IF="${rep_for_port[${ports[0]}]}" + [[ -z "$SERVER_IF" ]] && SERVER_IF="${rep_for_port[${ports[1]}]}" fi [[ -z "$CLIENT_RDMA" ]] && CLIENT_RDMA="$(rdma_for_netdev "$CLIENT_IF" || true)" [[ -z "$SERVER_RDMA" ]] && SERVER_RDMA="$(rdma_for_netdev "$SERVER_IF" || true)" From 57203aebdb9c0f9e27cfd853beb74fbfc0b35acf Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Fri, 10 Jul 2026 10:32:13 -0400 Subject: [PATCH 11/18] #15 - Add --workload-fft-len flag to pin the FFT transform length The FFT workload's 1-D C2C transform length was a hardcoded constant (1024), while GEMM already had --workload-gemm-dim. Add a symmetric --workload-fft-len so the FFT compute size is a benchmark knob, independent of the GEMM dimension: - parse_workload_fft_len() mirrors parse_workload_gemm_dim; GpuWorkload::init() takes an fft_len arg (default 1024) and threads it into cufftPlan1d, replacing the kFftLen constant. The working set still fans out across as many batched length-N transforms as fit. - Thread fft_len through every bench main (rdma, socket, raw gpudirect/hds) and rx_count_worker, exactly as gemm_dim is. - run_spark_bench.sh gains an FFT_LEN env (default 1024) passed as --workload-fft-len for all backends. Verified in-container: FFT_LEN=512 logs "fft shape = 1953 x C2C length-512". Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- AGENTS.md | 2 +- examples/bench_workload.cu | 31 +++++++++++++++++++++++-------- examples/bench_workload.h | 22 +++++++++++++++++----- examples/raw_bench_common.cpp | 5 +++-- examples/raw_bench_common.h | 3 ++- examples/raw_gpudirect_bench.cpp | 3 ++- examples/raw_hds_bench.cpp | 3 ++- examples/rdma_bench.cpp | 10 ++++++---- examples/run_spark_bench.sh | 10 ++++++++++ examples/socket_bench.cpp | 9 +++++---- 10 files changed, 71 insertions(+), 27 deletions(-) diff --git a/AGENTS.md b/AGENTS.md index 3eea151..7448afc 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -44,7 +44,7 @@ There is no unit test suite. Verification is done via the benchmark executables | `daqiri_bench_socket` | `socket_bench.cpp` | `daqiri_bench_socket_{udp,tcp}_tx_rx.yaml`, `daqiri_bench_socket_{udp,tcp}_tx_rx_spark_netns.yaml` (combined-role netns bases) | | `daqiri_pool_ring_bench` | `pool_ring_bench.cpp` | none — microbenchmark comparing `daqiri::Ring`/`daqiri::ObjectPool` vs DPDK `rte_ring`/`rte_mempool` (SPSC/MPMC, single/bulk, thread sweep). The `rte_*` comparison arm compiles only in a DPDK-enabled build; takes no YAML/CLI args | -The four `raw_*` benches share `raw_bench_common.{cpp,h}` and accept `--seconds N`. `daqiri_bench_rdma` and `daqiri_bench_socket` also take `--mode {tx,rx,both}`. `daqiri_bench_raw_gpudirect`, `daqiri_bench_raw_hds`, `daqiri_bench_rdma`, and `daqiri_bench_socket` additionally accept `--workload none|fft|gemm|gemm_fp16` — a reusable representative GPU workload (`examples/bench_workload.{h,cu}`, cuFFT/cuBLAS) run once per received reorder window on the **actual received payload**. Each backend first assembles the burst's payloads into one contiguous GPU buffer via `examples/bench_pipeline.{h,cu}` (`ReorderPipeline`): a sequence-number reorder kernel for the out-of-order transports (DPDK raw, UDP) and an arrival-order gather for the in-order ones (RoCE RC, TCP); sockets stage host→device first since their payloads land in pageable host memory. The reorder/gather kernels (`packet_reorder_copy_payload_by_sequence`, `packet_gather_copy_payload`) live in `src/kernels.cu`. `gemm` is FP32 `cublasSgemm`; `gemm_fp16` is the same-size mixed-precision FP16/tensor-core `cublasGemmEx` (inference-style); the contiguous buffer supplies the FFT input / GEMM A operand. `--workload-gemm-dim N` pins the square GEMM side length (default 1024), so the FLOP count per call (2·n³) is FIXED and the compute working set is exactly n·n·elem_size, read from the front of each received I/O unit (the unit must be at least that large). Used by `run_spark_bench.sh`'s `WORKLOAD` / `GEMM_DIM` env (all backends) to fill the CSV `post_process` / `post_process_gemm_dim` columns (issue #15). +The four `raw_*` benches share `raw_bench_common.{cpp,h}` and accept `--seconds N`. `daqiri_bench_rdma` and `daqiri_bench_socket` also take `--mode {tx,rx,both}`. `daqiri_bench_raw_gpudirect`, `daqiri_bench_raw_hds`, `daqiri_bench_rdma`, and `daqiri_bench_socket` additionally accept `--workload none|fft|gemm|gemm_fp16` — a reusable representative GPU workload (`examples/bench_workload.{h,cu}`, cuFFT/cuBLAS) run once per received reorder window on the **actual received payload**. Each backend first assembles the burst's payloads into one contiguous GPU buffer via `examples/bench_pipeline.{h,cu}` (`ReorderPipeline`): a sequence-number reorder kernel for the out-of-order transports (DPDK raw, UDP) and an arrival-order gather for the in-order ones (RoCE RC, TCP); sockets stage host→device first since their payloads land in pageable host memory. The reorder/gather kernels (`packet_reorder_copy_payload_by_sequence`, `packet_gather_copy_payload`) live in `src/kernels.cu`. `gemm` is FP32 `cublasSgemm`; `gemm_fp16` is the same-size mixed-precision FP16/tensor-core `cublasGemmEx` (inference-style); the contiguous buffer supplies the FFT input / GEMM A operand. `--workload-gemm-dim N` pins the square GEMM side length (default 1024), so the FLOP count per call (2·n³) is FIXED and the compute working set is exactly n·n·elem_size, read from the front of each received I/O unit (the unit must be at least that large). `--workload-fft-len N` pins the 1-D C2C transform length for `fft` (default 1024; independent of the GEMM dimension); the working set is fanned out across as many batched length-N transforms as fit. Used by `run_spark_bench.sh`'s `WORKLOAD` / `GEMM_DIM` / `FFT_LEN` env (all backends) to fill the CSV `post_process` / `post_process_gemm_dim` columns (issue #15). ```bash ./build/examples/daqiri_bench_raw_gpudirect ./build/examples/daqiri_bench_raw_tx_rx.yaml --seconds 10 diff --git a/examples/bench_workload.cu b/examples/bench_workload.cu index 2f055d1..2c71611 100644 --- a/examples/bench_workload.cu +++ b/examples/bench_workload.cu @@ -95,6 +95,18 @@ int parse_workload_gemm_dim(int argc, char** argv) { return 1024; // default square GEMM side length } +int parse_workload_fft_len(int argc, char** argv) { + for (int i = 2; i + 1 < argc; i += 2) { + if (std::string(argv[i]) == "--workload-fft-len") { + const long long v = std::atoll(argv[i + 1]); + if (v > 0) { + return static_cast(v); + } + } + } + return kFftLen; // default 1-D C2C transform length +} + int parse_workload_sync_interval(int argc, char** argv) { for (int i = 2; i + 1 < argc; i += 2) { if (std::string(argv[i]) == "--workload-sync-interval") { @@ -171,7 +183,7 @@ void GpuWorkload::destroy() { } bool GpuWorkload::init(BenchWorkload kind, size_t input_bytes, int sync_interval, - int gemm_dim) { + int gemm_dim, int fft_len) { kind_ = kind; sync_interval_ = sync_interval > 0 ? sync_interval : 1; run_count_ = 0; @@ -191,12 +203,15 @@ bool GpuWorkload::init(BenchWorkload kind, size_t input_bytes, int sync_interval stream_ = stream; if (kind_ == BenchWorkload::Fft) { - // Fan the working set out across batched length-kFftLen C2C transforms. The + // 1-D C2C transform length, pinned via --workload-fft-len (default kFftLen); + // held fixed while the I/O unit is swept, like the GEMM side length. + fft_len_ = std::max(2, fft_len > 0 ? fft_len : kFftLen); + // Fan the working set out across batched length-fft_len_ C2C transforms. The // transform reads the caller's input buffer and writes the owned fft_out_; // total*sizeof(cufftComplex) <= bytes so the input always holds enough data. - const size_t n_complex = std::max(kFftLen, bytes / sizeof(cufftComplex)); - const int batch = std::max(1, static_cast(n_complex / kFftLen)); - const size_t total = static_cast(kFftLen) * batch; + const size_t n_complex = std::max(fft_len_, bytes / sizeof(cufftComplex)); + const int batch = std::max(1, static_cast(n_complex / fft_len_)); + const size_t total = static_cast(fft_len_) * batch; if (cudaMalloc(&fft_out_, total * sizeof(cufftComplex)) != cudaSuccess || cudaMemset(fft_out_, 0, total * sizeof(cufftComplex)) != cudaSuccess) { @@ -205,7 +220,7 @@ bool GpuWorkload::init(BenchWorkload kind, size_t input_bytes, int sync_interval return false; } cufftHandle plan = 0; - if (cufftPlan1d(&plan, kFftLen, CUFFT_C2C, batch) != CUFFT_SUCCESS) { + if (cufftPlan1d(&plan, fft_len_, CUFFT_C2C, batch) != CUFFT_SUCCESS) { std::cerr << "GpuWorkload: cufftPlan1d failed\n"; destroy(); return false; @@ -218,8 +233,8 @@ bool GpuWorkload::init(BenchWorkload kind, size_t input_bytes, int sync_interval } // Explicit shape so every published FFT number carries its compute size. // ~5·N·log2(N) flops per length-N C2C transform, times `batch` transforms. - const double flops = 5.0 * kFftLen * std::log2(static_cast(kFftLen)) * batch; - std::cerr << "GpuWorkload: fft shape = " << batch << " x C2C length-" << kFftLen + const double flops = 5.0 * fft_len_ * std::log2(static_cast(fft_len_)) * batch; + std::cerr << "GpuWorkload: fft shape = " << batch << " x C2C length-" << fft_len_ << " (working set " << (bytes >> 10) << " KiB, " << (flops / 1e6) << " MFLOP/call)\n"; } else { // Gemm or GemmFp16 // Square matmul with the side length n pinned directly (FP32 and FP16 use the diff --git a/examples/bench_workload.h b/examples/bench_workload.h index d4f8ed0..7bcdeeb 100644 --- a/examples/bench_workload.h +++ b/examples/bench_workload.h @@ -41,6 +41,13 @@ BenchWorkload parse_workload(int argc, char** argv); // Returns 1024 (the default) if unset. Mirrors parse_workload's stride. int parse_workload_gemm_dim(int argc, char** argv); +// Parse "--workload-fft-len N" from argv: the 1-D C2C transform length. The burst +// working set is fanned out across as many batched length-N transforms as fit, so +// N sets the per-transform cost (~5*N*log2(N) flops) while the batch count tracks +// the I/O unit. Independent of --workload-gemm-dim. Returns 1024 (the default) if +// unset. Mirrors parse_workload's stride. +int parse_workload_fft_len(int argc, char** argv); + // Parse "--workload-sync-interval N" from argv: drain the GPU stream every N compute // calls (bounds outstanding GPU work). Larger N = deeper async queue, fewer CPU // stalls waiting on the GPU; N=1 is fully synchronous. Used to characterize how much @@ -98,11 +105,14 @@ class GpuWorkload { // GEMM; <=0 falls back to the 1024 default). kind == None leaves the object an // inert no-op. sync_interval bounds outstanding GPU work (sync every N runs) for // the run()+sync() callers (DPDK/socket); the RoCE path bounds work with events - // instead and ignores it. Logs the chosen problem shape (GEMM n / FFT - // length+batch) and FLOP count so every published benchmark number carries its - // explicit compute size. Returns false on CUDA / library error; the caller may - // warn and continue with the workload disabled (enabled() will report false). - bool init(BenchWorkload kind, size_t input_bytes, int sync_interval = 2, int gemm_dim = 1024); + // instead and ignores it. `fft_len` is the 1-D FFT transform length for the FFT + // kind (>0; <=0 falls back to the 1024 default), ignored by the GEMM kinds. + // Logs the chosen problem shape (GEMM n / FFT length+batch) and FLOP count so + // every published benchmark number carries its explicit compute size. Returns + // false on CUDA / library error; the caller may warn and continue with the + // workload disabled (enabled() will report false). + bool init(BenchWorkload kind, size_t input_bytes, int sync_interval = 2, int gemm_dim = 1024, + int fft_len = 1024); // Enqueue one representative FFT/SGEMM on the internal stream, reading `input` // (a device pointer to >= the GEMM working set of valid bytes). No-op unless @@ -169,6 +179,8 @@ class GpuWorkload { void* cublas_ = nullptr; // cublasHandle_t int fft_plan_ = -1; // cufftHandle (-1 == unset) + int fft_len_ = 0; // 1-D C2C transform length (set in init) + // Device scratch (operands NOT sourced from received data). void* fft_out_ = nullptr; // cufftComplex[fft_total_] (FFT output) void* gemm_b_ = nullptr; // B operand diff --git a/examples/raw_bench_common.cpp b/examples/raw_bench_common.cpp index 28db189..7ca75b3 100644 --- a/examples/raw_bench_common.cpp +++ b/examples/raw_bench_common.cpp @@ -550,7 +550,7 @@ void print_queue_stats(const char *direction, const std::string &interface_name, void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, BenchWorkload workload, const ReorderGeometry& geom, int workload_gemm_dim, - int workload_sync_interval) { + int workload_sync_interval, int workload_fft_len) { if (!set_current_thread_affinity(cfg.cpu_core, "bench_rx")) { stop.store(true); return; @@ -565,7 +565,8 @@ void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, Bench GpuWorkload gpu_workload; ReorderPipeline pipeline; if (workload != BenchWorkload::None) { - if (!gpu_workload.init(workload, batch_bytes, workload_sync_interval, workload_gemm_dim) || + if (!gpu_workload.init(workload, batch_bytes, workload_sync_interval, workload_gemm_dim, + workload_fft_len) || !pipeline.init(ReorderMode::SeqReorder, geom.packets_per_batch, out_payload_len, geom.payload_byte_offset, geom.seq_bit_offset, geom.seq_bit_width, /*staging_needed=*/false, gpu_workload.stream())) { diff --git a/examples/raw_bench_common.h b/examples/raw_bench_common.h index bc006dd..e57cf41 100644 --- a/examples/raw_bench_common.h +++ b/examples/raw_bench_common.h @@ -161,6 +161,7 @@ struct ReorderGeometry { // workload == None leaves the bare-loopback count-only path untouched. void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, BenchWorkload workload = BenchWorkload::None, const ReorderGeometry& geom = {}, - int workload_gemm_dim = 1024, int workload_sync_interval = 2); + int workload_gemm_dim = 1024, int workload_sync_interval = 2, + int workload_fft_len = 1024); } // namespace daqiri::bench diff --git a/examples/raw_gpudirect_bench.cpp b/examples/raw_gpudirect_bench.cpp index 06d26b4..c536f41 100644 --- a/examples/raw_gpudirect_bench.cpp +++ b/examples/raw_gpudirect_bench.cpp @@ -173,6 +173,7 @@ int main(int argc, char **argv) { const double target_gbps = daqiri::bench::parse_target_gbps(argc, argv); const auto workload = daqiri::bench::parse_workload(argc, argv); const int workload_gemm_dim = daqiri::bench::parse_workload_gemm_dim(argc, argv); + const int workload_fft_len = daqiri::bench::parse_workload_fft_len(argc, argv); const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); @@ -222,7 +223,7 @@ int main(int argc, char **argv) { rx_threads.reserve(rx_configs.size()); for (const auto &cfg : rx_configs) { rx_threads.emplace_back(daqiri::bench::rx_count_worker, cfg, std::ref(stop), workload, geom, - workload_gemm_dim, workload_sync_interval); + workload_gemm_dim, workload_sync_interval, workload_fft_len); } tx_threads.reserve(tx_configs.size()); for (const auto &cfg : tx_configs) { diff --git a/examples/raw_hds_bench.cpp b/examples/raw_hds_bench.cpp index 8d936ee..1706157 100644 --- a/examples/raw_hds_bench.cpp +++ b/examples/raw_hds_bench.cpp @@ -148,6 +148,7 @@ int main(int argc, char **argv) { const int run_seconds = daqiri::bench::parse_run_seconds(argc, argv); const auto workload = daqiri::bench::parse_workload(argc, argv); const int workload_gemm_dim = daqiri::bench::parse_workload_gemm_dim(argc, argv); + const int workload_fft_len = daqiri::bench::parse_workload_fft_len(argc, argv); const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); if (daqiri::daqiri_init(argv[1]) != daqiri::Status::SUCCESS) { @@ -188,7 +189,7 @@ int main(int argc, char **argv) { } rx_thread = std::thread(daqiri::bench::rx_count_worker, daqiri::bench::parse_rx(root), std::ref(stop), - workload, geom, workload_gemm_dim, workload_sync_interval); + workload, geom, workload_gemm_dim, workload_sync_interval, workload_fft_len); } if (has_tx) { tx_thread = diff --git a/examples/rdma_bench.cpp b/examples/rdma_bench.cpp index fe4a19b..a0d6cea 100644 --- a/examples/rdma_bench.cpp +++ b/examples/rdma_bench.cpp @@ -83,7 +83,7 @@ RdmaBenchConfig parse_rdma_cfg(const YAML::Node& node) { void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pacer, std::atomic& stop, RdmaWorkerStats& stats, daqiri::bench::BenchWorkload workload, int workload_gemm_dim, - int workload_sync_interval, int workload_max_inflight) { + int workload_sync_interval, int workload_max_inflight, int workload_fft_len) { const char *thread_name = cfg.server ? "rdma_bench_server" : "rdma_bench_client"; if (!daqiri::bench::set_current_thread_affinity(cfg.cpu_core, thread_name)) { @@ -104,7 +104,8 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa daqiri::bench::GpuWorkload gpu_workload; daqiri::bench::ReorderPipeline pipeline; if (workload != daqiri::bench::BenchWorkload::None) { - if (!gpu_workload.init(workload, msg, workload_sync_interval, workload_gemm_dim) || + if (!gpu_workload.init(workload, msg, workload_sync_interval, workload_gemm_dim, + workload_fft_len) || !pipeline.init(daqiri::bench::ReorderMode::GatherOnly, /*packets_per_batch=*/1, msg, /*payload_byte_offset=*/0, /*seq_bit_offset=*/0, @@ -370,6 +371,7 @@ int main(int argc, char** argv) { } const auto workload = daqiri::bench::parse_workload(argc, argv); const int workload_gemm_dim = daqiri::bench::parse_workload_gemm_dim(argc, argv); + const int workload_fft_len = daqiri::bench::parse_workload_fft_len(argc, argv); const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const int workload_max_inflight = daqiri::bench::parse_workload_max_inflight(argc, argv); @@ -409,12 +411,12 @@ int main(int argc, char** argv) { if (run_server) { server_thread = std::thread(rdma_worker, server_cfg, std::ref(server_pacer), std::ref(stop), std::ref(server_stats), workload, workload_gemm_dim, - workload_sync_interval, workload_max_inflight); + workload_sync_interval, workload_max_inflight, workload_fft_len); } if (run_client) { client_thread = std::thread(rdma_worker, client_cfg, std::ref(client_pacer), std::ref(stop), std::ref(client_stats), workload, workload_gemm_dim, - workload_sync_interval, workload_max_inflight); + workload_sync_interval, workload_max_inflight, workload_fft_len); } if (!server_thread.joinable() && !client_thread.joinable()) { diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index a1603db..6444ba6 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -33,6 +33,9 @@ # 1024), held fixed so FLOPs/call (2·n³) is constant. The # compute working set is n·n·elem_size, read from the front of # each received I/O unit. Recorded in post_process_gemm_dim. +# FFT_LEN — the 1-D C2C transform length (--workload-fft-len; default +# 1024) for WORKLOAD=fft, held fixed while the I/O unit is +# swept. Independent of GEMM_DIM. # SYNC_INTERVAL — drain the GPU stream every N compute calls # (--workload-sync-interval; default 2). Sweep it (1 2 4 8 16 32) # to see how much of the receive+compute ceiling is single-thread @@ -110,6 +113,12 @@ GEMM_DIM="${GEMM_DIM:-1024}" if [[ ! "$GEMM_DIM" =~ ^[0-9]+$ ]]; then echo "Invalid GEMM_DIM '$GEMM_DIM' (expected a positive integer)" >&2; exit 1 fi +# FFT_LEN: the 1-D C2C transform length (--workload-fft-len) for WORKLOAD=fft, held +# FIXED while the I/O unit is swept. Independent of GEMM_DIM. Default 1024. +FFT_LEN="${FFT_LEN:-1024}" +if [[ ! "$FFT_LEN" =~ ^[0-9]+$ ]]; then + echo "Invalid FFT_LEN '$FFT_LEN' (expected a positive integer)" >&2; exit 1 +fi # SYNC_INTERVAL: drain the GPU stream every N compute calls (--workload-sync-interval). # Larger N lets more GEMMs queue asynchronously before the single receive+compute # thread blocks on the GPU, so sweeping it (e.g. 1 2 4 8 16 32) shows how much of the @@ -472,6 +481,7 @@ run_cell() { if [[ "$WORKLOAD_EFF" != "none" ]]; then bench_extra+=(--workload "$WORKLOAD_EFF") bench_extra+=(--workload-gemm-dim "$GEMM_DIM") + bench_extra+=(--workload-fft-len "$FFT_LEN") [[ -n "$SYNC_INTERVAL" ]] && bench_extra+=(--workload-sync-interval "$SYNC_INTERVAL") [[ -n "$MAX_INFLIGHT" ]] && bench_extra+=(--workload-max-inflight "$MAX_INFLIGHT") fi diff --git a/examples/socket_bench.cpp b/examples/socket_bench.cpp index 40048de..948307b 100644 --- a/examples/socket_bench.cpp +++ b/examples/socket_bench.cpp @@ -103,7 +103,7 @@ bool socket_transport_is_tcp(const YAML::Node& root) { void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pacer, std::atomic& stop, SocketWorkerStats& stats, daqiri::bench::BenchWorkload workload, bool is_tcp, int workload_gemm_dim, - int workload_sync_interval) { + int workload_sync_interval, int workload_fft_len) { const char *thread_name = cfg.server ? "socket_bench_server" : "socket_bench_client"; if (!daqiri::bench::set_current_thread_affinity(cfg.cpu_core, thread_name)) { @@ -126,7 +126,7 @@ void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer daqiri::bench::ReorderPipeline pipeline; if (workload != daqiri::bench::BenchWorkload::None) { if (!gpu_workload.init(workload, static_cast(packets_per_batch) * msg, - workload_sync_interval, workload_gemm_dim) || + workload_sync_interval, workload_gemm_dim, workload_fft_len) || !pipeline.init(is_tcp ? daqiri::bench::ReorderMode::GatherOnly : daqiri::bench::ReorderMode::SeqReorder, packets_per_batch, msg, /*payload_byte_offset=*/0, @@ -259,6 +259,7 @@ int main(int argc, char** argv) { } const auto workload = daqiri::bench::parse_workload(argc, argv); const int workload_gemm_dim = daqiri::bench::parse_workload_gemm_dim(argc, argv); + const int workload_fft_len = daqiri::bench::parse_workload_fft_len(argc, argv); const int workload_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); @@ -298,12 +299,12 @@ int main(int argc, char** argv) { if (run_server) { server_thread = std::thread(socket_worker, server_cfg, std::ref(server_pacer), std::ref(stop), std::ref(server_stats), workload, is_tcp, workload_gemm_dim, - workload_sync_interval); + workload_sync_interval, workload_fft_len); } if (run_client) { client_thread = std::thread(socket_worker, client_cfg, std::ref(client_pacer), std::ref(stop), std::ref(client_stats), workload, is_tcp, workload_gemm_dim, - workload_sync_interval); + workload_sync_interval, workload_fft_len); } if (!server_thread.joinable() && !client_thread.joinable()) { From 4d0df9a94d30f53ad27e40fd178de0315ce93a9b Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Fri, 10 Jul 2026 10:36:24 -0400 Subject: [PATCH 12/18] #15 - Re-measure one-way RoCE results in the DGX Spark perf report Now that the RoCE bench is one-way, re-measure and refresh the perf report so all three transports are apples-to-apples: - Results Summary RoCE row and message-size sweep updated with one-way numbers (best case 106.0 Gb/s at 64 KB; 8 MB 105.0). Large messages rose vs the old bidirectional run (reverse-direction contention removed); the op-rate-bound small cells are the honest single-direction rate. - CPU-utilization table: fix the receiver core. For rdma, cpu_rx measured the idle client RX core (18); the one-way receiver is the SERVER, whose RX poller + bench worker share core 19. Point CPU_RX at 19 and re-measure -- the receiver core is ~1.1% at the bare-loopback 8 MB cell, confirming the RoCE RC "HCA DMAs straight to memory" signature with the correct core. - Reproduce: collapse the two redundant GPU-workload blocks into one (a workload-active sweep auto-pins to the summary-table cell, so it equals smoke); add a general smoke-vs-sweep explanation; document the FFT default and the new FFT_LEN / --workload-fft-len knob; drop the redundant explicit GEMM_DIM=1024. - Rename "headline cell" -> "summary-table cell" throughout. Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- docs/benchmarks/performance-dgx-spark.md | 91 ++++++++++++------------ examples/run_spark_bench.sh | 7 +- 2 files changed, 50 insertions(+), 48 deletions(-) diff --git a/docs/benchmarks/performance-dgx-spark.md b/docs/benchmarks/performance-dgx-spark.md index 78c8255..1df3677 100644 --- a/docs/benchmarks/performance-dgx-spark.md +++ b/docs/benchmarks/performance-dgx-spark.md @@ -38,7 +38,7 @@ four-pair aggregate is shown. | Stream / Protocol | Best case | Throughput | Drops | Notes | | ----------------- | --------- | ---------: | ----- | ----- | | Raw Ethernet / GPUDirect | 4 KB packet | **105.5 ±0.9 Gb/s** | 0 | 98.5 Gb/s single-queue at 8 KB | -| Socket / RoCE (SEND) | 8 MB message | **102.2 ±0.3 Gb/s** | 0 | Single QP, batch 1 | +| Socket / RoCE (SEND) | 64 KB message | **106.0 ±0.1 Gb/s** | 0 | Single QP, batch 1 | | Socket / TCP | 8 KB × 4 pairs | **97.2 ±2.8 Gb/s** | ~0 | Flow-controlled (App TX = App RX) | | Socket / UDP | 8 KB × 4 pairs | **29.8 ±0.2 Gb/s** | ~51% loss | Receiver goodput; unpaced sender | @@ -85,7 +85,7 @@ pps ≈ Gb/s ÷ ((payload + 64) × 8). These small-payload cells are flat across size and stable run-to-run. Because every cell is drop-free, the achieved rate is also the no-drop rate: pacing the sender below it hits the target with zero drops. -**CPU utilization** (headline cell, 8000 B / batch 10240, unpaced): +**CPU utilization** (summary-table cell, 8000 B / batch 10240, unpaced): | Core | Busy% | Note | | ------------------------ | ----: | ------------------------------------- | @@ -147,36 +147,37 @@ saturate the wire; the smallest messages are bound by per-operation software overhead. **Message-size sweep (single QP, batch 1, 0 drops).** Mean of 3 reps; run-to-run -spread ≤0.8 Gb/s (<2%) in every cell. +spread ≤0.7 Gb/s (<1%) in every cell. | Message size | Gb/s | | ------------ | ---: | -| 8 MB | **102.2** | -| 1 MB | 101.3 | -| 64 KB | 101.6 | -| 8 KB | 60.7 | -| 4 KB | 38.0 | +| 8 MB | 105.0 | +| 1 MB | 105.1 | +| 64 KB | **106.0** | +| 8 KB | 51.5 | +| 4 KB | 28.6 | -Messages ≥64 KB hold ~101–102 Gb/s at the wire ceiling. Below that the path is +Messages ≥64 KB hold ~105–106 Gb/s at the wire ceiling. Below that the path is operation-rate-bound (per-operation software overhead, not a stall) rather than -wire-bound, and every cell is drop-free. At 8 KB (60.7 Gb/s) and 4 KB (38.0 Gb/s) +wire-bound, and every cell is drop-free. At 8 KB (51.5 Gb/s) and 4 KB (28.6 Gb/s) a dedicated bench-worker core, separate from the RoCE engine thread, sustains the operation rate, as it does for small DPDK packets. A per-message flow-control window keeps enough operations in flight to amortize that overhead: it pre-posts `rx_depth` receives before sending and caps the transmit side at `tx_depth`, each sized to the message so the in-flight window stays full. -**CPU utilization** (headline cell, 8 MB message, batch 1, unpaced): +**CPU utilization** (summary-table cell, 8 MB message, batch 1, unpaced): -| Core | Busy% | Note | -| ------------------- | ----: | ----------------------------------------------- | -| Master (CPU 8) | 5.9% | Orchestration only | -| Client TX (CPU 17) | 74.9% | Post-and-poll spin; rate-independent | -| Server RX (CPU 18) | 0.0% | HCA writes straight to memory; CPU uninvolved | +| Core | Busy% | Note | +| -------------------- | ----: | ----------------------------------------------- | +| Master (CPU 8) | 0.7% | Orchestration only | +| Client TX (CPU 17) | 74.8% | Post-and-poll spin; rate-independent | +| Server RX (CPU 19) | 1.1% | HCA DMAs straight to memory; worker only reaps completions | -The idle RX core is the expected RoCE RC signature — the HCA places incoming -data directly into registered memory with no CPU involvement. The GPU stays -idle here too (SM and memory-controller ~0%; DMA target, not a compute engine). +The near-idle RX core is the expected RoCE RC signature — the HCA places incoming +data directly into registered memory, so the receive worker only reaps +completions and reposts (~1% at this message rate). The GPU stays idle here too +(SM and memory-controller ~0%; DMA target, not a compute engine). ## Socket / TCP @@ -334,6 +335,12 @@ apt-get install -y iproute2 iputils-ping ethtool iperf3 rdma-core ibverbs-utils These provide `ip`/`nstat` (`iproute2`), `ethtool`, and `ib_send_bw` (`perftest`). +Each `run_spark_bench.sh ` invocation takes a **mode** that sets +which cells run: `sweep` runs the full payload × batch × pairs matrix (the +per-transport message-size tables above), while `smoke` runs just the single +summary-table cell — one payload/batch/pairs operating point. `REPEATS=N` repeats +every cell N times for error bars. + **Raw Ethernet / GPUDirect (DPDK)** drives the two physical ports directly, so the `dq_wire_*` namespaces must **not** be up — they capture the ports and hide them from DPDK. Tear them down first (no-op if they were never created). @@ -371,40 +378,30 @@ before running; tear it down when finished: ./scripts/setup_spark_wire_loopback_netns.sh down # tear down when done ``` -**GPU workload (FFT / GEMM)** re-runs a backend with a representative GPU -workload in the receive path by exporting `WORKLOAD`. It composes with the same -modes and netns setup as above (dpdk in the default namespace, rdma in the -`dq_wire_*` namespaces): +**GPU workload (FFT / GEMM)** re-runs a backend with a representative GPU workload +in the receive path by exporting `WORKLOAD` (`none` | `fft` | `gemm` | +`gemm_fp16`), run once per received I/O unit on the real payload. Each call is a +fixed **1024³ GEMM** (override with `GEMM_DIM` / `--workload-gemm-dim`) or a batched +**length-1024 FFT** (override with `FFT_LEN` / `--workload-fft-len`) — both compute +sizes are held constant while the message size varies, so the FLOP count per call +is fixed. It composes with the same netns setup as above (dpdk in the default +namespace, rdma in the `dq_wire_*` namespaces). Use `smoke` — the single +summary-table cell that the fixed-n table reports — and run all three workloads +with error bars: ```bash -# Raw / GPUDirect (netns down, ETH_DST_ADDR exported as above) -WORKLOAD=fft ./examples/run_spark_bench.sh dpdk smoke -WORKLOAD=gemm ./examples/run_spark_bench.sh dpdk smoke -# Socket / RoCE (netns up) -WORKLOAD=fft ./examples/run_spark_bench.sh rdma smoke -WORKLOAD=gemm ./examples/run_spark_bench.sh rdma smoke -``` - -The chosen workload lands in the CSV `post_process` column; compare `gbps` / -`gpu_sm_pct` against the matching `WORKLOAD=none` baseline. - -**Fixed-size GEMM comparison** pins the matrix dimension with `GEMM_DIM` -(`--workload-gemm-dim`, default 1024) so the FLOP count is constant across transports; -each side runs one fixed 1024³ GEMM per received I/O unit, reading its first 4 MB. Lands -in the CSV `post_process_gemm_dim` column (the sweep is restricted to the headline -message size automatically when a workload is active): - -```bash -# RoCE (netns up) -for WL in none fft gemm; do - WORKLOAD=$WL GEMM_DIM=1024 REPEATS=3 ./examples/run_spark_bench.sh rdma sweep -done -# Raw (netns down, ETH_DST_ADDR exported) +# RoCE (netns up); Raw is identical with `dpdk`, netns down, ETH_DST_ADDR exported. for WL in none fft gemm; do - WORKLOAD=$WL GEMM_DIM=1024 REPEATS=3 ./examples/run_spark_bench.sh dpdk sweep + WORKLOAD=$WL REPEATS=3 ./examples/run_spark_bench.sh rdma smoke done ``` +In the workload case the payload size is fixed per backend (8 KB for DPDK, 8 MB +message for RoCE), so a `sweep` only steps through batch size (DPDK) or +client/server pairs (sockets). The workload lands in the CSV `post_process` column +(with the GEMM dimension in `post_process_gemm_dim`); compare each `gbps` / +`gpu_sm_pct` against the `WORKLOAD=none` baseline from the same loop. + Each run writes `bench-results/--/runs.csv`. See [Socket and RDMA Benchmarking](socket_benchmarking.md) and [Raw Ethernet Benchmarking](raw_benchmarking.md) for the namespace setup and diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index 6444ba6..e172b0f 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -201,7 +201,12 @@ case "$BACKEND" in # the kernel's local routing table. BASE_YAML="$SCRIPT_DIR/daqiri_bench_rdma_tx_rx_spark_netns.yaml" BENCH_BIN="$BUILD_DIR/examples/daqiri_bench_rdma" - CPU_MASTER=8; CPU_TX=17; CPU_RX=18 + # One-way roles: the client (send-only) drives the TX-queue core 17; the server + # (receive-only) runs its RX-queue poller AND bench worker on core 19. Measure + # the sender core as cpu_tx and the SERVER RECEIVE core (19, not the idle + # client-RX core 18) as cpu_rx, so cpu_rx_pct reflects the true RoCE RC + # receiver cost. + CPU_MASTER=8; CPU_TX=17; CPU_RX=19 ;; # Single-frame UDP sizes (<= the ~8972 B MTU payload, so no IP fragmentation). # 65507 is intentionally excluded: it fragments into ~8 packets and, under From 1598b71ddae3d23c52dc26d1958ecd3ecbe7bd26 Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Fri, 10 Jul 2026 11:25:57 -0400 Subject: [PATCH 13/18] #15 - Sync bench usage strings and docs with the new workload flags Address Greptile doc-sync review on PR #223: - Bench usage strings listed only --workload-gemm-dim. Add the other workload flags each main actually parses: --workload-fft-len and --workload-sync-interval to all four, plus --workload-max-inflight to the rdma bench (RoCE event path). - raw_benchmarking.md: note that --workload-gemm-dim / --workload-fft-len pin the GEMM side length / FFT transform length so the per-call FLOP count is fixed across the I/O sweep. - run_spark_bench.sh: GEMM_DIM / FFT_LEN validation accepted 0 (^[0-9]+$) despite the "positive integer" message; tighten to ^[1-9][0-9]*$. (The review's claim that the docs still document --workload-batch-bytes is stale -- that flag was already removed; no references remain in the tree.) Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- docs/benchmarks/raw_benchmarking.md | 2 +- examples/raw_gpudirect_bench.cpp | 3 ++- examples/raw_hds_bench.cpp | 3 ++- examples/rdma_bench.cpp | 3 ++- examples/run_spark_bench.sh | 4 ++-- examples/socket_bench.cpp | 3 ++- 6 files changed, 11 insertions(+), 7 deletions(-) diff --git a/docs/benchmarks/raw_benchmarking.md b/docs/benchmarks/raw_benchmarking.md index fa10a2a..59c4c57 100644 --- a/docs/benchmarks/raw_benchmarking.md +++ b/docs/benchmarks/raw_benchmarking.md @@ -238,7 +238,7 @@ After having modified the configuration file, ensure you have connected an SFP c By default the application runs for 10 seconds and then exits. You can change the duration by passing `--seconds ` after the YAML path, or stop it gracefully at any time with `Ctrl-C`. -`daqiri_bench_raw_gpudirect` and `daqiri_bench_raw_hds` also accept `--workload none|fft|gemm|gemm_fp16`, which runs a representative GPU workload once per received reorder window on the **actual received packet data**: `fft` (batched cuFFT C2C transform), `gemm` (FP32 `cublasSgemm`), or `gemm_fp16` (the same-size mixed-precision FP16/tensor-core matmul that models inference). Each received burst's payloads are first reordered by sequence number into a contiguous GPU buffer (`examples/bench_pipeline.{h,cu}`) that the compute then consumes. The same `--workload` flag is honoured by the RoCE bench (`daqiri_bench_rdma`, in-order gather) and the socket bench (`daqiri_bench_socket`, host→device stage then UDP reorder / TCP gather). See the [DGX Spark GPU-workload results](performance-dgx-spark.md#gpu-workloads-in-the-receive-path). +`daqiri_bench_raw_gpudirect` and `daqiri_bench_raw_hds` also accept `--workload none|fft|gemm|gemm_fp16`, which runs a representative GPU workload once per received reorder window on the **actual received packet data**: `fft` (batched cuFFT C2C transform), `gemm` (FP32 `cublasSgemm`), or `gemm_fp16` (the same-size mixed-precision FP16/tensor-core matmul that models inference). Each received burst's payloads are first reordered by sequence number into a contiguous GPU buffer (`examples/bench_pipeline.{h,cu}`) that the compute then consumes. `--workload-gemm-dim N` (default 1024) pins the square GEMM side length and `--workload-fft-len N` (default 1024) the 1-D FFT transform length, so the FLOP count per call stays constant as the I/O unit is swept. The same flags are honoured by the RoCE bench (`daqiri_bench_rdma`, in-order gather) and the socket bench (`daqiri_bench_socket`, host→device stage then UDP reorder / TCP gather). See the [DGX Spark GPU-workload results](performance-dgx-spark.md#gpu-workloads-in-the-receive-path). ## Flow programming smoke test diff --git a/examples/raw_gpudirect_bench.cpp b/examples/raw_gpudirect_bench.cpp index c536f41..c6eda0b 100644 --- a/examples/raw_gpudirect_bench.cpp +++ b/examples/raw_gpudirect_bench.cpp @@ -163,7 +163,8 @@ int main(int argc, char **argv) { if (argc < 2) { std::cerr << "Usage: " << argv[0] << " [--seconds N] [--target-gbps G] " - "[--workload none|fft|gemm|gemm_fp16] [--workload-gemm-dim N]\n"; + "[--workload none|fft|gemm|gemm_fp16] [--workload-gemm-dim N] " + "[--workload-fft-len N] [--workload-sync-interval N]\n"; return 1; } diff --git a/examples/raw_hds_bench.cpp b/examples/raw_hds_bench.cpp index 1706157..b9e4c6a 100644 --- a/examples/raw_hds_bench.cpp +++ b/examples/raw_hds_bench.cpp @@ -141,7 +141,8 @@ int main(int argc, char **argv) { if (argc < 2) { std::cerr << "Usage: " << argv[0] << " [--seconds N] " - "[--workload none|fft|gemm|gemm_fp16] [--workload-gemm-dim N]\n"; + "[--workload none|fft|gemm|gemm_fp16] [--workload-gemm-dim N] " + "[--workload-fft-len N] [--workload-sync-interval N]\n"; return 1; } diff --git a/examples/rdma_bench.cpp b/examples/rdma_bench.cpp index a0d6cea..0b6dcce 100644 --- a/examples/rdma_bench.cpp +++ b/examples/rdma_bench.cpp @@ -353,7 +353,8 @@ int main(int argc, char** argv) { std::cerr << "Usage: " << argv[0] << " [--seconds N] [--mode server|client|both] " "[--target-gbps G] [--workload none|fft|gemm|gemm_fp16] " - "[--workload-gemm-dim N]\n"; + "[--workload-gemm-dim N] [--workload-fft-len N] " + "[--workload-sync-interval N] [--workload-max-inflight N]\n"; return 1; } diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index e172b0f..d88f4ba 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -110,13 +110,13 @@ esac # so that unit must be at least that large. Default 1024. Recorded in the CSV # post_process_gemm_dim column. GEMM_DIM="${GEMM_DIM:-1024}" -if [[ ! "$GEMM_DIM" =~ ^[0-9]+$ ]]; then +if [[ ! "$GEMM_DIM" =~ ^[1-9][0-9]*$ ]]; then echo "Invalid GEMM_DIM '$GEMM_DIM' (expected a positive integer)" >&2; exit 1 fi # FFT_LEN: the 1-D C2C transform length (--workload-fft-len) for WORKLOAD=fft, held # FIXED while the I/O unit is swept. Independent of GEMM_DIM. Default 1024. FFT_LEN="${FFT_LEN:-1024}" -if [[ ! "$FFT_LEN" =~ ^[0-9]+$ ]]; then +if [[ ! "$FFT_LEN" =~ ^[1-9][0-9]*$ ]]; then echo "Invalid FFT_LEN '$FFT_LEN' (expected a positive integer)" >&2; exit 1 fi # SYNC_INTERVAL: drain the GPU stream every N compute calls (--workload-sync-interval). diff --git a/examples/socket_bench.cpp b/examples/socket_bench.cpp index 948307b..1812b62 100644 --- a/examples/socket_bench.cpp +++ b/examples/socket_bench.cpp @@ -241,7 +241,8 @@ int main(int argc, char** argv) { std::cerr << "Usage: " << argv[0] << " [--seconds N] [--mode server|client|both] " "[--target-gbps G] [--workload none|fft|gemm|gemm_fp16] " - "[--workload-gemm-dim N]\n"; + "[--workload-gemm-dim N] [--workload-fft-len N] " + "[--workload-sync-interval N]\n"; return 1; } From dd71cae98e0872be92cdbe0cbd0dbc0c8d876c02 Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Fri, 10 Jul 2026 16:48:45 -0400 Subject: [PATCH 14/18] #15 - Address PR #223 review: drop FFT flop estimate, clarify CPU notes - examples/bench_workload.h: the --workload-fft-len comment gave the FFT cost as ~5*N*log2(N) flops -- the radix-2 estimate, valid only for power-of-two N; cuFFT uses different algorithms (and costs) for other lengths. Drop the equation so the comment holds for any N (r3560531760). - performance-dgx-spark.md: "rate-independent" was undefined jargon (r3560468420). Reword the RoCE TX and DPDK poller CPU notes to say the busy% comes from the post-and-poll / poll-mode spin loop rather than the throughput, and add a sentence spelling out that the core stays near this level whether the link runs at 10 or 100 Gb/s. Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- docs/benchmarks/performance-dgx-spark.md | 17 ++++++++++------- examples/bench_workload.h | 4 ++-- 2 files changed, 12 insertions(+), 9 deletions(-) diff --git a/docs/benchmarks/performance-dgx-spark.md b/docs/benchmarks/performance-dgx-spark.md index 1df3677..e4a51d3 100644 --- a/docs/benchmarks/performance-dgx-spark.md +++ b/docs/benchmarks/performance-dgx-spark.md @@ -90,8 +90,8 @@ also the no-drop rate: pacing the sender below it hits the target with zero drop | Core | Busy% | Note | | ------------------------ | ----: | ------------------------------------- | | Master (CPU 8) | 3.7% | Orchestration only; mostly idle | -| TX queue poller (CPU 17) | ~92% | Poll-mode spin; rate-independent | -| RX queue poller (CPU 18) | ~92% | Poll-mode spin; rate-independent | +| TX queue poller (CPU 17) | ~92% | Poll-mode busy-spin; ~constant vs load | +| RX queue poller (CPU 18) | ~92% | Poll-mode busy-spin; ~constant vs load | The benchmark app workers run on their own cores (TX 16, RX 19) alongside these pollers; this run sampled only the poller cores. @@ -171,13 +171,16 @@ sized to the message so the in-flight window stays full. | Core | Busy% | Note | | -------------------- | ----: | ----------------------------------------------- | | Master (CPU 8) | 0.7% | Orchestration only | -| Client TX (CPU 17) | 74.8% | Post-and-poll spin; rate-independent | +| Client TX (CPU 17) | 74.8% | Busy-spins posting sends and polling completions | | Server RX (CPU 19) | 1.1% | HCA DMAs straight to memory; worker only reaps completions | -The near-idle RX core is the expected RoCE RC signature — the HCA places incoming -data directly into registered memory, so the receive worker only reaps -completions and reposts (~1% at this message rate). The GPU stays idle here too -(SM and memory-controller ~0%; DMA target, not a compute engine). +The TX core busy-spins in a post-and-poll loop, so its ~75% busy time is set by +that spin, not by the throughput: it stays near this level whether the link runs +at 10 or 100 Gb/s (the same reason the DPDK pollers sit near 92% regardless of +offered load). The near-idle RX core is the expected RoCE RC signature — the HCA +places incoming data directly into registered memory, so the receive worker only +reaps completions and reposts (~1% at this message rate). The GPU stays idle here +too (SM and memory-controller ~0%; DMA target, not a compute engine). ## Socket / TCP diff --git a/examples/bench_workload.h b/examples/bench_workload.h index 7bcdeeb..0eb18b6 100644 --- a/examples/bench_workload.h +++ b/examples/bench_workload.h @@ -43,8 +43,8 @@ int parse_workload_gemm_dim(int argc, char** argv); // Parse "--workload-fft-len N" from argv: the 1-D C2C transform length. The burst // working set is fanned out across as many batched length-N transforms as fit, so -// N sets the per-transform cost (~5*N*log2(N) flops) while the batch count tracks -// the I/O unit. Independent of --workload-gemm-dim. Returns 1024 (the default) if +// N sets the per-transform cost while the batch count tracks the I/O unit. +// Independent of --workload-gemm-dim. Returns 1024 (the default) if // unset. Mirrors parse_workload's stride. int parse_workload_fft_len(int argc, char** argv); From c662a9fb8f11d4b9372ef270c5517b06a2ca563c Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Fri, 10 Jul 2026 16:48:55 -0400 Subject: [PATCH 15/18] #15 - Self-certify RoCE wire transit and harden the phy-delta check The DPDK bench path already asserted per-cell wire transit via *_phy SerDes counters; the RoCE path did not, so an on-chip eswitch loopback could inflate a RoCE number above the 100GbE line rate undetected. Add the same check to the RoCE path: a netns-aware phy_counter reads the server rx_packets_phy and client tx_packets_phy inside the dq_wire_{server,client} loopback namespaces around each cell. Also harden the threshold on all three checks (RoCE, DPDK, multi-queue). A bare ">0" let a stray background/control packet pass an on-chip run as "wire OK" -- a +22 phy delta slipped through on a 23M-message run. Require the phy delta to reach ~half the message/packet count before certifying wire transit. Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- examples/run_spark_bench.sh | 63 +++++++++++++++++++++++++++++----- examples/run_spark_mq_bench.sh | 12 ++++--- 2 files changed, 62 insertions(+), 13 deletions(-) diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index d88f4ba..88a08bb 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -207,6 +207,17 @@ case "$BACKEND" in # client-RX core 18) as cpu_rx, so cpu_rx_pct reflects the true RoCE RC # receiver cost. CPU_MASTER=8; CPU_TX=17; CPU_RX=19 + # Resolve the server (RX) / client (TX) netdevs inside the wire-loopback + # namespaces so each cell can assert wire transit via *_phy SerDes counters, + # exactly like the dpdk path. RoCE loops over the SAME cable, so a non-advancing + # server rx_packets_phy flags the on-chip eswitch short-cut instead of a true + # over-the-cable loopback -- the tell for a RoCE number above the ~99 Gb/s + # 100GbE line-rate ceiling. Empty if the netns is not up yet (the per-cell check + # then just skips with a WARN). Override RDMA_{SERVER,CLIENT}_NETDEV if needed. + RDMA_SERVER_NS="${RDMA_SERVER_NS:-dq_wire_server}" + RDMA_CLIENT_NS="${RDMA_CLIENT_NS:-dq_wire_client}" + RDMA_SERVER_NETDEV="${RDMA_SERVER_NETDEV:-$(ip netns exec "$RDMA_SERVER_NS" ls /sys/class/net 2>/dev/null | grep -vx lo | head -n1 || true)}" + RDMA_CLIENT_NETDEV="${RDMA_CLIENT_NETDEV:-$(ip netns exec "$RDMA_CLIENT_NS" ls /sys/class/net 2>/dev/null | grep -vx lo | head -n1 || true)}" ;; # Single-frame UDP sizes (<= the ~8972 B MTU payload, so no IP fragmentation). # 65507 is intentionally excluded: it fragments into ~8 packets and, under @@ -355,10 +366,13 @@ cpu_busy_pct() { # Sum a NIC *_phy SerDes counter (proves traffic crossed the cable, not an on-chip # eswitch short-cut). Empty netdev -> 0. Mirrors run_spark_mq_bench.sh's phy check. +# Optional third arg is a network namespace (the RoCE/socket wire loopback moves the +# cabled netdevs into dq_wire_{server,client}); empty reads the current namespace. phy_counter() { - local netdev="$1" key="$2" + local netdev="$1" key="$2" ns="${3:-}" [[ -z "$netdev" ]] && { echo 0; return; } - ethtool -S "$netdev" 2>/dev/null \ + { if [[ -n "$ns" ]]; then ip netns exec "$ns" ethtool -S "$netdev" 2>/dev/null + else ethtool -S "$netdev" 2>/dev/null; fi; } \ | awk -F'[: ]+' -v k="$key" '$2 == k { s += $3 } END { printf "%d", s+0 }' } @@ -552,6 +566,12 @@ run_cell() { nsys_pre=("$NSYS_BIN" profile --trace=cuda,osrt --sample=cpu --cpuctxsw=process-tree --force-overwrite=true --output "$cell_dir/roce_server") fi + # Snapshot the netns *_phy counters around the run to assert the RoCE traffic + # actually crossed the cable (client tx -> server rx over the wire), not the + # on-chip eswitch short-cut that lets a loopback exceed the 100GbE line rate. + local phy_tx_before phy_rx_before + phy_tx_before="$(phy_counter "$RDMA_CLIENT_NETDEV" tx_packets_phy "$RDMA_CLIENT_NS")" + phy_rx_before="$(phy_counter "$RDMA_SERVER_NETDEV" rx_packets_phy "$RDMA_SERVER_NS")" ip netns exec dq_wire_server "${nsys_pre[@]}" "$BENCH_BIN" "$yaml" \ --seconds "$server_seconds" "${bench_extra[@]}" --mode server \ > "$cell_dir/server_stdout.txt" 2> "$cell_dir/server_stderr.txt" & @@ -572,11 +592,30 @@ run_cell() { "$cell_dir/roce_server.nsys-rep" >&2 || true echo "nsys report: $cell_dir/roce_server.nsys-rep" >&2 fi + local phy_tx_delta phy_rx_delta + phy_tx_delta=$(( $(phy_counter "$RDMA_CLIENT_NETDEV" tx_packets_phy "$RDMA_CLIENT_NS") - phy_tx_before )) + phy_rx_delta=$(( $(phy_counter "$RDMA_SERVER_NETDEV" rx_packets_phy "$RDMA_SERVER_NS") - phy_rx_before )) # RDMA prints "Client complete: ... send_completions=N send_bytes=N seconds=S". pkts="$(extract_field 'Client complete' send_completions "$stdout")" bytes="$(extract_field 'Client complete' send_bytes "$stdout")" secs="$(extract_field 'Client complete' seconds "$stdout")" rx_bytes="$bytes" + # Wire-transit check. On the cable the server's rx_*_phy advances by at least one + # SerDes packet per RDMA message -- many more once a message exceeds the MTU and + # segments -- so a genuine over-the-wire run shows phy_rx_delta >= the message + # count. An on-chip eswitch short-cut leaves rx_*_phy essentially FLAT (only a + # handful of stray background/control packets), so a bare ">0" is not enough: a + # +22 on a 23M-message run passed the old check while never touching the cable. + # Require the delta to reach half the message count (generous margin for counter + # slack) before certifying wire transit. + local phy_min=$(( ${pkts:-0} / 2 )) + if [[ -z "$RDMA_SERVER_NETDEV" ]]; then + echo "WARN: $cell could not resolve server netdev in $RDMA_SERVER_NS (netns up?); skipped wire (*_phy) check" >&2 + elif [[ "$phy_rx_delta" -lt "$phy_min" || "$phy_rx_delta" -le 0 ]]; then + echo "WARN: $cell rx_packets_phy advanced only +$phy_rx_delta on $RDMA_SERVER_NETDEV (expected >= ~$phy_min, one per message) -- traffic likely did NOT cross the wire (on-chip eswitch short-cut?)" >&2 + else + echo "INFO: $cell wire OK -- client tx_packets_phy +$phy_tx_delta, server rx_packets_phy +$phy_rx_delta (>= $phy_min msgs)" >&2 + fi else local yaml="$cell_dir/config.yaml" generate_yaml "$yaml" "$payload" "$batch" @@ -590,13 +629,6 @@ run_cell() { local phy_tx_delta phy_rx_delta phy_tx_delta=$(( $(phy_counter "$DPDK_TX_NETDEV" tx_packets_phy) - phy_tx_before )) phy_rx_delta=$(( $(phy_counter "$DPDK_RX_NETDEV" rx_packets_phy) - phy_rx_before )) - if [[ -z "$DPDK_RX_NETDEV" ]]; then - echo "WARN: $cell could not resolve rx_port netdev ($DPDK_RX_PCI); skipped wire (*_phy) check" >&2 - elif [[ "$phy_rx_delta" -le 0 ]]; then - echo "WARN: $cell rx_packets_phy did not advance ($phy_rx_delta) on $DPDK_RX_NETDEV -- traffic may NOT have crossed the wire (on-chip eswitch short-cut?)" >&2 - else - echo "INFO: $cell wire OK -- p0 tx_packets_phy +$phy_tx_delta, p1 rx_packets_phy +$phy_rx_delta" >&2 - fi # For RX-bearing benches "RX complete" is authoritative; fall back to "TX complete". pkts="$(extract_field 'RX complete' packets "$stdout")" bytes="$(extract_field 'RX complete' bytes "$stdout")" @@ -607,6 +639,19 @@ run_cell() { secs="$(extract_field 'TX complete' seconds "$stdout")" fi rx_bytes="$bytes" + # Wire-transit check. Raw Ethernet is one SerDes packet per app packet, so a + # genuine over-the-wire run advances rx_*_phy by ~the received packet count; an + # on-chip eswitch short-cut leaves it near-flat. Require the delta to reach half + # the packet count -- a bare ">0" would let a stray background packet false-pass + # (the same hole the RoCE check had). + local phy_min=$(( ${pkts:-0} / 2 )) + if [[ -z "$DPDK_RX_NETDEV" ]]; then + echo "WARN: $cell could not resolve rx_port netdev ($DPDK_RX_PCI); skipped wire (*_phy) check" >&2 + elif [[ "$phy_rx_delta" -lt "$phy_min" || "$phy_rx_delta" -le 0 ]]; then + echo "WARN: $cell rx_packets_phy advanced only +$phy_rx_delta on $DPDK_RX_NETDEV (expected >= ~$phy_min) -- traffic likely did NOT cross the wire (on-chip eswitch short-cut?)" >&2 + else + echo "INFO: $cell wire OK -- p0 tx_packets_phy +$phy_tx_delta, p1 rx_packets_phy +$phy_rx_delta (>= $phy_min)" >&2 + fi fi cp "$stderr" "$DRIVER_LOG" diff --git a/examples/run_spark_mq_bench.sh b/examples/run_spark_mq_bench.sh index 4787678..920d250 100755 --- a/examples/run_spark_mq_bench.sh +++ b/examples/run_spark_mq_bench.sh @@ -265,14 +265,18 @@ run_cell() { local drops; drops="$(parse_dpdk_drops "$stderr")" - # Wire-traffic confirmation: rx_packets_phy must advance over the run. + # Wire-traffic confirmation: rx_packets_phy must advance by ~the received packet + # count (raw Ethernet is one SerDes packet per app packet). A bare ">0" would let a + # stray background packet false-pass an on-chip eswitch short-cut, so require the + # delta to reach half the packet count. local phy_delta=$(( phy_after - phy_before )) + local phy_min=$(( ${pkts:-0} / 2 )) if [[ -z "$RX_NETDEV" ]]; then echo "WARN: $cell p$payload could not resolve RX netdev for $RX_PCI; skipped wire (*_phy) check" >&2 - elif [[ "$phy_delta" -le 0 ]]; then - echo "WARN: $cell p$payload rx_packets_phy did not advance ($phy_delta) -- traffic may not have crossed the wire" >&2 + elif [[ "$phy_delta" -lt "$phy_min" || "$phy_delta" -le 0 ]]; then + echo "WARN: $cell p$payload rx_packets_phy advanced only +$phy_delta (expected >= ~$phy_min) -- traffic likely did NOT cross the wire (on-chip eswitch short-cut?)" >&2 else - echo "INFO: $cell p$payload wire OK -- rx_packets_phy +$phy_delta" >&2 + echo "INFO: $cell p$payload wire OK -- rx_packets_phy +$phy_delta (>= $phy_min)" >&2 fi # Per-core busy% over the bench window, in CSV column order (see CPU_CORES). From 654c70f7758cc2322aaa72279eca0f0149c3f3c5 Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Sun, 12 Jul 2026 22:28:06 -0400 Subject: [PATCH 16/18] #15 - Fix socket port collision, mq lib path, and mq plot rep-averaging Three bench-harness fixes surfaced while re-measuring the DGX Spark sockets and multi-queue cells over the wire: - run_spark_bench.sh: generate_socket_yaml's local_addr/remote_addr port rewrite required a quoted "tcp://..." URI, but gen_spark_netns_config.py emits it unquoted (PyYAML strips the quotes), so the substitution silently missed. Every multi-pair TCP/UDP cell then reused one port -- both servers bound the same port -> 'Address already in use', a dead pair, and a wedged sweep. Make the quote optional in all three URI sed patterns. (Overlaps the fix/socket-netns-rx work.) - run_spark_mq_bench.sh: LD_LIBRARY_PATH pointed at $BUILD_DIR, but the per-engine sub-libs (libdaqiri_dpdk.so.0 etc.) live in $BUILD_DIR/src, so the bench failed with 'cannot open libdaqiri_dpdk.so.0'. Put $BUILD_DIR/src first. - plot_mq_payload_sweep.py: predated REPEATS>1 and plotted every rep as its own point; average reps per (cell, payload) so the line has one point per x. Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- examples/run_spark_bench.sh | 6 +++--- examples/run_spark_mq_bench.sh | 8 +++++--- scripts/plot_mq_payload_sweep.py | 24 +++++++++++++++--------- 3 files changed, 23 insertions(+), 15 deletions(-) diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index 88a08bb..5b6b2f0 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -450,15 +450,15 @@ generate_socket_yaml() { python3 "$NETNS_GEN" "$BASE_YAML" --role server | \ sed -E \ -e "s|^( *message_size: ).*|\1$payload|g" \ - -e "s|^( *local_addr: \"[a-z]+://[0-9.]+:)[0-9]+(\")|\1$srv_port\2|" \ + -e "s|^( *local_addr: \"?[a-z]+://[0-9.]+:)[0-9]+(\"?)|\1$srv_port\2|" \ -e "s|^( *server_port: ).*|\1$srv_port|" \ -e "s|^( *cpu_core: ).*|\1$core|" \ > "$server_out" python3 "$NETNS_GEN" "$BASE_YAML" --role client | \ sed -E \ -e "s|^( *message_size: ).*|\1$payload|g" \ - -e "s|^( *local_addr: \"[a-z]+://[0-9.]+:)[0-9]+(\")|\1$cli_port\2|" \ - -e "s|^( *remote_addr: \"[a-z]+://[0-9.]+:)[0-9]+(\")|\1$srv_port\2|" \ + -e "s|^( *local_addr: \"?[a-z]+://[0-9.]+:)[0-9]+(\"?)|\1$cli_port\2|" \ + -e "s|^( *remote_addr: \"?[a-z]+://[0-9.]+:)[0-9]+(\"?)|\1$srv_port\2|" \ -e "s|^( *server_port: ).*|\1$srv_port|" \ -e "s|^( *cpu_core: ).*|\1$core|" \ > "$client_out" diff --git a/examples/run_spark_mq_bench.sh b/examples/run_spark_mq_bench.sh index 920d250..88a61f9 100755 --- a/examples/run_spark_mq_bench.sh +++ b/examples/run_spark_mq_bench.sh @@ -83,9 +83,11 @@ REPEATS="${REPEATS:-1}" MQ_BASE="$SCRIPT_DIR/daqiri_bench_raw_tx_rx_spark_mq.yaml" MQ_GEN="$SCRIPT_DIR/../scripts/gen_spark_mq_config.py" -# Match run_spark_bench.sh: prefer the installed shared libs, falling back to -# the build tree. Keep both so a fresh build dir still resolves. -export LD_LIBRARY_PATH="/opt/daqiri/lib:$BUILD_DIR:${LD_LIBRARY_PATH:-}" +# Resolve the DAQIRI shared libs from the build tree first. The per-engine +# sub-libraries (libdaqiri_dpdk.so.0 etc.) live in $BUILD_DIR/src, not $BUILD_DIR, +# so that directory must be on the path or the bench fails to load them; keep +# $BUILD_DIR and the installed /opt/daqiri/lib as fallbacks. +export LD_LIBRARY_PATH="$BUILD_DIR/src:$BUILD_DIR:/opt/daqiri/lib:${LD_LIBRARY_PATH:-}" TS="$(date -u +%Y%m%dT%H%M%SZ)" OUT_DIR="$SCRIPT_DIR/../bench-results/$TS-dpdk-mq" diff --git a/scripts/plot_mq_payload_sweep.py b/scripts/plot_mq_payload_sweep.py index 173c3fb..1ecbf4f 100644 --- a/scripts/plot_mq_payload_sweep.py +++ b/scripts/plot_mq_payload_sweep.py @@ -4,8 +4,9 @@ """Render the DGX Spark DPDK multi-queue payload sweep as a line plot. Reads a runs.csv produced by examples/run_spark_mq_bench.sh (columns: -cell,tx_cores,rx_cores,payload,gbps,pps,drops,cpu8,cpu16,cpu17,cpu18,cpu19) -and plots achieved Gb/s vs payload size, one line per (TX,RX) core cell. +cell,tx_cores,rx_cores,payload,rep,gbps,pps,drops,cpu...) and plots achieved +Gb/s vs payload size, one line per (TX,RX) core cell. Multiple reps per +(cell, payload) are averaged. Usage: scripts/plot_mq_payload_sweep.py [output.svg] @@ -36,16 +37,21 @@ def load(csv_path): - """cell -> list of (payload_bytes, gbps), sorted by payload.""" - series = OrderedDict((c, []) for c in CELL_ORDER) + """cell -> list of (payload_bytes, mean gbps over reps), sorted by payload.""" + from collections import defaultdict + + acc = OrderedDict((c, defaultdict(list)) for c in CELL_ORDER) with open(csv_path, newline="") as fh: for row in csv.DictReader(fh): cell = row["cell"].strip() - if cell not in series: - series[cell] = [] # tolerate unexpected cells - series[cell].append((int(row["payload"]), float(row["gbps"]))) - for cell in series: - series[cell].sort(key=lambda pt: pt[0]) + if cell not in acc: + acc[cell] = defaultdict(list) # tolerate unexpected cells + acc[cell][int(row["payload"])].append(float(row["gbps"])) + series = OrderedDict() + for cell, by_payload in acc.items(): + series[cell] = sorted( + (payload, sum(vals) / len(vals)) for payload, vals in by_payload.items() + ) return series From 9660bc178bab48c095a5d2a56c8017ef41c0407e Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Sun, 12 Jul 2026 22:28:19 -0400 Subject: [PATCH 17/18] #15 - Refresh DGX Spark perf report to phy-verified 100 GbE numbers The loopback cable is 100 GbE, so the report's >100 Gb/s figures could not have crossed the wire -- they were on-chip eswitch hairpin (and/or a since-swapped 200 Gbps cable). Re-measured every cell over the phy-verified p0<->p1 wire loop (rx_packets_phy confirmed to advance) and corrected the numbers: - RoCE SEND: 106.0 -> 97.6 peak (8MB/1MB/64KB all ~97; small cells unchanged). - DPDK sweep: 105.5@4KB -> ~98.8 peak; 1KB 92 -> 97 (stale low); best case now ~98.8 @ 4-8KB. - Multi-queue: the '2nd TX core scales to 108.8' story was pure on-chip -- over the wire one queue already saturates (~97-99), and cores only help small, packet-rate-bound payloads via RX (64B 20->27, 256B 50->66). Section rewritten, payload-sweep plot re-rendered. - GPU workload: RoCE none 108.0 -> 96.6 (now ~= DPDK 98.7); table + prose redone. - Sockets (re-run after the port-collision fix): TCP 4-pair 97.2 -> 90.6, UDP 4-pair 29.8 -> 28.5. - SUT table now states the 100 GbE link rate. Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- docs/benchmarks/performance-dgx-spark.md | 134 ++-- docs/images/spark-mq-payload-sweep.svg | 812 +++++++++++------------ 2 files changed, 477 insertions(+), 469 deletions(-) diff --git a/docs/benchmarks/performance-dgx-spark.md b/docs/benchmarks/performance-dgx-spark.md index e4a51d3..70fbf3a 100644 --- a/docs/benchmarks/performance-dgx-spark.md +++ b/docs/benchmarks/performance-dgx-spark.md @@ -24,7 +24,7 @@ loopback). The exact commands are collected under [Reproduce](#reproduce) below. | Component | Detail | | --------- | ------ | | Platform | DGX Spark (GB10), 20 cores, isolcpus `16-19` (the multi-queue sweep expands this; see [Multi-queue core scaling](#multi-queue-core-scaling)) | -| NIC | ConnectX-7, ports p0 ↔ p1 cross-cabled (single-host loopback), MTU 9000 | +| NIC | ConnectX-7, ports p0 ↔ p1 cross-cabled with a **100 GbE QSFP28** loopback cable (single-host loopback), MTU 9000 | | Build | Release (`-DCMAKE_BUILD_TYPE=Release`), `DAQIRI_ENGINE="dpdk ibverbs"` | | Loopback | Raw/DPDK uses the two physical ports directly; socket/RoCE use the `dq_wire_*` network-namespace wire loopback | | Core pinning | Each direction has a busy-spin queue poller and an app worker on separate isolated X925 cores (PR #149). Single-queue: DPDK pollers 17/18, workers 16/19. Multi-queue: TX pollers 16/19, RX pollers 18/9, each with its own worker core, master 8 (with `isolcpus=5-9,15-19`). | @@ -35,12 +35,17 @@ Each transport at its best-case **operation size**. Raw/RoCE are single-stream; socket TCP/UDP scale with the number of client/server pairs, so the four-pair aggregate is shown. +The **100 GbE QSFP28 loopback cable sets the maximum data rate** here — not the +ConnectX-7 (which is rated for higher line rates) or the software path. A 100 GbE +link tops out near ~99.6 Gb/s of payload, so every large-transfer result saturates +just under that ceiling. + | Stream / Protocol | Best case | Throughput | Drops | Notes | | ----------------- | --------- | ---------: | ----- | ----- | -| Raw Ethernet / GPUDirect | 4 KB packet | **105.5 ±0.9 Gb/s** | 0 | 98.5 Gb/s single-queue at 8 KB | -| Socket / RoCE (SEND) | 64 KB message | **106.0 ±0.1 Gb/s** | 0 | Single QP, batch 1 | -| Socket / TCP | 8 KB × 4 pairs | **97.2 ±2.8 Gb/s** | ~0 | Flow-controlled (App TX = App RX) | -| Socket / UDP | 8 KB × 4 pairs | **29.8 ±0.2 Gb/s** | ~51% loss | Receiver goodput; unpaced sender | +| Raw Ethernet / GPUDirect | 8 KB packet | **98.8 ±0.1 Gb/s** | 0 | Flat ~98.7 across 4–8 KB, all batch sizes | +| Socket / RoCE (SEND) | 64 KB message | **97.6 ±0.1 Gb/s** | 0 | Single QP, batch 1 | +| Socket / TCP | 1 MiB × 4 pairs | **90.6 ±1.5 Gb/s** | ~0 | Flow-controlled (App TX = App RX) | +| Socket / UDP | 8 KB × 4 pairs | **28.5 ±0.6 Gb/s** | ~54% loss | Receiver goodput; unpaced sender | Each transport is best read at its own best-case operation size (see the per-transport tables below); a single cross-transport unit of work isn't meaningful here, since @@ -49,14 +54,14 @@ operation boundary. ## Raw Ethernet / GPUDirect (DPDK) -Physical port-to-port loopback, GPU-resident payloads. Throughput peaks at -**105.5 Gb/s** at a 4 KB payload; 8 KB packets hold **98.5 Gb/s** drop-free -across all batch sizes. Packet handling is CPU-bound (see the CPU -utilization table below). Throughput is flat across batch size and stable -run-to-run (3 reps per cell, ≤1% spread). +Physical port-to-port loopback, GPU-resident payloads. Throughput saturates the +100 GbE line at **~98.8 Gb/s** for 4–8 KB payloads, drop-free across all batch +sizes. Packet handling is CPU-bound (see the CPU utilization table below). +Throughput is flat across batch size and stable run-to-run (3 reps per cell, +≤1% spread). Achieved Gb/s measured at App RX (equal to App TX, since every cell is -drop-free), unpaced, mean of 3 reps; run-to-run spread ≤0.9 Gb/s (<1%): +drop-free), unpaced, mean of 3 reps; run-to-run spread ≤0.5 Gb/s (<1%): @@ -69,17 +74,17 @@ drop-free), unpaced, mean of 3 reps; run-to-run spread ≤0.9 Gb/s (<1%): - - - - - + + + + +
8000 B98.598.098.098.5
4096 B105.2105.3105.5105.1
1024 B92.191.891.791.7
256 B50.049.849.849.8
64 B20.420.520.520.4
8000 B98.898.898.898.7
4096 B98.698.898.798.6
1024 B97.197.297.297.1
256 B49.749.649.649.5
64 B20.220.220.420.2
-At ≥4 KB the link saturates (~98–105 Gb/s) regardless of batch. Below that the -path is packet-rate-bound: 1 KB ~92 Gb/s (10.5 M pps), 256 B ~50 Gb/s (19.5 M pps), -64 B ~20 Gb/s (20 M pps) — a ~20 M pps single-queue ceiling (the multi-queue +At ≥1 KB the link saturates near line rate (~97–99 Gb/s) regardless of batch. +Below that the path is packet-rate-bound: 256 B ~50 Gb/s (19.5 M pps), 64 B +~20 Gb/s (20 M pps) — a ~20 M pps single-queue ceiling (the multi-queue section lifts it). Gb/s here is the L2 frame rate including the 64 B header, so pps ≈ Gb/s ÷ ((payload + 64) × 8). These small-payload cells are flat across batch size and stable run-to-run. Because every cell is drop-free, the achieved rate is @@ -90,8 +95,8 @@ also the no-drop rate: pacing the sender below it hits the target with zero drop | Core | Busy% | Note | | ------------------------ | ----: | ------------------------------------- | | Master (CPU 8) | 3.7% | Orchestration only; mostly idle | -| TX queue poller (CPU 17) | ~92% | Poll-mode busy-spin; ~constant vs load | -| RX queue poller (CPU 18) | ~92% | Poll-mode busy-spin; ~constant vs load | +| TX queue poller (CPU 17) | ~92% | Poll-mode busy-spin | +| RX queue poller (CPU 18) | ~92% | Poll-mode busy-spin | The benchmark app workers run on their own cores (TX 16, RX 19) alongside these pollers; this run sampled only the poller cores. @@ -102,10 +107,10 @@ not a compute engine. ### Multi-queue core scaling -Each packet-handling core spins in poll-mode. At 8 KB a single -TX core caps throughput near ~98 Gb/s while a single RX core already drains the -line, so the second TX core is the lever: it scales `(1,1)` to `(2,1)` from -97.5 to 108.8 Gb/s, while a second RX core adds little. The matrix sweeps +Each packet-handling core spins in poll-mode. At large payloads (≥1 KB) a single +queue already saturates the 100 GbE line (~97–99 Gb/s), so extra cores add +nothing there — the multi-queue win is confined to the small, +packet-rate-bound payloads, where **RX cores** are the lever. The matrix sweeps (TX cores, RX cores) over `(1,1)`, `(1,2)`, `(2,1)`, `(2,2)`. Each queue is served by a poll-mode driver core plus a separate bench-worker @@ -117,26 +122,29 @@ base `daqiri_bench_raw_tx_rx_spark_mq.yaml` (the balanced 2,2 superset) by `scripts/gen_spark_mq_config.py`; generated by `examples/run_spark_mq_bench.sh`, 30 s per cell, 0 drops. +Achieved Gb/s at a **256 B payload** (the packet-rate-bound regime where core +count matters); at ≥1 KB every cell converges at the wire ceiling regardless: + | Cell | TX pollers | RX pollers | Achieved Gb/s | | ---- | ---------- | ---------- | ------------: | -| (1,1) | 16 | 18 | 97.5 | -| (1,2) | 16 | 18,9 | 99.0 | -| (2,1) | 16,19 | 18 | **108.8** | -| (2,2) | 16,19 | 18,9 | **108.8** | +| (1,1) | 16 | 18 | 50.0 | +| (1,2) | 16 | 18,9 | **66.4** | +| (2,1) | 16,19 | 18 | 49.0 | +| (2,2) | 16,19 | 18,9 | 64.7 | -Which core is the bottleneck flips with payload size. Sweeping each cell from -64 B to 8 KB: +A second **RX** core lifts 256 B from 50.0 to 66.4 Gb/s; a second **TX** core does +nothing (49.0 ≈ 50.0). The full payload sweep shows why — the bottleneck depends +on payload size: ![DPDK multi-queue throughput vs UDP payload size on DGX Spark, one line per (TX,RX) core count](../images/spark-mq-payload-sweep.svg) At small payloads the path is packet-rate-bound, so **RX cores** are the lever: -at 64 B a second RX core lifts throughput from 20.3 to 26.8 Gb/s (~20 M → ~26 M -pps) while a second TX core does nothing. At large payloads it inverts to the -byte/line-rate-bound regime where **TX cores** are the lever: at 8 KB the second -TX core takes 97.5 to 108.8 Gb/s (the 8 KB result above) while a second -RX core adds nothing. The curves cross around 1–4 KB, where the link saturates -and all four cells converge near ~104–109 Gb/s. Every cell is drop-free. -Generated by `examples/run_spark_mq_bench.sh` (30 s per point) and +a second RX core lifts 64 B from 20.3 to 26.9 Gb/s (~20 M → ~27 M pps) and 256 B +from 50.0 to 66.4 Gb/s, while a second TX core does nothing. At large payloads a +single queue already saturates the wire, so all four cells converge near +~97–99 Gb/s at ≥1 KB and neither extra core helps. Every cell is drop-free. +Generated by +`examples/run_spark_mq_bench.sh` (30 s per point) and `scripts/plot_mq_payload_sweep.py`. ## Socket / RoCE @@ -146,20 +154,20 @@ is App RX goodput, equal to App TX with 0 drops. Large messages up to 64 KB saturate the wire; the smallest messages are bound by per-operation software overhead. -**Message-size sweep (single QP, batch 1, 0 drops).** Mean of 3 reps; run-to-run -spread ≤0.7 Gb/s (<1%) in every cell. +**Message-size sweep (single QP, batch 1, 0 drops).** Mean ± sample std over 3 +reps; run-to-run spread <1% in every cell. | Message size | Gb/s | | ------------ | ---: | -| 8 MB | 105.0 | -| 1 MB | 105.1 | -| 64 KB | **106.0** | -| 8 KB | 51.5 | -| 4 KB | 28.6 | - -Messages ≥64 KB hold ~105–106 Gb/s at the wire ceiling. Below that the path is -operation-rate-bound (per-operation software overhead, not a stall) rather than -wire-bound, and every cell is drop-free. At 8 KB (51.5 Gb/s) and 4 KB (28.6 Gb/s) +| 8 MB | 96.8 ±0.3 | +| 1 MB | 96.9 ±0.2 | +| 64 KB | **97.6 ±0.1** | +| 8 KB | 51.8 ±0.2 | +| 4 KB | 28.5 ±0.0 | + +Messages ≥64 KB hold ~97 Gb/s at the 100 GbE wire ceiling (line rate). Below that +the path is operation-rate-bound (per-operation software overhead, not a stall) +rather than wire-bound, and every cell is drop-free. At 8 KB (51.8 Gb/s) and 4 KB (28.5 Gb/s) a dedicated bench-worker core, separate from the RoCE engine thread, sustains the operation rate, as it does for small DPDK packets. A per-message flow-control window keeps enough operations in flight to amortize that overhead: it pre-posts @@ -202,9 +210,9 @@ Throughput in Gb/s (App TX = App RX), mean ± std over 3 reps: - 1000 B13.5±0.227.3±0.254.8±0.2 - 8000 B30.8±5.868.0±9.797.2±2.8 - 1 MiB31.6±0.158.7±0.393.7±1.3 + 1000 B13.7±0.427.2±0.155.0±0.4 + 8000 B25.3±8.646.4±4.287.2±1.6 + 1 MiB31.5±1.351.9±2.290.6±1.5 @@ -232,8 +240,8 @@ Each cell shows **receiver goodput in Gb/s** (mean ± std over 3 reps) with the - 1000 B3.6 ±0.011% loss7.7 ±0.016% loss13.5 ±0.16% loss - 8000 B9.6 ±0.556% loss20.9 ±1.357% loss29.8 ±0.251% loss + 1000 B4.5 ±0.215% loss8.2 ±0.114% loss14.1 ±0.315% loss + 8000 B9.8 ±1.764% loss18.9 ±0.464% loss28.5 ±0.654% loss @@ -304,17 +312,17 @@ at an **8 KB payload** (~8 MB reorder window, 1024 packets × 8000 B), matched t | Workload | DPDK (Raw / GPUDirect) | RoCE (RC) | | -------- | ---------------------: | --------: | -| none (baseline) | 98.4 ±0.2 | 108.0 ±0.1 | -| FFT | 94.4 ±0.3 | 107.2 ±0.2 | -| GEMM (FP32) | 95.4 ±0.3 | 103.4 ±0.1 | +| none (baseline) | 98.7 ±0.0 | 96.6 ±0.3 | +| FFT | 95.7 ±0.8 | 95.6 ±0.1 | +| GEMM (FP32) | 96.6 ±0.2 | 90.2 ±1.1 | -Throughput in Gb/s. The two `none` baselines differ (98.4 vs 108.0) because each transport -runs at its own operation size — DPDK's 8 KB packets carry more per-packet overhead than -RoCE's large messages — not because of the workload. +Throughput in Gb/s. Both `none` baselines sit at the ~97–99 Gb/s wire ceiling +(DPDK 98.7, RoCE 96.6), as expected for two line-rate transports. -**GPU compute barely dents line rate on either transport:** both the GEMM and FFT leave -GPU headroom (SM well under 100%), so throughput stays wire-limited, not compute-limited. -The reorder/gather step assembles each unit's payload into one contiguous GPU buffer. +**GPU compute dents line rate only modestly, and stays wire-limited, not +compute-limited** (SM well under 100% throughout). FFT is nearly free on both +(SM ~6–17%, ≤1 Gb/s off baseline). The reorder/gather step assembles +each unit's payload into one contiguous GPU buffer. ## Reproduce diff --git a/docs/images/spark-mq-payload-sweep.svg b/docs/images/spark-mq-payload-sweep.svg index 8d6b309..998761d 100644 --- a/docs/images/spark-mq-payload-sweep.svg +++ b/docs/images/spark-mq-payload-sweep.svg @@ -6,11 +6,11 @@ - 2026-06-11T23:20:48.401040 + 2026-07-10T18:39:19.745269 image/svg+xml - Matplotlib v3.10.9, https://matplotlib.org/ + Matplotlib v3.11.0, https://matplotlib.org/ @@ -30,35 +30,35 @@ z - - + - - + - + - - - - + + - + - + - + - - - - - + + + - + - + - + - - - - - - + + + + - + - + - + - - - - - + + + + - + - + - + - - - - - + + + + - + - - - - - + - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + - + - - + - - + + - + - + - - - + + + - + - + - - - + + + - + - + - - - + + + - + - + - - - + + + - + - + - - - - + + + + - + - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + - + - - - - - - - + + + + + + - + - - - - - - - + + + + + + - + - - - - - - - + + + + + + - + - - - - - - - + + + + + + - - - - - + - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + - + - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + - - + - + - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + - - + - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + - - + - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + - - + - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + - - + + From a3a4232b9ac29e66e153927a1b01d3ecba9673c7 Mon Sep 17 00:00:00 2001 From: rgurunathan Date: Mon, 13 Jul 2026 10:59:25 -0400 Subject: [PATCH 18/18] #15 - Pin socket send/receive to separate cores for stable single streams The socket sweep co-pinned a pair's send (client) and receive (server) processes on one core. A shared send/receive core ping-pongs a single stream and can wedge it at half rate for an entire run -- the 8 KB (~1 MSS) cells were bimodal (~15 vs ~30 Gb/s, std +-8.6). Pin the two sides to SEPARATE isolated cores, kept in the SAME CPU cluster (pairs 0-1 in 15-19, pairs 2-3 in 5-9): straddling GB10's two big-core clusters makes every payload pay cross-cluster cache-coherence latency and regresses large messages, so intra-cluster placement is required. Single-stream variance drops sharply; the only cost is the 4-pair cells, whose eight cores must span both clusters (the 5-9 pairs sit farther from the NIC), so they scale slightly sub-linearly. Add SOCKET_NOPIN=1 (cpu_core -1, unpinned) for the pinned-vs-scheduler comparison; not used in the report. Refresh the TCP/UDP tables, summary rows, methodology prose, and the SUT core-pinning row in performance-dgx-spark.md to the re-measured split numbers. Assisted-by: Claude Opus 4.8 Signed-off-by: rgurunathan --- docs/benchmarks/performance-dgx-spark.md | 40 ++++++++++++--------- examples/run_spark_bench.sh | 44 ++++++++++++++++++------ 2 files changed, 56 insertions(+), 28 deletions(-) diff --git a/docs/benchmarks/performance-dgx-spark.md b/docs/benchmarks/performance-dgx-spark.md index 70fbf3a..fdc0777 100644 --- a/docs/benchmarks/performance-dgx-spark.md +++ b/docs/benchmarks/performance-dgx-spark.md @@ -27,7 +27,7 @@ loopback). The exact commands are collected under [Reproduce](#reproduce) below. | NIC | ConnectX-7, ports p0 ↔ p1 cross-cabled with a **100 GbE QSFP28** loopback cable (single-host loopback), MTU 9000 | | Build | Release (`-DCMAKE_BUILD_TYPE=Release`), `DAQIRI_ENGINE="dpdk ibverbs"` | | Loopback | Raw/DPDK uses the two physical ports directly; socket/RoCE use the `dq_wire_*` network-namespace wire loopback | -| Core pinning | Each direction has a busy-spin queue poller and an app worker on separate isolated X925 cores (PR #149). Single-queue: DPDK pollers 17/18, workers 16/19. Multi-queue: TX pollers 16/19, RX pollers 18/9, each with its own worker core, master 8 (with `isolcpus=5-9,15-19`). | +| Core pinning | Each direction has a busy-spin queue poller and an app worker on separate isolated X925 cores (PR #149). Single-queue: DPDK pollers 17/18, workers 16/19. Multi-queue: TX pollers 16/19, RX pollers 18/9, each with its own worker core, master 8. Sockets pin each pair's send and receive to separate cores in the same CPU cluster (all with `isolcpus=5-9,15-19`). | ## Results Summary (C++ loopback) @@ -44,8 +44,8 @@ just under that ceiling. | ----------------- | --------- | ---------: | ----- | ----- | | Raw Ethernet / GPUDirect | 8 KB packet | **98.8 ±0.1 Gb/s** | 0 | Flat ~98.7 across 4–8 KB, all batch sizes | | Socket / RoCE (SEND) | 64 KB message | **97.6 ±0.1 Gb/s** | 0 | Single QP, batch 1 | -| Socket / TCP | 1 MiB × 4 pairs | **90.6 ±1.5 Gb/s** | ~0 | Flow-controlled (App TX = App RX) | -| Socket / UDP | 8 KB × 4 pairs | **28.5 ±0.6 Gb/s** | ~54% loss | Receiver goodput; unpaced sender | +| Socket / TCP | 8 KB × 4 pairs | **87.3 ±2.2 Gb/s** | ~0 | Flow-controlled (App TX = App RX) | +| Socket / UDP | 8 KB × 4 pairs | **34.5 ±0.6 Gb/s** | ~48% loss | Receiver goodput; unpaced sender | Each transport is best read at its own best-case operation size (see the per-transport tables below); a single cross-transport unit of work isn't meaningful here, since @@ -192,10 +192,14 @@ too (SM and memory-controller ~0%; DMA target, not a compute engine). ## Socket / TCP -Four one-way TCP client/server pairs over the netns wire loopback, each pair -pinned to one isolated core (16–19). TCP self-paces via flow control, so App TX -equals App RX with effectively no app-level loss. `message_size` is the per-send -byte count of a stream (no datagram boundary, no fragmentation). +Four one-way TCP client/server pairs over the netns wire loopback. Each pair's +send (client) and receive (server) sides pin to **separate** isolated cores in one +CPU cluster (pairs 1–2 in 15–19, pairs 3–4 in 5–9). A shared send/receive core +ping-pongs a single stream and can wedge it at half rate for a whole run, so +splitting the two sides keeps single-stream throughput stable. TCP self-paces via +flow control, so App TX equals App RX with effectively no app-level loss. +`message_size` is the per-send byte count of a stream (no datagram boundary, no +fragmentation). Throughput in Gb/s (App TX = App RX), mean ± std over 3 reps: @@ -210,21 +214,23 @@ Throughput in Gb/s (App TX = App RX), mean ± std over 3 reps: - 1000 B13.7±0.427.2±0.155.0±0.4 - 8000 B25.3±8.646.4±4.287.2±1.6 - 1 MiB31.5±1.351.9±2.290.6±1.5 + 1000 B14.2±0.427.6±0.445.2±0.1 + 8000 B28.9±2.942.4±2.787.3±2.2 + 1 MiB32.1±2.251.5±2.483.7±0.4 Throughput scales with the pair count; retransmits stay negligible over the run. +At four pairs the eight cores span both clusters, and the pairs in 5–9 sit farther +from the NIC, so the 4-pair cells scale slightly sub-linearly. ## Socket / UDP -Four one-way UDP client/server pairs, same one-core-per-pair pinning. UDP has no -flow control, so each sender runs flat-out and the receiver drops whatever it -cannot drain — the loss column is an inherent property of unpaced UDP, not a -fault. App RX is the delivered goodput; App-level loss is `(App TX − App RX) / -App TX`. +Four one-way UDP client/server pairs, same per-side pinning (send and receive on +separate cores). UDP has no flow control, so each sender runs flat-out and the +receiver drops whatever it cannot drain — the loss column is an inherent property +of unpaced UDP. App RX is the delivered goodput; App-level loss is +`(App TX − App RX) / App TX`. Each cell shows **receiver goodput in Gb/s** (mean ± std over 3 reps) with the **app-level loss %** dimmed beneath it: @@ -240,8 +246,8 @@ Each cell shows **receiver goodput in Gb/s** (mean ± std over 3 reps) with the - 1000 B4.5 ±0.215% loss8.2 ±0.114% loss14.1 ±0.315% loss - 8000 B9.8 ±1.764% loss18.9 ±0.464% loss28.5 ±0.654% loss + 1000 B4.0 ±0.144% loss6.8 ±0.852% loss15.2 ±0.340% loss + 8000 B12.2 ±1.269% loss18.9 ±0.259% loss34.5 ±0.648% loss diff --git a/examples/run_spark_bench.sh b/examples/run_spark_bench.sh index 5b6b2f0..25358e3 100755 --- a/examples/run_spark_bench.sh +++ b/examples/run_spark_bench.sh @@ -431,28 +431,50 @@ generate_yaml() { esac } -# One isolated core per socket pair: 16 + (idx % 4) -> cores 16-19 for pairs 0-3 (Spark -# isolcpus), wrapping (oversubscribing) beyond four. Both the server and the client of a -# pair pin to this same core. Sharing one CPU across the send and receive sides couples -# the send rate to the receive rate (the sender cannot outrun the receiver it time-slices -# with), so each pair sits near the per-core ceiling and N pairs scale to ~N x that. This -# matches the reference four-pair methodology and keeps App TX ~= App RX with low loss. -pair_core() { echo $(( 16 + ($1 % 4) )); } +# Separate isolated cores for the send (client) and receive (server) sides of each +# socket pair, so a pair's two processes never time-slice one CPU. Co-locating them +# (the old "one core per pair" setup) forces a send/recv ping-pong on a single core +# that, at mid-size messages (e.g. 8000 B), makes a single TCP stream metastable -- +# it locks into an efficient (~30 Gb/s) or a serialized (~15 Gb/s) mode for a whole +# run, producing bimodal, high-variance results. +# +# CRITICAL: keep each pair's two cores in the SAME CPU cluster. GB10's ten big X925 +# cores are two clusters of five (5-9 and 15-19); straddling a pair across clusters +# makes every payload pay cross-cluster cache-coherence latency and *regresses* large +# messages (1 MiB 4-pair dropped ~90 -> ~77 Gb/s in testing). So pairs 0-1 sit in the +# 15-19 cluster and pairs 2-3 in the 5-9 cluster, each send/recv intra-cluster +# (master 8, spare 15). NOTE: separating send/recv also lets the UDP sender outrun the +# receiver (no shared-core self-pacing), so UDP loss rises vs the co-pinned setup -- +# the more representative behaviour for an unpaced sender. +SRV_PIN_CORES=(16 18 5 7) +CLI_PIN_CORES=(17 19 6 9) +pair_server_core() { echo "${SRV_PIN_CORES[$(( $1 % 4 ))]}"; } +pair_client_core() { echo "${CLI_PIN_CORES[$(( $1 % 4 ))]}"; } # Write the server/client YAML pair for socket pair `idx`: split the combined base # per role, then substitute message_size, unique ports (SRV/CLI_PORT_BASE + idx), -# and pin every queue to the pair's core. +# and pin the server (receive) side and client (send) side to DIFFERENT isolated +# cores so the pair does not time-slice one CPU (see pair_server_core comment). generate_socket_yaml() { local idx="$1" payload="$2" server_out="$3" client_out="$4" local srv_port=$(( SRV_PORT_BASE + idx )) local cli_port=$(( CLI_PORT_BASE + idx )) - local core; core="$(pair_core "$idx")" + # SOCKET_NOPIN=1 runs the bench workers unpinned (cpu_core -1 -> no affinity, the + # scheduler places them). For gathering the pinned-vs-non-pinned comparison only; + # the published report stays pinned like the other backends. + local server_core client_core + if [[ -n "${SOCKET_NOPIN:-}" ]]; then + server_core=-1; client_core=-1 + else + server_core="$(pair_server_core "$idx")" + client_core="$(pair_client_core "$idx")" + fi python3 "$NETNS_GEN" "$BASE_YAML" --role server | \ sed -E \ -e "s|^( *message_size: ).*|\1$payload|g" \ -e "s|^( *local_addr: \"?[a-z]+://[0-9.]+:)[0-9]+(\"?)|\1$srv_port\2|" \ -e "s|^( *server_port: ).*|\1$srv_port|" \ - -e "s|^( *cpu_core: ).*|\1$core|" \ + -e "s|^( *cpu_core: ).*|\1$server_core|" \ > "$server_out" python3 "$NETNS_GEN" "$BASE_YAML" --role client | \ sed -E \ @@ -460,7 +482,7 @@ generate_socket_yaml() { -e "s|^( *local_addr: \"?[a-z]+://[0-9.]+:)[0-9]+(\"?)|\1$cli_port\2|" \ -e "s|^( *remote_addr: \"?[a-z]+://[0-9.]+:)[0-9]+(\"?)|\1$srv_port\2|" \ -e "s|^( *server_port: ).*|\1$srv_port|" \ - -e "s|^( *cpu_core: ).*|\1$core|" \ + -e "s|^( *cpu_core: ).*|\1$client_core|" \ > "$client_out" }