HPC-X, MPI, PMIx & NCCL: A Full Dissection of the GPU Cluster Communication Stack

In the context of distributed GPU training, terms like HPC-X, OpenMPI, UCX, PMIx, and NCCL appear frequently, but their lines of responsibility are often blurred. The most common misconception is that HPC-X or UCX is involved in the data transfer of an all-reduce operation. In reality, the division of labor in this stack is crystal clear.

The Big Picture: Division of Labor

The entire GPU cluster communication stack can be broken down into five distinct layers:

Layer	Component	Responsibility
User Layer	User Code / nccl-tests	Invokes MPI_* and ncclAllReduce
MPI Layer	OpenMPI (from HPC-X)	Process management, PMIx rendezvous, control messages
Transport Layer	UCX → RC/DC QP	MPI point-to-point transport backend (does not touch tensor data)
Data Plane	NCCL → GPU kernel → IB QP	Tensor data plane, GPUDirect RDMA, the hot path
Hardware Layer	InfiniBand / NVLink / PCIe	Physical transport medium

All of HPC-X's work is finished and off the critical path once ncclCommInitRank completes.

§ 01 — HPC-X: A Software Distribution, Not a Component

HPC-X is a high-performance MPI software distribution packaged by NVIDIA. It bundles several independent components, deeply optimized for InfiniBand/RoCE networks.

Component	Function	Role in nccl-tests
OpenMPI	MPI runtime and API	Process management, control message routing
UCX	Point-to-point transport engine	Transfers the 128B ncclUniqueId, handles Barriers
HCOLL	MPI collective communication accelerator	Accelerates MPI_Allreduce (NCCL does not use this)
SHARP	In-network computing	Only active if NCCL CollNet is used
XPMEM	Cross-process shared memory	Intra-node MPI communication

[⚠ Common Misconception] HCOLL accelerates MPI_Allreduce, but the hot path in nccl-tests is ncclAllReduce. The two are completely independent. HCOLL has zero impact on NCCL performance numbers.

When is HPC-X Actually Called?

MPI_Init: UCX creates endpoints, an RC/DC QP is established for each pair of ranks, and the PMIx fence completes process discovery.

The Sole Intersection: MPI_Bcast(ncclUniqueId, 128B) — a one-time UCX eager send. After this, the two paths never cross again.

MPI_Barrier (before/after timing): A small UCX control message, not within the timed window.

ncclAllReduce Hot Path: UCX is not involved at all. NCCL exclusively uses its own IB QPs, and the GPU kernel drives RDMA directly.

MPI_Finalize: UCX endpoints are destroyed, and QPs are closed.

§ 02 — The Relationship Between UCX and OpenMPI

UCX (Unified Communication X) is the default point-to-point transport backend for OpenMPI within HPC-X. Their relationship is like a kernel and a distro: OpenMPI provides the MPI API, and UCX is the engine that actually moves the data via the PML (Point-to-point Management Layer) plugin.

The call chain for UCX within MPI_Init:

ucp_init()                       ← Process-level, once
  └── ucp_worker_create()         ← Thread-level worker, manages progress
          └── ucp_ep_create() × (N-1)    ← Create endpoint for each peer
                  └── ibv_create_qp()    ← Create RC or DC QP
                      ibv_modify_qp(INIT → RTR → RTS)

[Key Fact] The IB QPs used by UCX in nccl-tests and those used by NCCL are two completely separate sets of QPs. On large-scale clusters, the total number of QPs from both can approach the IB device limit (typically ~131,072).

Key Environment Variable Boundaries

UCX / MPI Control Plane: UCX_TLS=rc,shm,self, UCX_NET_DEVICES=mlx5_0:1, OMPI_MCA_pml=ucx.

NCCL Data Plane: NCCL_IB_HCA=mlx5_0,mlx5_1, NCCL_NET_GDR_LEVEL=SYS, NCCL_ALGO=Ring/Tree.

§ 03 — The Only Intersection of HPC-X and NCCL

NCCL has no built-in mechanism for process discovery. It relies on an external entity to broadcast a 128-byte rendezvous token (ncclUniqueId). MPI's MPI_Bcast is used for exactly this purpose—it is their only true data intersection.

The ncclUniqueId contains the address of a TCP listening socket opened by rank 0:

struct ncclUniqueId {
  char internal[128];
  // [0..5]   = Rank 0's IP:port (bootstrap TCP listen socket)
  // [6..127] = Random nonce (to prevent misconnections)
};

The code structure of nccl-tests clearly shows this division:

// ── Part handled by MPI ──────────────────────────
MPI_Init(&argc, &argv);
MPI_Comm_rank(MPI_COMM_WORLD, &rank);

