From d588003afed6f09f0d16832895ecb70780a60082 Mon Sep 17 00:00:00 2001 From: Chloe Crozier Date: Fri, 12 Jun 2026 09:05:26 +0000 Subject: [PATCH 1/7] #17 - Add RTX PRO 6000 benchmark YAMLs and baseline Partial #17: raw GPUDirect configs for RTX PRO 6000 Blackwell plus measured baseline on the dev box. No HDS, RoCE, FFT, or GEMM yet. Signed-off-by: Chloe Crozier --- AGENTS.md | 2 +- docs/tutorials/configuration-walkthrough.md | 4 + examples/CMakeLists.txt | 4 + ...ri_bench_raw_sw_loopback_rtx_pro_6000.yaml | 77 +++++++++++++++ .../daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml | 93 +++++++++++++++++++ ...qiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml | 93 +++++++++++++++++++ ..._raw_tx_rx_rtx_pro_6000_nic_same_port.yaml | 75 +++++++++++++++ examples/rtx_pro_6000_baseline.md | 60 ++++++++++++ 8 files changed, 407 insertions(+), 1 deletion(-) create mode 100644 examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml create mode 100644 examples/daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml create mode 100644 examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml create mode 100644 examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml create mode 100644 examples/rtx_pro_6000_baseline.md diff --git a/AGENTS.md b/AGENTS.md index 7ef7b25..3cb274d 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -32,7 +32,7 @@ There is no unit test suite. Verification is done via the benchmark executables | Executable | Source | Typical config | |---|---|---| -| `daqiri_bench_raw_gpudirect` | `raw_gpudirect_bench.cpp` | `daqiri_bench_raw_tx_rx.yaml`, `daqiri_bench_raw_tx_rx_4q.yaml`, `daqiri_bench_raw_tx_rx_spark.yaml`, `daqiri_bench_raw_{tx,rx}_spark_xhost.yaml`, `daqiri_bench_raw_sw_loopback.yaml`, `daqiri_bench_raw_rx_multi_q.yaml`, `daqiri_bench_raw_tx_rx_spark_mq.yaml` (mq base; `run_spark_mq_bench.sh` derives the 4 cells via `scripts/gen_spark_mq_config.py`) | +| `daqiri_bench_raw_gpudirect` | `raw_gpudirect_bench.cpp` | `daqiri_bench_raw_tx_rx.yaml`, `daqiri_bench_raw_tx_rx_4q.yaml`, `daqiri_bench_raw_tx_rx_spark.yaml`, `daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml`, `daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml`, `daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml`, `daqiri_bench_raw_{tx,rx}_spark_xhost.yaml`, `daqiri_bench_raw_sw_loopback.yaml`, `daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml`, `daqiri_bench_raw_rx_multi_q.yaml`, `daqiri_bench_raw_tx_rx_spark_mq.yaml` (mq base; `run_spark_mq_bench.sh` derives the 4 cells via `scripts/gen_spark_mq_config.py`) | | `daqiri_bench_raw_hds` | `raw_hds_bench.cpp` | `daqiri_bench_raw_tx_rx_hds.yaml` | | `daqiri_bench_raw_reorder_seq` | `raw_reorder_seq_bench.cpp` | `daqiri_bench_raw_tx_rx_reorder_seq_1024*.yaml`, `daqiri_bench_raw_rx_reorder_seq_*.yaml` | | `daqiri_bench_raw_reorder_quantize` | `raw_reorder_quantize_bench.cpp` | `daqiri_bench_raw_tx_rx_reorder_quantize_seq_batch.yaml` | diff --git a/docs/tutorials/configuration-walkthrough.md b/docs/tutorials/configuration-walkthrough.md index b2232bc..9e0957e 100644 --- a/docs/tutorials/configuration-walkthrough.md +++ b/docs/tutorials/configuration-walkthrough.md @@ -32,6 +32,10 @@ For a shorter selection guide, start with the [Benchmarking overview](../benchma - **DGX Spark multi-queue core-scaling matrix** (prefilled) — one base config [`daqiri_bench_raw_tx_rx_spark_mq.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_spark_mq.yaml) (the balanced TX=2/RX=2 superset; cores TX → 16,17, RX → 18,19) from which `examples/run_spark_mq_bench.sh` (via `scripts/gen_spark_mq_config.py`) derives the four `(TX, RX)` cells — (1,1), (1,2) (RX scaling), (2,1) (TX scaling), (2,2) (balanced) — by pruning queues/flows. All run on `daqiri_bench_raw_gpudirect` at the native 8 KB shape. - **DGX Spark cross-host** (prefilled, runs on two Sparks) — [`daqiri_bench_raw_tx_spark_xhost.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_spark_xhost.yaml) on the TX host and [`daqiri_bench_raw_rx_spark_xhost.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_rx_spark_xhost.yaml) on the RX host. Each host runs `daqiri_bench_raw_gpudirect` against its own half; cables connect p0↔p0 between the two boxes. See the [Cross-host two-DGX-Spark loopback](../benchmarks/raw_benchmarking.md#cross-host-two-dgx-spark-loopback) section for run details. - **No physical NIC available** — [`daqiri_bench_raw_sw_loopback.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_sw_loopback.yaml). `loopback: "sw"`, no NIC required. Useful for first-time build verification, not representative of production performance. + - **RTX PRO 6000 Blackwell — no cable** — [`daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml). `kind: device`, `affinity: 0`; build with [`CMAKE_CUDA_ARCHITECTURES=120`](../tutorials/bare-metal-cmake-build.md). SW loopback smoke test only. + - **RTX PRO 6000 Blackwell — real NIC, dual-port on one card** (prefilled dev box) — [`daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml). `61:00.0` p0 → `61:00.1` p1, GPU 0 TX / GPU 1 RX; needs L2 link between ports (not SW loopback). See [`rtx_pro_6000_baseline.md`](https://github.com/nvidia/daqiri/blob/main/examples/rtx_pro_6000_baseline.md) for hardware limits and measured baseline. + - **RTX PRO 6000 — same-PF NIC attempt** (experimental) — [`daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml). Single port `61:00.0` TX+RX; failed `daqiri_init` on reference box — kept for follow-up. + - **RTX PRO 6000 Blackwell — dual-NIC loopback** (generic template) — [`daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml). Placeholders for Cliff's 800 Gbps cross-card target; fill PCIe BDFs and MACs. To watch the same raw loopback benchmark with live Prometheus and Grafana counters, use the Grafana compose stack described in diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 829d88e..2f4d123 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -34,6 +34,10 @@ set(DAQIRI_BENCH_CONFIGS daqiri_bench_raw_rx_reorder_seq_batch.yaml daqiri_bench_raw_rx_multi_q.yaml daqiri_bench_raw_sw_loopback.yaml + daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml + daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml + daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml + daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml daqiri_example_gds_write_sw_loopback.yaml daqiri_example_gds_write_tx_rx.yaml daqiri_example_pcap_writer_sw_loopback.yaml diff --git a/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml b/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml new file mode 100644 index 0000000..df02770 --- /dev/null +++ b/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml @@ -0,0 +1,77 @@ +# RTX PRO 6000 Blackwell (discrete dGPU) software-loopback smoke test. +# No NIC or cable required — validates build + GPUDirect on one GPU. +# Not representative of wire-speed performance; use the hardware template +# (daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml) once a QSFP loopback cable is installed. +# +# Build (native sm_120): +# cmake -S . -B build -DBUILD_SHARED_LIBS=ON -DDAQIRI_BUILD_PYTHON=OFF \ +# -DDAQIRI_MGR="dpdk socket rdma" -DCMAKE_CUDA_ARCHITECTURES=120 +# cmake --build build -j +# +# Run: +# ./build/examples/daqiri_bench_raw_gpudirect \ +# ./build/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml --seconds 30 +# +# memory_regions[].affinity is the CUDA device index (not nvidia-smi GPU id). +# Change affinity to target a different GPU, e.g. affinity: 1 for CUDA device 1. +# +%YAML 1.2 +--- +daqiri: + cfg: + version: 1 + stream_type: "raw" + master_core: 3 + debug: false + log_level: "info" + loopback: "sw" + + memory_regions: + - name: "Data_TX_GPU" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + - name: "Data_RX_GPU" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + + interfaces: + - name: "loopback_ports" + address: "loopback" + tx: + queues: + - name: "tx_q_0" + id: 0 + batch_size: 10240 + cpu_core: 11 + timeout_us: 1000 + memory_regions: + - "Data_TX_GPU" + offloads: + - "tx_eth_src" + rx: + queues: + - name: "rq_q_0" + id: 0 + cpu_core: 9 + timeout_us: 1000 + batch_size: 10240 + memory_regions: + - "Data_RX_GPU" + +bench_rx: + interface_name: "loopback_ports" + +bench_tx: + interface_name: "loopback_ports" + batch_size: 10240 + payload_size: 8000 + header_size: 64 + eth_dst_addr: 00:00:00:00:00:00 + ip_src_addr: 000.000.000.0 + ip_dst_addr: 000.000.000.0 + udp_src_port: 4096 + udp_dst_port: 4096 diff --git a/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml new file mode 100644 index 0000000..4af275e --- /dev/null +++ b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml @@ -0,0 +1,93 @@ +# RTX PRO 6000 Blackwell dual-NIC hardware loopback TEMPLATE (generic placeholders). +# For a prefilled single-card dual-port run on the reference dev box, see +# daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml instead. +# +# Requires L2 connectivity between tx_port and rx_port (QSFP cable, loopback optic, or switch). +# NOT runnable until are replaced for your system. +# +# Target topology (Cliff's 800 Gbps vision — two cards or two ports at line rate): +# GPU 0 (CUDA affinity 0) --TX buffers--> NIC0 --link--> NIC1 --RX buffers--> GPU 1 +# +# Hardware limits: +# - 800 Gbps aggregate needs two ~400G ports with an active link each; no cable = no wire test. +# - This server lacks Spark's dual-PF-per-port on-chip eswitch shortcut (see baseline doc). +# - SW loopback (daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml) never touches the NIC. +# +# Discovery helpers: +# lspci -d 15b3: +# ibdev2netdev +# cat /sys/class/net//address # eth_dst_addr = rx_port MAC +# nvidia-smi topo -m # NUMA / PCIe proximity for cpu_core picks +# +# Build: same as daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml header. +# +%YAML 1.2 +--- +daqiri: + cfg: + version: 1 + stream_type: "raw" + master_core: 3 + debug: false + log_level: "info" + loopback: "" + + memory_regions: + - name: "Data_TX_GPU" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + - name: "Data_RX_GPU" + kind: "device" + affinity: 1 + num_bufs: 51200 + buf_size: 8064 + + interfaces: + - name: "tx_port" + address: <0000:00:00.0> + tx: + queues: + - name: "tx_q_0" + id: 0 + batch_size: 10240 + cpu_core: 11 + memory_regions: + - "Data_TX_GPU" + offloads: + - "tx_eth_src" + - name: "rx_port" + address: <0000:00:00.1> + rx: + flow_isolation: true + queues: + - name: "rq_q_0" + id: 0 + cpu_core: 9 + batch_size: 10240 + memory_regions: + - "Data_RX_GPU" + flows: + - name: "flow_0" + id: 0 + action: + type: queue + id: 0 + match: + udp_src: 4096 + udp_dst: 4096 + +bench_rx: +- interface_name: "rx_port" + +bench_tx: +- interface_name: "tx_port" + batch_size: 10240 + payload_size: 8000 + header_size: 64 + eth_dst_addr: <00:00:00:00:00:00> + ip_src_addr: <1.2.3.4> + ip_dst_addr: <5.6.7.8> + udp_src_port: 4096 + udp_dst_port: 4096 diff --git a/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml new file mode 100644 index 0000000..f9cb21f --- /dev/null +++ b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml @@ -0,0 +1,93 @@ +# RTX PRO 6000 Blackwell — real NIC loopback on this dev box (no external cable required +# if p0/p1 already have link). Uses one ConnectX-7 / BF-3 dual-port card: +# tx_port 0000:61:00.0 (ens1f0np0, p0) GPU CUDA 0 +# rx_port 0000:61:00.1 (ens1f1np1, p1) GPU CUDA 1 +# +# Hardware limitations (read before comparing numbers): +# - NOT Cliff's 800 Gbps dual-card test. This is one ASIC, two ports (~400G class each). +# - Requires L2 connectivity between p0 and p1 (QSFP cable, passive loopback optic, or +# switch). If carrier=0 on either port, this config will not pass traffic. +# - Unlike DGX Spark, this server has one PF per physical port — no on-chip eswitch +# loopback between PFs on the same port without a link. +# - SW loopback (daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml) does not use the NIC. +# +# Verify link: cat /sys/class/net/ens1f{0,1}np*/carrier +# eth_dst_addr: rx_port MAC — cat /sys/class/net/ens1f1np1/address on this dev box +# Verify wire vs on-chip after a run: tx_phy_packets / rx_phy_packets near zero = on-chip; +# rising with TX/RX = packets crossed the SerDes. +# +# Build: cmake ... -DCMAKE_CUDA_ARCHITECTURES=120 && cmake --build build -j +# Run: +# sudo ./build/examples/daqiri_bench_raw_gpudirect \ +# ./build/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml --seconds 30 +# +%YAML 1.2 +--- +daqiri: + cfg: + version: 1 + stream_type: "raw" + master_core: 3 + debug: false + log_level: "info" + loopback: "" + + memory_regions: + - name: "Data_TX_GPU" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + - name: "Data_RX_GPU" + kind: "device" + affinity: 1 + num_bufs: 51200 + buf_size: 8064 + + interfaces: + - name: "tx_port" + address: 0000:61:00.0 + tx: + queues: + - name: "tx_q_0" + id: 0 + batch_size: 10240 + cpu_core: 11 + memory_regions: + - "Data_TX_GPU" + offloads: + - "tx_eth_src" + - name: "rx_port" + address: 0000:61:00.1 + rx: + flow_isolation: true + queues: + - name: "rq_q_0" + id: 0 + cpu_core: 9 + batch_size: 10240 + memory_regions: + - "Data_RX_GPU" + flows: + - name: "flow_0" + id: 0 + action: + type: queue + id: 0 + match: + udp_src: 4096 + udp_dst: 4096 + +bench_rx: +- interface_name: "rx_port" + +bench_tx: +- interface_name: "tx_port" + batch_size: 10240 + payload_size: 8000 + header_size: 64 + eth_dst_addr: <00:00:00:00:00:00> + ip_src_addr: 1.2.3.4 + ip_dst_addr: 5.6.7.8 + udp_src_port: 4096 + udp_dst_port: 4096 diff --git a/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml new file mode 100644 index 0000000..77ec6cc --- /dev/null +++ b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml @@ -0,0 +1,75 @@ +# Same-PF TX+RX on 0000:61:00.0 (ens1f0np0). Alternative when p0<->p1 are not cabled. +# eth_dst_addr: this port's own MAC (hairpin / L2 loopback if supported) — +# cat /sys/class/net//address +# See daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml for dual-port p0->p1 attempt. +# +%YAML 1.2 +--- +daqiri: + cfg: + version: 1 + stream_type: "raw" + master_core: 3 + debug: false + log_level: "info" + loopback: "" + + memory_regions: + - name: "Data_TX_GPU" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + - name: "Data_RX_GPU" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + + interfaces: + - name: "tx_port" + address: 0000:61:00.0 + tx: + queues: + - name: "tx_q_0" + id: 0 + batch_size: 10240 + cpu_core: 11 + memory_regions: + - "Data_TX_GPU" + offloads: + - "tx_eth_src" + - name: "rx_port" + address: 0000:61:00.0 + rx: + flow_isolation: true + queues: + - name: "rq_q_0" + id: 0 + cpu_core: 9 + batch_size: 10240 + memory_regions: + - "Data_RX_GPU" + flows: + - name: "flow_0" + id: 0 + action: + type: queue + id: 0 + match: + udp_src: 4096 + udp_dst: 4096 + +bench_rx: +- interface_name: "rx_port" + +bench_tx: +- interface_name: "tx_port" + batch_size: 10240 + payload_size: 8000 + header_size: 64 + eth_dst_addr: <00:00:00:00:00:00> + ip_src_addr: 1.2.3.4 + ip_dst_addr: 5.6.7.8 + udp_src_port: 4096 + udp_dst_port: 4096 diff --git a/examples/rtx_pro_6000_baseline.md b/examples/rtx_pro_6000_baseline.md new file mode 100644 index 0000000..bd2babc --- /dev/null +++ b/examples/rtx_pro_6000_baseline.md @@ -0,0 +1,60 @@ +# RTX PRO 6000 benchmark baseline + +Host: x86_64 EPYC · GPU: NVIDIA RTX PRO 6000 Blackwell Server Edition · `CMAKE_CUDA_ARCHITECTURES=120` · branch `ccrozier-rtx-pro-6000-bench` · 2026-06-12 + +## Hardware limitations (read first) + +| What | Limit | +|---|---| +| **Cliff 800 Gbps target** | Two ~400G ports with an **active L2 link** (QSFP cable, loopback optic, or switch). Not achievable without that link. | +| **This server vs DGX Spark** | One PF per physical port. **No** Spark-style on-chip eswitch loopback (two PFs on the same port). | +| **`loopback: "sw"`** | Does **not** use the NIC. Measures in-process DPDK path only; can exceed line rate (not comparable to Gbps on the wire). | +| **`carrier=1`** | Link up on a port ≠ p0 and p1 are cabled **to each other**. Our dual-port run proved this: TX on p0, RX 0 on p1. | +| **Starting point** | SW smoke test validates GPUDirect build. Real NIC numbers need a completed L2 loop; use `tx_phy_packets` / `rx_phy_packets` to confirm wire vs internal. | + +## Results (this dev box) + +| Config | GPU (CUDA) | Mode | Duration | TX Gbps | RX Gbps | Notes | +|---|---|---|---|---|---|---| +| `daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml` | 0 | `loopback: sw` | 30s | 1579.5 | 1579.5 | No NIC; GPUDirect smoke baseline | +| `daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml` | 0 TX / 1 RX | real NIC `61:00.0`→`61:00.1` | 30s | 381.5 | 0 | NIC TX path works; **no RX** — p0/p1 not looped (phy_pkts ≪ vport) | +| `daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml` | 0 | same PF `61:00.0` | — | — | — | `daqiri_init` failed (extmem pool on single port); not a baseline | +| `daqiri_bench_raw_sw_loopback_reorder_seq_1024.yaml` | CPU huge | `loopback: sw` + reorder | 30s | — | 160.2 | GPU reorder kernel path; CPU buffers not GPUDirect | +| `daqiri_bench_socket_udp_tx_rx.yaml` (`iterations: 0`) | host | kernel UDP `127.0.0.1` | 15s | 4.7 | 4.7 | Kernel socket baseline; stock YAML uses `iterations: 1000` (not a perf test) | + +### NIC dual-port run detail + +- Card: `0000:61:00.0` (ens1f0np0, p0) TX → `0000:61:00.1` (ens1f1np1, p1) RX +- `eth_dst_addr`: `c4:70:bd:c2:8a:93` (p1 MAC) +- DPDK `tx_phy_packets` / `rx_phy_packets` on port 0: **2** / **76** vs **177M** TX vport packets → traffic did not cross SerDes between ports + +## Commands + +```bash +# SW smoke (no NIC) +sudo ./build/examples/daqiri_bench_raw_gpudirect \ + ./build/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml --seconds 30 + +# Real NIC dual-port (needs p0↔p1 L2 loop) +sudo ./build/examples/daqiri_bench_raw_gpudirect \ + ./build/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml --seconds 30 +``` + +## What we can still run now (#17, partial) + +| Test | Status | Notes | +|---|---|---| +| Raw SW loopback GPUDirect | Done | Best no-cable perf number | +| Raw NIC TX (one port) | Done | Proves mlx5 + GPUDirect TX | +| Reorder SW loopback | Done | ~160 Gbps RX, CPU-side buffers | +| Socket UDP (`iterations: 0`) | Done | ~5 Gbps kernel baseline | +| Socket TCP (`iterations: 0`) | Easy | Same yaml tweak | +| HDS / RoCE / NIC closed-loop | Blocked | Need filled YAML + L2 loop | +| FFT / GEMM workloads | Not in repo | #17 asks for these | + +## Follow-ups + +- Install QSFP cable (or passive loopback) between p0 and p1 on `61:00.x`, re-run `_nic.yaml` +- Scale to cross-card 800 Gbps using `daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml` template + second card +- Tune CPU cores from `nvidia-smi topo -m` once RX path is up +- Add RTX socket YAML with `iterations: 0` for repeatable kernel baseline From 694e52774516261c2cdbda83312af5b7b40f86f0 Mon Sep 17 00:00:00 2001 From: Chloe Crozier Date: Fri, 12 Jun 2026 09:19:34 +0000 Subject: [PATCH 2/7] #17 - Doc sync and YAML fixes for RTX PRO 6000 benchmarks Adjust sw loopback IP placeholders; expand raw_benchmarking tip block. Signed-off-by: Chloe Crozier --- docs/benchmarks/raw_benchmarking.md | 20 +++++++++++++++++++ ...ri_bench_raw_sw_loopback_rtx_pro_6000.yaml | 4 ++-- 2 files changed, 22 insertions(+), 2 deletions(-) diff --git a/docs/benchmarks/raw_benchmarking.md b/docs/benchmarks/raw_benchmarking.md index 20264b2..04b527d 100644 --- a/docs/benchmarks/raw_benchmarking.md +++ b/docs/benchmarks/raw_benchmarking.md @@ -59,6 +59,26 @@ docker run --rm -it --privileged \ The Spark configs also pin the benchmark application's `bench_tx.cpu_core` / `bench_rx.cpu_core` fields to the high-frequency Cortex-X925 cores. Keep both the DAQIRI queue cores and the application worker cores on cores 16-19 unless you intentionally want a lower-power core in the measurement. +!!! tip "RTX PRO 6000 Blackwell (x86_64 workstation / server)" + + For discrete Blackwell RTX PRO 6000 systems, build with [`CMAKE_CUDA_ARCHITECTURES=120`](../tutorials/bare-metal-cmake-build.md) and use these configs: + + - [`daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml) — software loopback, no NIC required. Validates the GPUDirect build path; throughput is not wire-rate. See measured numbers in [`examples/rtx_pro_6000_baseline.md`](https://github.com/nvidia/daqiri/blob/main/examples/rtx_pro_6000_baseline.md). + - [`daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml) — prefilled dual-port NIC run (dev-box PCIe BDFs; fill `eth_dst_addr` from the rx_port MAC). Port 0 TX on GPU CUDA 0, port 1 RX on GPU CUDA 1. Requires an L2 link between the two ports (QSFP cable, passive loopback optic, or switch). `carrier=1` on both ports does not guarantee they are looped to each other. + - [`daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml) — generic `` template for cross-card or custom topology (800 Gbps target once cabled). + - [`daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml) — experimental same-port TX+RX; failed `daqiri_init` on the reference box. + + Unlike DGX Spark, typical RTX Pro servers expose one PF per physical port — there is no on-chip eswitch shortcut between two PFs on the same port without a link. After a NIC run, confirm whether traffic crossed the wire using `tx_phy_packets` / `rx_phy_packets` in the DPDK extended stats (near zero = on-chip or no wire loop; rising with vport counts = over-the-wire). Full constraints and baseline results: [`rtx_pro_6000_baseline.md`](https://github.com/nvidia/daqiri/blob/main/examples/rtx_pro_6000_baseline.md). + + ```bash + sudo ./daqiri_bench_raw_gpudirect \ + ./examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml --seconds 30 + + # Once p0 and p1 are cabled: + sudo ./daqiri_bench_raw_gpudirect \ + ./examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml --seconds 30 + ``` + #### Cross-host two-DGX-Spark loopback If you have two DGX Sparks cross-cabled p0↔p0 instead of a chassis QSFP loop on one machine, use the `_xhost` configs. Each host runs only its own role, so the YAML on each side configures one port instead of two. Both hosts must already be set up per the [DGX Spark profile](../tutorials/system_configuration.md#dgx-spark-profile), with one adjustment: the `daqiri-tx` (`1.1.1.1/24`) and `daqiri-rx` (`2.2.2.2/24`) nmcli profiles are *split across* the two hosts — bring up `daqiri-tx` on the TX host's p0 and `daqiri-rx` on the RX host's p0, instead of both on one box. diff --git a/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml b/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml index df02770..b514bb3 100644 --- a/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml +++ b/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml @@ -71,7 +71,7 @@ bench_tx: payload_size: 8000 header_size: 64 eth_dst_addr: 00:00:00:00:00:00 - ip_src_addr: 000.000.000.0 - ip_dst_addr: 000.000.000.0 + ip_src_addr: 0.0.0.0 + ip_dst_addr: 0.0.0.0 udp_src_port: 4096 udp_dst_port: 4096 From ad58d11d2d76ef7d014c55a901aff78ee717be54 Mon Sep 17 00:00:00 2001 From: Chloe Crozier Date: Fri, 10 Jul 2026 17:25:29 +0000 Subject: [PATCH 3/7] #17 - Port RTX PRO 6000 benchmark runner Signed-off-by: Chloe Crozier --- examples/CMakeLists.txt | 23 +- examples/bench_pipeline.cu | 204 +++++ examples/bench_pipeline.h | 123 +++ examples/bench_workload.cu | 330 ++++++++ examples/bench_workload.h | 144 ++++ ...iri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml | 52 ++ ...ri_bench_raw_sw_loopback_rtx_pro_6000.yaml | 2 +- ...ri_bench_raw_tx_only_rtx_pro_6000_nic.yaml | 46 ++ ...qiri_bench_raw_tx_rx_hds_rtx_pro_6000.yaml | 93 +++ .../daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml | 104 +++ examples/raw_bench_common.cpp | 105 ++- examples/raw_bench_common.h | 26 +- examples/raw_gpudirect_bench.cpp | 41 +- examples/raw_hds_bench.cpp | 48 +- examples/rdma_bench.cpp | 134 ++- examples/run_rtx_pro_bench.sh | 774 ++++++++++++++++++ examples/socket_bench.cpp | 113 ++- scripts/discover_rtx_pro_topology.sh | 149 ++++ src/engines/socket/daqiri_socket_engine.cpp | 15 +- 19 files changed, 2487 insertions(+), 39 deletions(-) create mode 100644 examples/bench_pipeline.cu create mode 100644 examples/bench_pipeline.h create mode 100644 examples/bench_workload.cu create mode 100644 examples/bench_workload.h create mode 100644 examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml create mode 100644 examples/daqiri_bench_raw_tx_only_rtx_pro_6000_nic.yaml create mode 100644 examples/daqiri_bench_raw_tx_rx_hds_rtx_pro_6000.yaml create mode 100644 examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml create mode 100755 examples/run_rtx_pro_bench.sh create mode 100755 scripts/discover_rtx_pro_topology.sh diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 2f4d123..a0b6fce 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -26,6 +26,8 @@ set(DAQIRI_BENCH_CONFIGS daqiri_bench_raw_tx_rx.yaml daqiri_bench_raw_tx_rx_4q.yaml daqiri_bench_raw_tx_rx_hds.yaml + daqiri_bench_raw_tx_rx_hds_rtx_pro_6000.yaml + daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml daqiri_bench_raw_tx_rx_reorder_seq_1024.yaml daqiri_bench_raw_tx_rx_reorder_seq_1024_cpu.yaml daqiri_bench_raw_sw_loopback_reorder_seq_1024.yaml @@ -38,11 +40,13 @@ set(DAQIRI_BENCH_CONFIGS daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml + daqiri_bench_raw_tx_only_rtx_pro_6000_nic.yaml daqiri_example_gds_write_sw_loopback.yaml daqiri_example_gds_write_tx_rx.yaml daqiri_example_pcap_writer_sw_loopback.yaml daqiri_example_pcap_writer_tx_rx.yaml daqiri_bench_rdma_tx_rx.yaml + daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml daqiri_bench_socket_udp_tx_rx.yaml daqiri_bench_socket_tcp_tx_rx.yaml ) @@ -55,7 +59,7 @@ function(link_daqiri_bench target) target_compile_features(${target} PRIVATE cxx_std_17) if(NOT BUILD_SHARED_LIBS AND UNIX AND NOT APPLE) set(DAQIRI_EXAMPLE_STATIC_LIBS daqiri::daqiri) - foreach(engine IN ITEMS dpdk socket rdma) + foreach(engine IN ITEMS dpdk socket rdma ibverbs) if(TARGET daqiri_${engine}) list(APPEND DAQIRI_EXAMPLE_STATIC_LIBS daqiri_${engine}) endif() @@ -73,9 +77,10 @@ function(link_daqiri_bench target) endfunction() function(add_daqiri_raw_bench target source) - add_executable(${target} ${source} raw_bench_common.cpp raw_bench_common_cuda.cu) + add_executable(${target} ${source} raw_bench_common.cpp raw_bench_common_cuda.cu + bench_workload.cu bench_pipeline.cu) link_daqiri_bench(${target}) - target_link_libraries(${target} PRIVATE CUDA::cudart) + target_link_libraries(${target} PRIVATE CUDA::cudart CUDA::cufft CUDA::cublas) set_target_properties(${target} PROPERTIES CUDA_ARCHITECTURES "${DAQIRI_EXAMPLES_CUDA_ARCHITECTURES}" BUILD_RPATH "$ORIGIN/../src;$ORIGIN/../src/third_party/yaml-cpp" @@ -102,17 +107,21 @@ add_daqiri_raw_bench(daqiri_bench_raw_reorder_quantize raw_reorder_quantize_benc add_daqiri_raw_bench(daqiri_example_gds_write gds_write_example.cpp) add_daqiri_raw_bench(daqiri_example_pcap_writer pcap_writer_example.cpp) -add_executable(daqiri_bench_rdma rdma_bench.cpp raw_bench_common.cpp) +add_executable(daqiri_bench_rdma rdma_bench.cpp raw_bench_common.cpp + bench_workload.cu bench_pipeline.cu) link_daqiri_bench(daqiri_bench_rdma) -target_link_libraries(daqiri_bench_rdma PRIVATE CUDA::cudart) +target_link_libraries(daqiri_bench_rdma PRIVATE CUDA::cudart CUDA::cufft CUDA::cublas) set_target_properties(daqiri_bench_rdma PROPERTIES + CUDA_ARCHITECTURES "${DAQIRI_EXAMPLES_CUDA_ARCHITECTURES}" BUILD_RPATH "$ORIGIN/../src;$ORIGIN/../src/third_party/yaml-cpp" ) -add_executable(daqiri_bench_socket socket_bench.cpp raw_bench_common.cpp) +add_executable(daqiri_bench_socket socket_bench.cpp raw_bench_common.cpp + bench_workload.cu bench_pipeline.cu) link_daqiri_bench(daqiri_bench_socket) -target_link_libraries(daqiri_bench_socket PRIVATE CUDA::cudart) +target_link_libraries(daqiri_bench_socket PRIVATE CUDA::cudart CUDA::cufft CUDA::cublas) set_target_properties(daqiri_bench_socket PROPERTIES + CUDA_ARCHITECTURES "${DAQIRI_EXAMPLES_CUDA_ARCHITECTURES}" BUILD_RPATH "$ORIGIN/../src;$ORIGIN/../src/third_party/yaml-cpp" ) diff --git a/examples/bench_pipeline.cu b/examples/bench_pipeline.cu new file mode 100644 index 0000000..62c0128 --- /dev/null +++ b/examples/bench_pipeline.cu @@ -0,0 +1,204 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. + * All rights reserved. SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "bench_pipeline.h" + +#include + +#include +#include +#include + +#include "../src/kernels.h" + +namespace daqiri::bench { + +namespace { +// kReorderDataTypeSame / kReorderEndiannessNetwork from src/kernels.cu -- a pure +// vectorized copy (no quantization), reading the seq number network-byte-order. +constexpr uint8_t kDataTypeSame = 0; +constexpr uint8_t kEndianNetwork = 1; + +cudaStream_t as_stream(void* s) { + return static_cast(s); +} + +__global__ void packet_gather_copy_payload_kernel(void* out, const void* const* in, + uint32_t payload_len, + uint32_t payload_byte_offset, + uint32_t num_pkts) { + const uint32_t pkt_idx = blockIdx.x; + if (pkt_idx >= num_pkts) { + return; + } + + const auto* src_pkt = static_cast(in[pkt_idx]); + if (src_pkt == nullptr) { + return; + } + const auto* src = src_pkt + payload_byte_offset; + auto* dst = static_cast(out) + (static_cast(pkt_idx) * payload_len); + for (uint32_t offset = threadIdx.x; offset < payload_len; offset += blockDim.x) { + dst[offset] = src[offset]; + } +} + +void packet_gather_copy_payload(void* out, const void* const* in, uint32_t payload_len, + uint32_t payload_byte_offset, uint32_t num_pkts, + cudaStream_t stream) { + if (out == nullptr || in == nullptr || payload_len == 0 || num_pkts == 0) { + return; + } + packet_gather_copy_payload_kernel<<>>( + out, in, payload_len, payload_byte_offset, num_pkts); +} +} // namespace + +ReorderPipeline::~ReorderPipeline() { + destroy(); +} + +void ReorderPipeline::destroy() { + if (ordered_ != nullptr) { + cudaFree(ordered_); + ordered_ = nullptr; + } + if (staging_ != nullptr) { + cudaFree(staging_); + staging_ = nullptr; + } + if (dev_ptrs_ != nullptr) { + cudaFree(dev_ptrs_); + dev_ptrs_ = nullptr; + } + if (host_ptrs_ != nullptr) { + delete[] static_cast(host_ptrs_); + host_ptrs_ = nullptr; + } + stream_ = nullptr; // not owned + ok_ = false; +} + +bool ReorderPipeline::init(ReorderMode mode, uint32_t packets_per_batch, uint32_t out_payload_len, + uint32_t payload_byte_offset, uint16_t seq_bit_offset, + uint8_t seq_bit_width, bool staging_needed, void* stream) { + mode_ = mode; + staging_needed_ = staging_needed; + packets_per_batch_ = packets_per_batch; + out_payload_len_ = out_payload_len; + payload_byte_offset_ = payload_byte_offset; + seq_bit_offset_ = seq_bit_offset; + seq_bit_width_ = seq_bit_width; + collected_ = 0; + staged_ = 0; + + if (stream == nullptr || packets_per_batch == 0 || out_payload_len == 0) { + ok_ = false; // disabled (e.g. workload == none): inert no-op, not an error + return true; + } + stream_ = stream; + + batch_bytes_ = static_cast(packets_per_batch) * out_payload_len; + stage_slot_bytes_ = static_cast(payload_byte_offset) + out_payload_len; + + if (cudaMalloc(&ordered_, batch_bytes_) != cudaSuccess || + cudaMemset(ordered_, 0, batch_bytes_) != cudaSuccess) { + std::cerr << "ReorderPipeline: ordered buffer alloc failed\n"; + destroy(); + return false; + } + if (cudaMalloc(&dev_ptrs_, static_cast(packets_per_batch) * sizeof(const void*)) != + cudaSuccess) { + std::cerr << "ReorderPipeline: pointer array alloc failed\n"; + destroy(); + return false; + } + host_ptrs_ = new (std::nothrow) const void*[packets_per_batch]; + if (host_ptrs_ == nullptr) { + std::cerr << "ReorderPipeline: host pointer array alloc failed\n"; + destroy(); + return false; + } + if (staging_needed_) { + if (cudaMalloc(&staging_, static_cast(packets_per_batch) * stage_slot_bytes_) != + cudaSuccess) { + std::cerr << "ReorderPipeline: staging buffer alloc failed\n"; + destroy(); + return false; + } + } + + ok_ = true; + return true; +} + +void ReorderPipeline::reset_batch() { + collected_ = 0; + staged_ = 0; +} + +void ReorderPipeline::add_device_packet(const void* dptr) { + if (!ok_ || dptr == nullptr || collected_ >= packets_per_batch_) { + return; + } + static_cast(host_ptrs_)[collected_++] = dptr; +} + +void* ReorderPipeline::stage_host_packet(const void* hptr, uint32_t len) { + if (!ok_ || !staging_needed_ || hptr == nullptr || staged_ >= packets_per_batch_) { + return nullptr; + } + const uint32_t copy_len = + len < stage_slot_bytes_ ? len : static_cast(stage_slot_bytes_); + auto* dst = static_cast(staging_) + static_cast(staged_) * stage_slot_bytes_; + cudaMemcpyAsync(dst, hptr, copy_len, cudaMemcpyHostToDevice, as_stream(stream_)); + ++staged_; + return dst; +} + +const void* ReorderPipeline::finish_batch() { + if (!ok_ || collected_ == 0) { + return nullptr; + } + + // Single-packet in-order case (one large RoCE/TCP message): the source buffer + // is already contiguous and device-resident, so feed it straight to the + // workload with no kernel and no copy. + if (mode_ == ReorderMode::GatherOnly && collected_ == 1) { + const auto* src = static_cast(static_cast(host_ptrs_)[0]); + return src + payload_byte_offset_; + } + + cudaMemcpyAsync(dev_ptrs_, host_ptrs_, static_cast(collected_) * sizeof(const void*), + cudaMemcpyHostToDevice, as_stream(stream_)); + + const auto* const* in = static_cast(dev_ptrs_); + if (mode_ == ReorderMode::SeqReorder) { + packet_reorder_copy_payload_by_sequence( + ordered_, in, out_payload_len_, out_payload_len_, payload_byte_offset_, collected_, + seq_bit_offset_, seq_bit_width_, /*batch_bit_offset=*/0, + /*batch_bit_width=*/0, /*has_batch_number=*/0, packets_per_batch_, + /*max_slot_idx=*/packets_per_batch_ - 1, kDataTypeSame, kDataTypeSame, kEndianNetwork, + /*batch_id_out=*/nullptr, as_stream(stream_)); + } else { + packet_gather_copy_payload(ordered_, in, out_payload_len_, payload_byte_offset_, collected_, + as_stream(stream_)); + } + return ordered_; +} + +} // namespace daqiri::bench diff --git a/examples/bench_pipeline.h b/examples/bench_pipeline.h new file mode 100644 index 0000000..9393d48 --- /dev/null +++ b/examples/bench_pipeline.h @@ -0,0 +1,123 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. + * All rights reserved. SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include + +namespace daqiri::bench { + +// How a burst's packets are assembled into the contiguous device buffer that the +// GpuWorkload then consumes. +enum class ReorderMode { + // Out-of-order capable transports (DPDK raw, UDP): read a per-packet sequence + // number and place each payload at slot = seq % packets_per_batch. + SeqReorder, + // Reliable / in-order transports (RoCE RC, TCP): place each payload at its + // arrival index (no reshuffle). When a single already-device-resident buffer + // makes up the batch (one large message), this is a zero-copy pass-through. + GatherOnly, +}; + +// Per-RX-thread helper that turns a burst of received packets into one +// contiguous device buffer for the GpuWorkload, doing the reorder/gather (and, +// for sockets, the host->device staging) on the workload's CUDA stream so the +// kernel orders before the workload with no extra synchronization. +// +// CUDA types are hidden behind opaque void* so this header is includable from +// plain .cpp bench mains (same pattern as bench_workload.h). +// +// Usage per burst: +// pipe.reset_batch(); +// for each packet p: +// const void* d = staging_needed ? pipe.stage_host_packet(host_ptr, len) +// : device_ptr; // GPU-accessible src +// pipe.add_device_packet(d); +// const void* ordered = pipe.finish_batch(); // launches kernel +// workload.run(ordered); +class ReorderPipeline { + public: + ReorderPipeline() = default; + ~ReorderPipeline(); + ReorderPipeline(const ReorderPipeline&) = delete; + ReorderPipeline& operator=(const ReorderPipeline&) = delete; + + // packets_per_batch slots of out_payload_len bytes form the contiguous output + // buffer. payload_byte_offset / seq_bit_offset / seq_bit_width describe the + // packet layout (seq fields ignored in GatherOnly). staging_needed allocates a + // device staging buffer for host sources (sockets). stream is the + // GpuWorkload's cudaStream_t (share it); pass nullptr to disable the pipeline. + // Returns false on CUDA error (caller may warn and continue disabled). + bool init(ReorderMode mode, uint32_t packets_per_batch, uint32_t out_payload_len, + uint32_t payload_byte_offset, uint16_t seq_bit_offset, uint8_t seq_bit_width, + bool staging_needed, void* stream); + + // Begin accumulating a new batch. + void reset_batch(); + + // Record one GPU-accessible source pointer (packet base, before + // payload_byte_offset). Ignored once packets_per_batch is reached. + void add_device_packet(const void* dptr); + + // Copy `len` host bytes into the device staging buffer and return the device + // pointer to feed add_device_packet(). Returns nullptr if disabled/full. + void* stage_host_packet(const void* hptr, uint32_t len); + + uint32_t collected() const { + return collected_; + } + + // Launch the reorder/gather kernel for the collected packets and return the + // contiguous ordered device buffer (packets_per_batch * out_payload_len) to + // hand to GpuWorkload::run(). For the single-packet GatherOnly pass-through it + // returns the source pointer directly (no kernel, no copy). Returns nullptr if + // disabled or nothing was collected. + const void* finish_batch(); + + size_t batch_bytes() const { + return batch_bytes_; + } + bool enabled() const { + return ok_; + } + + private: + void destroy(); + + ReorderMode mode_ = ReorderMode::GatherOnly; + bool ok_ = false; + bool staging_needed_ = false; + uint32_t packets_per_batch_ = 0; + uint32_t out_payload_len_ = 0; + uint32_t payload_byte_offset_ = 0; + uint16_t seq_bit_offset_ = 0; + uint8_t seq_bit_width_ = 0; + size_t batch_bytes_ = 0; + size_t stage_slot_bytes_ = 0; + + uint32_t collected_ = 0; + uint32_t staged_ = 0; + + void* stream_ = nullptr; // cudaStream_t (shared, not owned) + void* ordered_ = nullptr; // device output buffer (owned) + void* staging_ = nullptr; // device staging buffer (owned, sockets only) + void* dev_ptrs_ = nullptr; // device void*[packets_per_batch] (owned) + void* host_ptrs_ = nullptr; // host const void*[packets_per_batch] (owned) +}; + +} // namespace daqiri::bench diff --git a/examples/bench_workload.cu b/examples/bench_workload.cu new file mode 100644 index 0000000..e86c70e --- /dev/null +++ b/examples/bench_workload.cu @@ -0,0 +1,330 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. + * All rights reserved. SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "bench_workload.h" + +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +namespace daqiri::bench { +namespace { + +// 1D FFT length; the burst working set is fanned out across as many batched +// transforms of this length as fit. +constexpr int kFftLen = 1024; +// Default working-set size when the caller passes bytes_per_burst == 0. +constexpr size_t kDefaultBytes = 1u << 16; // 64 KiB + +cufftHandle as_fft_plan(int p) { + return static_cast(p); +} +cudaStream_t as_stream(void* s) { + return static_cast(s); +} +cublasHandle_t as_cublas(void* h) { + return static_cast(h); +} + +} // namespace + +BenchWorkload parse_workload(int argc, char** argv) { + BenchWorkload workload = BenchWorkload::None; + for (int i = 2; i + 1 < argc; i += 2) { + if (std::string(argv[i]) == "--workload") { + const std::string val = argv[i + 1]; + if (val == "fft") { + workload = BenchWorkload::Fft; + } else if (val == "gemm") { + workload = BenchWorkload::Gemm; + } else if (val == "gemm_fp16" || val == "gemm-fp16") { + workload = BenchWorkload::GemmFp16; + } else if (val == "none") { + workload = BenchWorkload::None; + } else { + std::cerr << "Unknown --workload value '" << val + << "' (expected none|fft|gemm|gemm_fp16); using none\n"; + workload = BenchWorkload::None; + } + } + } + return workload; +} + +size_t parse_workload_batch_bytes(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") { + 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: + return "fft"; + case BenchWorkload::Gemm: + return "gemm"; + case BenchWorkload::GemmFp16: + return "gemm_fp16"; + case BenchWorkload::None: + default: + return "none"; + } +} + +GpuWorkload::~GpuWorkload() { + destroy(); +} + +void GpuWorkload::destroy() { + if (fft_plan_ >= 0) { + cufftDestroy(as_fft_plan(fft_plan_)); + fft_plan_ = -1; + } + if (cublas_ != nullptr) { + cublasDestroy(as_cublas(cublas_)); + cublas_ = nullptr; + } + if (fft_out_ != nullptr) { + cudaFree(fft_out_); + fft_out_ = nullptr; + } + if (gemm_b_ != nullptr) { + cudaFree(gemm_b_); + gemm_b_ = nullptr; + } + if (gemm_c_ != nullptr) { + cudaFree(gemm_c_); + gemm_c_ = nullptr; + } + if (stream_ != nullptr) { + cudaStreamDestroy(as_stream(stream_)); + stream_ = nullptr; + } + ok_ = false; +} + +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; + if (kind_ == BenchWorkload::None) { + ok_ = false; + return true; // inert no-op object; not an error + } + + const size_t bytes = batch_bytes > 0 ? batch_bytes : kDefaultBytes; + + cudaStream_t stream = nullptr; + if (cudaStreamCreate(&stream) != cudaSuccess) { + std::cerr << "GpuWorkload: cudaStreamCreate failed\n"; + destroy(); + return false; + } + stream_ = stream; + + if (kind_ == BenchWorkload::Fft) { + // Fan the working set out across batched length-kFftLen 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; + + if (cudaMalloc(&fft_out_, total * sizeof(cufftComplex)) != cudaSuccess || + cudaMemset(fft_out_, 0, total * sizeof(cufftComplex)) != cudaSuccess) { + std::cerr << "GpuWorkload: FFT buffer alloc failed\n"; + destroy(); + return false; + } + cufftHandle plan = 0; + if (cufftPlan1d(&plan, kFftLen, CUFFT_C2C, batch) != CUFFT_SUCCESS) { + std::cerr << "GpuWorkload: cufftPlan1d failed\n"; + destroy(); + return false; + } + fft_plan_ = static_cast(plan); + if (cufftSetStream(as_fft_plan(fft_plan_), stream) != CUFFT_SUCCESS) { + std::cerr << "GpuWorkload: cufftSetStream failed\n"; + 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. 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) { + std::cerr << "GpuWorkload: GEMM buffer alloc failed\n"; + destroy(); + return false; + } + cublasHandle_t handle = nullptr; + if (cublasCreate(&handle) != CUBLAS_STATUS_SUCCESS) { + std::cerr << "GpuWorkload: cublasCreate failed\n"; + destroy(); + return false; + } + cublas_ = handle; + if (cublasSetStream(handle, stream) != CUBLAS_STATUS_SUCCESS) { + std::cerr << "GpuWorkload: cublasSetStream failed\n"; + destroy(); + return false; + } + // Explicit shape so every published GEMM number carries its compute size. A + // square nxnxn matmul is 2*n^3 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 + // buffer: this surfaces a misconfigured cuFFT/cuBLAS call (which would + // otherwise no-op silently on the hot path and look like a free workload) and + // excludes one-time library setup from the measured run. The warmup buffer is + // freed immediately; the measured runs read the caller's received-data buffer. + void* warmup_in = nullptr; + const bool warmup_ok = cudaMalloc(&warmup_in, bytes) == cudaSuccess && + cudaMemset(warmup_in, 0, bytes) == cudaSuccess && issue_op(warmup_in) && + cudaStreamSynchronize(stream) == cudaSuccess; + if (warmup_in != nullptr) { + cudaFree(warmup_in); + } + if (!warmup_ok) { + std::cerr << "GpuWorkload: warmup of '" << workload_name(kind_) + << "' failed (cuda=" << cudaGetErrorString(cudaGetLastError()) << ")\n"; + destroy(); + return false; + } + + ok_ = true; + return true; +} + +bool GpuWorkload::issue_op(const void* input) { + if (kind_ == BenchWorkload::Fft) { + return cufftExecC2C(as_fft_plan(fft_plan_), + const_cast(static_cast(input)), + static_cast(fft_out_), CUFFT_FORWARD) == CUFFT_SUCCESS; + } + const float alpha = 1.0f; + const float beta = 0.0f; + if (kind_ == BenchWorkload::Gemm) { + return cublasSgemm(as_cublas(cublas_), CUBLAS_OP_N, CUBLAS_OP_N, gemm_n_, gemm_n_, gemm_n_, + &alpha, static_cast(input), gemm_n_, + static_cast(gemm_b_), gemm_n_, &beta, + static_cast(gemm_c_), gemm_n_) == CUBLAS_STATUS_SUCCESS; + } + // GemmFp16: FP16 inputs, FP32 accumulate, on the tensor cores. + return cublasGemmEx(as_cublas(cublas_), CUBLAS_OP_N, CUBLAS_OP_N, gemm_n_, gemm_n_, gemm_n_, + &alpha, input, CUDA_R_16F, gemm_n_, gemm_b_, CUDA_R_16F, gemm_n_, &beta, + gemm_c_, CUDA_R_16F, gemm_n_, CUBLAS_COMPUTE_32F, + CUBLAS_GEMM_DEFAULT_TENSOR_OP) == CUBLAS_STATUS_SUCCESS; +} + +void GpuWorkload::run(const void* input) { + if (!enabled() || input == nullptr) { + return; + } + issue_op(input); + ++run_count_; + maybe_sync(); +} + +void GpuWorkload::maybe_sync() { + if (!enabled()) { + return; + } + if (run_count_ % static_cast(sync_interval_) == 0) { + cudaStreamSynchronize(as_stream(stream_)); + } +} + +void GpuWorkload::sync() { + if (stream_ != nullptr) { + cudaStreamSynchronize(as_stream(stream_)); + } +} + +} // namespace daqiri::bench diff --git a/examples/bench_workload.h b/examples/bench_workload.h new file mode 100644 index 0000000..6a1fd0c --- /dev/null +++ b/examples/bench_workload.h @@ -0,0 +1,144 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. + * All rights reserved. SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +namespace daqiri::bench { + +// Representative GPU workloads that can be dropped into any benchmark's receive +// path to model downstream GPU compute. Bare loopback (None) vs loopback + FFT +// vs loopback + GEMM. GemmFp16 is the same square matmul as Gemm but in +// mixed-precision (FP16 inputs, FP32 accumulate) on the tensor cores -- the core +// op of GPU inference, and far faster than the FP32 path on tensor-core GPUs. +enum class BenchWorkload { None, Fft, Gemm, GemmFp16 }; + +// Parse "--workload none|fft|gemm|gemm_fp16" from argv (default None). Mirrors +// 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-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); + +// Engine-agnostic representative GPU compute, run once per received burst on the +// ACTUAL received packet data. +// +// The caller hands run() a contiguous device buffer holding the reordered / +// gathered payload of one burst (see ReorderPipeline in bench_pipeline.h). The +// op reads that buffer as its input operand: FFT transforms it in place of the +// scratch input; GEMM uses it as the A matrix. The bytes are arbitrary packet +// data, which is fine for a throughput benchmark -- the FLOP profile and memory +// footprint are unchanged; only the input source differs. This makes the +// measurement an honest end-to-end "receive then process the data" cost. +// +// Owns its own CUDA stream, cuFFT plan, and cuBLAS handle. cuFFT plans and +// cuBLAS handles are not safe to share across threads, so construct one +// GpuWorkload per RX worker thread (each multi-queue RX thread gets its own). +// The stream is shared with the ReorderPipeline so the reorder/gather kernel and +// the workload are serialized on the same stream without an explicit sync. +class GpuWorkload { + public: + GpuWorkload() = default; + ~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); + + // Enqueue one representative FFT/SGEMM on the internal stream, reading `input` + // (a device pointer to >= batch_bytes 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 + // the critical path without unbounded queueing. + void maybe_sync(); + + // Drain any remaining queued work (call once on shutdown). + void sync(); + + bool enabled() const { + return kind_ != BenchWorkload::None && ok_; + } + BenchWorkload kind() const { + return kind_; + } + + // CUDA stream (cudaStream_t) this workload runs on; share it with the + // ReorderPipeline so the reorder/gather kernel orders before run(). null when + // the workload is disabled. + void* stream() const { + return stream_; + } + + private: + void destroy(); + // Enqueue one op on the stream reading `input`; returns false if the + // cuFFT/cuBLAS call reports an error at enqueue. Shared by run() (hot path) + // and init() (warmup/validate, against an internal zeroed buffer). + bool issue_op(const void* input); + + BenchWorkload kind_ = BenchWorkload::None; + bool ok_ = false; + int sync_interval_ = 16; + unsigned long long run_count_ = 0; + + // Opaque handles, cast in the .cu so this header stays free of CUDA library + // includes (it is included from plain .cpp bench mains). + void* stream_ = nullptr; // cudaStream_t + void* cublas_ = nullptr; // cublasHandle_t + int fft_plan_ = -1; // cufftHandle (-1 == unset) + + // Device scratch (operands NOT sourced from received data). + void* fft_out_ = nullptr; // cufftComplex[fft_total_] (FFT output) + void* gemm_b_ = nullptr; // B operand + void* gemm_c_ = nullptr; // C output + int gemm_n_ = 0; +}; + +} // namespace daqiri::bench diff --git a/examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml b/examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml new file mode 100644 index 0000000..103ddd1 --- /dev/null +++ b/examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml @@ -0,0 +1,52 @@ +# RTX PRO 6000 Blackwell -- ibverbs MPRQ RX-only (host-staged). +# +# Selected via `engine: "ibverbs"` under `stream_type: "raw"`. Uses HOST hugepage +# memory (GPUDirect device registration is not yet available on this backend). +# FFT/GEMM workloads therefore include host->device staging cost -- label results +# as `host-staged` when comparing against DPDK HDS. +# +# RX-only: pair with a DPDK TX sender (run_rtx_pro_bench.sh ibverbs mode starts +# daqiri_bench_raw_gpudirect on tx_port in the background). +# +# Run (manual): +# sudo ./build/examples/daqiri_bench_raw_gpudirect \ +# ./build/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml --seconds 30 & +# sudo ./build/examples/daqiri_bench_raw_gpudirect \ +# ./build/examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml \ +# --seconds 30 --workload fft +# +%YAML 1.2 +--- +daqiri: + cfg: + version: 1 + stream_type: "raw" + engine: "ibverbs" + master_core: 3 + debug: false + log_level: "info" + loopback: "" + + memory_regions: + - name: "Data_RX_HOST" + kind: "huge" + affinity: 1 + num_bufs: 4096 + buf_size: 8064 + + interfaces: + - name: "rx_port" + address: 0000:61:00.1 + rx: + queues: + - name: "rx_q_0" + id: 0 + cpu_core: 9 + batch_size: 1024 + timeout_us: 2000 + memory_regions: + - "Data_RX_HOST" + +bench_rx: + interface_name: "rx_port" + cpu_core: 8 diff --git a/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml b/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml index b514bb3..89a8ce9 100644 --- a/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml +++ b/examples/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml @@ -5,7 +5,7 @@ # # Build (native sm_120): # cmake -S . -B build -DBUILD_SHARED_LIBS=ON -DDAQIRI_BUILD_PYTHON=OFF \ -# -DDAQIRI_MGR="dpdk socket rdma" -DCMAKE_CUDA_ARCHITECTURES=120 +# -DDAQIRI_ENGINE="dpdk ibverbs" -DCMAKE_CUDA_ARCHITECTURES=120 # cmake --build build -j # # Run: diff --git a/examples/daqiri_bench_raw_tx_only_rtx_pro_6000_nic.yaml b/examples/daqiri_bench_raw_tx_only_rtx_pro_6000_nic.yaml new file mode 100644 index 0000000..0d74cac --- /dev/null +++ b/examples/daqiri_bench_raw_tx_only_rtx_pro_6000_nic.yaml @@ -0,0 +1,46 @@ +# RTX PRO 6000 -- DPDK GPUDirect TX-only sender for ibverbs RX pairing. +# Binds a single TX port; the RX port stays kernel-bound for ibverbs MPRQ. +# +%YAML 1.2 +--- +daqiri: + cfg: + version: 1 + stream_type: "raw" + master_core: 3 + debug: false + log_level: "info" + loopback: "" + + memory_regions: + - name: "Data_TX_GPU" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + + interfaces: + - name: "tx_port" + address: 0000:61:00.0 + tx: + queues: + - name: "tx_q_0" + id: 0 + batch_size: 10240 + cpu_core: 11 + memory_regions: + - "Data_TX_GPU" + offloads: + - "tx_eth_src" + +bench_tx: + interface_name: "tx_port" + cpu_core: 10 + batch_size: 10240 + payload_size: 8000 + header_size: 64 + eth_dst_addr: <00:00:00:00:00:00> + ip_src_addr: 1.2.3.4 + ip_dst_addr: 5.6.7.8 + udp_src_port: 4096 + udp_dst_port: 4096 diff --git a/examples/daqiri_bench_raw_tx_rx_hds_rtx_pro_6000.yaml b/examples/daqiri_bench_raw_tx_rx_hds_rtx_pro_6000.yaml new file mode 100644 index 0000000..ce6c12f --- /dev/null +++ b/examples/daqiri_bench_raw_tx_rx_hds_rtx_pro_6000.yaml @@ -0,0 +1,93 @@ +# RTX PRO 6000 -- Header-Data Split wire closed-loop (dual-port L2 loop). +# TX port + RX port must have L2 connectivity. Override BDFs/MAC via discovery +# or run_rtx_pro_bench.sh --tx-bdf / --rx-bdf / --rx-mac. +# +# Run: +# sudo ./build/examples/daqiri_bench_raw_hds --seconds 30 --workload fft +# sudo ./examples/run_rtx_pro_bench.sh dpdk-hds issue17 --seconds 30 +# +%YAML 1.2 +--- +daqiri: + cfg: + version: 1 + stream_type: "raw" + master_core: 3 + debug: false + log_level: "info" + loopback: "" + + memory_regions: + - name: "Data_TX_CPU" + kind: "huge" + affinity: 0 + num_bufs: 51200 + buf_size: 64 + - name: "Data_TX_GPU" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + - name: "Data_RX_CPU" + kind: "huge" + affinity: 1 + num_bufs: 51200 + buf_size: 64 + - name: "Data_RX_GPU" + kind: "device" + affinity: 1 + num_bufs: 51200 + buf_size: 8000 + + interfaces: + - name: "tx_port" + address: 0000:61:00.0 + tx: + queues: + - name: "tx_q_0" + id: 0 + batch_size: 4096 + cpu_core: 11 + memory_regions: + - "Data_TX_CPU" + - "Data_TX_GPU" + offloads: + - "tx_eth_src" + - name: "rx_port" + address: 0000:61:00.1 + rx: + flow_isolation: true + queues: + - name: "rq_q_0" + id: 0 + cpu_core: 9 + batch_size: 4096 + memory_regions: + - "Data_RX_CPU" + - "Data_RX_GPU" + flows: + - name: "flow_0" + id: 0 + action: + type: queue + id: 0 + match: + udp_src: 4096 + udp_dst: 4096 + ipv4_len: 8050 + +bench_rx: + interface_name: "rx_port" + cpu_core: 8 + +bench_tx: + interface_name: "tx_port" + cpu_core: 10 + batch_size: 4096 + payload_size: 8000 + header_size: 64 + eth_dst_addr: <00:00:00:00:00:00> + ip_src_addr: 1.2.3.4 + ip_dst_addr: 5.6.7.8 + udp_src_port: 4096 + udp_dst_port: 4096 diff --git a/examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml b/examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml new file mode 100644 index 0000000..0d01db1 --- /dev/null +++ b/examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml @@ -0,0 +1,104 @@ +# RTX PRO 6000 Blackwell -- RoCE closed-loop (client + server in one process). +# Uses host_pinned memory by default. Device memory RoCE requires a working GPUDirect RDMA MR path (nvidia-peermem or dma-buf support); use host_pinned as the portable fallback. +# +# Replace client/server IPs with addresses on your RoCE-capable interfaces. +# Default uses link-local style addresses on the bench host's mlx5 ports. +# +# Run: +# sudo ./build/examples/daqiri_bench_rdma \ +# ./build/examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml \ +# --seconds 30 --mode both --workload gemm +# +%YAML 1.2 +--- +daqiri: + cfg: + version: 1 + stream_type: "socket" + protocol: "roce" + master_core: 3 + debug: false + log_level: "info" + + memory_regions: + - name: "DATA_RX_GPU_SERVER" + kind: "host_pinned" + affinity: 1 + num_bufs: 20 + buf_size: 9000000 + - name: "DATA_TX_GPU_SERVER" + kind: "host_pinned" + affinity: 1 + num_bufs: 20 + buf_size: 9000000 + - name: "DATA_TX_GPU_CLIENT" + kind: "host_pinned" + affinity: 0 + num_bufs: 20 + buf_size: 9000000 + - name: "DATA_RX_GPU_CLIENT" + kind: "host_pinned" + affinity: 0 + num_bufs: 20 + buf_size: 9000000 + + interfaces: + - name: my_client + address: 10.100.1.1 + socket_config: + mode: client + remote_ip: 10.100.3.1 + remote_port: 4096 + roce_config: + transport_mode: RC + tx: + queues: + - name: "Client_TX_Queue" + id: 0 + batch_size: 1 + cpu_core: 7 + rx: + queues: + - name: "Client_RX_Queue" + id: 0 + cpu_core: 7 + batch_size: 1 + - name: my_server + address: 10.100.3.1 + socket_config: + mode: server + local_ip: 10.100.3.1 + local_port: 4096 + roce_config: + transport_mode: RC + rx: + queues: + - name: "Server_RX_Queue" + id: 0 + cpu_core: 8 + batch_size: 1 + tx: + queues: + - name: "Server_TX_Queue" + id: 0 + cpu_core: 8 + batch_size: 1 + +rdma_bench_server: + server_address: 10.100.3.1 + server_port: 4096 + message_size: 8000000 + send: true + receive: true + cpu_core: 9 + server: true + +rdma_bench_client: + message_size: 8000000 + client_address: 10.100.1.1 + server_address: 10.100.3.1 + server_port: 4096 + receive: true + send: true + cpu_core: 10 + server: false diff --git a/examples/raw_bench_common.cpp b/examples/raw_bench_common.cpp index e46847b..61fa8d7 100644 --- a/examples/raw_bench_common.cpp +++ b/examples/raw_bench_common.cpp @@ -22,6 +22,7 @@ #include #include +#include #include #include #include @@ -548,12 +549,70 @@ void print_queue_stats(const char *direction, const std::string &interface_name, << std::endl; } -void rx_count_worker(const RawBenchRxConfig &cfg, std::atomic &stop) { +void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, BenchWorkload workload, + 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; } + // Per-thread representative GPU workload + double-buffered reorder slots + // (both inert unless --workload is set). Each slot owns its ordered output + // buffer and holds the source DAQIRI burst until the slot's stream work is + // complete, so RX can enqueue GPU work without a per-burst synchronization. + constexpr size_t kWorkloadSlots = 2; + const uint32_t out_payload_len = geom.out_payload_len > 0 ? geom.out_payload_len : 1; + const size_t batch_bytes = static_cast(geom.packets_per_batch) * out_payload_len; + GpuWorkload gpu_workload; + struct WorkloadSlot { + ReorderPipeline pipeline; + cudaEvent_t done = nullptr; + daqiri::BurstParams *burst = nullptr; + bool busy = false; + }; + std::array workload_slots; + bool run_workload = false; + if (workload != BenchWorkload::None) { + bool ok = gpu_workload.init(workload, batch_bytes, workload_sync_interval, workload_gemm_n); + for (auto &slot : workload_slots) { + ok = ok && slot.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()) && + cudaEventCreateWithFlags(&slot.done, cudaEventDisableTiming) == cudaSuccess; + } + if (!ok) { + std::cerr << "RX workload init failed; continuing without GPU workload\n"; + for (auto &slot : workload_slots) { + if (slot.done != nullptr) { + cudaEventDestroy(slot.done); + slot.done = nullptr; + } + } + gpu_workload.sync(); + } else { + run_workload = true; + } + } + + auto release_slot = [](WorkloadSlot &slot, bool wait) { + if (!slot.busy) { + return false; + } + const cudaError_t status = wait ? cudaEventSynchronize(slot.done) : cudaEventQuery(slot.done); + if (status == cudaSuccess) { + daqiri::free_all_packets_and_burst_rx(slot.burst); + slot.burst = nullptr; + slot.busy = false; + return true; + } + if (status != cudaErrorNotReady) { + std::cerr << "RX workload CUDA event error: " << cudaGetErrorString(status) << "\n"; + } + return false; + }; + size_t next_slot = 0; + const int port_id = daqiri::get_port_id(cfg.interface_name); if (port_id < 0) { std::cerr << "Invalid RX interface_name: " << cfg.interface_name << "\n"; @@ -591,15 +650,55 @@ void rx_count_worker(const RawBenchRxConfig &cfg, std::atomic &stop) { } got_any = true; auto &stats = queue_stats[static_cast(q)]; - stats.packets += static_cast(daqiri::get_num_packets(burst)); + const int num_pkts = static_cast(daqiri::get_num_packets(burst)); + stats.packets += static_cast(num_pkts); stats.bytes += daqiri::get_burst_tot_byte(burst); ++stats.bursts; - daqiri::free_all_packets_and_burst_rx(burst); + + if (run_workload) { + for (auto &slot : workload_slots) { + release_slot(slot, false); + } + + auto &slot = workload_slots[next_slot]; + release_slot(slot, true); + next_slot = (next_slot + 1) % workload_slots.size(); + + slot.pipeline.reset_batch(); + bool enqueued_work = false; + for (int i = 0; i < num_pkts; ++i) { + slot.pipeline.add_device_packet( + daqiri::get_segment_packet_ptr(burst, geom.payload_segment, i)); + if (slot.pipeline.collected() == geom.packets_per_batch) { + gpu_workload.run(slot.pipeline.finish_batch()); + slot.pipeline.reset_batch(); + enqueued_work = true; + } + } + + if (enqueued_work) { + cudaEventRecord(slot.done, static_cast(gpu_workload.stream())); + slot.burst = burst; + slot.busy = true; + } else { + daqiri::free_all_packets_and_burst_rx(burst); + } + } else { + daqiri::free_all_packets_and_burst_rx(burst); + } } if (!got_any) { std::this_thread::sleep_for(std::chrono::microseconds(100)); } } + gpu_workload.sync(); + for (auto &slot : workload_slots) { + release_slot(slot, true); + if (slot.done != nullptr) { + cudaEventDestroy(slot.done); + slot.done = nullptr; + } + } const double secs = std::chrono::duration(std::chrono::steady_clock::now() - t0) .count(); diff --git a/examples/raw_bench_common.h b/examples/raw_bench_common.h index 17e9705..496a2c7 100644 --- a/examples/raw_bench_common.h +++ b/examples/raw_bench_common.h @@ -20,6 +20,9 @@ #include #include +#include "bench_pipeline.h" +#include "bench_workload.h" + #include #include #include @@ -137,6 +140,27 @@ void wait_for_stop(int run_seconds, std::atomic &stop); void print_queue_stats(const char *direction, const std::string &interface_name, int queue_id, const RawBenchQueueStats &stats, double seconds); -void rx_count_worker(const RawBenchRxConfig &cfg, std::atomic &stop); +// Describes how received packets are reordered/gathered into the contiguous +// device buffer the GpuWorkload consumes. Defaults match the raw native shape; +// the bench main fills it from the TX config. +struct ReorderGeometry { + uint32_t packets_per_batch = 1024; // slots in the ordered buffer + uint32_t out_payload_len = 0; // bytes/slot (0 => use payload_size) + uint32_t payload_byte_offset = 0; // offset of payload within the packet + uint16_t seq_bit_offset = 0; // bit offset of the seq number + uint8_t seq_bit_width = 32; // seq number width (bits) + int payload_segment = 0; // packet segment carrying payload (1 = HDS) +}; + +// When `workload != None`, each RX thread builds its own GpuWorkload (cuFFT/ +// cuBLAS handles are not thread-safe to share) and its own ReorderPipeline, then +// for each received burst reorders the real packet payloads into a contiguous +// device buffer (geom) and runs one representative compute on THAT buffer. The +// stream is shared so the reorder kernel orders before the workload; the burst +// 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 = {}, + 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 4240c05..0181454 100644 --- a/examples/raw_gpudirect_bench.cpp +++ b/examples/raw_gpudirect_bench.cpp @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -112,6 +113,14 @@ void tx_worker(const daqiri::bench::RawBenchTxConfig &cfg, eth_dst, ip_src, ip_dst, src_port, dst_port); std::memcpy(packet_template.data() + cfg.header_size, payload_template.data(), cfg.payload_size); + // Inject a per-packet sequence number (network byte order) at the start + // of the payload so the RX-side reorder workload has a real seq to sort + // by. The buffer is initialized once and reused, so seq == the packet's + // fixed index within the batch. + if (cfg.payload_size >= sizeof(uint32_t)) { + const uint32_t seq = htonl(static_cast(i)); + std::memcpy(packet_template.data() + cfg.header_size, &seq, sizeof(seq)); + } daqiri::bench::finalize_udp_ipv4_checksums(packet_template.data()); if (cudaMemcpy(gpu_pkt, packet_template.data(), packet_template.size(), cudaMemcpyHostToDevice) != cudaSuccess) { @@ -153,7 +162,8 @@ void tx_worker(const daqiri::bench::RawBenchTxConfig &cfg, int main(int argc, char **argv) { if (argc < 2) { std::cerr << "Usage: " << argv[0] - << " [--seconds N] [--target-gbps G]\n"; + << " [--seconds N] [--target-gbps G] " + "[--workload none|fft|gemm|gemm_fp16] [--workload-batch-bytes N]\n"; return 1; } @@ -161,6 +171,10 @@ int main(int argc, char **argv) { daqiri::bench::grafana::init_prometheus_metrics_from_env(); 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_sync_interval = daqiri::bench::parse_workload_sync_interval(argc, argv); const auto root = YAML::LoadFile(argv[1]); std::vector rx_configs; @@ -188,10 +202,31 @@ int main(int argc, char **argv) { std::vector rx_threads; daqiri::bench::TokenBucketPacer tx_pacer(target_gbps); + // Reorder geometry for the real-data workload: the seq number sits at the + // payload start (just after the UDP/IP header), and a window of up to 1024 + // packets is reordered into a contiguous ~8 MB device buffer (matching the + // RoCE 8 MB working set) that the compute then consumes. + daqiri::bench::ReorderGeometry geom; + if (!tx_configs.empty()) { + const auto& tx = tx_configs.front(); + geom.payload_segment = 0; + geom.payload_byte_offset = tx.header_size; + 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); + } 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)); + 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 250b57b..6de5298 100644 --- a/examples/raw_hds_bench.cpp +++ b/examples/raw_hds_bench.cpp @@ -21,6 +21,7 @@ #include #include +#include #include #include #include @@ -69,6 +70,8 @@ void tx_worker(const daqiri::bench::RawBenchTxConfig &cfg, } std::unordered_set initialized_payload_buffers; + daqiri::bench::RawBenchQueueStats stats; + const auto t0 = std::chrono::steady_clock::now(); while (!stop.load()) { auto *msg = daqiri::create_tx_burst_params(); @@ -130,19 +133,36 @@ void tx_worker(const daqiri::bench::RawBenchTxConfig &cfg, daqiri::free_all_packets_and_burst_tx(msg); continue; } - daqiri::send_tx_burst(msg); + if (daqiri::send_tx_burst(msg) == daqiri::Status::SUCCESS) { + stats.packets += static_cast(num_pkts); + stats.bytes += static_cast(num_pkts) * + static_cast(cfg.header_size + cfg.payload_size); + ++stats.bursts; + } } + + const double secs = + std::chrono::duration(std::chrono::steady_clock::now() - t0) + .count(); + daqiri::bench::print_queue_stats("TX", cfg.interface_name, cfg.queue_id, stats, + secs); } } // namespace int main(int argc, char **argv) { if (argc < 2) { - std::cerr << "Usage: " << argv[0] << " [--seconds N]\n"; + std::cerr << "Usage: " << argv[0] + << " [--seconds N] " + "[--workload none|fft|gemm|gemm_fp16] [--workload-batch-bytes 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_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"; @@ -162,8 +182,28 @@ int main(int argc, char **argv) { std::thread rx_thread; if (has_rx) { - rx_thread = std::thread(daqiri::bench::rx_count_worker, - daqiri::bench::parse_rx(root), std::ref(stop)); + // HDS reorder geometry: the payload lives in segment 1 (GPU memory), so the + // reorder reads it from offset 0 of that segment. The HDS TX does not inject + // a per-packet sequence number, so the seq-based reorder still exercises the + // kernel but mostly collides on slot 0 -- fine for a throughput benchmark + // (the FLOP/copy volume is unchanged). + daqiri::bench::ReorderGeometry geom; + if (has_tx) { + const auto tx = daqiri::bench::parse_tx(root); + geom.payload_segment = 1; + geom.payload_byte_offset = 0; + 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); + } + 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 b1beb4a..00d7266 100644 --- a/examples/rdma_bench.cpp +++ b/examples/rdma_bench.cpp @@ -17,6 +17,7 @@ #include +#include #include #include #include @@ -25,6 +26,7 @@ #include #include #include +#include #include "raw_bench_common.h" #include @@ -78,7 +80,9 @@ RdmaBenchConfig parse_rdma_cfg(const YAML::Node& node) { } void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pacer, - std::atomic& stop, RdmaWorkerStats& stats) { + std::atomic& stop, RdmaWorkerStats& stats, + 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)) { @@ -86,6 +90,55 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa return; } + // Representative GPU workload run on each received message's REAL data (no-op + // unless --workload set). RoCE/RC delivers one large in-order message, so the + // recv buffer is already contiguous and (host_pinned on the integrated GB10) + // 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. + 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) || + !pipeline.init(daqiri::bench::ReorderMode::GatherOnly, + /*packets_per_batch=*/1, msg, + /*payload_byte_offset=*/0, /*seq_bit_offset=*/0, + /*seq_bit_width=*/0, /*staging_needed=*/false, gpu_workload.stream())) { + std::cerr << "RDMA workload init failed; continuing without GPU workload\n"; + } + } + 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 = + run_workload ? std::max(1, std::min(cfg.rx_depth / 2, 8)) : 0; + auto flush_held_recv = [&]() { + if (held_recv.empty()) { + return; + } + gpu_workload.sync(); + for (auto* held : held_recv) { + daqiri::free_tx_burst(held); + } + held_recv.clear(); + }; + int outstanding_send = 0; int outstanding_recv = 0; uint64_t send_wr_id = 0x1234; @@ -93,8 +146,24 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa uintptr_t conn_id = 0; std::string send_mr = cfg.server ? "DATA_TX_GPU_SERVER" : "DATA_TX_GPU_CLIENT"; std::string recv_mr = cfg.server ? "DATA_RX_GPU_SERVER" : "DATA_RX_GPU_CLIENT"; + bool recv_primed = !cfg.receive; + + // After `stop` is signalled, drain in-flight completions for a short window so + // pending SENDs land cleanly instead of getting WR_FLUSH_ERR on disconnect. + auto drain_deadline = std::chrono::steady_clock::time_point::max(); + const auto drain_window = std::chrono::milliseconds(500); - while (!stop.load()) { + while (true) { + const bool stopped = stop.load(); + if (stopped && drain_deadline == std::chrono::steady_clock::time_point::max()) { + drain_deadline = std::chrono::steady_clock::now() + drain_window; + } + if (stopped) { + if (outstanding_send == 0 && outstanding_recv == 0) { break; } + if (std::chrono::steady_clock::now() >= drain_deadline) { break; } + } + + if (conn_id == 0 && stopped) { break; } if (conn_id == 0) { daqiri::Status s = daqiri::Status::GENERIC_FAILURE; if (cfg.server) { @@ -107,6 +176,7 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa std::this_thread::sleep_for(std::chrono::milliseconds(200)); continue; } + recv_primed = !cfg.receive; } auto post_req = [&](int& outstanding, int depth, uint64_t& wr_id, daqiri::RDMAOpCode op, @@ -156,7 +226,20 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa return posted; }; - bool posted_work = refill_receives(); + if (!recv_primed && !stopped) { + int idle_spins = 0; + while (!stop.load() && outstanding_recv < cfg.rx_depth) { + if (refill_receives()) { + idle_spins = 0; + } else { + if (++idle_spins > 100) { break; } + std::this_thread::sleep_for(std::chrono::microseconds(100)); + } + } + recv_primed = !cfg.receive || outstanding_recv > 0; + } + + bool posted_work = stopped ? false : refill_receives(); bool got_completion = false; while (true) { daqiri::BurstParams* completion = nullptr; @@ -173,13 +256,31 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa outstanding_recv--; stats.recv_completions++; stats.recv_bytes += static_cast(cfg.message_size); + 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); + } + // 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(); + } + continue; // do not free yet; freed by flush_held_recv() + } } daqiri::free_tx_burst(completion); } - posted_work = refill_receives() || posted_work; + if (!stopped) { posted_work = refill_receives() || posted_work; } const bool local_receive_window_ready = !cfg.receive || outstanding_recv > 0; - if (cfg.send && local_receive_window_ready) { + if (!stopped && cfg.send && local_receive_window_ready) { while (outstanding_send < cfg.tx_depth && post_req(outstanding_send, cfg.tx_depth, send_wr_id, daqiri::RDMAOpCode::SEND, send_mr)) { @@ -188,9 +289,14 @@ 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(); std::this_thread::sleep_for(std::chrono::microseconds(100)); } } + flush_held_recv(); + gpu_workload.sync(); } } // namespace @@ -198,7 +304,9 @@ void rdma_worker(const RdmaBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pa int main(int argc, char** argv) { if (argc < 2) { std::cerr << "Usage: " << argv[0] - << " [--seconds N] [--mode server|client|both] [--target-gbps G]\n"; + << " [--seconds N] [--mode server|client|both] " + "[--target-gbps G] [--workload none|fft|gemm|gemm_fp16] " + "[--workload-batch-bytes N]\n"; return 1; } @@ -214,6 +322,10 @@ int main(int argc, char** argv) { target_gbps = std::stod(argv[i + 1]); } } + 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) { @@ -249,12 +361,14 @@ 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)); + 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); } if (run_client) { - client_thread = std::thread(rdma_worker, client_cfg, std::ref(client_pacer), - std::ref(stop), std::ref(client_stats)); + 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); } if (!server_thread.joinable() && !client_thread.joinable()) { diff --git a/examples/run_rtx_pro_bench.sh b/examples/run_rtx_pro_bench.sh new file mode 100755 index 0000000..faa278a --- /dev/null +++ b/examples/run_rtx_pro_bench.sh @@ -0,0 +1,774 @@ +#!/usr/bin/env bash +# +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# +# RTX PRO 6000 benchmark runner (discrete dGPU, arch 120). +# Primary: dual-port L2 wire closed-loop (real NIC traffic, any host with a link). +# Optional: sw-smoke (no NIC, build sanity only). +# Adapts examples/run_spark_bench.sh; override topology via discovery or CLI. +# +# Usage: +# ./run_rtx_pro_bench.sh [options] +# +# backend: dpdk | dpdk-hds | rdma | roce | ibverbs | socket-udp | socket-tcp +# mode: nic-smoke | issue17 | sweep | drop-curve | drop-curve-matrix | +# sw-smoke | build-only +# +# nic-smoke : wire closed-loop sanity (TX port -> RX port over L2) +# issue17 : none/fft/gemm on wire +# sw-smoke : optional build sanity (no NIC, non-wire) +# +# Options: +# --seconds N +# --workload none|fft|gemm|gemm_fp16 +# --build-targets t1,t2,... (build-only) +# --tx-bdf, --rx-bdf, --rx-mac, --tx-gpu, --rx-gpu +# +# Environment: +# DAQIRI_BUILD_DIR, ETH_DST_ADDR (required for dpdk wire modes) +# WORKLOAD, WORKLOAD_BATCH, GEMM_N, SYNC_INTERVAL (GPU post-process) +# RTX_TX_BDF, RTX_RX_BDF (defaults 61:00.0 / 61:00.1) +# +# Run inside the project container as root (per AGENTS.md). + +set -euo pipefail + +SCRIPT_DIR="$(cd "$(dirname "$0")" && pwd)" +REPO_ROOT="$(cd "$SCRIPT_DIR/.." && pwd)" +BUILD_DIR="${DAQIRI_BUILD_DIR:-$REPO_ROOT/build}" +DISCOVER="$REPO_ROOT/scripts/discover_rtx_pro_topology.sh" + +BACKEND="${1:-}" +MODE="${2:-}" +shift 2 2>/dev/null || true + +RUN_SECONDS=30 +BUILD_TARGETS="raw_gpudirect,raw_hds,raw_reorder_seq,raw_reorder_quantize,rdma,socket" +TX_BDF="" +RX_BDF="" +RX_MAC="" +TX_GPU="" +RX_GPU="" +CLI_WORKLOAD="" + +while [[ $# -gt 0 ]]; do + case "$1" in + --seconds) RUN_SECONDS="$2"; shift 2 ;; + --workload) CLI_WORKLOAD="$2"; shift 2 ;; + --build-targets) BUILD_TARGETS="$2"; shift 2 ;; + --tx-bdf) TX_BDF="$2"; shift 2 ;; + --rx-bdf) RX_BDF="$2"; shift 2 ;; + --rx-mac) RX_MAC="$2"; shift 2 ;; + --tx-gpu) TX_GPU="$2"; shift 2 ;; + --rx-gpu) RX_GPU="$2"; shift 2 ;; + -h|--help) + sed -n '1,40p' "$0" + exit 0 + ;; + *) + echo "Unknown option: $1" >&2 + exit 1 + ;; + esac +done + +if [[ -z "$BACKEND" || -z "$MODE" ]]; then + echo "Usage: $0 [--seconds N] [--workload none|fft|gemm] ..." >&2 + exit 1 +fi + +# roce is an alias for rdma; smoke is an alias for nic-smoke. +[[ "$BACKEND" == "roce" ]] && BACKEND="rdma" +if [[ "$MODE" == "smoke" ]]; then + MODE="nic-smoke" +fi +[[ "$MODE" == "workload-smoke" ]] && MODE="issue17" +# sw-workload-smoke is the no-cable analogue of workload-smoke/issue17. +[[ "$MODE" == "sw-workload-smoke" ]] && MODE="sw-issue17" + +# Software-loopback modes need no NIC/cable; all "sw-*" modes share this path. +# Detecting it from the mode (not a single literal) keeps the wire-only guards +# below correct for every current and future SW mode. +IS_SW=0 +[[ "$MODE" == sw-* ]] && IS_SW=1 + +WORKLOAD="${CLI_WORKLOAD:-${WORKLOAD:-none}}" +case "$WORKLOAD" in + none|fft|gemm|gemm_fp16) ;; + *) + echo "Invalid WORKLOAD '$WORKLOAD' (expected none|fft|gemm|gemm_fp16)" >&2 + exit 1 + ;; +esac +WORKLOAD_BATCH="${WORKLOAD_BATCH:-}" +GEMM_N="${GEMM_N:-}" +SYNC_INTERVAL="${SYNC_INTERVAL:-}" + +# -------------------------------------------------------------------------- +# build-only +# -------------------------------------------------------------------------- +if [[ "$MODE" == "build-only" ]]; then + IFS=',' read -r -a targets <<< "$BUILD_TARGETS" + cmake_args=( + -S "$REPO_ROOT" -B "$BUILD_DIR" + -DBUILD_SHARED_LIBS=ON + -DDAQIRI_BUILD_PYTHON=OFF + -DDAQIRI_BUILD_EXAMPLES=ON + -DDAQIRI_ENGINE="dpdk ibverbs" + -DCMAKE_CUDA_ARCHITECTURES=120 + ) + echo "Configuring: cmake ${cmake_args[*]}" + cmake "${cmake_args[@]}" + + declare -A target_map=( + [raw_gpudirect]=daqiri_bench_raw_gpudirect + [raw_hds]=daqiri_bench_raw_hds + [raw_reorder_seq]=daqiri_bench_raw_reorder_seq + [raw_reorder_quantize]=daqiri_bench_raw_reorder_quantize + [rdma]=daqiri_bench_rdma + [socket]=daqiri_bench_socket + ) + build_list=() + for t in "${targets[@]}"; do + t="$(echo "$t" | xargs)" + [[ -z "$t" ]] && continue + if [[ -n "${target_map[$t]:-}" ]]; then + build_list+=("${target_map[$t]}") + else + echo "Unknown build target: $t" >&2 + exit 1 + fi + done + echo "Building: ${build_list[*]}" + cmake --build "$BUILD_DIR" --target "${build_list[@]}" -j"$(nproc)" + exit $? +fi + +# -------------------------------------------------------------------------- +# Discovery + wire preflight +# -------------------------------------------------------------------------- +if [[ -x "$DISCOVER" ]]; then + # shellcheck disable=SC1090 + source "$DISCOVER" +fi +[[ -n "$TX_BDF" ]] && export RTX_TX_BDF="$TX_BDF" +[[ -n "$RX_BDF" ]] && export RTX_RX_BDF="$RX_BDF" +[[ -n "$RX_MAC" ]] && export ETH_DST_ADDR="$RX_MAC" + +TS="$(date -u +%Y%m%dT%H%M%SZ)" +OUT_ROOT="${DAQIRI_BENCH_RESULTS_DIR:-$REPO_ROOT/bench-results}" +if ! mkdir -p "$OUT_ROOT" 2>/dev/null || [[ ! -w "$OUT_ROOT" ]]; then + if [[ -n "${DAQIRI_BENCH_RESULTS_DIR:-}" ]]; then + echo "ERROR: cannot write DAQIRI_BENCH_RESULTS_DIR=$OUT_ROOT" >&2 + exit 1 + fi + OUT_ROOT="/tmp/daqiri-bench-results" + mkdir -p "$OUT_ROOT" +fi +OUT_DIR="$OUT_ROOT/${TS}-rtx-pro-${BACKEND}-${MODE}" +mkdir -p "$OUT_DIR" + +CSV="$OUT_DIR/runs.csv" +SUMMARY_CSV="$OUT_DIR/summary.csv" +echo "lang,backend,post_process,payload,batch,target_gbps,seconds,tx_packets,rx_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 "backend,workload,config,seconds,total_data_rate_gbps,total_packets_sent,total_packets_received,total_dropped,drop_source,notes" > "$SUMMARY_CSV" + +"$SCRIPT_DIR/bench_capture_environment.sh" "$OUT_DIR" + +DRIVER_LOG="$OUT_DIR/last_run.stderr" +FAILURES=0 + +TX_BDF="${RTX_TX_BDF:-0000:61:00.0}" +RX_BDF="${RTX_RX_BDF:-0000:61:00.1}" +TX_GPU="${TX_GPU:-0}" +RX_GPU="${RX_GPU:-1}" + +wire_preflight() { + if [[ "$IS_SW" == "1" ]]; then + return 0 + fi + case "$BACKEND" in + dpdk|dpdk-hds|ibverbs) + : "${ETH_DST_ADDR:?ETH_DST_ADDR must be set for wire modes (source scripts/discover_rtx_pro_topology.sh)}" + local tx_if="${RTX_TX_IFACE:-}" + local rx_if="${RTX_RX_IFACE:-}" + [[ -z "$tx_if" ]] && tx_if="$(resolve_iface_for_bdf "$TX_BDF" 2>/dev/null || true)" + [[ -z "$rx_if" ]] && rx_if="$(resolve_iface_for_bdf "$RX_BDF" 2>/dev/null || true)" + if [[ -n "$tx_if" && -n "$rx_if" ]]; then + local c0 c1 + c0="$(cat "/sys/class/net/${tx_if}/carrier" 2>/dev/null || echo 0)" + c1="$(cat "/sys/class/net/${rx_if}/carrier" 2>/dev/null || echo 0)" + if [[ "$c0" != "1" || "$c1" != "1" ]]; then + echo "WARNING: carrier not 1 on both ports ($tx_if=$c0 $rx_if=$c1); wire test may fail" >&2 + fi + fi + ;; + esac +} + +wire_preflight + +# -------------------------------------------------------------------------- +# Backend configuration +# -------------------------------------------------------------------------- +TX_SENDER_BIN="$BUILD_DIR/examples/daqiri_bench_raw_gpudirect" +case "$BACKEND" in + dpdk) + PAYLOADS_SWEEP=(8000 4096 1024 256 64) + BATCHES_SWEEP=(10240 4096 1024 256) + PAYLOADS_HEADLINE=(8000) + BATCHES_HEADLINE=(10240) + case "$MODE" in + sw-smoke|sw-issue17) + BASE_YAML="$SCRIPT_DIR/daqiri_bench_raw_sw_loopback_rtx_pro_6000.yaml" + ;; + nic-smoke|sweep|drop-curve|drop-curve-matrix|issue17) + BASE_YAML="$SCRIPT_DIR/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml" + ;; + *) + echo "Unknown mode for dpdk: $MODE" >&2 + exit 1 + ;; + esac + BENCH_BIN="$BUILD_DIR/examples/daqiri_bench_raw_gpudirect" + CPU_MASTER=3 + CPU_TX=11 + CPU_RX=9 + ;; + dpdk-hds) + PAYLOADS_SWEEP=(8000 4096 2048) + BATCHES_SWEEP=(4096 2048 1024) + PAYLOADS_HEADLINE=(8000) + BATCHES_HEADLINE=(4096) + BASE_YAML="$SCRIPT_DIR/daqiri_bench_raw_tx_rx_hds_rtx_pro_6000.yaml" + BENCH_BIN="$BUILD_DIR/examples/daqiri_bench_raw_hds" + CPU_MASTER=3 + CPU_TX=10 + CPU_RX=8 + ;; + rdma) + PAYLOADS_SWEEP=(8000000 1048576 65536 8192 4096) + BATCHES_SWEEP=(1) + PAYLOADS_HEADLINE=(8000000) + BATCHES_HEADLINE=(1) + BASE_YAML="$SCRIPT_DIR/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml" + BENCH_BIN="$BUILD_DIR/examples/daqiri_bench_rdma" + CPU_MASTER=3 + CPU_TX=7 + CPU_RX=8 + ;; + ibverbs) + PAYLOADS_SWEEP=(8000) + BATCHES_SWEEP=(1024) + PAYLOADS_HEADLINE=(8000) + BATCHES_HEADLINE=(1024) + BASE_YAML="$SCRIPT_DIR/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml" + TX_YAML="$SCRIPT_DIR/daqiri_bench_raw_tx_only_rtx_pro_6000_nic.yaml" + BENCH_BIN="$BUILD_DIR/examples/daqiri_bench_raw_gpudirect" + CPU_MASTER=3 + CPU_TX=10 + CPU_RX=8 + ;; + socket-udp) + PAYLOADS_SWEEP=(1472 1024 256 64) + BATCHES_SWEEP=(1) + PAYLOADS_HEADLINE=(1472) + BATCHES_HEADLINE=(1) + BASE_YAML="$SCRIPT_DIR/daqiri_bench_socket_udp_tx_rx.yaml" + BENCH_BIN="$BUILD_DIR/examples/daqiri_bench_socket" + CPU_MASTER=3 + CPU_TX=7 + CPU_RX=8 + ;; + socket-tcp) + PAYLOADS_SWEEP=(1048576 65536 1024) + BATCHES_SWEEP=(1) + PAYLOADS_HEADLINE=(65536) + BATCHES_HEADLINE=(1) + BASE_YAML="$SCRIPT_DIR/daqiri_bench_socket_tcp_tx_rx.yaml" + BENCH_BIN="$BUILD_DIR/examples/daqiri_bench_socket" + CPU_MASTER=3 + CPU_TX=7 + CPU_RX=8 + ;; + *) + echo "Unknown backend: $BACKEND" >&2 + exit 1 + ;; +esac + +DROP_CURVE_TARGETS=(1 5 10 25 50 75 100 0) + +# SW loopback rides the single-segment GPUDirect ring in the dpdk engine. HDS +# (segment split), RoCE, ibverbs, and sockets have no software-loopback path, +# so fail early with a clear pointer rather than emitting a misleading run. +if [[ "$IS_SW" == "1" && "$BACKEND" != "dpdk" ]]; then + echo "SW-loopback modes (sw-smoke / sw-issue17) are only supported on the 'dpdk' backend." >&2 + echo "Use: $0 dpdk sw-issue17 --workload fft (no cable needed)" >&2 + echo "For $BACKEND, use a wire mode (nic-smoke / issue17) with a cabled link." >&2 + exit 1 +fi + +if [[ ! -x "$BENCH_BIN" ]]; then + echo "Benchmark binary not found: $BENCH_BIN (run build-only mode first)" >&2 + exit 1 +fi + +# -------------------------------------------------------------------------- +# Helpers +# -------------------------------------------------------------------------- +extract_field() { + local prefix="$1" field="$2" file="$3" + grep -E "^$prefix" "$file" | tail -n1 | grep -oE " $field=[^ ]+" | head -n1 | sed -E "s/.*$field=//" +} + +parse_dpdk_drops() { + local log="$1" + local sum=0 v + for key in imissed ierrors rx_nombuf; do + v="$(grep -oE "$key=[0-9]+" "$log" 2>/dev/null | grep -oE '[0-9]+$' | sort -n | tail -n1 || true)" + [[ -n "${v:-}" ]] && sum=$((sum + v)) + done + # DAQIRI_LOG_ERROR lines: "Dropped N packets ... (total: T)" + v="$(grep -oE 'Dropped [0-9]+ packets.*\(total: [0-9]+\)' "$log" 2>/dev/null \ + | grep -oE 'total: [0-9]+' | grep -oE '[0-9]+$' | sort -n | tail -n1 || true)" + [[ -n "${v:-}" && "$v" -gt "$sum" ]] && sum="$v" + echo "$sum" +} + +parse_rdma_drops() { + local log="$1" + grep -c 'CQ error' "$log" 2>/dev/null || echo 0 +} + +snapshot_proc_net_udp() { + awk 'NR>1 { sum += $13 } END { print sum+0 }' /proc/net/udp 2>/dev/null || echo 0 +} + +snapshot_nstat() { + nstat -a 2>/dev/null | awk '/TcpExtTCPLostRetransmit|TcpRetransSegs|TcpInErrs/ { s += $2 } END { print s+0 }' || echo 0 +} + +snapshot_cpu_stat() { + awk '/^cpu[0-9]+/ { + total = $2+$3+$4+$5+$6+$7+$8 + busy = total - $5 - $6 + print $1, total, busy + }' /proc/stat > "$1" +} + +cpu_busy_pct() { + local before="$1" after="$2" cpu_idx="$3" + awk -v cpu="cpu$cpu_idx" ' + NR == FNR { b_total[$1] = $2; b_busy[$1] = $3; next } + { a_total[$1] = $2; a_busy[$1] = $3 } + END { + dt = a_total[cpu] - b_total[cpu] + db = a_busy[cpu] - b_busy[cpu] + if (dt > 0) printf "%.1f", (db * 100.0) / dt + else printf "0.0" + } + ' "$before" "$after" +} + +max_phy_counter() { + local key="$1" file="$2" + # DPDK logs use "tx_phy_packets:\t123" (colon), not key=value. + grep -oE "${key}:[[:space:]]*[0-9]+" "$file" 2>/dev/null \ + | grep -oE '[0-9]+$' \ + | sort -n \ + | tail -n1 +} + +log_phy_counters() { + local stderr_file="$1" + local tx_phy rx_phy + tx_phy="$(max_phy_counter 'tx_phy_packets' "$stderr_file")" + rx_phy="$(max_phy_counter 'rx_phy_packets' "$stderr_file")" + tx_phy="${tx_phy:-0}" + rx_phy="${rx_phy:-0}" + echo "post-run phy: tx_phy_packets=${tx_phy} rx_phy_packets=${rx_phy}" >> "$OUT_DIR/wire_validation.txt" + echo "${tx_phy} ${rx_phy}" +} + +wire_phy_ok() { + local stderr_file="$1" rx_pkts="$2" + local counts tx_phy rx_phy + counts="$(log_phy_counters "$stderr_file")" + read -r tx_phy rx_phy <<< "$counts" + # Dual-port closed-loop: TX port drives tx_phy, RX port drives rx_phy. Take + # the per-port maxima from extended stats (each counter appears twice). + if [[ "${rx_pkts:-0}" -gt 1000 && "${rx_phy:-0}" -lt 100 ]]; then + echo "ERROR: rx_phy_packets=${rx_phy:-0} flat while rx_pkts=$rx_pkts -- traffic may not have crossed the link" >&2 + return 1 + fi + if [[ "${rx_pkts:-0}" -gt 1000 && "${tx_phy:-0}" -lt 100 ]]; then + echo "ERROR: tx_phy_packets=${tx_phy:-0} flat while rx_pkts=$rx_pkts -- TX port may not be driving the wire" >&2 + return 1 + fi + return 0 +} + +generate_yaml() { + local out="$1" payload="$2" batch="$3" + case "$BACKEND" in + dpdk) + # SW loopback has no NIC MAC/BDF to substitute (address is "loopback" and + # eth_dst is a don't-care), so only rewrite the destination MAC on the + # wire path where discovery populated ETH_DST_ADDR. + local sed_args=( + -e "s|^( *payload_size: ).*|\1$payload|" + -e "s|^( *batch_size: ).*|\1$batch|g" + -e "s|address: 0000:61:00.0|address: $TX_BDF|g" + -e "s|address: 0000:61:00.1|address: $RX_BDF|g" + ) + if [[ "$IS_SW" != "1" && -n "${ETH_DST_ADDR:-}" ]]; then + sed_args+=( + -e "s|<00:00:00:00:00:00>|$ETH_DST_ADDR|g" + -e "s|^( *eth_dst_addr: ).*|\1$ETH_DST_ADDR|" + ) + fi + sed -E "${sed_args[@]}" "$BASE_YAML" > "$out" + awk -v txg="$TX_GPU" -v rxg="$RX_GPU" ' + /^ - name: "Data_TX_GPU"/ { in_tx=1; in_rx=0 } + /^ - name: "Data_RX_GPU"/ { in_rx=1; in_tx=0 } + in_tx && /^ affinity:/ { sub(/[0-9]+$/, txg); in_tx=0 } + in_rx && /^ affinity:/ { sub(/[0-9]+$/, rxg); in_rx=0 } + { print } + ' "$out" > "${out}.tmp" && mv "${out}.tmp" "$out" + ;; + dpdk-hds) + local ipv4_len=$((payload + 50)) + sed -E \ + -e "s|^( *payload_size: ).*|\1$payload|" \ + -e "s|^( *batch_size: ).*|\1$batch|g" \ + -e "s|<00:00:00:00:00:00>|$ETH_DST_ADDR|g" \ + -e "s|^( *eth_dst_addr: ).*|\1$ETH_DST_ADDR|" \ + -e "s|^( *ipv4_len: ).*|\1$ipv4_len|" \ + -e "s|address: 0000:61:00.0|address: $TX_BDF|g" \ + -e "s|address: 0000:61:00.1|address: $RX_BDF|g" \ + "$BASE_YAML" > "$out" + ;; + ibverbs) + sed -E \ + -e "s|address: 0000:61:00.1|address: $RX_BDF|g" \ + "$BASE_YAML" > "$out" + ;; + rdma) + sed -E "s|^( *message_size: ).*|\1$payload|g" "$BASE_YAML" > "$out" + ;; + socket-udp|socket-tcp) + sed -E \ + -e "s|^( *message_size: ).*|\1$payload|g" \ + -e "s|^( *max_payload_size: ).*|\1$payload|g" \ + -e "s|^( *buf_size: ).*|\1$payload|g" \ + "$BASE_YAML" > "$out" + ;; + esac +} + +generate_tx_yaml() { + local out="$1" payload="${2:-8000}" batch="${3:-10240}" + sed -E \ + -e "s|<00:00:00:00:00:00>|$ETH_DST_ADDR|g" \ + -e "s|^( *eth_dst_addr: ).*|\1$ETH_DST_ADDR|" \ + -e "s|address: 0000:61:00.0|address: $TX_BDF|g" \ + -e "s|^( *payload_size: ).*|\1$payload|" \ + -e "s|^( *batch_size: ).*|\1$batch|g" \ + "$TX_YAML" > "$out" + awk -v txg="$TX_GPU" ' + /^ - name: "Data_TX_GPU"/ { in_tx=1 } + in_tx && /^ affinity:/ { sub(/[0-9]+$/, txg); in_tx=0 } + { print } + ' "$out" > "${out}.tmp" && mv "${out}.tmp" "$out" +} + +append_summary_row() { + local workload="$1" config="$2" secs="$3" data_rate="$4" tx_pkts="$5" rx_pkts="$6" drops="$7" drop_kind="$8" notes="$9" + echo "$BACKEND,$workload,$config,$secs,$data_rate,$tx_pkts,$rx_pkts,$drops,$drop_kind,$notes" >> "$SUMMARY_CSV" +} + +run_cell() { + local lang="$1" payload="$2" batch="$3" target_gbps="$4" workload="${5:-$WORKLOAD}" + local cell="$lang-$BACKEND-w${workload}-p$payload-b$batch-g$target_gbps" + local cell_dir="$OUT_DIR/$cell" + mkdir -p "$cell_dir" + + local yaml="$cell_dir/config.yaml" + generate_yaml "$yaml" "$payload" "$batch" + + local udp_before tcp_before + udp_before="$(snapshot_proc_net_udp)" + tcp_before="$(snapshot_nstat)" + + snapshot_cpu_stat "$cell_dir/cpu_stat.before" + + ( nvidia-smi dmon -s pucvmet -c "$RUN_SECONDS" > "$cell_dir/nvidia_smi_dmon.txt" 2>&1 ) & + local dmon_pid=$! + + local stdout="$cell_dir/stdout.txt" + local stderr="$cell_dir/stderr.txt" + local args=("$yaml" --seconds "$RUN_SECONDS") + [[ "$target_gbps" != "0" ]] && args+=(--target-gbps "$target_gbps") + [[ "$BACKEND" == "rdma" || "$BACKEND" =~ ^socket- ]] && args+=(--mode both) + if [[ "$workload" != "none" ]]; then + args+=(--workload "$workload") + [[ -n "$WORKLOAD_BATCH" ]] && args+=(--workload-batch-bytes "$WORKLOAD_BATCH") + [[ -n "$GEMM_N" ]] && args+=(--workload-gemm-n "$GEMM_N") + [[ -n "$SYNC_INTERVAL" ]] && args+=(--workload-sync-interval "$SYNC_INTERVAL") + fi + + local tx_pid="" + local rx_pid="" + local bench_rc=0 + if [[ "$BACKEND" == "ibverbs" ]]; then + # ibverbs RX init needs hugepages; start it before the DPDK TX sender so both + # processes are not competing for the same 1G pages during EAL bring-up. + "$BENCH_BIN" "${args[@]}" > "$stdout" 2> "$stderr" & + rx_pid=$! + sleep 5 + local tx_yaml="$cell_dir/tx_sender.yaml" + generate_tx_yaml "$tx_yaml" "$payload" "$batch" + "$TX_SENDER_BIN" "$tx_yaml" --seconds "$((RUN_SECONDS + 5))" > "$cell_dir/tx_stdout.txt" 2> "$cell_dir/tx_stderr.txt" & + tx_pid=$! + wait "$rx_pid" || bench_rc=$? + wait "$tx_pid" 2>/dev/null || true + cp "$stderr" "$DRIVER_LOG" + else + "$BENCH_BIN" "${args[@]}" > "$stdout" 2> "$stderr" || bench_rc=$? + cp "$stderr" "$DRIVER_LOG" + fi + + snapshot_cpu_stat "$cell_dir/cpu_stat.after" + wait "$dmon_pid" 2>/dev/null || true + + local tx_pkts rx_pkts bytes secs tx_bytes rx_bytes + tx_pkts="$(extract_field 'TX complete' packets "$stdout")" + tx_bytes="$(extract_field 'TX complete' bytes "$stdout")" + rx_pkts="$(extract_field 'RX complete' packets "$stdout")" + rx_bytes="$(extract_field 'RX complete' bytes "$stdout")" + secs="$(extract_field 'RX complete' seconds "$stdout")" + [[ -z "$secs" ]] && secs="$(extract_field 'TX complete' seconds "$stdout")" + + if [[ -z "$tx_pkts" && -z "$rx_pkts" ]]; then + case "$BACKEND" in + rdma) + tx_pkts="$(extract_field 'Client complete' send_completions "$stdout")" + tx_bytes="$(extract_field 'Client complete' send_bytes "$stdout")" + rx_pkts="$(extract_field 'Client complete' recv_completions "$stdout")" + rx_bytes="$(extract_field 'Client complete' recv_bytes "$stdout")" + secs="$(extract_field 'Client complete' seconds "$stdout")" + ;; + socket-udp|socket-tcp) + tx_pkts="$(extract_field 'Client complete' sent_packets "$stdout")" + tx_bytes="$(extract_field 'Client complete' sent_bytes "$stdout")" + rx_pkts="$(extract_field 'Client complete' recv_packets "$stdout")" + rx_bytes="$(extract_field 'Client complete' recv_bytes "$stdout")" + secs="$(extract_field 'Client complete' seconds "$stdout")" + ;; + esac + fi + + bytes="${rx_bytes:-${tx_bytes:-0}}" + local stats_missing=0 + if [[ -z "${secs:-}" ]]; then + stats_missing=1 + fi + if [[ "$bench_rc" -ne 0 || "$stats_missing" -ne 0 ]]; then + echo "ERROR: $cell failed (rc=$bench_rc stats_missing=$stats_missing)" >&2 + echo " stdout: $stdout" >&2 + echo " stderr: $stderr" >&2 + return 1 + fi + + if [[ "$BACKEND" =~ ^(dpdk|dpdk-hds)$ && "$IS_SW" != "1" ]]; then + if [[ -z "${rx_pkts:-}" || "$rx_pkts" == "0" ]]; then + echo "ERROR: $cell NIC RX packets zero -- check cable/MAC/flow" >&2 + return 1 + fi + wire_phy_ok "$stderr" "$rx_pkts" || return 1 + fi + + if [[ "$BACKEND" == "ibverbs" && "$IS_SW" != "1" ]]; then + local tx_sender_pkts tx_sender_rc=0 + tx_sender_pkts="$(extract_field 'TX complete' packets "$cell_dir/tx_stdout.txt")" + if grep -q 'daqiri_init failed' "$cell_dir/tx_stderr.txt" 2>/dev/null; then + echo "ERROR: $cell ibverbs TX sender failed (daqiri_init) -- see $cell_dir/tx_stderr.txt" >&2 + return 1 + fi + if [[ -z "${tx_sender_pkts:-}" || "$tx_sender_pkts" == "0" ]]; then + echo "ERROR: $cell ibverbs TX sender sent zero packets -- RX port may still be kernel-bound to ibverbs only" >&2 + return 1 + fi + tx_pkts="${tx_sender_pkts}" + local min_wire_rx=$((RUN_SECONDS * 50000)) + [[ "$min_wire_rx" -lt 500000 ]] && min_wire_rx=500000 + if [[ "${rx_pkts:-0}" -lt "$min_wire_rx" ]]; then + echo "ERROR: $cell ibverbs RX too low ($rx_pkts pkts, need >=$min_wire_rx) -- check flow/MAC/cable" >&2 + return 1 + fi + fi + + if [[ "$workload" != "none" && "$IS_SW" != "1" && "${rx_pkts:-0}" -lt 1000000 ]]; then + echo "ERROR: $cell workload=$workload RX too low ($rx_pkts pkts) -- GPU post-process path likely stalled" >&2 + return 1 + fi + + tx_pkts="${tx_pkts:-0}" + rx_pkts="${rx_pkts:-0}" + local pkts="$rx_pkts" + [[ "$pkts" == "0" && "$tx_pkts" != "0" ]] && pkts="$tx_pkts" + + local pps gbps rx_gbps + pps="$(awk -v p="$pkts" -v s="$secs" 'BEGIN { if (s+0>0) printf "%.0f", p/s; else print 0 }')" + gbps="$(awk -v b="${tx_bytes:-0}" -v s="$secs" 'BEGIN { if (s+0>0) printf "%.3f", (b*8.0)/s/1e9; else print 0 }')" + rx_gbps="$(awk -v b="${rx_bytes:-0}" -v s="$secs" 'BEGIN { if (s+0>0) printf "%.3f", (b*8.0)/s/1e9; else print 0 }')" + local data_rate_gbps="$rx_gbps" + [[ "$data_rate_gbps" == "0.000" || -z "$data_rate_gbps" ]] && data_rate_gbps="$gbps" + + local drops drops_kind notes="" + case "$BACKEND" in + dpdk|dpdk-hds|ibverbs) + drops="$(parse_dpdk_drops "$stderr")" + drops_kind="dpdk-imissed+ierrors+nombuf" + ;; + rdma) + drops="$(parse_rdma_drops "$stderr")" + drops_kind="rdma-cqe-error" + ;; + socket-udp) + local udp_after + udp_after="$(snapshot_proc_net_udp)" + drops="$((udp_after - udp_before))" + drops_kind="udp-proc-net-udp-drops" + ;; + socket-tcp) + local tcp_after + tcp_after="$(snapshot_nstat)" + drops="$((tcp_after - tcp_before))" + drops_kind="tcp-nstat-retrans+inerrs" + ;; + esac + + if [[ "$IS_SW" == "1" ]]; then + notes="sw-loopback-nonwire" + fi + + if [[ "$BACKEND" == "ibverbs" ]]; then + notes="host-staged" + if [[ -n "${tx_pkts:-}" && "$tx_pkts" -gt 0 && "$rx_pkts" -lt "$tx_pkts" ]]; then + local external_loss=$((tx_pkts - rx_pkts)) + drops=$((drops + external_loss)) + drops_kind="${drops_kind}+sender-minus-receiver" + fi + fi + + if [[ "$IS_SW" != "1" && "${drops:-0}" -gt 1000000 ]]; then + local nic_drops="$drops" + local drop_base="${tx_pkts:-0}" + if [[ "$BACKEND" == "ibverbs" ]]; then + # Paired DPDK TX + ibverbs RX: sender-minus-receiver is RX backpressure, not a + # NIC drop. Fail only on DPDK/imissed counters from the RX process log. + nic_drops="$(parse_dpdk_drops "$stderr")" + drop_base="${rx_pkts:-0}" + fi + if [[ "$drop_base" -gt 0 && "$nic_drops" -gt 1000000 ]]; then + local drop_pct + drop_pct="$(awk -v d="$nic_drops" -v t="$drop_base" 'BEGIN { printf "%.0f", (d*100.0)/t }')" + if [[ "$drop_pct" -gt 5 ]]; then + echo "ERROR: $cell drop rate ${drop_pct}% (${nic_drops} drops / ${drop_base} pkts) exceeds 5% wire smoke threshold" >&2 + return 1 + fi + fi + fi + + local cpu_master_pct cpu_tx_pct cpu_rx_pct + cpu_master_pct="$(cpu_busy_pct "$cell_dir/cpu_stat.before" "$cell_dir/cpu_stat.after" "$CPU_MASTER")" + cpu_tx_pct="$(cpu_busy_pct "$cell_dir/cpu_stat.before" "$cell_dir/cpu_stat.after" "$CPU_TX")" + cpu_rx_pct="$(cpu_busy_pct "$cell_dir/cpu_stat.before" "$cell_dir/cpu_stat.after" "$CPU_RX")" + + local gpu_sm gpu_mem + gpu_sm="$(awk '/^ *[0-9]/ { count++; sum += $5 } END { if (count) printf "%.1f", sum/count; else print 0 }' \ + "$cell_dir/nvidia_smi_dmon.txt" 2>/dev/null || echo 0)" + 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}" + local pp_gemm_n="${GEMM_N:-derived}" + local pp_sync="${SYNC_INTERVAL:-default}" + + echo "$lang,$BACKEND,$workload,$payload,$batch,$target_gbps,$secs,$tx_pkts,$rx_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" + + append_summary_row "$workload" "$(basename "$yaml")" "$secs" "$data_rate_gbps" "$tx_pkts" "$rx_pkts" "$drops" "$drops_kind" "$notes" +} + +run_cell_or_record_failure() { + run_cell "$@" || FAILURES=$((FAILURES + 1)) +} + +# -------------------------------------------------------------------------- +# Driver +# -------------------------------------------------------------------------- +case "$MODE" in + issue17|sw-issue17) + # Same none/fft/gemm workload matrix on the wire (issue17) or over the + # software loopback (sw-issue17). The SW variant validates the FFT/GEMM + # post-process path end-to-end with no cable, so results are labelled + # non-wire in the summary via the SW config name. + for wl in none fft gemm; do + for p in "${PAYLOADS_HEADLINE[@]}"; do + for b in "${BATCHES_HEADLINE[@]}"; do + run_cell_or_record_failure cpp "$p" "$b" 0 "$wl" + done + done + done + ;; + sw-smoke|nic-smoke) + for p in "${PAYLOADS_HEADLINE[@]}"; do + for b in "${BATCHES_HEADLINE[@]}"; do + run_cell_or_record_failure cpp "$p" "$b" 0 + done + done + ;; + sweep) + for p in "${PAYLOADS_SWEEP[@]}"; do + for b in "${BATCHES_SWEEP[@]}"; do + run_cell_or_record_failure cpp "$p" "$b" 0 + done + done + ;; + drop-curve) + for p in "${PAYLOADS_HEADLINE[@]}"; do + for b in "${BATCHES_HEADLINE[@]}"; do + for g in "${DROP_CURVE_TARGETS[@]}"; do + run_cell_or_record_failure cpp "$p" "$b" "$g" + done + done + done + ;; + drop-curve-matrix) + for p in "${PAYLOADS_SWEEP[@]}"; do + for b in "${BATCHES_HEADLINE[@]}"; do + for g in "${DROP_CURVE_TARGETS[@]}"; do + run_cell_or_record_failure cpp "$p" "$b" "$g" + done + done + done + ;; + *) + echo "Unknown mode: $MODE" >&2 + exit 1 + ;; +esac + +echo +echo "Results in: $OUT_DIR" +echo "CSV: $CSV" +echo "Summary: $SUMMARY_CSV" + +if [[ "$FAILURES" -ne 0 ]]; then + echo "Failed cells: $FAILURES" >&2 + exit 1 +fi diff --git a/examples/socket_bench.cpp b/examples/socket_bench.cpp index c3aaaa1..9369d56 100644 --- a/examples/socket_bench.cpp +++ b/examples/socket_bench.cpp @@ -17,6 +17,9 @@ #include +#include + +#include #include #include #include @@ -74,8 +77,34 @@ SocketBenchConfig parse_socket_cfg(const YAML::Node& node) { return cfg; } +// Detect the socket transport from the YAML: TCP is an in-order byte stream +// (gather-only), UDP datagrams can reorder (seq-based reorder). Scans the +// interfaces' socket_config.local_addr scheme. +bool socket_transport_is_tcp(const YAML::Node& root) { + const auto cfg = root["daqiri"] && root["daqiri"]["cfg"] ? root["daqiri"]["cfg"] : root; + const auto ifaces = cfg["interfaces"]; + if (ifaces && ifaces.IsSequence()) { + for (const auto& iface : ifaces) { + const auto sc = iface["socket_config"]; + if (!sc) { + continue; + } + const auto addr = sc["local_addr"].as(""); + if (addr.rfind("tcp://", 0) == 0) { + return true; + } + if (addr.rfind("udp://", 0) == 0) { + return false; + } + } + } + return false; // default to UDP semantics +} + void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer& pacer, - std::atomic& stop, SocketWorkerStats& stats) { + 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) { const char *thread_name = cfg.server ? "socket_bench_server" : "socket_bench_client"; if (!daqiri::bench::set_current_thread_affinity(cfg.cpu_core, thread_name)) { @@ -83,6 +112,33 @@ void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer return; } + // Representative GPU workload on the REAL received data. Socket data lands in + // pageable host memory, so the pipeline stages it to the GPU (a host->device + // 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; + 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) || + !pipeline.init(is_tcp ? daqiri::bench::ReorderMode::GatherOnly + : daqiri::bench::ReorderMode::SeqReorder, + packets_per_batch, msg, /*payload_byte_offset=*/0, + /*seq_bit_offset=*/0, /*seq_bit_width=*/is_tcp ? 0 : 32, + /*staging_needed=*/true, gpu_workload.stream())) { + std::cerr << "Socket workload init failed; continuing without GPU workload\n"; + } + } + const bool run_workload = gpu_workload.enabled() && pipeline.enabled(); + uintptr_t conn_id = 0; uint16_t port = 0; uint16_t queue = 0; @@ -124,9 +180,20 @@ void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer daqiri::set_header(msg, port, queue, 1, 1); if (daqiri::get_tx_packet_burst(msg) == daqiri::Status::SUCCESS) { + if (daqiri::set_packet_lengths(msg, 0, {cfg.message_size}) != daqiri::Status::SUCCESS) { + std::cerr << "socket_bench message_size exceeds configured socket buffer size\n"; + daqiri::free_all_packets_and_burst_tx(msg); + stop.store(true); + continue; + } auto* payload = reinterpret_cast(daqiri::get_packet_ptr(msg, 0)); std::memset(payload, static_cast(stats.sent_packets & 0xff), cfg.message_size); - daqiri::set_packet_lengths(msg, 0, {cfg.message_size}); + // Inject a sequence number (network byte order) at the payload start so + // the UDP RX-side reorder has a real seq to sort by. Harmless for TCP. + if (cfg.message_size >= static_cast(sizeof(uint32_t))) { + const uint32_t seq = htonl(static_cast(stats.sent_packets)); + std::memcpy(payload, &seq, sizeof(seq)); + } daqiri::set_connection_id(msg, conn_id); @@ -144,9 +211,28 @@ void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer daqiri::BurstParams* burst = nullptr; if (daqiri::get_rx_burst(&burst, conn_id, cfg.server) == daqiri::Status::SUCCESS && burst != nullptr) { - const uint64_t rx_pkts = static_cast(daqiri::get_num_packets(burst)); - stats.received_packets += rx_pkts; + const int num_pkts = static_cast(daqiri::get_num_packets(burst)); + stats.received_packets += static_cast(num_pkts); stats.received_bytes += daqiri::get_burst_tot_byte(burst); + + if (run_workload) { + // Stage each received (host) payload to the GPU; the copy persists in + // the pipeline's device buffer, so a batch can accumulate across bursts + // (UDP bursts are one datagram). When a batch fills, reorder/gather it + // and run the compute. TCP (packets_per_batch == 1) flushes per chunk. + for (int i = 0; i < num_pkts; ++i) { + const auto len = daqiri::get_packet_length(burst, i); + void* d = pipeline.stage_host_packet(daqiri::get_packet_ptr(burst, i), len); + pipeline.add_device_packet(d); + if (pipeline.collected() == packets_per_batch) { + gpu_workload.run(pipeline.finish_batch()); + pipeline.reset_batch(); + } + } + // Drain so the host->device staging copies (which read this burst's + // memory) complete before the burst is freed/recycled. + gpu_workload.sync(); + } daqiri::free_all_packets_and_burst_rx(burst); } else { std::this_thread::sleep_for(std::chrono::microseconds(100)); @@ -160,7 +246,9 @@ void socket_worker(const SocketBenchConfig& cfg, daqiri::bench::TokenBucketPacer int main(int argc, char** argv) { if (argc < 2) { std::cerr << "Usage: " << argv[0] - << " [--seconds N] [--mode server|client|both] [--target-gbps G]\n"; + << " [--seconds N] [--mode server|client|both] " + "[--target-gbps G] [--workload none|fft|gemm|gemm_fp16] " + "[--workload-batch-bytes N]\n"; return 1; } @@ -176,8 +264,13 @@ int main(int argc, char** argv) { target_gbps = std::stod(argv[i + 1]); } } + 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); if (daqiri::daqiri_init(argv[1]) != daqiri::Status::SUCCESS) { std::cerr << "daqiri_init failed\n"; return 1; @@ -211,12 +304,14 @@ 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)); + 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); } if (run_client) { - client_thread = std::thread(socket_worker, client_cfg, std::ref(client_pacer), - std::ref(stop), std::ref(client_stats)); + 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); } if (!server_thread.joinable() && !client_thread.joinable()) { diff --git a/scripts/discover_rtx_pro_topology.sh b/scripts/discover_rtx_pro_topology.sh new file mode 100755 index 0000000..b464121 --- /dev/null +++ b/scripts/discover_rtx_pro_topology.sh @@ -0,0 +1,149 @@ +#!/usr/bin/env bash +# +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# +# Emit RTX PRO 6000 topology for benchmark YAML generation. +# Run on the host while NICs are kernel-bound (before DPDK devbind). +# +# Usage: +# ./scripts/discover_rtx_pro_topology.sh # print summary +# source ./scripts/discover_rtx_pro_topology.sh # export ETH_DST_ADDR, etc. +# +# Override when auto-detection is wrong: +# RTX_TX_BDF RTX_RX_BDF RTX_TX_IFACE RTX_RX_IFACE ETH_DST_ADDR + +set -u + +read_mac() { + local iface="$1" + local path="/sys/class/net/${iface}/address" + [[ -r "$path" ]] || return 1 + cat "$path" +} + +read_carrier() { + local iface="$1" + local path="/sys/class/net/${iface}/carrier" + [[ -r "$path" ]] || return 1 + tr -d '\n' < "$path" +} + +bdf_for_iface() { + local iface="$1" + local path="/sys/class/net/${iface}/device" + [[ -L "$path" ]] || return 1 + basename "$(readlink -f "$path")" +} + +# Collect kernel mlx5 netdevs with link up, sorted for stable defaults. +list_mlx_ifaces() { + if command -v ibdev2netdev >/dev/null 2>&1; then + ibdev2netdev 2>/dev/null | awk '/\(Up\)/ { + for (i=1;i<=NF;i++) if ($i ~ /^ens/) { print $i; break } + }' + return 0 + fi + ls /sys/class/net 2>/dev/null | awk '/^ens/ { print }' +} + +resolve_iface_for_bdf() { + local bdf="$1" + if command -v ibdev2netdev >/dev/null 2>&1; then + ibdev2netdev 2>/dev/null | awk -v b="${bdf#0000:}" ' + $0 ~ ("0000:" b) { for (i=1;i<=NF;i++) if ($i ~ /^ens/) { print $i; exit } } + ' | head -n1 + fi +} + +pick_default_pair() { + mapfile -t ifaces < <(list_mlx_ifaces) + local up=() + local iface carrier + for iface in "${ifaces[@]}"; do + carrier="$(read_carrier "$iface" 2>/dev/null || echo 0)" + [[ "$carrier" == "1" ]] && up+=("$iface") + done + if [[ ${#up[@]} -ge 2 ]]; then + echo "${up[0]} ${up[1]}" + return 0 + fi + if [[ ${#ifaces[@]} -ge 2 ]]; then + echo "${ifaces[0]} ${ifaces[1]}" + return 0 + fi + return 1 +} + +TX_BDF="${RTX_TX_BDF:-}" +RX_BDF="${RTX_RX_BDF:-}" +TX_IFACE="${RTX_TX_IFACE:-}" +RX_IFACE="${RTX_RX_IFACE:-}" + +# Legacy reference-box defaults when those devices exist. +if [[ -z "$TX_BDF" && -d /sys/bus/pci/devices/0000:61:00.0 ]]; then + TX_BDF="0000:61:00.0" + RX_BDF="${RX_BDF:-0000:61:00.1}" + TX_IFACE="${TX_IFACE:-ens1f0np0}" + RX_IFACE="${RX_IFACE:-ens1f1np1}" +fi + +if [[ -z "$TX_IFACE" || ! -d "/sys/class/net/${TX_IFACE}" ]]; then + if [[ -n "$TX_BDF" ]]; then + resolved="$(resolve_iface_for_bdf "$TX_BDF" || true)" + [[ -n "$resolved" ]] && TX_IFACE="$resolved" + fi +fi +if [[ -z "$RX_IFACE" || ! -d "/sys/class/net/${RX_IFACE}" ]]; then + if [[ -n "$RX_BDF" ]]; then + resolved="$(resolve_iface_for_bdf "$RX_BDF" || true)" + [[ -n "$resolved" ]] && RX_IFACE="$resolved" + fi +fi + +if [[ -z "$TX_IFACE" || -z "$RX_IFACE" || "$TX_IFACE" == "$RX_IFACE" ]]; then + if pair="$(pick_default_pair)"; then + read -r _tx _rx <<< "$pair" + TX_IFACE="${TX_IFACE:-$_tx}" + RX_IFACE="${RX_IFACE:-$_rx}" + fi +fi + +if [[ -z "$TX_BDF" && -n "$TX_IFACE" ]]; then + TX_BDF="$(bdf_for_iface "$TX_IFACE" || true)" +fi +if [[ -z "$RX_BDF" && -n "$RX_IFACE" ]]; then + RX_BDF="$(bdf_for_iface "$RX_IFACE" || true)" +fi + +RX_MAC="$(read_mac "$RX_IFACE" 2>/dev/null || true)" +TX_MAC="$(read_mac "$TX_IFACE" 2>/dev/null || true)" +P0_CARRIER="$(read_carrier "$TX_IFACE" 2>/dev/null || echo "?")" +P1_CARRIER="$(read_carrier "$RX_IFACE" 2>/dev/null || echo "?")" + +export RTX_TX_BDF="${TX_BDF:-}" +export RTX_RX_BDF="${RX_BDF:-}" +export RTX_TX_IFACE="${TX_IFACE:-}" +export RTX_RX_IFACE="${RX_IFACE:-}" +export RTX_TX_MAC="${TX_MAC:-}" +export RTX_RX_MAC="${RX_MAC:-}" +export ETH_DST_ADDR="${ETH_DST_ADDR:-$RX_MAC}" +export RX_IFACE="$RX_IFACE" + +echo "RTX PRO 6000 topology" +echo " TX BDF: ${TX_BDF:-unknown} iface=${TX_IFACE:-unknown} mac=${TX_MAC:-unknown} carrier=$P0_CARRIER" +echo " RX BDF: ${RX_BDF:-unknown} iface=${RX_IFACE:-unknown} mac=${RX_MAC:-unknown} carrier=$P1_CARRIER" +echo " ETH_DST_ADDR=${ETH_DST_ADDR:-}" + +if [[ -z "${RX_MAC:-}" ]]; then + echo "WARNING: could not read RX MAC; set ETH_DST_ADDR manually" >&2 + return 0 2>/dev/null || exit 0 +fi + +if [[ "$P0_CARRIER" != "1" || "$P1_CARRIER" != "1" ]]; then + echo "WARNING: one or both ports not carrier=1; wire tests may fail" >&2 +fi + +if [[ -z "${TX_BDF:-}" || -z "${RX_BDF:-}" ]]; then + echo "WARNING: could not resolve PCIe BDFs; pass --tx-bdf/--rx-bdf to run_rtx_pro_bench.sh" >&2 +fi diff --git a/src/engines/socket/daqiri_socket_engine.cpp b/src/engines/socket/daqiri_socket_engine.cpp index f6d04fb..a8357fe 100644 --- a/src/engines/socket/daqiri_socket_engine.cpp +++ b/src/engines/socket/daqiri_socket_engine.cpp @@ -355,7 +355,20 @@ Status SocketEngine::set_packet_lengths(BurstParams* burst, int idx, return Status::INVALID_PARAMETER; } - burst->pkt_lens[0][idx] = static_cast(*(lens.begin())); + const int len = *(lens.begin()); + if (len < 0) { + return Status::INVALID_PARAMETER; + } + const auto* ep = endpoint_for_port(burst->hdr.hdr.port_id); + if (ep != nullptr && static_cast(len) > ep->max_packet_size) { + DAQIRI_LOG_ERROR("Socket packet length {} exceeds configured max_payload_size/buf_size {} on port {}", + len, + ep->max_packet_size, + burst->hdr.hdr.port_id); + return Status::NO_SPACE_AVAILABLE; + } + + burst->pkt_lens[0][idx] = static_cast(len); return Status::SUCCESS; } From 9f01b732c9b0c9d3e5c5843b3212dabf58b717ba Mon Sep 17 00:00:00 2001 From: Chloe Crozier Date: Fri, 10 Jul 2026 18:37:05 +0000 Subject: [PATCH 4/7] #17 - Fix RTX PRO 6000 wire benchmarks and add core-scaling sweep Fix three failure classes surfaced by wire closed-loop testing and add a multi-queue core/payload sweep with plotting. Signed-off-by: Chloe Crozier --- .gitignore | 8 + ...aqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml | 137 +++++++ .../daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml | 26 +- examples/raw_bench_common.cpp | 26 +- examples/run_rtx_pro_bench.sh | 71 +++- examples/run_rtx_pro_mq_bench.sh | 272 ++++++++++++++ examples/run_rtx_pro_push_400g.sh | 245 +++++++++++++ scripts/gen_rtx_pro_mq_config.py | 117 ++++++ scripts/plot_rtx_pro_bench.py | 345 ++++++++++++++++++ 9 files changed, 1229 insertions(+), 18 deletions(-) create mode 100644 examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml create mode 100755 examples/run_rtx_pro_mq_bench.sh create mode 100755 examples/run_rtx_pro_push_400g.sh create mode 100755 scripts/gen_rtx_pro_mq_config.py create mode 100755 scripts/plot_rtx_pro_bench.py diff --git a/.gitignore b/.gitignore index c23e9a4..33e9ffb 100644 --- a/.gitignore +++ b/.gitignore @@ -2,6 +2,14 @@ build*/ site/ bench-results/ +# Benchmark plot generation (scripts/plot_rtx_pro_bench.py, plot_mq_payload_sweep.py) +bench-results/**/plots/ +bench-results/**/*.svg +bench-results/**/*.png +bench-results/*.log +docs/images/rtx-pro-*.svg +docs/images/rtx-pro-*.png + # tune_system.py default output pcie_schematic.png diff --git a/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml new file mode 100644 index 0000000..3b6b3d3 --- /dev/null +++ b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml @@ -0,0 +1,137 @@ +# RTX PRO 6000 — multi-queue core-scaling base (TX=2, RX=2 superset). +# +# Matrix cells (TX,RX) = (1,1),(1,2),(2,1),(2,2) are derived at sweep time by +# scripts/gen_rtx_pro_mq_config.py (invoked by run_rtx_pro_mq_bench.sh), which +# prunes queues/flows/memory-regions/bench entries down to each cell. +# +# Wire closed-loop: tx_port -> rx_port over L2 (cable/optic/switch). Fill +# eth_dst_addr from discovery (ETH_DST_ADDR) or: +# cat /sys/class/net//address +# +# Core map (pollers are DPDK EAL lcores; workers are bench pthreads): +# master_core 3 +# TX q0: poller 11 / worker 10 +# TX q1: poller 12 / worker 13 +# RX q0: poller 9 / worker 8 +# RX q1: poller 14 / worker 15 +# +%YAML 1.2 +--- +daqiri: + cfg: + version: 1 + stream_type: "raw" + master_core: 3 + debug: false + log_level: "info" + loopback: "" + + memory_regions: + - name: "Data_TX_GPU_0" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + - name: "Data_TX_GPU_1" + kind: "device" + affinity: 0 + num_bufs: 51200 + buf_size: 8064 + - name: "Data_RX_GPU_0" + kind: "device" + affinity: 1 + num_bufs: 51200 + buf_size: 8064 + - name: "Data_RX_GPU_1" + kind: "device" + affinity: 1 + num_bufs: 51200 + buf_size: 8064 + + interfaces: + - name: "tx_port" + address: "0000:75:00.0" + tx: + queues: + - name: "tx_q_0" + id: 0 + batch_size: 10240 + cpu_core: 11 + memory_regions: + - "Data_TX_GPU_0" + offloads: + - "tx_eth_src" + - name: "tx_q_1" + id: 1 + batch_size: 10240 + cpu_core: 12 + memory_regions: + - "Data_TX_GPU_1" + offloads: + - "tx_eth_src" + - name: "rx_port" + address: "0000:05:00.0" + rx: + flow_isolation: true + queues: + - name: "rx_q_0" + id: 0 + cpu_core: 9 + batch_size: 10240 + memory_regions: + - "Data_RX_GPU_0" + - name: "rx_q_1" + id: 1 + cpu_core: 14 + batch_size: 10240 + memory_regions: + - "Data_RX_GPU_1" + flows: + - name: "flow_0" + id: 0 + action: + type: queue + id: 0 + match: + udp_src: 4096 + udp_dst: 4096 + - name: "flow_1" + id: 1 + action: + type: queue + id: 1 + match: + udp_src: 4097 + udp_dst: 4097 + +bench_rx: +- interface_name: "rx_port" + queue_id: 0 + cpu_core: 8 +- interface_name: "rx_port" + queue_id: 1 + cpu_core: 15 + +bench_tx: +- interface_name: "tx_port" + queue_id: 0 + cpu_core: 10 + batch_size: 10240 + payload_size: 8000 + header_size: 64 + eth_dst_addr: <00:00:00:00:00:00> + ip_src_addr: 1.2.3.4 + ip_dst_addr: 5.6.7.8 + udp_src_port: 4096 + udp_dst_port: 4096 +- interface_name: "tx_port" + queue_id: 1 + cpu_core: 13 + batch_size: 10240 + payload_size: 8000 + header_size: 64 + eth_dst_addr: <00:00:00:00:00:00> + ip_src_addr: 1.2.3.4 + ip_dst_addr: 5.6.7.8 + udp_src_port: 4097 + udp_dst_port: 4097 diff --git a/examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml b/examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml index 0d01db1..6bd74ae 100644 --- a/examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml +++ b/examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml @@ -1,8 +1,13 @@ # RTX PRO 6000 Blackwell -- RoCE closed-loop (client + server in one process). # Uses host_pinned memory by default. Device memory RoCE requires a working GPUDirect RDMA MR path (nvidia-peermem or dma-buf support); use host_pinned as the portable fallback. # -# Replace client/server IPs with addresses on your RoCE-capable interfaces. -# Default uses link-local style addresses on the bench host's mlx5 ports. +# The client (10.100.0.1) and server (10.100.0.2) live on ONE connected subnet +# so a single-host dual-port loop resolves the peer over the cable. Assign the +# two addresses to the two RoCE-capable cabled ports before running, e.g.: +# sudo ip addr add 10.100.0.1/24 dev +# sudo ip addr add 10.100.0.2/24 dev +# sudo sysctl -w net.ipv4.conf.all.accept_local=1 net.ipv4.conf.all.rp_filter=0 +# (run_rtx_pro_bench.sh rdma does this automatically from discovered topology). # # Run: # sudo ./build/examples/daqiri_bench_rdma \ @@ -15,7 +20,6 @@ daqiri: cfg: version: 1 stream_type: "socket" - protocol: "roce" master_core: 3 debug: false log_level: "info" @@ -44,11 +48,10 @@ daqiri: interfaces: - name: my_client - address: 10.100.1.1 + address: 10.100.0.1 socket_config: mode: client - remote_ip: 10.100.3.1 - remote_port: 4096 + local_addr: "roce://10.100.0.1" roce_config: transport_mode: RC tx: @@ -64,11 +67,10 @@ daqiri: cpu_core: 7 batch_size: 1 - name: my_server - address: 10.100.3.1 + address: 10.100.0.2 socket_config: mode: server - local_ip: 10.100.3.1 - local_port: 4096 + local_addr: "roce://10.100.0.2:4096" roce_config: transport_mode: RC rx: @@ -85,7 +87,7 @@ daqiri: batch_size: 1 rdma_bench_server: - server_address: 10.100.3.1 + server_address: 10.100.0.2 server_port: 4096 message_size: 8000000 send: true @@ -95,8 +97,8 @@ rdma_bench_server: rdma_bench_client: message_size: 8000000 - client_address: 10.100.1.1 - server_address: 10.100.3.1 + client_address: 10.100.0.1 + server_address: 10.100.0.2 server_port: 4096 receive: true send: true diff --git a/examples/raw_bench_common.cpp b/examples/raw_bench_common.cpp index 61fa8d7..54f8408 100644 --- a/examples/raw_bench_common.cpp +++ b/examples/raw_bench_common.cpp @@ -572,7 +572,25 @@ void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, Bench }; std::array workload_slots; bool run_workload = false; - if (workload != BenchWorkload::None) { + bool workload_init_done = (workload == BenchWorkload::None); + + // The GPU workload is initialized lazily on the first received burst so the + // cuFFT/cuBLAS handles, reorder buffers, and CUDA stream are created on the + // SAME device that holds the RX packet buffers. The RX memory region may live + // on a non-default GPU (e.g. affinity: 1 on the dual-port wire path); binding + // the workload to the default device would make every reorder/compute kernel + // dereference foreign device pointers and fault with an illegal access. + auto init_workload = [&](daqiri::BurstParams *burst) { + workload_init_done = true; + const void *sample = + daqiri::get_segment_packet_ptr(burst, geom.payload_segment, 0); + cudaPointerAttributes attrs{}; + if (sample != nullptr && + cudaPointerGetAttributes(&attrs, sample) == cudaSuccess && + attrs.type == cudaMemoryTypeDevice) { + cudaSetDevice(attrs.device); + } + cudaGetLastError(); // clear any benign query error before real work bool ok = gpu_workload.init(workload, batch_bytes, workload_sync_interval, workload_gemm_n); for (auto &slot : workload_slots) { ok = ok && slot.pipeline.init(ReorderMode::SeqReorder, geom.packets_per_batch, @@ -593,7 +611,7 @@ void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, Bench } else { run_workload = true; } - } + }; auto release_slot = [](WorkloadSlot &slot, bool wait) { if (!slot.busy) { @@ -655,6 +673,10 @@ void rx_count_worker(const RawBenchRxConfig& cfg, std::atomic& stop, Bench stats.bytes += daqiri::get_burst_tot_byte(burst); ++stats.bursts; + if (!workload_init_done) { + init_workload(burst); + } + if (run_workload) { for (auto &slot : workload_slots) { release_slot(slot, false); diff --git a/examples/run_rtx_pro_bench.sh b/examples/run_rtx_pro_bench.sh index faa278a..2283f33 100755 --- a/examples/run_rtx_pro_bench.sh +++ b/examples/run_rtx_pro_bench.sh @@ -209,6 +209,54 @@ wire_preflight() { wire_preflight +# -------------------------------------------------------------------------- +# RoCE (rdma) single-host wire loop: the client and server IPs must sit on ONE +# connected subnet across the two cabled RoCE ports so rdma_cm can resolve the +# peer over the cable. Assign them to the discovered TX/RX interfaces (and relax +# rp_filter / enable accept_local for same-host cross-NIC delivery), then remove +# the addresses we added on exit so we don't disturb the host's base config. +# -------------------------------------------------------------------------- +RDMA_IP_CLIENT="10.100.0.1" +RDMA_IP_SERVER="10.100.0.2" +RDMA_NET_CIDR="24" +RDMA_ADDED_CLIENT_IF="" +RDMA_ADDED_SERVER_IF="" + +rdma_net_teardown() { + [[ -n "$RDMA_ADDED_CLIENT_IF" ]] && \ + ip addr del "${RDMA_IP_CLIENT}/${RDMA_NET_CIDR}" dev "$RDMA_ADDED_CLIENT_IF" 2>/dev/null || true + [[ -n "$RDMA_ADDED_SERVER_IF" ]] && \ + ip addr del "${RDMA_IP_SERVER}/${RDMA_NET_CIDR}" dev "$RDMA_ADDED_SERVER_IF" 2>/dev/null || true +} + +rdma_net_setup() { + [[ "$BACKEND" == "rdma" && "$IS_SW" != "1" ]] || return 0 + local tx_if="${RTX_TX_IFACE:-}" rx_if="${RTX_RX_IFACE:-}" + [[ -z "$tx_if" ]] && tx_if="$(resolve_iface_for_bdf "$TX_BDF" 2>/dev/null || true)" + [[ -z "$rx_if" ]] && rx_if="$(resolve_iface_for_bdf "$RX_BDF" 2>/dev/null || true)" + if [[ -z "$tx_if" || -z "$rx_if" ]]; then + echo "WARNING: could not resolve RoCE interfaces for auto IP setup; assign ${RDMA_IP_CLIENT}/${RDMA_IP_SERVER} manually" >&2 + return 0 + fi + if ! ip -br addr show "$tx_if" 2>/dev/null | grep -q "$RDMA_IP_CLIENT"; then + ip addr add "${RDMA_IP_CLIENT}/${RDMA_NET_CIDR}" dev "$tx_if" 2>/dev/null && RDMA_ADDED_CLIENT_IF="$tx_if" + fi + if ! ip -br addr show "$rx_if" 2>/dev/null | grep -q "$RDMA_IP_SERVER"; then + ip addr add "${RDMA_IP_SERVER}/${RDMA_NET_CIDR}" dev "$rx_if" 2>/dev/null && RDMA_ADDED_SERVER_IF="$rx_if" + fi + sysctl -q -w net.ipv4.conf.all.accept_local=1 2>/dev/null || true + sysctl -q -w net.ipv4.conf.all.rp_filter=0 2>/dev/null || true + for i in "$tx_if" "$rx_if"; do + sysctl -q -w "net.ipv4.conf.${i}.rp_filter=0" 2>/dev/null || true + sysctl -q -w "net.ipv4.conf.${i}.arp_ignore=1" 2>/dev/null || true + sysctl -q -w "net.ipv4.conf.${i}.arp_announce=2" 2>/dev/null || true + done + trap rdma_net_teardown EXIT + echo "RoCE net: ${RDMA_IP_CLIENT} on ${tx_if}, ${RDMA_IP_SERVER} on ${rx_if}" >&2 +} + +rdma_net_setup + # -------------------------------------------------------------------------- # Backend configuration # -------------------------------------------------------------------------- @@ -339,7 +387,11 @@ parse_dpdk_drops() { parse_rdma_drops() { local log="$1" - grep -c 'CQ error' "$log" 2>/dev/null || echo 0 + local n + # grep -c prints "0" AND exits 1 on no match; capture the count and ignore + # the exit status so we return a single clean integer. + n="$(grep -c 'CQ error' "$log" 2>/dev/null)" || true + echo "${n:-0}" } snapshot_proc_net_udp() { @@ -610,9 +662,20 @@ run_cell() { fi fi - if [[ "$workload" != "none" && "$IS_SW" != "1" && "${rx_pkts:-0}" -lt 1000000 ]]; then - echo "ERROR: $cell workload=$workload RX too low ($rx_pkts pkts) -- GPU post-process path likely stalled" >&2 - return 1 + # Stalled-GPU guard for the raw small-packet backends only. rdma/socket move + # large messages (e.g. 8 MB RDMA writes), so a handful of thousand "packets" + # is a full line rate there and a compute-heavy GEMM legitimately lowers the + # message count -- a fixed 1M-packet floor would false-positive on those. + if [[ "$workload" != "none" && "$IS_SW" != "1" ]]; then + local min_workload_rx=0 + case "$BACKEND" in + dpdk|dpdk-hds|ibverbs) min_workload_rx=1000000 ;; + *) min_workload_rx=1 ;; + esac + if [[ "${rx_pkts:-0}" -lt "$min_workload_rx" ]]; then + echo "ERROR: $cell workload=$workload RX too low ($rx_pkts pkts, need >=$min_workload_rx) -- GPU post-process path likely stalled" >&2 + return 1 + fi fi tx_pkts="${tx_pkts:-0}" diff --git a/examples/run_rtx_pro_mq_bench.sh b/examples/run_rtx_pro_mq_bench.sh new file mode 100755 index 0000000..7cc6e36 --- /dev/null +++ b/examples/run_rtx_pro_mq_bench.sh @@ -0,0 +1,272 @@ +#!/usr/bin/env bash +# +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# +# RTX PRO 6000 DPDK multi-queue core-scaling + payload sweep. +# +# Sweeps the four (TX, RX) poller-count cells at multiple payload sizes to show +# how wire throughput scales as cores are added and how that varies with packet +# size. Cells: +# +# cell TX pollers RX pollers +# 1t1r 11 9 +# 1t2r 11 9,14 +# 2t1r 11,12 9 +# 2t2r 11,12 9,14 +# +# Usage (as root, privileged container): +# source scripts/discover_rtx_pro_topology.sh +# ./examples/run_rtx_pro_mq_bench.sh +# +# Environment: +# RUN_SECONDS per cell (default 20) +# PAYLOADS space-separated payload bytes (default "64 256 1024 4096 8000") +# REPEATS repeats per (cell, payload) (default 1) +# RTX_TX_BDF, RTX_RX_BDF, RTX_RX_IFACE, ETH_DST_ADDR from discovery +# +# Plot afterwards: +# scripts/plot_rtx_pro_bench.py bench-results/-rtx-pro-mq/runs.csv + +set -u +set -o pipefail + +SCRIPT_DIR="$(cd "$(dirname "$0")" && pwd)" +REPO_ROOT="$(cd "$SCRIPT_DIR/.." && pwd)" +BUILD_DIR="${DAQIRI_BUILD_DIR:-$REPO_ROOT/build}" +DISCOVER="$REPO_ROOT/scripts/discover_rtx_pro_topology.sh" +BENCH_BIN="$BUILD_DIR/examples/daqiri_bench_raw_gpudirect" +RUN_SECONDS="${RUN_SECONDS:-20}" +PAYLOADS="${PAYLOADS:-64 256 1024 4096 8000}" +REPEATS="${REPEATS:-1}" + +MQ_BASE="$SCRIPT_DIR/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml" +MQ_GEN="$REPO_ROOT/scripts/gen_rtx_pro_mq_config.py" +PLOT_SCRIPT="$REPO_ROOT/scripts/plot_rtx_pro_bench.py" + +export LD_LIBRARY_PATH="/opt/daqiri/lib:$BUILD_DIR:${LD_LIBRARY_PATH:-}" + +if [[ -x "$DISCOVER" ]]; then + # shellcheck disable=SC1090 + source "$DISCOVER" +fi + +TX_BDF="${RTX_TX_BDF:-0000:75:00.0}" +RX_BDF="${RTX_RX_BDF:-0000:05:00.0}" +TX_GPU="${TX_GPU:-0}" +RX_GPU="${RX_GPU:-1}" +RX_NETDEV="${RTX_RX_IFACE:-}" + +TS="$(date -u +%Y%m%dT%H%M%SZ)" +OUT_DIR="${DAQIRI_BENCH_RESULTS_DIR:-$REPO_ROOT/bench-results}/$TS-rtx-pro-mq" +mkdir -p "$OUT_DIR" + +CSV="$OUT_DIR/runs.csv" +echo "cell,tx_cores,rx_cores,payload,rep,gbps,rx_gbps,pps,drops,cpu_master,cpu_tx_q0_poll,cpu_tx_q0_work,cpu_tx_q1_poll,cpu_tx_q1_work,cpu_rx_q0_poll,cpu_rx_q0_work,cpu_rx_q1_poll,cpu_rx_q1_work,gpu_sm,gpu_mem" > "$CSV" + +if [[ -x "$SCRIPT_DIR/bench_capture_environment.sh" ]]; then + "$SCRIPT_DIR/bench_capture_environment.sh" "$OUT_DIR" || true +fi + +if [[ ! -x "$BENCH_BIN" ]]; then + echo "ERROR: bench binary not found: $BENCH_BIN" >&2 + exit 1 +fi +if [[ ! -f "$MQ_BASE" || ! -f "$MQ_GEN" ]]; then + echo "ERROR: MQ base/generator missing: $MQ_BASE / $MQ_GEN" >&2 + exit 1 +fi +if [[ -z "${ETH_DST_ADDR:-}" ]]; then + echo "ERROR: ETH_DST_ADDR must be set (source scripts/discover_rtx_pro_topology.sh)" >&2 + exit 1 +fi + +CELLS=( + "1t1r 1 1" + "1t2r 1 2" + "2t1r 2 1" + "2t2r 2 2" +) + +# master, TX q0 poll/work, TX q1 poll/work, RX q0 poll/work, RX q1 poll/work +CPU_CORES=(3 11 10 12 13 9 8 14 15) + +FAILURES=0 + +sum_field() { + local prefix="$1" field="$2" file="$3" + grep -E "^$prefix" "$file" 2>/dev/null \ + | grep -oE " $field=[0-9]+" \ + | sed -E "s/.* $field=//" \ + | awk '{ s += $1 } END { printf "%d", s+0 }' +} + +max_field() { + local prefix="$1" field="$2" file="$3" + grep -E "^$prefix" "$file" 2>/dev/null \ + | grep -oE " $field=[0-9.]+" \ + | sed -E "s/.* $field=//" \ + | awk '{ if ($1+0 > m+0) m = $1 } END { printf "%s", (m == "" ? 0 : m) }' +} + +parse_dpdk_drops() { + local log="$1" + local sum=0 v + for key in imissed ierrors rx_nombuf; do + v="$(grep -oE "$key=[0-9]+" "$log" 2>/dev/null | tail -n1 | sed -E "s/.*=//" || true)" + [[ -n "${v:-}" ]] && sum=$((sum + v)) + done + v="$(grep -oE 'total: [0-9]+' "$log" 2>/dev/null | tail -n1 | grep -oE '[0-9]+$' || true)" + [[ -n "${v:-}" ]] && sum=$((sum + v)) + echo "$sum" +} + +max_phy_counter() { + local key="$1" file="$2" + grep -oE "${key}:[[:space:]]*[0-9]+" "$file" 2>/dev/null \ + | grep -oE '[0-9]+$' \ + | sort -n \ + | tail -n1 +} + +snapshot_cpu_stat() { + awk '/^cpu[0-9]+/ { + total = $2+$3+$4+$5+$6+$7+$8 + busy = total - $5 - $6 + print $1, total, busy + }' /proc/stat > "$1" +} + +cpu_busy_pct() { + local before="$1" after="$2" cpu_idx="$3" + awk -v cpu="cpu$cpu_idx" ' + NR == FNR { b_total[$1] = $2; b_busy[$1] = $3; next } + { a_total[$1] = $2; a_busy[$1] = $3 } + END { + dt = a_total[cpu] - b_total[cpu] + db = a_busy[cpu] - b_busy[cpu] + if (dt > 0) printf "%.1f", (db * 100.0) / dt + else printf "0.0" + } + ' "$before" "$after" +} + +gpu_sample() { + local out="$1" + nvidia-smi --query-gpu=utilization.gpu,utilization.memory \ + --format=csv,noheader,nounits -i "$TX_GPU,$RX_GPU" 2>/dev/null \ + | awk -F', ' '{ + sm += $1; mem += $2; n++ + } END { + if (n > 0) printf "%.1f,%.1f", sm/n, mem/n + else printf "0.0,0.0" + }' > "$out" || echo "0.0,0.0" > "$out" +} + +run_cell() { + local cell="$1" tx_count="$2" rx_count="$3" payload="$4" rep="${5:-1}" + local tx_cores rx_cores + [[ "$tx_count" == 2 ]] && tx_cores="11|12" || tx_cores="11" + [[ "$rx_count" == 2 ]] && rx_cores="9|14" || rx_cores="9" + local run_dir="$OUT_DIR/$cell/p$payload/r$rep" + mkdir -p "$run_dir" + + local tmp_cfg="$run_dir/config.yaml" + if ! python3 "$MQ_GEN" "$MQ_BASE" --tx "$tx_count" --rx "$rx_count" \ + --payload "$payload" --eth-dst "$ETH_DST_ADDR" \ + --tx-bdf "$TX_BDF" --rx-bdf "$RX_BDF" \ + --tx-gpu "$TX_GPU" --rx-gpu "$RX_GPU" \ + > "$tmp_cfg" 2> "$run_dir/gen.err"; then + echo "ERROR: $cell p$payload config generation failed" >&2 + cat "$run_dir/gen.err" >&2 + return 1 + fi + + local stdout="$run_dir/stdout.txt" + local stderr="$run_dir/stderr.txt" + + snapshot_cpu_stat "$run_dir/cpu_stat.before" + local phy_before phy_after + phy_before="$(max_phy_counter 'rx_phy_packets' /dev/null)" + phy_before="${phy_before:-0}" + + local bench_rc=0 + "$BENCH_BIN" "$tmp_cfg" --seconds "$RUN_SECONDS" > "$stdout" 2> "$stderr" || bench_rc=$? + + phy_after="$(max_phy_counter 'rx_phy_packets' "$stderr")" + phy_after="${phy_after:-0}" + snapshot_cpu_stat "$run_dir/cpu_stat.after" + + local tx_pkts tx_bytes rx_pkts rx_bytes secs + tx_pkts="$(sum_field 'TX complete' packets "$stdout")" + tx_bytes="$(sum_field 'TX complete' bytes "$stdout")" + rx_pkts="$(sum_field 'RX complete' packets "$stdout")" + rx_bytes="$(sum_field 'RX complete' bytes "$stdout")" + secs="$(max_field 'RX complete' seconds "$stdout")" + + if [[ "$bench_rc" -ne 0 || "${rx_pkts:-0}" -eq 0 || -z "${secs:-}" || "$secs" == "0" ]]; then + echo "ERROR: $cell p$payload failed (rc=$bench_rc rx_pkts=${rx_pkts:-0})" >&2 + return 1 + fi + + local pps tx_gbps rx_gbps drops + pps="$(awk -v p="$rx_pkts" -v s="$secs" 'BEGIN { if (s+0>0) printf "%.0f", p/s; else print 0 }')" + tx_gbps="$(awk -v b="$tx_bytes" -v s="$secs" 'BEGIN { if (s+0>0) printf "%.3f", (b*8.0)/s/1e9; else print 0 }')" + rx_gbps="$(awk -v b="$rx_bytes" -v s="$secs" 'BEGIN { if (s+0>0) printf "%.3f", (b*8.0)/s/1e9; else print 0 }')" + drops="$(parse_dpdk_drops "$stderr")" + + local phy_delta=$(( phy_after - phy_before )) + if [[ "${rx_pkts:-0}" -gt 1000 && "${phy_after:-0}" -lt 100 ]]; then + echo "WARN: $cell p$payload rx_phy_packets=${phy_after:-0} flat -- check cable" >&2 + elif [[ "$phy_delta" -gt 0 ]]; then + echo "INFO: $cell p$payload wire OK -- rx_phy_packets +$phy_delta" >&2 + fi + + local cpu_vals=() + local c + for c in "${CPU_CORES[@]}"; do + cpu_vals+=("$(cpu_busy_pct "$run_dir/cpu_stat.before" "$run_dir/cpu_stat.after" "$c")") + done + local cpu_csv gpu_csv + cpu_csv="$(IFS=,; echo "${cpu_vals[*]}")" + gpu_csv="$(gpu_sample "$run_dir/gpu.txt")" + + echo "$cell,$tx_cores,$rx_cores,$payload,$rep,$tx_gbps,$rx_gbps,$pps,$drops,$cpu_csv,$gpu_csv" | tee -a "$CSV" +} + +echo "RTX PRO 6000 multi-queue sweep -- ${RUN_SECONDS}s per (cell, payload)" +echo "TX BDF=$TX_BDF RX BDF=$RX_BDF ETH_DST=$ETH_DST_ADDR" +echo "Payloads: $PAYLOADS" +echo "Output: $OUT_DIR" +echo + +for entry in "${CELLS[@]}"; do + read -r cell tx_count rx_count <<< "$entry" + for payload in $PAYLOADS; do + for rep in $(seq 1 "$REPEATS"); do + run_cell "$cell" "$tx_count" "$rx_count" "$payload" "$rep" \ + || FAILURES=$((FAILURES + 1)) + done + done +done + +echo +echo "==================== RTX PRO multi-queue sweep ====================" +printf "%-6s %-9s %-9s %8s %10s %10s %8s\n" "cell" "tx_cores" "rx_cores" "payload" "tx_Gbps" "rx_Gbps" "drops" +printf "%-6s %-9s %-9s %8s %10s %10s %8s\n" "----" "--------" "--------" "-------" "-------" "-------" "-----" +while IFS=, read -r cell tx_cores rx_cores payload rep tx_gbps rx_gbps pps drops _rest; do + printf "%-6s %-9s %-9s %8s %10s %10s %8s\n" "$cell" "$tx_cores" "$rx_cores" "$payload" "$tx_gbps" "$rx_gbps" "$drops" +done < <(tail -n +2 "$CSV") +echo "===================================================================" +echo +echo "Results in: $OUT_DIR" +echo "CSV: $CSV" + +if [[ -f "$PLOT_SCRIPT" ]]; then + python3 "$PLOT_SCRIPT" "$CSV" "$OUT_DIR/plots" && echo "Plots: $OUT_DIR/plots/" +fi + +if [[ "$FAILURES" -ne 0 ]]; then + echo "Failed cells: $FAILURES" >&2 + exit 1 +fi diff --git a/examples/run_rtx_pro_push_400g.sh b/examples/run_rtx_pro_push_400g.sh new file mode 100755 index 0000000..44ce34a --- /dev/null +++ b/examples/run_rtx_pro_push_400g.sh @@ -0,0 +1,245 @@ +#!/usr/bin/env bash +# +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# +# RTX PRO 6000 — push toward 400 Gbps using the documented levers: +# 1) MQ 2t2r + larger TX/RX batches +# 2) Dual parallel wire loops (two bench processes, different EAL prefixes) +# 3) Baseline 1t1r for comparison +# +# Usage (root, privileged container): +# source scripts/discover_rtx_pro_topology.sh +# sudo -E env RTX_TX_BDF="$RTX_TX_BDF" RTX_RX_BDF="$RTX_RX_BDF" \ +# ETH_DST_ADDR="$ETH_DST_ADDR" RUN_SECONDS=20 \ +# ./examples/run_rtx_pro_push_400g.sh +# +# Output: bench-results/-rtx-pro-push400g/{runs.csv,plots/} + +set -u +set -o pipefail + +SCRIPT_DIR="$(cd "$(dirname "$0")" && pwd)" +REPO_ROOT="$(cd "$SCRIPT_DIR/.." && pwd)" +BUILD_DIR="${DAQIRI_BUILD_DIR:-$REPO_ROOT/build}" +DISCOVER="$REPO_ROOT/scripts/discover_rtx_pro_topology.sh" +BENCH_BIN="$BUILD_DIR/examples/daqiri_bench_raw_gpudirect" +MQ_BASE="$SCRIPT_DIR/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml" +MQ_GEN="$REPO_ROOT/scripts/gen_rtx_pro_mq_config.py" +NIC_BASE="$SCRIPT_DIR/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml" +PLOT_SCRIPT="$REPO_ROOT/scripts/plot_rtx_pro_bench.py" + +RUN_SECONDS="${RUN_SECONDS:-20}" +TX_GPU="${TX_GPU:-0}" +RX_GPU="${RX_GPU:-1}" + +export LD_LIBRARY_PATH="/opt/daqiri/lib:$BUILD_DIR:${LD_LIBRARY_PATH:-}" + +if [[ -x "$DISCOVER" ]]; then + # shellcheck disable=SC1090 + source "$DISCOVER" +fi + +TX_BDF="${RTX_TX_BDF:-0000:75:00.0}" +RX_BDF="${RTX_RX_BDF:-0000:05:00.0}" +TX_IF="${RTX_TX_IFACE:-}" +RX_IF="${RTX_RX_IFACE:-}" + +[[ -z "${ETH_DST_ADDR:-}" ]] && { echo "ERROR: ETH_DST_ADDR required" >&2; exit 1; } + +TS="$(date -u +%Y%m%dT%H%M%SZ)" +OUT_DIR="${DAQIRI_BENCH_RESULTS_DIR:-$REPO_ROOT/bench-results}/$TS-rtx-pro-push400g" +mkdir -p "$OUT_DIR" + +CSV="$OUT_DIR/runs.csv" +echo "experiment,cell,tx_cores,rx_cores,payload,batch,rep,tx_gbps,rx_gbps,pps,drops,notes" > "$CSV" + +FAILURES=0 + +sum_field() { + local prefix="$1" field="$2" file="$3" + grep -E "^$prefix" "$file" 2>/dev/null \ + | grep -oE " $field=[0-9]+" | sed -E "s/.* $field=//" \ + | awk '{ s += $1 } END { printf "%d", s+0 }' +} + +max_field() { + local prefix="$1" field="$2" file="$3" + grep -E "^$prefix" "$file" 2>/dev/null \ + | grep -oE " $field=[0-9.]+" | sed -E "s/.* $field=//" \ + | awk '{ if ($1+0 > m+0) m = $1 } END { printf "%s", (m == "" ? 0 : m) }' +} + +parse_dpdk_drops() { + local log="$1" sum=0 v + for key in imissed ierrors rx_nombuf; do + v="$(grep -oE "$key=[0-9]+" "$log" 2>/dev/null | tail -n1 | sed -E 's/.*=//' || true)" + [[ -n "${v:-}" ]] && sum=$((sum + v)) + done + echo "$sum" +} + +max_phy() { + grep -oE "${1}:[[:space:]]*[0-9]+" "$2" 2>/dev/null \ + | grep -oE '[0-9]+$' | sort -n | tail -n1 +} + +gen_nic_yaml() { + local out="$1" txbdf="$2" rxbdf="$3" rxmac="$4" payload="${5:-8000}" batch="${6:-10240}" + sed -E \ + -e "s|address: 0000:61:00.0|address: $txbdf|g" \ + -e "s|address: 0000:61:00.1|address: $rxbdf|g" \ + -e "s|<00:00:00:00:00:00>|$rxmac|g" \ + -e "s|^( *payload_size: ).*|\1$payload|" \ + -e "s|^( *batch_size: ).*|\1$batch|g" \ + "$NIC_BASE" > "$out" + awk -v txg="$TX_GPU" -v rxg="$RX_GPU" ' + /^ - name: "Data_TX_GPU"/ { in_tx=1; in_rx=0 } + /^ - name: "Data_RX_GPU"/ { in_rx=1; in_tx=0 } + in_tx && /^ affinity:/ { sub(/[0-9]+$/, txg); in_tx=0 } + in_rx && /^ affinity:/ { sub(/[0-9]+$/, rxg); in_rx=0 } + { print } + ' "$out" > "${out}.tmp" && mv "${out}.tmp" "$out" +} + +record_run() { + local exp="$1" cell="$2" tx_cores="$3" rx_cores="$4" payload="$5" batch="$6" + local stdout="$7" stderr="$8" notes="$9" + local tx_pkts rx_pkts tx_bytes rx_bytes secs tx_gbps rx_gbps pps drops + tx_pkts="$(sum_field 'TX complete' packets "$stdout")" + rx_pkts="$(sum_field 'RX complete' packets "$stdout")" + tx_bytes="$(sum_field 'TX complete' bytes "$stdout")" + rx_bytes="$(sum_field 'RX complete' bytes "$stdout")" + secs="$(max_field 'RX complete' seconds "$stdout")" + if [[ "${rx_pkts:-0}" -eq 0 || -z "${secs:-}" || "$secs" == "0" ]]; then + echo "FAIL: $exp $cell p$payload b$batch — rx_pkts=$rx_pkts" >&2 + FAILURES=$((FAILURES + 1)) + return 1 + fi + tx_gbps="$(awk -v b="$tx_bytes" -v s="$secs" 'BEGIN { printf "%.3f", (b*8.0)/s/1e9 }')" + rx_gbps="$(awk -v b="$rx_bytes" -v s="$secs" 'BEGIN { printf "%.3f", (b*8.0)/s/1e9 }')" + pps="$(awk -v p="$rx_pkts" -v s="$secs" 'BEGIN { printf "%.0f", p/s }')" + drops="$(parse_dpdk_drops "$stderr")" + echo "$exp,$cell,$tx_cores,$rx_cores,$payload,$batch,1,$tx_gbps,$rx_gbps,$pps,$drops,$notes" | tee -a "$CSV" +} + +run_mq_cell() { + local exp="$1" tx_count="$2" rx_count="$3" payload="$4" batch="$5" + local cell notes tx_cores rx_cores + if [[ "$tx_count" == 2 && "$rx_count" == 2 ]]; then cell="2t2r" + elif [[ "$tx_count" == 1 && "$rx_count" == 2 ]]; then cell="1t2r" + elif [[ "$tx_count" == 2 && "$rx_count" == 1 ]]; then cell="2t1r" + else cell="1t1r"; fi + if [[ "$tx_count" == 2 ]]; then tx_cores="11|12"; else tx_cores="11"; fi + if [[ "$rx_count" == 2 ]]; then rx_cores="9|14"; else rx_cores="9"; fi + notes="batch=$batch mq" + local run_dir="$OUT_DIR/mq-${cell}-p${payload}-b${batch}" + mkdir -p "$run_dir" + local cfg="$run_dir/config.yaml" + python3 "$MQ_GEN" "$MQ_BASE" --tx "$tx_count" --rx "$rx_count" \ + --payload "$payload" --batch "$batch" --eth-dst "$ETH_DST_ADDR" \ + --tx-bdf "$TX_BDF" --rx-bdf "$RX_BDF" \ + --tx-gpu "$TX_GPU" --rx-gpu "$RX_GPU" > "$cfg" 2>"$run_dir/gen.err" || return 1 + # Patch queue batch_size in daqiri cfg (generator sets bench_tx only). + awk -v b="$batch" ' + /^ batch_size:/ { sub(/[0-9]+$/, b); print; next } + { print } + ' "$cfg" > "${cfg}.tmp" && mv "${cfg}.tmp" "$cfg" + "$BENCH_BIN" "$cfg" --seconds "$RUN_SECONDS" > "$run_dir/stdout.txt" 2>"$run_dir/stderr.txt" || true + record_run "$exp" "$cell" "$tx_cores" "$rx_cores" "$payload" "$batch" \ + "$run_dir/stdout.txt" "$run_dir/stderr.txt" "$notes" +} + +run_nic_loop() { + local exp="$1" txbdf="$2" rxbdf="$3" rxmac="$4" label="$5" + local payload="${6:-8000}" batch="${7:-10240}" + local run_dir="$OUT_DIR/${exp}-${label}" + mkdir -p "$run_dir" + gen_nic_yaml "$run_dir/config.yaml" "$txbdf" "$rxbdf" "$rxmac" "$payload" "$batch" + "$BENCH_BIN" "$run_dir/config.yaml" --seconds "$RUN_SECONDS" \ + > "$run_dir/stdout.txt" 2>"$run_dir/stderr.txt" || true + record_run "$exp" "1t1r" "11" "9" "$payload" "$batch" \ + "$run_dir/stdout.txt" "$run_dir/stderr.txt" "$label" +} + +echo "=== RTX PRO push-400G experiments (${RUN_SECONDS}s each) ===" +echo "Primary loop: TX $TX_BDF -> RX $RX_BDF" +echo "Output: $OUT_DIR" +echo + +# Lever 1+5: MQ core scaling + batch tuning at 8 KB +for batch in 10240 16384 20480; do + run_mq_cell "mq-tune" 2 2 8000 "$batch" || true +done +run_mq_cell "mq-tune" 1 1 8000 10240 || true + +# Lever 4: single-queue batch sweep +for batch in 10240 16384 20480 32768; do + run_nic_loop "batch-tune" "$TX_BDF" "$RX_BDF" "$ETH_DST_ADDR" "loop1-b${batch}" 8000 "$batch" || true +done + +# Lever 2: small-packet MQ 2t2r (pps stress) +run_mq_cell "small-pkt" 2 2 64 10240 || true +run_mq_cell "small-pkt" 2 2 256 10240 || true + +# Lever 3: second wire pair probe (ens21 <-> ens20) +PAIR2_TX_BDF="0000:f5:00.0" +PAIR2_RX_BDF="0000:85:00.0" +PAIR2_RX_MAC="cc:40:f3:c6:ec:d8" +run_nic_loop "dual-loop-probe" "$PAIR2_TX_BDF" "$PAIR2_RX_BDF" "$PAIR2_RX_MAC" \ + "ens21-to-ens20" 8000 10240 || true +run_nic_loop "dual-loop-probe" "$PAIR2_RX_BDF" "$PAIR2_TX_BDF" \ + "$(cat /sys/class/net/ens21f0np0/address 2>/dev/null || echo cc:40:f3:c6:e5:30)" \ + "ens20-to-ens21" 8000 10240 || true + +# Lever 3b: dual parallel loops (primary + pair2 if pair2 works) +PAIR2_DIR="$OUT_DIR/dual-parallel" +mkdir -p "$PAIR2_DIR" +gen_nic_yaml "$PAIR2_DIR/loop1.yaml" "$TX_BDF" "$RX_BDF" "$ETH_DST_ADDR" 8000 16384 +gen_nic_yaml "$PAIR2_DIR/loop2.yaml" "$PAIR2_TX_BDF" "$PAIR2_RX_BDF" "$PAIR2_RX_MAC" 8000 16384 + +echo "Starting dual parallel run (loop1 + loop2)..." >&2 +"$BENCH_BIN" "$PAIR2_DIR/loop1.yaml" --seconds "$RUN_SECONDS" > "$PAIR2_DIR/l1.stdout" 2>"$PAIR2_DIR/l1.stderr" & +pid1=$! +"$BENCH_BIN" "$PAIR2_DIR/loop2.yaml" --seconds "$RUN_SECONDS" > "$PAIR2_DIR/l2.stdout" 2>"$PAIR2_DIR/l2.stderr" & +pid2=$! +wait "$pid1" || true +wait "$pid2" || true + +l1_rx="$(sum_field 'RX complete' bytes "$PAIR2_DIR/l1.stdout")" +l2_rx="$(sum_field 'RX complete' bytes "$PAIR2_DIR/l2.stdout")" +l1_secs="$(max_field 'RX complete' seconds "$PAIR2_DIR/l1.stdout")" +l2_secs="$(max_field 'RX complete' seconds "$PAIR2_DIR/l2.stdout")" +agg_gbps="$(awk -v b1="$l1_rx" -v b2="$l2_rx" -v s1="$l1_secs" -v s2="$l2_secs" ' + BEGIN { + g1 = (s1+0>0) ? (b1*8.0)/s1/1e9 : 0 + g2 = (s2+0>0) ? (b2*8.0)/s2/1e9 : 0 + printf "%.3f", g1+g2 + }')" +l1_gbps="$(awk -v b="$l1_rx" -v s="$l1_secs" 'BEGIN { if (s+0>0) printf "%.3f", (b*8.0)/s/1e9; else print 0 }')" +l2_gbps="$(awk -v b="$l2_rx" -v s="$l2_secs" 'BEGIN { if (s+0>0) printf "%.3f", (b*8.0)/s/1e9; else print 0 }')" +echo "dual-parallel,aggregate,|,|,8000,16384,1,$l1_gbps,$l2_gbps,0,0,loop1=${l1_gbps}G loop2=${l2_gbps}G sum=${agg_gbps}G" | tee -a "$CSV" +echo "dual-parallel,sum,,,8000,16384,1,0,${agg_gbps},0,0,aggregate_rx_gbps" | tee -a "$CSV" + +echo +echo "==================== Results ====================" +column -t -s, "$CSV" 2>/dev/null || cat "$CSV" +echo +best="$(awk -F, 'NR>1 && $9+0>max {max=$9+0; line=$0} END {print line}' "$CSV")" +echo "Best single RX Gbps: $best" +best_agg="$(awk -F, '$1=="dual-parallel" && $2=="sum" {print $9; exit}' "$CSV")" +echo "Dual-loop aggregate RX Gbps: ${best_agg:-n/a}" + +# Plot MQ-style if we have enough rows +if [[ -f "$PLOT_SCRIPT" ]]; then + awk -F, 'NR==1 {print "cell,tx_cores,rx_cores,payload,rep,gbps,rx_gbps,pps,drops"; next} + $1 ~ /^mq-tune/ {print $2","$3","$4","$5",1,"$7","$8","$9","$10}' "$CSV" \ + > "$OUT_DIR/mq_subset.csv" + if [[ $(wc -l < "$OUT_DIR/mq_subset.csv") -gt 2 ]]; then + python3 "$PLOT_SCRIPT" "$OUT_DIR/mq_subset.csv" "$OUT_DIR/plots" || true + fi +fi + +echo "Results: $OUT_DIR" +[[ "$FAILURES" -gt 0 ]] && echo "Failures: $FAILURES" >&2 +exit 0 diff --git a/scripts/gen_rtx_pro_mq_config.py b/scripts/gen_rtx_pro_mq_config.py new file mode 100755 index 0000000..c402290 --- /dev/null +++ b/scripts/gen_rtx_pro_mq_config.py @@ -0,0 +1,117 @@ +#!/usr/bin/env python3 +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +"""Derive an RTX PRO 6000 multi-queue cell from the (TX=2, RX=2) base. + +Same pruning rules as scripts/gen_spark_mq_config.py, plus RTX-specific +substitutions for PCIe BDFs, GPU affinities, and the RX-port MAC. +""" + +from __future__ import annotations + +import argparse +import sys + +import yaml + +TX_PORT_BASE = 4096 + + +def prune(base: dict, tx: int, rx: int, payload: int, batch: int | None, + eth_dst: str | None, tx_bdf: str | None, rx_bdf: str | None, + tx_gpu: int, rx_gpu: int) -> dict: + cfg = base["daqiri"]["cfg"] + n_ports = max(tx, rx) + + def keep_region(name: str) -> bool: + if name.endswith("_1"): + return tx == 2 if "_TX_" in name else rx == 2 + return True + + cfg["memory_regions"] = [ + m for m in cfg["memory_regions"] if keep_region(m["name"]) + ] + for m in cfg["memory_regions"]: + if "_TX_" in m["name"]: + m["affinity"] = tx_gpu + elif "_RX_" in m["name"]: + m["affinity"] = rx_gpu + + for iface in cfg["interfaces"]: + if iface["name"] == "tx_port" and tx_bdf: + iface["address"] = tx_bdf + if iface["name"] == "rx_port" and rx_bdf: + iface["address"] = rx_bdf + if "tx" in iface: + iface["tx"]["queues"] = [ + q for q in iface["tx"]["queues"] if q["id"] == 0 or tx == 2 + ] + if "rx" in iface: + iface["rx"]["queues"] = [ + q for q in iface["rx"]["queues"] if q["id"] == 0 or rx == 2 + ] + flows = [] + for f in iface["rx"]["flows"]: + if f["id"] == 0: + flows.append(f) + elif n_ports == 2: + f = dict(f) + f["action"] = dict(f["action"]) + f["action"]["id"] = 1 if rx == 2 else 0 + flows.append(f) + iface["rx"]["flows"] = flows + + base["bench_rx"] = [ + b for b in base["bench_rx"] if b["queue_id"] == 0 or rx == 2 + ] + + bench_tx = [b for b in base["bench_tx"] if b["queue_id"] == 0 or tx == 2] + for b in bench_tx: + b["payload_size"] = payload + if batch is not None: + b["batch_size"] = batch + if eth_dst is not None: + b["eth_dst_addr"] = eth_dst + if b["queue_id"] == 0 and tx == 1 and rx == 2: + b["udp_src_port"] = f"{TX_PORT_BASE}-{TX_PORT_BASE + 1}" + b["udp_dst_port"] = f"{TX_PORT_BASE}-{TX_PORT_BASE + 1}" + base["bench_tx"] = bench_tx + + return base + + +def main() -> int: + ap = argparse.ArgumentParser(description=__doc__) + ap.add_argument("base", help="path to the (2,2) superset base YAML") + ap.add_argument("--tx", type=int, choices=(1, 2), required=True) + ap.add_argument("--rx", type=int, choices=(1, 2), required=True) + ap.add_argument("--payload", type=int, required=True) + ap.add_argument("--batch", type=int, default=None) + ap.add_argument("--eth-dst", default=None) + ap.add_argument("--tx-bdf", default=None) + ap.add_argument("--rx-bdf", default=None) + ap.add_argument("--tx-gpu", type=int, default=0) + ap.add_argument("--rx-gpu", type=int, default=1) + args = ap.parse_args() + + with open(args.base, encoding="utf-8") as fh: + base = yaml.safe_load(fh) + + out = prune( + base, + args.tx, + args.rx, + args.payload, + args.batch, + args.eth_dst, + args.tx_bdf, + args.rx_bdf, + args.tx_gpu, + args.rx_gpu, + ) + yaml.safe_dump(out, sys.stdout, sort_keys=False, default_flow_style=False) + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/scripts/plot_rtx_pro_bench.py b/scripts/plot_rtx_pro_bench.py new file mode 100755 index 0000000..82eb2be --- /dev/null +++ b/scripts/plot_rtx_pro_bench.py @@ -0,0 +1,345 @@ +#!/usr/bin/env python3 +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +"""Render RTX PRO 6000 multi-queue sweep plots (stdlib only, writes SVG). + +Reads runs.csv from examples/run_rtx_pro_mq_bench.sh and writes: + gbps_vs_payload.svg, core_scaling.svg, cpu_utilization.svg, drops_vs_payload.svg + +If matplotlib is installed, also writes PNG copies for convenience. +""" + +from __future__ import annotations + +import csv +import math +import sys +from collections import OrderedDict +from pathlib import Path +from xml.sax.saxutils import escape + +CELL_ORDER = ["1t1r", "1t2r", "2t1r", "2t2r"] +CELL_LABEL = { + "1t1r": "1 TX / 1 RX", + "1t2r": "1 TX / 2 RX", + "2t1r": "2 TX / 1 RX", + "2t2r": "2 TX / 2 RX", +} +COLORS = ["#4C72B0", "#55A868", "#C44E52", "#8172B2"] + +CPU_COLS = [ + ("cpu_master", "master"), + ("cpu_tx_q0_poll", "TX q0 poll"), + ("cpu_tx_q0_work", "TX q0 work"), + ("cpu_tx_q1_poll", "TX q1 poll"), + ("cpu_tx_q1_work", "TX q1 work"), + ("cpu_rx_q0_poll", "RX q0 poll"), + ("cpu_rx_q0_work", "RX q0 work"), + ("cpu_rx_q1_poll", "RX q1 poll"), + ("cpu_rx_q1_work", "RX q1 work"), +] + + +def load_rows(csv_path: Path) -> list[dict]: + with open(csv_path, newline="") as fh: + return list(csv.DictReader(fh)) + + +def series_by_cell(rows: list[dict], key: str) -> OrderedDict[str, list[tuple[int, float]]]: + out: OrderedDict[str, list[tuple[int, float]]] = OrderedDict((c, []) for c in CELL_ORDER) + for row in rows: + cell = row["cell"].strip() + if cell not in out: + out[cell] = [] + out[cell].append((int(row["payload"]), float(row[key]))) + for cell in out: + out[cell].sort(key=lambda pt: pt[0]) + return out + + +def largest_payload(rows: list[dict]) -> int: + return max(int(r["payload"]) for r in rows) + + +def rows_at_payload(rows: list[dict], payload: int) -> dict[str, dict]: + picked: dict[str, dict] = {} + for row in rows: + if int(row["payload"]) != payload: + continue + cell = row["cell"].strip() + if cell not in picked: + picked[cell] = row + return picked + + +class Svg: + def __init__(self, w: float, h: float, bg: str = "#ffffff"): + self.w, self.h = w, h + self.parts: list[str] = [] + if bg: + self.parts.append( + f'' + ) + + def text(self, x, y, s, size=12, anchor="start", weight="normal"): + self.parts.append( + f'{escape(s)}' + ) + + def line(self, x1, y1, x2, y2, color="#333", width=1): + self.parts.append( + f'' + ) + + def rect(self, x, y, w, h, fill="#4C72B0", stroke="#222", sw=0.6): + self.parts.append( + f'' + ) + + def circle(self, x, y, r=4, fill="#4C72B0"): + self.parts.append(f'') + + def polyline(self, pts, color="#4C72B0", width=2): + if len(pts) < 2: + return + d = " ".join(f"{x:.1f},{y:.1f}" for x, y in pts) + self.parts.append( + f'' + ) + + def save(self, path: Path): + body = "\n ".join(self.parts) + path.write_text( + f'\n' + f'\n {body}\n\n' + ) + + +def log_scale_x(payload: int, xmin: int, xmax: int, left: float, right: float) -> float: + if xmin <= 0 or xmax <= xmin: + return left + return left + (math.log2(payload) - math.log2(xmin)) / (math.log2(xmax) - math.log2(xmin)) * (right - left) + + +def plot_lines(title, xlabel, ylabel, series, out_path: Path, log_x: bool = True): + margin = dict(l=70, r=200, t=50, b=55) + W, H = 900, 520 + svg = Svg(W, H) + plot_l, plot_r = margin["l"], W - margin["r"] + plot_t, plot_b = margin["t"], H - margin["b"] + + all_y = [v for pts in series.values() for _, v in pts] + ymax = max(all_y) * 1.12 if all_y else 1 + payloads = sorted({p for pts in series.values() for p, _ in pts}) + xmin, xmax = payloads[0], payloads[-1] + + def x_map(p): + return log_scale_x(p, xmin, xmax, plot_l, plot_r) if log_x else plot_l + (p - xmin) / max(xmax - xmin, 1) * (plot_r - plot_l) + + def y_map(v): + return plot_b - (v / ymax) * (plot_b - plot_t) + + # grid + y ticks + for i in range(6): + yv = ymax * i / 5 + y = y_map(yv) + svg.line(plot_l, y, plot_r, y, color="#ddd") + svg.text(plot_l - 8, y + 4, f"{yv:.0f}", size=10, anchor="end") + + for p in payloads: + x = x_map(p) + svg.line(x, plot_t, x, plot_b, color="#eee") + svg.text(x, plot_b + 18, str(p), size=10, anchor="middle") + + svg.line(plot_l, plot_b, plot_r, plot_b, color="#333") + svg.line(plot_l, plot_t, plot_l, plot_b, color="#333") + + for idx, cell in enumerate(CELL_ORDER): + pts = series.get(cell) or [] + if not pts: + continue + color = COLORS[idx % len(COLORS)] + xy = [(x_map(p), y_map(v)) for p, v in pts] + svg.polyline(xy, color=color, width=2.5) + for x, y in xy: + svg.circle(x, y, r=4, fill=color) + + svg.text(W / 2, 24, title, size=15, anchor="middle", weight="bold") + svg.text(W / 2, H - 12, xlabel, size=12, anchor="middle") + svg.text(16, H / 2, ylabel, size=12, anchor="middle") + + ly = plot_t + 10 + for idx, cell in enumerate(CELL_ORDER): + if not series.get(cell): + continue + color = COLORS[idx % len(COLORS)] + lx = plot_r + 20 + svg.line(lx, ly, lx + 24, ly, color=color, width=3) + svg.circle(lx + 12, ly, r=4, fill=color) + svg.text(lx + 32, ly + 4, CELL_LABEL[cell], size=11) + ly += 22 + + svg.save(out_path) + + +def plot_bars(title, labels, values, out_path: Path): + W, H = 820, 480 + svg = Svg(W, H) + margin = dict(l=60, r=30, t=50, b=90) + plot_l, plot_r = margin["l"], W - margin["r"] + plot_t, plot_b = margin["t"], H - margin["b"] + ymax = max(values) * 1.15 if values else 1 + n = len(values) + gap = 18 + bar_w = (plot_r - plot_l - gap * (n - 1)) / max(n, 1) + + for i in range(6): + yv = ymax * i / 5 + y = plot_b - (yv / ymax) * (plot_b - plot_t) + svg.line(plot_l, y, plot_r, y, color="#ddd") + svg.text(plot_l - 8, y + 4, f"{yv:.0f}", size=10, anchor="end") + + for i, (lab, val) in enumerate(zip(labels, values)): + x = plot_l + i * (bar_w + gap) + h = (val / ymax) * (plot_b - plot_t) + y = plot_b - h + svg.rect(x, y, bar_w, h, fill=COLORS[i % len(COLORS)]) + svg.text(x + bar_w / 2, plot_b + 16, lab, size=10, anchor="middle") + svg.text(x + bar_w / 2, y - 6, f"{val:.0f}", size=10, anchor="middle") + + svg.line(plot_l, plot_b, plot_r, plot_b, color="#333") + svg.text(W / 2, 24, title, size=15, anchor="middle", weight="bold") + svg.text(16, H / 2, "RX Gb/s", size=12, anchor="middle") + svg.save(out_path) + + +def plot_cpu_panels(rows_at_max, max_payload: int, out_path: Path): + W, H = 1000, 620 + svg = Svg(W, H) + svg.text(W / 2, 24, f"CPU utilization at {max_payload} B payload", size=15, anchor="middle", weight="bold") + panel_w = (W - 60) / 2 + panel_h = (H - 80) / 2 + origins = [(30, 50), (30 + panel_w + 20, 50), (30, 50 + panel_h + 20), (30 + panel_w + 20, 50 + panel_h + 20)] + + for idx, cell in enumerate(CELL_ORDER): + row = rows_at_max.get(cell) + if not row: + continue + ox, oy = origins[idx] + labels, vals = [], [] + for col, label in CPU_COLS: + raw = row.get(col, "") + if not raw: + continue + v = float(raw) + if v <= 0.2 and "q1" in col: + continue + labels.append(label) + vals.append(v) + if not vals: + continue + svg.text(ox, oy - 8, CELL_LABEL[cell], size=12, weight="bold") + bar_h = min(22, (panel_h - 30) / max(len(vals), 1)) + for i, (lab, val) in enumerate(zip(labels, vals)): + y = oy + 10 + i * (bar_h + 4) + bw = (val / 100.0) * (panel_w - 120) + svg.rect(ox + 110, y, bw, bar_h - 2, fill=COLORS[idx % len(COLORS)]) + svg.text(ox + 4, y + bar_h - 5, lab, size=9) + svg.text(ox + 115 + bw + 4, y + bar_h - 5, f"{val:.0f}%", size=9) + svg.line(ox + 110, oy + panel_h - 10, ox + panel_w - 10, oy + panel_h - 10, color="#333") + svg.save(out_path) + + +def maybe_matplotlib_png(rows, out_dir: Path): + try: + import matplotlib + matplotlib.use("Agg") + import matplotlib.pyplot as plt + import numpy as np + except ImportError: + return + + gbps = series_by_cell(rows, "rx_gbps") + max_p = largest_payload(rows) + at_max = rows_at_payload(rows, max_p) + + fig, ax = plt.subplots(figsize=(8, 4.8)) + for idx, cell in enumerate(CELL_ORDER): + pts = gbps.get(cell) or [] + if not pts: + continue + xs, ys = zip(*pts) + ax.plot(xs, ys, marker="o", label=CELL_LABEL[cell], color=COLORS[idx]) + ax.set_xscale("log", base=2) + ax.set_xlabel("Payload (bytes)") + ax.set_ylabel("RX Gb/s") + ax.set_title("RTX PRO 6000 — throughput vs payload") + ax.grid(True, alpha=0.4) + ax.legend() + fig.tight_layout() + fig.savefig(out_dir / "gbps_vs_payload.png", dpi=150, facecolor="white") + plt.close(fig) + + labels = [CELL_LABEL[c] for c in CELL_ORDER if c in at_max] + values = [float(at_max[c]["rx_gbps"]) for c in CELL_ORDER if c in at_max] + fig, ax = plt.subplots(figsize=(7.5, 4.5)) + ax.bar(labels, values, color=COLORS[:len(values)]) + ax.set_ylabel("RX Gb/s") + ax.set_title(f"Core scaling at {max_p} B") + fig.tight_layout() + fig.savefig(out_dir / "core_scaling.png", dpi=150, facecolor="white") + plt.close(fig) + + +def main(argv: list[str]) -> int: + if len(argv) < 2: + print(__doc__, file=sys.stderr) + return 1 + csv_path = Path(argv[1]) + out_dir = Path(argv[2]) if len(argv) >= 3 else csv_path.parent / "plots" + out_dir.mkdir(parents=True, exist_ok=True) + + rows = load_rows(csv_path) + if not rows: + print(f"no rows in {csv_path}", file=sys.stderr) + return 1 + + max_p = largest_payload(rows) + at_max = rows_at_payload(rows, max_p) + gbps = series_by_cell(rows, "rx_gbps") + drops = series_by_cell(rows, "drops") + + plot_lines( + "RTX PRO 6000 — wire RX throughput vs payload", + "UDP payload size (bytes, log₂ scale)", + "RX Gb/s", + gbps, + out_dir / "gbps_vs_payload.svg", + ) + plot_bars( + f"Core scaling at {max_p} B payload", + [CELL_LABEL[c] for c in CELL_ORDER if c in at_max], + [float(at_max[c]["rx_gbps"]) for c in CELL_ORDER if c in at_max], + out_dir / "core_scaling.svg", + ) + plot_cpu_panels(at_max, max_p, out_dir / "cpu_utilization.svg") + plot_lines( + "RTX PRO 6000 — DPDK drops vs payload", + "UDP payload size (bytes, log₂ scale)", + "drops", + drops, + out_dir / "drops_vs_payload.svg", + ) + maybe_matplotlib_png(rows, out_dir) + + for p in sorted(out_dir.glob("*.svg")) + sorted(out_dir.glob("*.png")): + print(f"wrote {p}") + return 0 + + +if __name__ == "__main__": + sys.exit(main(sys.argv)) From 84a2b3d9aa3b6099a5c3696e3383c3a58100bbd9 Mon Sep 17 00:00:00 2001 From: Chloe Crozier Date: Fri, 10 Jul 2026 18:52:22 +0000 Subject: [PATCH 5/7] #17 - Add RTX PRO 6000 configs to docs decision tree The merge brought five RTX PRO 6000 example YAMLs (HDS, MQ base, ibverbs RX, TX-only, RoCE) into examples/ that lacked a leaf in the configuration walkthrough decision tree, failing scripts/check_doc_refs.py in the docs CI. Add a decision-tree entry for each so every shipped YAML is covered. Signed-off-by: Chloe Crozier --- docs/tutorials/configuration-walkthrough.md | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/docs/tutorials/configuration-walkthrough.md b/docs/tutorials/configuration-walkthrough.md index 51fe23b..e3b74bf 100644 --- a/docs/tutorials/configuration-walkthrough.md +++ b/docs/tutorials/configuration-walkthrough.md @@ -36,6 +36,9 @@ For a shorter selection guide, start with the [Benchmarking overview](../benchma - **RTX PRO 6000 Blackwell — real NIC, dual-port on one card** (prefilled dev box) — [`daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic.yaml). `61:00.0` p0 → `61:00.1` p1, GPU 0 TX / GPU 1 RX; needs L2 link between ports (not SW loopback). See [`rtx_pro_6000_baseline.md`](https://github.com/nvidia/daqiri/blob/main/examples/rtx_pro_6000_baseline.md) for hardware limits and measured baseline. - **RTX PRO 6000 — same-PF NIC attempt** (experimental) — [`daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_nic_same_port.yaml). Single port `61:00.0` TX+RX; failed `daqiri_init` on reference box — kept for follow-up. - **RTX PRO 6000 Blackwell — dual-NIC loopback** (generic template) — [`daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000.yaml). Placeholders for Cliff's 800 Gbps cross-card target; fill PCIe BDFs and MACs. + - **RTX PRO 6000 Blackwell — header-data split wire loopback** — [`daqiri_bench_raw_tx_rx_hds_rtx_pro_6000.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_hds_rtx_pro_6000.yaml). Runs on `daqiri_bench_raw_hds`; CPU header segment + GPU (`kind: device`) payload segment over the dual-port L2 link. + - **RTX PRO 6000 Blackwell — multi-queue core-scaling matrix** — one base config [`daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml) (the TX=2/RX=2 superset; GPU 0 TX / GPU 1 RX) from which `examples/run_rtx_pro_mq_bench.sh` (via `scripts/gen_rtx_pro_mq_config.py`) derives the four `(TX, RX)` cells — (1,1), (1,2), (2,1), (2,2) — by pruning queues/flows and sweeping payload sizes. All run on `daqiri_bench_raw_gpudirect`. + - **RTX PRO 6000 Blackwell — ibverbs (MPRQ) RX, paired DPDK TX** — the RX side uses [`daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml) (`engine: ibverbs`, runs on `daqiri_bench_raw_gpudirect` built with `ibverbs`), fed by a DPDK TX sender from [`daqiri_bench_raw_tx_only_rtx_pro_6000_nic.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_raw_tx_only_rtx_pro_6000_nic.yaml) on the other port over the L2 link. **Raw Ethernet hardware tunnel transforms** — run on `daqiri_bench_raw_gpudirect`; replace placeholders and use a raw DPDK or raw ibverbs build. @@ -51,6 +54,7 @@ For a shorter selection guide, start with the [Benchmarking overview](../benchma **Socket — RoCE (RDMA)** (`stream_type: "socket"`, `roce://` endpoints) — runs on `daqiri_bench_rdma` (use `--mode {tx,rx,both}`). Configs use `kind: host_pinned` regardless of platform. - **Generic** (template — replace IPs) — [`daqiri_bench_rdma_tx_rx.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_rdma_tx_rx.yaml). + - **RTX PRO 6000 Blackwell** (prefilled dual-port RoCE) — [`daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_rdma_tx_rx_rtx_pro_6000.yaml). Client/server on `10.100.0.1`/`10.100.0.2` across the two cabled ports; `examples/run_rtx_pro_bench.sh` auto-assigns the RoCE IPs and sysctls for a self-contained single-host loop. - **DGX Spark** (prefilled) — [`daqiri_bench_rdma_tx_rx_spark.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_rdma_tx_rx_spark.yaml). See [Socket and RDMA Benchmarking](../benchmarks/socket_benchmarking.md#run-the-rdma-roce-benchmark) for namespace and wire-counter run details. - **DGX Spark netns wire loopback** (prefilled, combined base) — [`daqiri_bench_rdma_tx_rx_spark_netns.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_rdma_tx_rx_spark_netns.yaml). Carries both roles; `examples/run_spark_bench.sh` (via `scripts/gen_spark_netns_config.py`) splits it per role and runs each in its own network namespace (`--mode server` / `--mode client`) so RDMA-CM resolves over the wire; see [Socket and RDMA Benchmarking](../benchmarks/socket_benchmarking.md#run-the-rdma-roce-benchmark). - **DGX Spark cross-host** (prefilled, runs on two Sparks) — [`daqiri_bench_rdma_tx_rx_spark_xhost.yaml`](https://github.com/nvidia/daqiri/blob/main/examples/daqiri_bench_rdma_tx_rx_spark_xhost.yaml). Run with `--mode server` on the RX host and `--mode client` on the TX host. Apply the [cross-host network setup](../tutorials/system_configuration.md#cross-host-variant-two-sparks) before running. See the [Cross-host two-DGX-Spark loopback](../benchmarks/raw_benchmarking.md#cross-host-two-dgx-spark-loopback) section for run details. From ac8acc25263d8f3f9d71a60802e8c564a1b2657c Mon Sep 17 00:00:00 2001 From: Chloe Crozier Date: Tue, 14 Jul 2026 13:38:09 +0000 Subject: [PATCH 6/7] #17 - Auto-discover RTX PRO GPU/CPU affinity for wire benchmarks Add rtx_pro_discover.py to pick PIX-linked CUDA ordinals per NIC and isolcpus-based poll cores from topology. Wire discovery into the MQ and single-backend runners, split multi-queue TX across two GPUs, default RoCE workload-max-inflight, and add run_rtx_pro_suite.sh for the standard RTX PRO sweep. Signed-off-by: ccrozier --- ...aqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml | 20 +- examples/run_rtx_pro_bench.sh | 22 +- examples/run_rtx_pro_mq_bench.sh | 46 ++- examples/run_rtx_pro_suite.sh | 50 +++ scripts/discover_rtx_pro_topology.sh | 19 ++ scripts/gen_rtx_pro_mq_config.py | 53 ++- scripts/rtx_pro_discover.py | 302 ++++++++++++++++++ 7 files changed, 483 insertions(+), 29 deletions(-) create mode 100755 examples/run_rtx_pro_suite.sh create mode 100755 scripts/rtx_pro_discover.py diff --git a/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml index 3b6b3d3..66a4c26 100644 --- a/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml +++ b/examples/daqiri_bench_raw_tx_rx_rtx_pro_6000_mq.yaml @@ -8,12 +8,14 @@ # eth_dst_addr from discovery (ETH_DST_ADDR) or: # cat /sys/class/net//address # -# Core map (pollers are DPDK EAL lcores; workers are bench pthreads): -# master_core 3 -# TX q0: poller 11 / worker 10 -# TX q1: poller 12 / worker 13 -# RX q0: poller 9 / worker 8 -# RX q1: poller 14 / worker 15 +# Core map (pollers are DPDK EAL lcores; workers are bench pthreads). +# Defaults below are placeholders; run_rtx_pro_mq_bench.sh replaces them via +# scripts/rtx_pro_discover.py (isolcpus on the NIC NUMA node + PIX GPU ordinals). +# master_core, TX q0 poll/work, TX q1 poll/work, RX q0 poll/work, RX q1 poll/work +# +# GPU affinity is the CUDA ordinal under CUDA_DEVICE_ORDER=PCI_BUS_ID. Discovery +# picks PIX-linked GPUs per NIC and spreads TX queue 1 onto a second PIX GPU when +# available. Override with RTX_TX_GPU / RTX_RX_GPU / RTX_TX_GPU2 / RTX_RX_GPU2. # %YAML 1.2 --- @@ -34,7 +36,7 @@ daqiri: buf_size: 8064 - name: "Data_TX_GPU_1" kind: "device" - affinity: 0 + affinity: 1 num_bufs: 51200 buf_size: 8064 - name: "Data_RX_GPU_0" @@ -50,7 +52,7 @@ daqiri: interfaces: - name: "tx_port" - address: "0000:75:00.0" + address: <0000:00:00.0> tx: queues: - name: "tx_q_0" @@ -70,7 +72,7 @@ daqiri: offloads: - "tx_eth_src" - name: "rx_port" - address: "0000:05:00.0" + address: <0000:00:00.1> rx: flow_isolation: true queues: diff --git a/examples/run_rtx_pro_bench.sh b/examples/run_rtx_pro_bench.sh index 2283f33..a9b4a99 100755 --- a/examples/run_rtx_pro_bench.sh +++ b/examples/run_rtx_pro_bench.sh @@ -104,6 +104,7 @@ esac WORKLOAD_BATCH="${WORKLOAD_BATCH:-}" GEMM_N="${GEMM_N:-}" SYNC_INTERVAL="${SYNC_INTERVAL:-}" +MAX_INFLIGHT="${MAX_INFLIGHT:-16}" # -------------------------------------------------------------------------- # build-only @@ -145,9 +146,12 @@ if [[ "$MODE" == "build-only" ]]; then exit $? fi +export LD_LIBRARY_PATH="/opt/daqiri/lib:$BUILD_DIR:${LD_LIBRARY_PATH:-}" 2>/dev/null || true + # -------------------------------------------------------------------------- # Discovery + wire preflight # -------------------------------------------------------------------------- +export CUDA_DEVICE_ORDER="${CUDA_DEVICE_ORDER:-PCI_BUS_ID}" if [[ -x "$DISCOVER" ]]; then # shellcheck disable=SC1090 source "$DISCOVER" @@ -181,8 +185,11 @@ FAILURES=0 TX_BDF="${RTX_TX_BDF:-0000:61:00.0}" RX_BDF="${RTX_RX_BDF:-0000:61:00.1}" -TX_GPU="${TX_GPU:-0}" -RX_GPU="${RX_GPU:-1}" +TX_GPU="${TX_GPU:-${RTX_TX_GPU:-0}}" +RX_GPU="${RX_GPU:-${RTX_RX_GPU:-1}}" +CPU_MASTER="${RTX_MASTER_CORE:-3}" +CPU_TX="${RTX_TX_Q0_POLL:-11}" +CPU_RX="${RTX_RX_Q0_POLL:-9}" wire_preflight() { if [[ "$IS_SW" == "1" ]]; then @@ -346,6 +353,11 @@ case "$BACKEND" in ;; esac +# Discovery overrides backend defaults when sourced from discover_rtx_pro_topology.sh. +CPU_MASTER="${RTX_MASTER_CORE:-$CPU_MASTER}" +CPU_TX="${RTX_TX_Q0_POLL:-$CPU_TX}" +CPU_RX="${RTX_RX_Q0_POLL:-$CPU_RX}" + DROP_CURVE_TARGETS=(1 5 10 25 50 75 100 0) # SW loopback rides the single-segment GPUDirect ring in the dpdk engine. HDS @@ -489,6 +501,11 @@ generate_yaml() { in_rx && /^ affinity:/ { sub(/[0-9]+$/, rxg); in_rx=0 } { print } ' "$out" > "${out}.tmp" && mv "${out}.tmp" "$out" + sed -E \ + -e "s|^( master_core:)[[:space:]]*[0-9]+|\1 $CPU_MASTER|" \ + -e "/- name: \"tx_port\"/,/- name: \"rx_port\"/ s|^( cpu_core:)[[:space:]]*[0-9]+|\1 $CPU_TX|" \ + -e "/- name: \"rx_port\"/,/^bench_/ s|^( cpu_core:)[[:space:]]*[0-9]+|\1 $CPU_RX|" \ + "$out" > "${out}.tmp" && mv "${out}.tmp" "$out" ;; dpdk-hds) local ipv4_len=$((payload + 50)) @@ -569,6 +586,7 @@ run_cell() { [[ -n "$WORKLOAD_BATCH" ]] && args+=(--workload-batch-bytes "$WORKLOAD_BATCH") [[ -n "$GEMM_N" ]] && args+=(--workload-gemm-n "$GEMM_N") [[ -n "$SYNC_INTERVAL" ]] && args+=(--workload-sync-interval "$SYNC_INTERVAL") + [[ -n "$MAX_INFLIGHT" && "$BACKEND" == "rdma" ]] && args+=(--workload-max-inflight "$MAX_INFLIGHT") fi local tx_pid="" diff --git a/examples/run_rtx_pro_mq_bench.sh b/examples/run_rtx_pro_mq_bench.sh index 7cc6e36..4fe404c 100755 --- a/examples/run_rtx_pro_mq_bench.sh +++ b/examples/run_rtx_pro_mq_bench.sh @@ -10,10 +10,10 @@ # size. Cells: # # cell TX pollers RX pollers -# 1t1r 11 9 -# 1t2r 11 9,14 -# 2t1r 11,12 9 -# 2t2r 11,12 9,14 +# 1t1r 1 1 +# 1t2r 1 2 +# 2t1r 2 1 +# 2t2r 2 2 # # Usage (as root, privileged container): # source scripts/discover_rtx_pro_topology.sh @@ -24,6 +24,7 @@ # PAYLOADS space-separated payload bytes (default "64 256 1024 4096 8000") # REPEATS repeats per (cell, payload) (default 1) # RTX_TX_BDF, RTX_RX_BDF, RTX_RX_IFACE, ETH_DST_ADDR from discovery +# RTX_TX_GPU, RTX_RX_GPU, RTX_TX_GPU2, RTX_RX_GPU2, RTX_CPU_CORES (discovery) # # Plot afterwards: # scripts/plot_rtx_pro_bench.py bench-results/-rtx-pro-mq/runs.csv @@ -45,17 +46,26 @@ MQ_GEN="$REPO_ROOT/scripts/gen_rtx_pro_mq_config.py" PLOT_SCRIPT="$REPO_ROOT/scripts/plot_rtx_pro_bench.py" export LD_LIBRARY_PATH="/opt/daqiri/lib:$BUILD_DIR:${LD_LIBRARY_PATH:-}" +export CUDA_DEVICE_ORDER="${CUDA_DEVICE_ORDER:-PCI_BUS_ID}" if [[ -x "$DISCOVER" ]]; then # shellcheck disable=SC1090 source "$DISCOVER" fi -TX_BDF="${RTX_TX_BDF:-0000:75:00.0}" -RX_BDF="${RTX_RX_BDF:-0000:05:00.0}" -TX_GPU="${TX_GPU:-0}" -RX_GPU="${RX_GPU:-1}" -RX_NETDEV="${RTX_RX_IFACE:-}" +TX_BDF="${RTX_TX_BDF:-}" +RX_BDF="${RTX_RX_BDF:-}" +TX_GPU="${TX_GPU:-${RTX_TX_GPU:-0}}" +RX_GPU="${RX_GPU:-${RTX_RX_GPU:-1}}" +TX_GPU2="${TX_GPU2:-${RTX_TX_GPU2:-$TX_GPU}}" +RX_GPU2="${RX_GPU2:-${RTX_RX_GPU2:-$RX_GPU}}" + +# master, TX q0 poll/work, TX q1 poll/work, RX q0 poll/work, RX q1 poll/work +if [[ -n "${RTX_CPU_CORES:-}" ]]; then + read -r -a CPU_CORES <<< "$RTX_CPU_CORES" +else + CPU_CORES=(3 11 10 12 13 9 8 14 15) +fi TS="$(date -u +%Y%m%dT%H%M%SZ)" OUT_DIR="${DAQIRI_BENCH_RESULTS_DIR:-$REPO_ROOT/bench-results}/$TS-rtx-pro-mq" @@ -80,6 +90,12 @@ if [[ -z "${ETH_DST_ADDR:-}" ]]; then echo "ERROR: ETH_DST_ADDR must be set (source scripts/discover_rtx_pro_topology.sh)" >&2 exit 1 fi +if [[ -z "$TX_BDF" || -z "$RX_BDF" ]]; then + echo "ERROR: RTX_TX_BDF and RTX_RX_BDF must be set (discovery or env)" >&2 + exit 1 +fi + +CPU_CORES_CSV="$(IFS=,; echo "${CPU_CORES[*]}")" CELLS=( "1t1r 1 1" @@ -88,9 +104,6 @@ CELLS=( "2t2r 2 2" ) -# master, TX q0 poll/work, TX q1 poll/work, RX q0 poll/work, RX q1 poll/work -CPU_CORES=(3 11 10 12 13 9 8 14 15) - FAILURES=0 sum_field() { @@ -154,7 +167,7 @@ cpu_busy_pct() { gpu_sample() { local out="$1" nvidia-smi --query-gpu=utilization.gpu,utilization.memory \ - --format=csv,noheader,nounits -i "$TX_GPU,$RX_GPU" 2>/dev/null \ + --format=csv,noheader,nounits -i "$TX_GPU,$TX_GPU2,$RX_GPU,$RX_GPU2" 2>/dev/null \ | awk -F', ' '{ sm += $1; mem += $2; n++ } END { @@ -166,8 +179,8 @@ gpu_sample() { run_cell() { local cell="$1" tx_count="$2" rx_count="$3" payload="$4" rep="${5:-1}" local tx_cores rx_cores - [[ "$tx_count" == 2 ]] && tx_cores="11|12" || tx_cores="11" - [[ "$rx_count" == 2 ]] && rx_cores="9|14" || rx_cores="9" + [[ "$tx_count" == 2 ]] && tx_cores="${CPU_CORES[1]}|${CPU_CORES[3]}" || tx_cores="${CPU_CORES[1]}" + [[ "$rx_count" == 2 ]] && rx_cores="${CPU_CORES[5]}|${CPU_CORES[7]}" || rx_cores="${CPU_CORES[5]}" local run_dir="$OUT_DIR/$cell/p$payload/r$rep" mkdir -p "$run_dir" @@ -176,6 +189,8 @@ run_cell() { --payload "$payload" --eth-dst "$ETH_DST_ADDR" \ --tx-bdf "$TX_BDF" --rx-bdf "$RX_BDF" \ --tx-gpu "$TX_GPU" --rx-gpu "$RX_GPU" \ + --tx-gpu2 "$TX_GPU2" --rx-gpu2 "$RX_GPU2" \ + --cpu-cores "$CPU_CORES_CSV" \ > "$tmp_cfg" 2> "$run_dir/gen.err"; then echo "ERROR: $cell p$payload config generation failed" >&2 cat "$run_dir/gen.err" >&2 @@ -236,6 +251,7 @@ run_cell() { echo "RTX PRO 6000 multi-queue sweep -- ${RUN_SECONDS}s per (cell, payload)" echo "TX BDF=$TX_BDF RX BDF=$RX_BDF ETH_DST=$ETH_DST_ADDR" +echo "TX GPU=$TX_GPU,$TX_GPU2 RX GPU=$RX_GPU,$RX_GPU2 cores=${CPU_CORES[*]}" echo "Payloads: $PAYLOADS" echo "Output: $OUT_DIR" echo diff --git a/examples/run_rtx_pro_suite.sh b/examples/run_rtx_pro_suite.sh new file mode 100755 index 0000000..2fb56de --- /dev/null +++ b/examples/run_rtx_pro_suite.sh @@ -0,0 +1,50 @@ +#!/usr/bin/env bash +# +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# +# Run the standard RTX PRO 6000 benchmark matrix and emit plots where available. +# +# Usage (root, privileged container): +# ./examples/run_rtx_pro_suite.sh +# +# Environment: +# RUN_SECONDS per MQ cell (default 20; nic/sw modes use 30) +# SKIP_MQ=1 skip multi-queue sweep +# SKIP_WIRE=1 skip wire-backed dpdk modes (nic-smoke, sweep) +# SKIP_SW=1 skip software loopback smoke + +set -euo pipefail + +SCRIPT_DIR="$(cd "$(dirname "$0")" && pwd)" +REPO_ROOT="$(cd "$SCRIPT_DIR/.." && pwd)" +RUN_SECONDS="${RUN_SECONDS:-20}" +WIRE_SECONDS="${WIRE_SECONDS:-30}" + +run_mq() { + echo "========== RTX PRO multi-queue sweep ==========" + RUN_SECONDS="$RUN_SECONDS" "$SCRIPT_DIR/run_rtx_pro_mq_bench.sh" +} + +run_wire_smoke() { + echo "========== RTX PRO dpdk nic-smoke ==========" + "$SCRIPT_DIR/run_rtx_pro_bench.sh" dpdk nic-smoke --seconds "$WIRE_SECONDS" +} + +run_sw_smoke() { + echo "========== RTX PRO dpdk sw-smoke ==========" + "$SCRIPT_DIR/run_rtx_pro_bench.sh" dpdk sw-smoke --seconds "$WIRE_SECONDS" +} + +run_sweep() { + echo "========== RTX PRO dpdk sweep ==========" + "$SCRIPT_DIR/run_rtx_pro_bench.sh" dpdk sweep --seconds "$WIRE_SECONDS" +} + +[[ "${SKIP_MQ:-0}" == "1" ]] || run_mq +[[ "${SKIP_WIRE:-0}" == "1" ]] || run_wire_smoke +[[ "${SKIP_SW:-0}" == "1" ]] || run_sw_smoke +[[ "${SKIP_SWEEP:-1}" == "1" ]] || run_sweep + +echo +echo "Suite complete. Results under: ${DAQIRI_BENCH_RESULTS_DIR:-$REPO_ROOT/bench-results}/" diff --git a/scripts/discover_rtx_pro_topology.sh b/scripts/discover_rtx_pro_topology.sh index b464121..bef4238 100755 --- a/scripts/discover_rtx_pro_topology.sh +++ b/scripts/discover_rtx_pro_topology.sh @@ -15,6 +15,8 @@ set -u +REPO_ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" + read_mac() { local iface="$1" local path="/sys/class/net/${iface}/address" @@ -147,3 +149,20 @@ fi if [[ -z "${TX_BDF:-}" || -z "${RX_BDF:-}" ]]; then echo "WARNING: could not resolve PCIe BDFs; pass --tx-bdf/--rx-bdf to run_rtx_pro_bench.sh" >&2 fi + +# GPU ordinals (PIX to each NIC) and poll-core layout for bench YAML generation. +# Requires CUDA (libcuda) and nvidia-smi; override with TX_GPU/RX_GPU/RTX_CPU_CORES. +export CUDA_DEVICE_ORDER="${CUDA_DEVICE_ORDER:-PCI_BUS_ID}" +DISCOVER_PY="$REPO_ROOT/scripts/rtx_pro_discover.py" +if [[ -f "$DISCOVER_PY" && -n "${TX_BDF:-}" && -n "${RX_BDF:-}" ]]; then + # shellcheck disable=SC1090 + eval "$(RTX_TX_BDF="$TX_BDF" RTX_RX_BDF="$RX_BDF" python3 "$DISCOVER_PY" 2>/dev/null)" || true +fi + +if [[ -n "${RTX_TX_GPU:-}" ]]; then + echo " TX GPU: ordinal ${RTX_TX_GPU} (+ queue2 ${RTX_TX_GPU2:-$RTX_TX_GPU})" + echo " RX GPU: ordinal ${RTX_RX_GPU} (+ queue2 ${RTX_RX_GPU2:-$RTX_RX_GPU})" +fi +if [[ -n "${RTX_CPU_CORES:-}" ]]; then + echo " poll cores: ${RTX_CPU_CORES}" +fi diff --git a/scripts/gen_rtx_pro_mq_config.py b/scripts/gen_rtx_pro_mq_config.py index c402290..28da52f 100755 --- a/scripts/gen_rtx_pro_mq_config.py +++ b/scripts/gen_rtx_pro_mq_config.py @@ -17,12 +17,33 @@ TX_PORT_BASE = 4096 +def apply_cpu_layout(cfg: dict, cores: list[int] | None) -> None: + if not cores or len(cores) < 9: + return + cfg["master_core"] = cores[0] + tx_poll0, tx_work0, tx_poll1, tx_work1 = cores[1], cores[2], cores[3], cores[4] + rx_poll0, rx_work0, rx_poll1, rx_work1 = cores[5], cores[6], cores[7], cores[8] + for iface in cfg["interfaces"]: + if "tx" in iface: + for q in iface["tx"]["queues"]: + q["cpu_core"] = tx_poll0 if q["id"] == 0 else tx_poll1 + if "rx" in iface: + for q in iface["rx"]["queues"]: + q["cpu_core"] = rx_poll0 if q["id"] == 0 else rx_poll1 + + def prune(base: dict, tx: int, rx: int, payload: int, batch: int | None, eth_dst: str | None, tx_bdf: str | None, rx_bdf: str | None, - tx_gpu: int, rx_gpu: int) -> dict: + tx_gpu: int, rx_gpu: int, + tx_gpu2: int | None = None, rx_gpu2: int | None = None, + cpu_cores: list[int] | None = None) -> dict: cfg = base["daqiri"]["cfg"] n_ports = max(tx, rx) + # Second-queue GPU defaults to the first-queue GPU (single-GPU behaviour). + tx_gpu2 = tx_gpu if tx_gpu2 is None else tx_gpu2 + rx_gpu2 = rx_gpu if rx_gpu2 is None else rx_gpu2 + def keep_region(name: str) -> bool: if name.endswith("_1"): return tx == 2 if "_TX_" in name else rx == 2 @@ -32,10 +53,13 @@ def keep_region(name: str) -> bool: m for m in cfg["memory_regions"] if keep_region(m["name"]) ] for m in cfg["memory_regions"]: + # Queue 1's memory region (name suffix "_1") can live on a different GPU + # than queue 0's so the two queues drive independent PCIe paths. + second = m["name"].endswith("_1") if "_TX_" in m["name"]: - m["affinity"] = tx_gpu + m["affinity"] = tx_gpu2 if second else tx_gpu elif "_RX_" in m["name"]: - m["affinity"] = rx_gpu + m["affinity"] = rx_gpu2 if second else rx_gpu for iface in cfg["interfaces"]: if iface["name"] == "tx_port" and tx_bdf: @@ -77,6 +101,15 @@ def keep_region(name: str) -> bool: b["udp_dst_port"] = f"{TX_PORT_BASE}-{TX_PORT_BASE + 1}" base["bench_tx"] = bench_tx + if cpu_cores and len(cpu_cores) >= 9: + apply_cpu_layout(cfg, cpu_cores) + tx_work0, tx_work1 = cpu_cores[2], cpu_cores[4] + rx_work0, rx_work1 = cpu_cores[6], cpu_cores[8] + for b in base["bench_rx"]: + b["cpu_core"] = rx_work0 if b["queue_id"] == 0 else rx_work1 + for b in base["bench_tx"]: + b["cpu_core"] = tx_work0 if b["queue_id"] == 0 else tx_work1 + return base @@ -92,8 +125,19 @@ def main() -> int: ap.add_argument("--rx-bdf", default=None) ap.add_argument("--tx-gpu", type=int, default=0) ap.add_argument("--rx-gpu", type=int, default=1) + ap.add_argument("--tx-gpu2", type=int, default=None, + help="GPU ordinal for the 2nd TX queue (default: --tx-gpu)") + ap.add_argument("--rx-gpu2", type=int, default=None, + help="GPU ordinal for the 2nd RX queue (default: --rx-gpu)") + ap.add_argument("--cpu-cores", default=None, + help="comma-separated poll layout: master,TX q0 poll/work," + " TX q1 poll/work, RX q0 poll/work, RX q1 poll/work") args = ap.parse_args() + cpu_cores = None + if args.cpu_cores: + cpu_cores = [int(x.strip()) for x in args.cpu_cores.split(",") if x.strip()] + with open(args.base, encoding="utf-8") as fh: base = yaml.safe_load(fh) @@ -108,6 +152,9 @@ def main() -> int: args.rx_bdf, args.tx_gpu, args.rx_gpu, + args.tx_gpu2, + args.rx_gpu2, + cpu_cores, ) yaml.safe_dump(out, sys.stdout, sort_keys=False, default_flow_style=False) return 0 diff --git a/scripts/rtx_pro_discover.py b/scripts/rtx_pro_discover.py new file mode 100755 index 0000000..279f4fe --- /dev/null +++ b/scripts/rtx_pro_discover.py @@ -0,0 +1,302 @@ +#!/usr/bin/env python3 +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +"""Discover RTX PRO 6000 bench affinity: CUDA ordinals, PIX GPUs per NIC, poll cores. + +Reads NIC PCIe BDFs (from env or CLI), matches CUDA-visible GPUs by PCI bus id, +uses ``nvidia-smi topo -m`` PIX links for GPUDirect placement, and picks poll +cores from ``isolcpus`` on the NIC's NUMA node (fallback: low cores on that node). + +Emits shell ``export`` lines for ``source`` from discover_rtx_pro_topology.sh. +""" + +from __future__ import annotations + +import argparse +import ctypes +import os +import re +import subprocess +import sys +from pathlib import Path + + +def normalize_bdf(bdf: str) -> str: + bdf = bdf.strip().lower() + if not bdf: + return bdf + # nvidia-smi uses 00000000:bb:dd.f; CUDA uses 0000:bb:dd.f + m = re.match(r"^(?:0000:)?(?:00000000:)?([0-9a-f]{2}:[0-9a-f]{2}\.[0-9a-f])$", bdf) + if m: + return f"0000:{m.group(1)}" + if not bdf.startswith("0000:"): + bdf = f"0000:{bdf}" + return bdf + + +def pci_numa_node(bdf: str) -> int: + path = Path(f"/sys/bus/pci/devices/{normalize_bdf(bdf)}/numa_node") + if path.is_file(): + return int(path.read_text().strip()) + return 0 + + +def parse_isolcpu_sets() -> list[set[int]]: + cmdline = Path("/proc/cmdline").read_text() + m = re.search(r"isolcpus=[^\s]+", cmdline) + if not m: + return [] + spec = m.group(0).split("=", 1)[1] + cores: set[int] = set() + for part in spec.split(","): + part = part.strip() + if not part or part in ("domain", "managed_irq", "nohz", "nohz-daemon"): + continue + if "-" in part: + lo, hi = part.split("-", 1) + cores.update(range(int(lo), int(hi) + 1)) + elif part.isdigit(): + cores.add(int(part)) + return [cores] if cores else [] + + +def numa_cpu_list(node: int) -> list[int]: + path = Path(f"/sys/devices/system/node/node{node}/cpulist") + if not path.is_file(): + return list(range(128)) + out: list[int] = [] + for part in path.read_text().strip().split(","): + if "-" in part: + lo, hi = part.split("-", 1) + out.extend(range(int(lo), int(hi) + 1)) + elif part: + out.append(int(part)) + return out + + +def pick_poll_cores(numa_node: int, count: int = 9) -> list[int]: + isol = parse_isolcpu_sets() + candidates = numa_cpu_list(numa_node) + if isol: + pool = sorted(c for c in candidates if any(c in s for s in isol)) + if len(pool) >= count: + return pool[:count] + # No isolcpus (or too few on this node): use the upper half of the node's + # core list so we stay away from core 0 / typical IRQ housekeeping. + if len(candidates) >= count + 4: + return candidates[4 : 4 + count] + return candidates[:count] + + +def load_cuda() -> ctypes.CDLL | None: + for name in ("libcuda.so.1", "libcuda.so"): + try: + return ctypes.CDLL(name) + except OSError: + continue + return None + + +def cuda_gpus_pci_order() -> list[tuple[int, str]]: + cuda = load_cuda() + if cuda is None: + return [] + cuda.cuInit(0) + n = ctypes.c_int() + if cuda.cuDeviceGetCount(ctypes.byref(n)) != 0: + return [] + buf = ctypes.create_string_buffer(64) + gpus: list[tuple[int, str]] = [] + for ordinal in range(n.value): + dev = ctypes.c_int() + if cuda.cuDeviceGet(ctypes.byref(dev), ordinal) != 0: + continue + if cuda.cuDeviceGetPCIBusId(buf, 64, dev) != 0: + continue + bdf = buf.value.decode().strip("\x00").lower() + gpus.append((ordinal, normalize_bdf(bdf))) + return gpus + + +def nvidia_smi_gpus() -> list[tuple[int, str]]: + try: + out = subprocess.check_output( + [ + "nvidia-smi", + "--query-gpu=index,pci.bus_id", + "--format=csv,noheader,nounits", + ], + text=True, + ) + except (subprocess.CalledProcessError, FileNotFoundError): + return [] + gpus: list[tuple[int, str]] = [] + for line in out.strip().splitlines(): + idx_s, bdf = [p.strip() for p in line.split(",", 1)] + gpus.append((int(idx_s), normalize_bdf(bdf))) + return gpus + + +def mlx5_nic_labels() -> dict[str, str]: + """Map normalized BDF -> topo label (NIC0, NIC1, ...).""" + labels: dict[str, str] = {} + try: + out = subprocess.check_output(["nvidia-smi", "topo", "-m"], text=True, stderr=subprocess.DEVNULL) + except (subprocess.CalledProcessError, FileNotFoundError): + return labels + out = re.sub(r"\x1b\[[0-9;]*m", "", out) + in_legend = False + for line in out.splitlines(): + if "NIC Legend" in line: + in_legend = True + continue + if not in_legend: + continue + m = re.match(r"\s*(NIC\d+):\s*(\S+)", line) + if not m: + continue + nic_label, mlx = m.group(1), m.group(2) + dev = Path(f"/sys/class/infiniband/{mlx}/device") + if dev.is_symlink(): + bdf = normalize_bdf(dev.resolve().name) + labels[bdf] = nic_label + return labels + + +def topo_pix_gpus(nic_label: str) -> list[str]: + """Return nvidia-smi GPU labels (GPU1, GPU2, ...) with PIX to nic_label.""" + try: + out = subprocess.check_output(["nvidia-smi", "topo", "-m"], text=True, stderr=subprocess.DEVNULL) + except (subprocess.CalledProcessError, FileNotFoundError): + return [] + out = re.sub(r"\x1b\[[0-9;]*m", "", out) + lines = [ln for ln in out.splitlines() if ln.strip()] + if not lines: + return [] + header = re.split(r"\t+", lines[0].strip()) + try: + nic_col = header.index(nic_label) + 1 # +1: row label occupies column 0 + except ValueError: + return [] + pix: list[str] = [] + for line in lines[1:]: + cols = re.split(r"\t+", line.strip()) + if not cols or not cols[0].startswith("GPU"): + continue + if len(cols) <= nic_col: + continue + if cols[nic_col].strip() == "PIX": + pix.append(cols[0].strip()) + return pix + + +def gpu_label_to_bdf(gpu_label: str) -> str | None: + # Topo labels GPU1..GPUN match nvidia-smi index 1..N (not 0-based). + m = re.match(r"GPU(\d+)", gpu_label) + if not m: + return None + smi_idx = int(m.group(1)) + for idx, bdf in nvidia_smi_gpus(): + if idx == smi_idx: + return bdf + return None + + +def ordinal_for_bdf(bdf: str, cuda_gpus: list[tuple[int, str]]) -> int | None: + bdf = normalize_bdf(bdf) + for ordinal, bus in cuda_gpus: + if bus == bdf: + return ordinal + return None + + +def pix_ordinals_for_nic(nic_bdf: str, cuda_gpus: list[tuple[int, str]]) -> list[int]: + labels = mlx5_nic_labels() + nic_label = labels.get(normalize_bdf(nic_bdf)) + if not nic_label: + return [] + ordinals: list[int] = [] + for gpu_label in topo_pix_gpus(nic_label): + bus = gpu_label_to_bdf(gpu_label) + if bus is None: + continue + ord_ = ordinal_for_bdf(bus, cuda_gpus) + if ord_ is not None: + ordinals.append(ord_) + return ordinals + + +def pick_gpu_pair(pix: list[int], fallback: int) -> tuple[int, int]: + if not pix: + return fallback, fallback + if len(pix) == 1: + return pix[0], pix[0] + return pix[0], pix[1] + + +def emit_exports(tx_bdf: str, rx_bdf: str) -> int: + cuda_gpus = cuda_gpus_pci_order() + if not cuda_gpus: + print("# WARNING: CUDA not available; GPU ordinals not discovered", file=sys.stderr) + return 1 + + tx_pix = pix_ordinals_for_nic(tx_bdf, cuda_gpus) + rx_pix = pix_ordinals_for_nic(rx_bdf, cuda_gpus) + + tx_gpu, tx_gpu2 = pick_gpu_pair(tx_pix, cuda_gpus[0][0]) + rx_gpu, rx_gpu2 = pick_gpu_pair(rx_pix, cuda_gpus[min(1, len(cuda_gpus) - 1)][0]) + + # Poll on the NUMA node of the TX port (both ports are usually co-located). + numa = pci_numa_node(tx_bdf) + cores = pick_poll_cores(numa, 9) + while len(cores) < 9: + cores.append(cores[-1] if cores else 0) + + # master, TX q0 poll/work, TX q1 poll/work, RX q0 poll/work, RX q1 poll/work + layout = { + "RTX_MASTER_CORE": cores[0], + "RTX_TX_Q0_POLL": cores[1], + "RTX_TX_Q0_WORK": cores[2], + "RTX_TX_Q1_POLL": cores[3], + "RTX_TX_Q1_WORK": cores[4], + "RTX_RX_Q0_POLL": cores[5], + "RTX_RX_Q0_WORK": cores[6], + "RTX_RX_Q1_POLL": cores[7], + "RTX_RX_Q1_WORK": cores[8], + } + + print(f"export RTX_TX_GPU={tx_gpu}") + print(f"export RTX_RX_GPU={rx_gpu}") + print(f"export RTX_TX_GPU2={tx_gpu2}") + print(f"export RTX_RX_GPU2={rx_gpu2}") + print(f'export RTX_CPU_CORES="{" ".join(str(c) for c in cores)}"') + for key, val in layout.items(): + print(f"export {key}={val}") + + bus_by_ord = {o: b for o, b in cuda_gpus} + print(f"# CUDA ordinals (PCI_BUS_ID): TX={tx_gpu}({bus_by_ord.get(tx_gpu,'?')}) " + f"TX2={tx_gpu2}({bus_by_ord.get(tx_gpu2,'?')}) " + f"RX={rx_gpu}({bus_by_ord.get(rx_gpu,'?')}) " + f"RX2={rx_gpu2}({bus_by_ord.get(rx_gpu2,'?')})", file=sys.stderr) + print(f"# poll cores NUMA {numa}: {' '.join(str(c) for c in cores)}", file=sys.stderr) + if not tx_pix: + print(f"# WARNING: no PIX GPU found for TX BDF {tx_bdf}; using fallback ordinal {tx_gpu}", file=sys.stderr) + if not rx_pix: + print(f"# WARNING: no PIX GPU found for RX BDF {rx_bdf}; using fallback ordinal {rx_gpu}", file=sys.stderr) + return 0 + + +def main() -> int: + ap = argparse.ArgumentParser(description=__doc__) + ap.add_argument("--tx-bdf", default=os.environ.get("RTX_TX_BDF", "")) + ap.add_argument("--rx-bdf", default=os.environ.get("RTX_RX_BDF", "")) + ap.add_argument("--emit-shell", action="store_true", default=True) + args = ap.parse_args() + if not args.tx_bdf or not args.rx_bdf: + print("ERROR: --tx-bdf and --rx-bdf required", file=sys.stderr) + return 1 + return emit_exports(args.tx_bdf, args.rx_bdf) + + +if __name__ == "__main__": + sys.exit(main()) From f18c875e47ff3e352a5d63a79f2c88748574e3c1 Mon Sep 17 00:00:00 2001 From: Chloe Crozier Date: Tue, 14 Jul 2026 16:55:02 +0000 Subject: [PATCH 7/7] #17 - Fix ibverbs catch-all flow steering on mlx5 Use priority 0 for the mlx5dv_dr catch-all matcher so the uint16_t priority does not truncate and rule creation no longer fails with ENOMEM during daqiri_init. Also pin the RTX PRO ibverbs RX hugepage region to NUMA node 0. Signed-off-by: ccrozier --- ...iri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml | 4 ++- src/engines/ibverbs/daqiri_ibverbs_engine.cpp | 27 +++++++++++-------- 2 files changed, 19 insertions(+), 12 deletions(-) diff --git a/examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml b/examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml index 103ddd1..da4a627 100644 --- a/examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml +++ b/examples/daqiri_bench_raw_rx_ibverbs_rtx_pro_6000.yaml @@ -27,10 +27,12 @@ daqiri: log_level: "info" loopback: "" + # affinity is the NUMA node for the huge-page RX region. Keep it on the RX + # port's node (discover with: cat /sys/class/net//device/numa_node). memory_regions: - name: "Data_RX_HOST" kind: "huge" - affinity: 1 + affinity: 0 num_bufs: 4096 buf_size: 8064 diff --git a/src/engines/ibverbs/daqiri_ibverbs_engine.cpp b/src/engines/ibverbs/daqiri_ibverbs_engine.cpp index 05ce401..630dd0c 100644 --- a/src/engines/ibverbs/daqiri_ibverbs_engine.cpp +++ b/src/engines/ibverbs/daqiri_ibverbs_engine.cpp @@ -51,7 +51,13 @@ namespace { // libdpdk: "cycles" are nanoseconds, so the timer frequency is 1e9. constexpr uint64_t ibv_timer_hz = 1'000'000'000ULL; constexpr FlowId kMaxIbverbsFlowTag = 0x00ffffffU; -constexpr int kIbverbsCatchAllPriority = 1'000'000; +// mlx5dv_dr_matcher_create() takes the priority as a uint16_t. A value that +// does not fit silently truncates, and on ConnectX-class NICs software steering +// then rejects the rule (mlx5dv_dr_rule_create fails with ENOMEM). The catch-all +// is only installed when it is the sole rule on a port (static flows and +// flow_isolation take other branches), so precedence is irrelevant and it uses +// priority 0. +constexpr int kIbverbsCatchAllPriority = 0; inline uint64_t ibv_now_ns() { return static_cast(std::chrono::duration_cast( std::chrono::steady_clock::now().time_since_epoch()) @@ -952,15 +958,10 @@ Status IbverbsEngine::devx_create_tir(IbvRxQueue& q) { return Status::SUCCESS; } -// mlx5dv_dr match buffer: 64 bytes interpreted as fte_match_set_lyr_2_4 when -// criteria_enable selects outer_headers. -namespace { -struct DrMatchBuf { - size_t match_sz; - uint64_t buf[8]; // == sizeof(fte_match_set_lyr_2_4) -}; // Full match parameter (512 B) for matches that reach misc_parameters_4 (the -// flex-parser sample registers), which sits past the first 64-byte section. +// flex-parser sample registers). The first 64 bytes alias fte_match_set_lyr_2_4 +// (outer_headers), so this buffer also serves L2/L3/L4 and catch-all rules. +namespace { struct DrMatchParam { size_t match_sz; uint64_t buf[64]; // == sizeof(fte_match_param) @@ -1930,8 +1931,12 @@ Status IbverbsEngine::install_port_flows() { "after initialization", port); } else if (installed == 0) { - // No per-flow rules: catch-all (criteria_enable = 0) -> first queue. - DrMatchBuf mask{}, val{}; + // No per-flow rules: wildcard catch-all -> first queue. This rule is only + // installed when it is the sole rule on the port (static flows take the + // branch above; flow_isolation takes the branch below), so it sits at + // priority 0 with an empty match (match_criteria_enable = 0) and matches + // every packet. + DrMatchParam mask{}, val{}; mask.match_sz = sizeof(mask.buf); val.match_sz = sizeof(val.buf); IbvRxQueue* dest = by_id.begin()->second;