nccl-tests/src/device_api/gin/README.md

GIN Device API P2P Performance Tests

These tests measure latency and bandwidth / message rate of NCCL GIN point-to-point device API operations (put, get, signal, and combined variants).

They require NCCL 2.30.7+, MPI, and must always run with exactly 2 ranks (one GPU each).

Build

From the root of the nccl-tests repository, build the tests with MPI enabled:

$ make -j MPI=1 MPI_HOME=/path/to/mpi NCCL_HOME=/path/to/nccl

If CUDA is not installed in /usr/local/cuda, set CUDA_HOME. Binaries are written under build/device_api/gin/.

You can also build a single binary, for example:

$ make -C src -j MPI=1 MPI_HOME=/path/to/mpi NCCL_HOME=/path/to/nccl \
    build/device_api/gin/throughput/ginPutBW_perf

With CMake, configure and build from the repository root:

$ cmake -S . -B build -DNCCL_HOME=/path/to/nccl
$ cmake --build build -j

Usage

Process count is managed by MPI and is not passed as a test argument. Use exactly two ranks:

$ mpirun -np 2 ./build/device_api/gin/<category>/<binary>_perf [OPTIONS]

Quick examples

Default put bandwidth size sweep (4B to 4M, factor 2):

$ mpirun -np 2 ./build/device_api/gin/throughput/ginPutBW_perf

put+signal bandwidth with 16 CTAs and 32 threads per CTA:

$ mpirun -np 2 ./build/device_api/gin/throughput/ginPutBW_perf \
    --gin_op put_signal -c 16 -t 32

put ping-pong latency from 4 B to 1 MiB:

$ mpirun -np 2 ./build/device_api/gin/latency/ginPutLatency_pingPong_perf \
    -b 4 -e 1M -f 2

Binaries

Category	Binary	Operations
Throughput	ginPutBW_perf	--gin_op put (default), put_signal, or put_counter
Throughput	ginGetBW_perf	get
Throughput	ginSignalBW_perf	signal
Latency	ginPutLatency_ping_perf	--gin_op put (default), put_signal, or put_counter
Latency	ginPutLatency_pingPong_perf	--gin_op put (default) or put_signal
Latency	ginGetLatency_ping_perf	get
Latency	ginSignalLatency_ping_perf	signal
Latency	ginSignalLatency_pingPong_perf	signal

put+counter is supported on ping latency and bandwidth only (not ping-pong): a counter reports local completion and cannot drive the peer's turn.

Output

Throughput tests print message size, message count, bandwidth (MiB/s), and message rate (MPPS).

Latency tests print size, iteration count, and one-way latency (us).

Rank 0 also prints the resolved configuration (RSM, CTA/thread counts for throughput, and any enabled optional flags) above the results table.

Arguments

All binaries share the same argument parser. Flags that a binary does not support are rejected rather than ignored.

• Sizes to scan
  • -b, --minbytes <size> minimum size in bytes. Default: sizeof(int). K/M/G suffixes OK.
  • -e, --maxbytes <size> maximum size / registered buffer size in bytes. Default: 4M. The size sweep runs through -e. For throughput, threads partition that buffer into maxbytes / size slots (threadIdx.x % slots); when size reaches -e, there is only one slot and threads reuse the same region.
  • -f <factor> multiplicative step between sizes. Default: 2.
  • -i <stepbytes> additive step in bytes; overrides -f when non-zero. Default: 0.
• Performance
  • -w <iters> warmup iterations per size (not timed). Default: 50; 0 disables.
  • -n, --iters <iters> measured iterations per size. Default: 500.
• Launch geometry (mainly for throughput)
  • -c, --num_ctas <n> number of CTAs. Default: 1.
  • -t, --num_threads <n> number of threads per CTA. Default: 1. Latency tests require -c 1 and -t 1 (the defaults); other values are rejected.
• GIN options
  • --gin_op <put|put_signal|put_counter> operation for put-family binaries. Default: put.
  • --gin_rsm <thread|cta|gpu> resource sharing mode. Default: gpu for all tests. Throughput rejects thread.
  • --gin_tx_depth <depth> GIN send queue depth. Default: 1024.
  • --gin_skip_credit_check skip GIN credit checks. Default: off.
  • --gin_strong_signal use Strong instead of Weak signals where supported. Default: Weak.
  • --gin_ag enable ncclGinOptFlagsAggregateRequests (throughput only). Default: off. Requires -t to be a multiple of 32.
  • --gin_bd bidirectional bandwidth: both ranks send; reported metric is the sum of each rank's average (throughput only). Default: off.

Optional optimizations

By default, optimization knobs above are off. Pass them explicitly when you want a more aggressive configuration, for example:

$ mpirun -np 2 ./build/device_api/gin/throughput/ginPutBW_perf \
    --gin_op put_signal -c 16 -t 32 --gin_rsm cta --gin_ag --gin_skip_credit_check