#15 - Add reusable FFT/GEMM GPU workloads to receive-path benchmarks#203
Conversation
Adds an engine-agnostic GPU-workload component (examples/bench_workload.{h,cu})
that runs a representative cuFFT C2C transform or cuBLAS SGEMM once per received
burst, sized to the burst's data volume (batch x payload). It operates on its own
GPU scratch buffers rather than the received bytes, so it drops into any
stream_type/engine without needing a payload device pointer (RoCE exposes none)
and is a verified no-op when --workload is none.
Wired --workload none|fft|gemm into the DPDK/HDS receive path (shared
rx_count_worker, used by daqiri_bench_raw_gpudirect and daqiri_bench_raw_hds) and
the RoCE path (rdma_bench rdma_worker, per RECEIVE completion). run_spark_bench.sh
gains a WORKLOAD env that fills the existing post_process CSV column. cuFFT/cuBLAS
are linked into the affected example targets.
DGX Spark results (Raw/GPUDirect, 8 KB native shape, batch 10240, single 30 s
run), added to docs/benchmarks/performance-dgx-spark.md: none ~98.5 Gb/s; FFT 94.2
Gb/s / 8.9% SM; GEMM 62.3 Gb/s / 82.4% SM. All three drop-free -- the receive
path backpressures against GPU load without dropping packets. RoCE workload rows
are left pending a netns RoCE sweep (blocked on a pre-existing RoCE-over-netns
connection issue, independent of this change).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
PR #156 moved the RoCE client peer endpoint into the application config and made daqiri_init reject socket_config.remote_addr on a RoCE client, but the Spark netns RDMA bench config was not updated. The generated client config failed daqiri_init ('RoCE client peer endpoints belong in application config'), so the client never connected and the server timed out with no peer. Switch the client interface to bind its own address via local_addr (matching the updated daqiri_bench_rdma_tx_rx.yaml); the server peer stays in rdma_bench_client. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
…y in DGX Spark perf doc Socket/RoCE 8 MB native message, single 30 s run: none 101.4 Gb/s; FFT 98.7 Gb/s / 26.9% SM; GEMM 79.3 Gb/s / 76.6% SM. All drop-free (0 CQ errors) -- same backpressure-not-loss behaviour as the DPDK path. Also expands the section to describe exactly what the FFT (batched 1-D cuFFT C2C) and GEMM (square cuBLAS SGEMM) steps compute and the real-world receiver patterns they model. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
Adds --workload gemm_fp16: the same square matmul as gemm but in mixed precision (FP16 inputs, FP32 accumulate) on the tensor cores via cublasGemmEx (CUBLAS_GEMM_DEFAULT_TENSOR_OP), the core op of GPU inference. Sized to the same matrix dimension as the FP32 gemm so the two isolate the precision/tensor-core effect. Wired through all three benches' --workload, run_spark_bench.sh WORKLOAD validation, and AGENTS.md. GpuWorkload::init now runs one warmup of the chosen op and checks the cuFFT/cuBLAS status + a stream sync, so a misconfigured call fails loudly at init (and the caller disables the workload) instead of silently no-op'ing on the hot path and looking like a free workload. issue_op() factors the dispatch shared by run() and the warmup. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
…ection polish Raw/GPUDirect 8 KB and Socket/RoCE 8 MB workload tables refreshed to mean +/- std over 3 reps, with the FP16 tensor-core GEMM result added for both backends: Raw 98.5 Gb/s @ 29% SM (vs FP32 61.5 @ 82%) and RoCE 100.1 Gb/s @ 35% SM (vs FP32 77.8 @ 77%) -- the same matmul, effectively free at inference precision, all drop-free. Renames the section to 'GPU workloads in the receive path', drops 'native-shape peak' from the Results Summary heading (clarified in text), removes the issue-#15 callout, and expands the methodology to describe what the FFT and GEMM/GEMM-FP16 steps compute. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
|
| Filename | Overview |
|---|---|
| examples/bench_workload.cu | New GpuWorkload class (cuFFT C2C + FP32/FP16-tensor-core GEMM); init-time warmup validates the op; stream ownership and handle lifetime are cleanly managed with destroy() covering every exit path. |
| examples/bench_workload.h | Public header for GpuWorkload and BenchWorkload enum; CUDA types hidden behind void* so it is includable from plain .cpp mains — clean design matching existing patterns. |
| examples/bench_pipeline.cu | ReorderPipeline correctly handles both SeqReorder and GatherOnly modes; the single-packet zero-copy pass-through avoids a redundant kernel launch for RoCE; all CUDA allocations freed in destroy(). |
| examples/raw_bench_common.cpp | Wires workload into rx_count_worker; per-burst gpu_workload.sync() before free_all_packets_and_burst_rx correctly prevents the reorder kernel from reading freed DMA buffers. |
| examples/rdma_bench.cpp | Batch-drain pipelining for RoCE is well-structured; flush_held_recv() is called on every exit (batch full, idle, shutdown), preserving the BurstParams free contract on all paths. |
| examples/socket_bench.cpp | Socket workload correctly stages host payloads to the GPU and syncs the stream before freeing the burst; per-burst sync is necessary here since the staging copy reads the burst's pageable memory. |
| src/kernels.cu | New packet_gather_copy_payload kernel follows the same vectorized-copy pattern as the existing reorder kernel; one block per packet, 128 threads, input null-check present. |
| examples/CMakeLists.txt | Adds bench_workload.cu + bench_pipeline.cu to all four bench targets and links CUDA::cufft + CUDA::cublas; also correctly adds CUDA_ARCHITECTURES to rdma and socket benches that were previously missing it. |
| examples/run_spark_bench.sh | WORKLOAD / WORKLOAD_BATCH env vars validated at startup with case/regex guards; CSV post_process column now records the actual workload instead of the hardcoded 'none'. |
| examples/raw_gpudirect_bench.cpp | Injects per-packet network-byte-order seq numbers into TX templates and wires ReorderGeometry into rx_count_worker; packets_per_batch capped to tx.batch_size. |
| examples/raw_hds_bench.cpp | HDS reorder uses segment 1 (GPU payload) with zero payload_byte_offset; no seq injection since HDS TX does not write seq numbers — acknowledged in comment as acceptable for a throughput benchmark. |
| examples/daqiri_bench_rdma_tx_rx_spark_netns.yaml | Fixes post-#156 RoCE client schema by replacing rejected remote_addr with local_addr; comment explains the reason for the change. |
| docs/benchmarks/performance-dgx-spark.md | Adds full GPU-workload results section with methodology, per-transport tables, and batch-size sweep; prose and numbers are internally consistent. |
Reviews (5): Last reviewed commit: "#15 - Parse --workload-batch-bytes in so..." | Re-trigger Greptile
…llback A typo like --workload ftt previously mapped silently to none, producing a baseline run that looked valid. Emit a stderr warning for unknown values and match the legitimate "none" explicitly so it stays quiet. (Greptile PR #203) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
…kload The post_process column emitted $WORKLOAD unconditionally, so a socket run with WORKLOAD=fft recorded a GPU workload that was never applied (socket ignores --workload). Reduce to an effective workload (none on socket backends) and use it for both --workload forwarding and the CSV column. (Greptile PR #203) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
…/gather
The --workload component previously ran cuFFT/cuBLAS on private GPU scratch
buffers, so the numbers measured GPU-load headroom, not the cost of processing
the data that actually arrived. Make every backend run the workload on the REAL
received payload, preceded by a packet-reorder step:
- GpuWorkload now consumes a caller-supplied contiguous device buffer
(run(const void* input)); FFT reads it, GEMM uses it as the A operand.
- New ReorderPipeline (examples/bench_pipeline.{h,cu}) assembles a burst's
payloads into one contiguous device buffer and shares the workload's stream:
a sequence-number reorder for out-of-order transports (DPDK raw, UDP) via the
existing kernel, an arrival-order gather for in-order ones (RoCE RC, TCP) via
a new packet_gather_copy_payload kernel; a single large device-resident
message is a zero-copy pass-through.
- DPDK raw + HDS: drive the pipeline from rx_count_worker; TX injects a
per-packet sequence number. RoCE: gather pass-through of the recv MR.
Sockets: gain --workload, stage host->device, UDP reorder / TCP gather, and
inject a UDP seq number.
- run_spark_bench.sh: all backends now honour WORKLOAD (socket included).
Validated on DGX Spark: DPDK raw (gemm_fp16, 96.8 Gb/s, drop-free, GPU active
on reordered data) and RoCE (fft, 92.7 Gb/s, drop-free, gather pass-through).
Socket netns RX is blocked by a pre-existing harness regression (tracked
separately); the socket wiring is complete and compiles.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
…tables) Rewrite the DGX Spark GPU-workloads section: drop the "representative compute, not a data transform" caveat (no longer true), document the per-backend reorder/gather methodology and the socket host->device staging cost, and replace the stale scratch-buffer result tables with placeholders to fill from re-run sweeps. Note that socket rows are blocked on a separate netns RX regression. Update raw_benchmarking.md's --workload paragraph and the AGENTS.md benchmark note (scratch -> real data + reorder; all backends honour --workload). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
Drop the pending/regression admonition and the socket (UDP/TCP) result rows from the GPU-workloads section; those numbers depend on a separate netns receive fix. Keep the DPDK and RoCE placeholder tables (the validated paths) for the 3-rep re-run, and split RoCE into its own table. The per-backend methodology (incl. the socket staging path) stays, since the capability is implemented. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
Populate the Raw/GPUDirect and RoCE tables with mean+/-std over 3 reps (30s each) on the real received data. All cells drop-free. Raw: FFT 95.5, FP32 GEMM 94.9 (56% SM), FP16 GEMM 97.7 Gb/s. RoCE: FFT 93.0, FP32 GEMM 32.6 (78% SM, serialized per 8MB message), FP16 GEMM 85.3 Gb/s. Add prose explaining the raw-vs-RoCE FP32 gap (raw pipelines ~10 windows/burst; RoCE drains the stream per message). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
Add --workload-batch-bytes (env WORKLOAD_BATCH): the GPU compute working set per call, independent of the I/O unit. RoCE sub-divides each 8 MB message into batch-sized slices (one compute per slice, one sync amortized over them); the raw benches size their reorder window from the batch instead of a fixed 1024-packet window. The GEMM dimension scales as n=sqrt(batch/4), so sweeping the batch traces the compute-intensity vs throughput curve and makes the RoCE per-message serialization legible. run_spark_bench.sh gains the WORKLOAD_BATCH env and records it in a new post_process_batch CSV column (appended last, so existing column positions and the mq DictReader consumer are unaffected). Doc gains a batch-size-sweep subsection (placeholder table + sweep loop), focused on gemm_fp16. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
…s with raw) Previously the RoCE worker drained the GPU stream after every received message (the gather pass-through reads the recv buffer directly, so it had to finish before the buffer was reposted), fully serializing compute against receive and making the FP32 GEMM look far heavier than the raw path (32.6 vs 94.9 Gb/s). Hold a small batch of RECEIVE completions instead: holding a completion keeps its recv buffer live, so the pass-through compute can read it while later messages arrive into other pool buffers. Enqueue the computes, drain the stream once per batch (bounded well under rx_depth so receives are not starved), then free them all. This mirrors the raw path, which holds a whole burst and drains once per burst -- now both overlap receive with compute and the comparison is apples-to-apples. Reset the RoCE result table to placeholders (the prior numbers were measured serially) and reword the prose. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
RoCE fixed-batch table (3-rep, post overlap fix): none 101.7, FFT 93.6, FP32 GEMM 35.0 (79% SM), FP16 GEMM 85.8 Gb/s. Batch-size sweep (gemm_fp16, RoCE, 1 rep): non-monotonic with a 2-4 MB sweet spot reaching ~92 Gb/s -- on par with the raw path -- confirming the RoCE/raw gap at the 8 MB default is pipelining depth (one GEMM/message vs ~10 back-to-back in a raw burst), not transport. Raw sweep column and a repeat-averaged RoCE sweep still to collect. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
…-sensitive) Raw (DPDK) gemm_fp16 batch sweep holds line rate (~98 Gb/s) across every batch -- a raw burst always packs ~10 reorder windows so the GPU stays fed regardless of window size. RoCE swings 37->92 Gb/s with a 2-4 MB sweet spot, confirming the RoCE/raw gap at the default batch is pipelining depth (one GEMM/message), not transport: a 2-4 MB RoCE batch packs several GEMMs/message and converges with raw. No TBDs remain in the perf doc. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
Keep the RoCE summary at the 8 MB native operating point (batch-matched to the raw path's ~8.19 MB reorder window) and add a compact 4 MB batch-tuned block (3-rep): FP32 GEMM 76.3 Gb/s (2.2x the 8 MB default), FP16 GEMM 92.3 Gb/s (within ~5% of raw). Makes the pipelining-recovery result explicit alongside the native numbers. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
git-clang-format fixups (include ordering / wrapping) on the files touched by the real-data workload + batch-sweep work. No functional change. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
run_spark_bench.sh forwards --workload-batch-bytes to every backend and records post_process_batch in the CSV, but socket_bench ignored it -- so a socket batch sweep produced identical numbers while the CSV claimed different batches. Parse it and size the UDP reorder window / working set from it (default ~8 MB). TCP still gathers one recv chunk per compute (batch sub-chunking for TCP left to the socket follow-up, with the netns RX fix). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: rgurunathan <rgurunathan@nvidia.com>
Closes #15
What
examples/bench_workload.{h,cu}— reusable, engine-agnostic GPU workload run per received burst on its own GPU scratch (no payload pointer → drops into any engine; no-op fornone, validated by an init-time warmup).--workload none|fft|gemm|gemm_fp16: batched cuFFT C2C; FP32cublasSgemm; and the same-size mixed-precision FP16/tensor-corecublasGemmEx(inference-style). Wired into DPDK/HDS (daqiri_bench_raw_gpudirect,daqiri_bench_raw_hds) and RoCE (daqiri_bench_rdma);run_spark_bench.shWORKLOADenv fills thepost_processCSV column.socket_config.remote_addris now rejected for RoCE clients).Results (DGX Spark, 3-rep mean ± std, all drop-free)
(Gb/s / GPU SM%.) The receive path backpressures against GPU load without dropping packets, and the same matmul at FP16/tensor-core precision returns to line rate at far lower SM than FP32.
How workloads run
Each backend's RX path now feeds the actual received payload to the GPU workload, assembled into one contiguous device buffer per batch by a shared
ReorderPipeline(examples/bench_pipeline.{h,cu}) that shares the workload's CUDA stream:packet_reorder_copy_payload_by_sequence) places each packet's payload atslot = seq % packets_per_batch, producing an ordered contiguous buffer. TX injects a per-packet seq number.packet_gather_copy_payload). A single large device-resident message is a zero-copy pass-through (the recv buffer is the contiguous buffer).The assembled buffer then feeds
GpuWorkload(examples/bench_workload.{h,cu}): cuFFT C2C, or cuBLAS FP32Sgemm/ FP16 tensor-coreGemmEx(the buffer is the FFT input / GEMMAoperand). To overlap compute with ingest, each path holds a batch of received data and drains the GPU stream once per batch (raw: a ~10-window burst; RoCE: a small batch of held completions whose recv buffers stay live during the pass-through), so the receiver backpressures against GPU load rather than dropping.--workload-batch-bytesdecouples the compute working set from the I/O unit (RoCE sub-chunks each message; raw sizes its reorder window), enabling the batch-size sweep.Socket (UDP/TCP) results are deferred to a follow-up PR: the netns wire-loopback harness has a pre-existing receive-path regression (server binds but receives nothing), independent of this change, so socket numbers will land with that fix on
fix/socket-netns-rx.