// ── The only intersection point ──────────────────
ncclUniqueId id;
if (rank == 0) ncclGetUniqueId(&id);       // Rank 0 opens a TCP listen socket
MPI_Bcast(&id, sizeof(id), MPI_BYTE, 0,    // UCX eagerly sends the 128B
          MPI_COMM_WORLD);

// ── From here on, it's all NCCL. MPI is off the critical path ──
ncclCommInitRank(&comm, nranks, id, rank); // NCCL builds its own QPs
// ...
ncclAllReduce(sendbuf, recvbuf, count,     // NCCL data plane
              ncclFloat, ncclSum, comm, stream);

What NCCL Does After MPI_Bcast

Once all ranks receive the same id, ncclCommInitRank executes a sequence of steps to establish its own communication fabric, completely independent of UCX.

[Design Insight] The NCCL bootstrap TCP socket is not closed after ncclCommInitRank returns, but it is never used for data transfer. The hot path exclusively uses the IB QPs. The bootstrap TCP socket remains idle for the entire duration of the training job.

§ 04 — The Full Call Chain of ncclAllReduce

Each time ncclAllReduce is called, the CPU performs three lightweight tasks, after which the GPU operates autonomously.

ncclAllReduce(...)
  ├─ ncclEnqueueCheck()        ← Selects algo/proto based on message size
  ├─ ncclSaveKernels()         ← Packs parameters into a work queue
  └─ cudaLaunchKernel()        ← CPU returns, GPU takes over
       └─ GPU kernel executes (CPU not involved):
          ├─ reduce-scatter: N-1 steps, RDMA Write to next rank
          └─ all-gather:     N-1 steps, RDMA Write to broadcast result

Why Small Messages Are Slow (in multi-node scenarios)

This has nothing to do with HPC-X or UCX. The bottleneck for small messages is the accumulation of (Number of Ring Steps × IB RTT).

Node Count	Ring Steps	Accumulated IB RTT (≈2μs/hop)
4 Nodes	6 steps	~12 μs
16 Nodes	30 steps	~60 μs
64 Nodes	126 steps	~252 μs

§ 05 — What Happens Inside MPI_Init

MPI_Init is the heaviest operation in the entire HPC-X stack, taking anywhere from hundreds of milliseconds to several seconds. It completes five main phases in sequence, with the heaviest being UCX initialization, which can take seconds on large clusters due to the O(N²) QP handshakes required for the RC transport.

[Important Tip] Seeing MPI_Init hang for a few seconds on a large-scale training job is normal—it's not a hang, but the inherent cost of UCX QP handshakes. Switching to the DC transport (UCX_TLS=dc,shm,self) can significantly reduce this startup time.

§ 06 — PMIx: The Nervous System Between Slurm and MPI

PMIx (Process Management Interface for Exascale) is the standard interface between a process manager and a runtime library. It defines how Slurm and MPI exchange information. Slurm acts as the PMIx server, while OpenMPI is the client.

§ 07 — Slurm vs. MPI vs. NCCL: A Full Comparison

Dimension	Slurm	MPI / HPC-X	NCCL
Essence	Resource Manager	Communication Library + Runtime	GPU Collective Comm. Library
Data Type	Metadata (node allocation)	Metadata (UCX addresses)	Tensors (float/bf16)
Transport	Management Plane	TCP / IB RC QP	NVLink / IB RDMA
Initiator	slurmctld	CPU	GPU kernel
GPU Aware	Allocates but unaware of topology	CUDA-aware optional	Deeply aware of NVLink/PCIe

§ 08 — Summary: One-Line Descriptions

Slurm / slurmd   →  "On which machines to run, and for how long" (Resource Manager)
PMIx server      →  "Where are these processes, how to contact them" (Process Discovery Bus)
OpenMPI          →  Implements the MPI API, routes to UCX/HCOLL
UCX              →  Transports MPI control messages, does not touch tensor data
NCCL bootstrap   →  Uses a TCP ring to exchange topology info, builds IB QPs
NCCL hot path    →  GPU kernel drives RDMA, approaching hardware line speed
─────────────────────────────────────────────────────────
The intersection of HPC-X and NCCL:
  MPI_Bcast(ncclUniqueId, 128B)
  Once these 128 bytes are transferred, the two paths never cross again.

[Core Conclusion] If you encounter performance issues in nccl-tests or any NCCL-based training job: High latency for small messages → Look at NCCL algorithm selection and RTT accumulation from the node count, unrelated to UCX. Low bandwidth for large messages → Look at the NCCL IB QP configuration (NCCL_IB_HCA, GDR, nvidia_peermem), unrelated to UCX. Slow startup → Look at UCX QP establishment (consider switching to DC transport) or PMIx Fence latency.