Generative AI workloads are quickly outgrowing the reminiscence and compute price range of single GPUs. For inference builders constructing media era pipelines, the problem is scaling throughout a number of units with out sacrificing the crucial optimizations—like kernel fusions, reminiscence planning, and quantization—that NVIDIA TensorRT delivers for manufacturing deployments.
Multi-device inference help, a brand new function launched in TensorRT 11.0, brings native high-performance multi-GPU inference to the TensorRT runtime, enabling multi-device manufacturing deployments focusing on edge units.
Combining the multi-device inference help in TensorRT with Torch-TensorRT, builders can convert and deploy large PyTorch fashions out-of-framework, shattering single-device reminiscence and compute limits.
Obtain TensorRT 11.0 with multi-device inference help from NVIDIA Developer Portal to unlock native, high-performance multi-device acceleration to your fashions.
NVIDIA NCCL: The transport layer for distributed inference
The NVIDIA Collective Communications Library (NCCL) gives high-performance multi-GPU and multi-node collective operations powering large-scale mannequin coaching throughout hundreds of GPUs. NCCL mechanically selects the optimum transport for a given topology, abstracting NVIDIA NVLink, NVIDIA NVSwitch, PCIe, and InfiniBand behind a uniform interface. By integrating straight with NCCL, TensorRT inherits this transport optimization for inference workloads, when operating multi-device inference. For extra info on NCCL, see https://developer.nvidia.com/nccl.
The brand new multi-device function covers the total set of NVIDIA NCCL distributed collectives: AllReduce, Broadcast, Cut back, AllGather, ReduceScatter, AlltoAll, Collect, and Scatter.
Parallelism methods for distributed inference
Distributed inference might be expressed utilizing a number of parallelism methods, every with totally different trade-offs between reminiscence financial savings, compute scaling, and communication overhead. The most typical methods are tensor parallelism and context parallelism.
Tensor parallelism
In tensor parallelism, the weights of a single layer are partitioned throughout GPUs. Every GPU computes a shard of the layer’s matrix multiplication after which combines partial outcomes via a collective to supply the total output. This reduces per-device reminiscence weight, making it the pure (and sometimes the one) selection when a person layer’s weights exceed the reminiscence of a single GPU, unbiased of the enter sequence size or batch dimension.
In a transformer block, column-parallel projections (for instance, QKV and the MLP up-projection) are paired with row-parallel projections (the eye output and the MLP down-projection) so that every block requires solely a single AllReduce, preserving communication overhead bounded.


Context parallelism
In context parallelism, the enter sequence is partitioned throughout GPUs alongside the sequence dimension. Every GPU processes solely a slice of the sequence, whereas collective operations make the worldwide sequence obtainable the place wanted, reminiscent of throughout consideration. Context parallelism is especially efficient for long-sequence workloads, the place consideration’s quadratic scaling with sequence size makes it the dominant client of compute and reminiscence.
It is usually an particularly pure match for diffusion and DiT fashions, whose bidirectional consideration sidesteps the load-imbalance points that come up with causal masks.
Learn the Context Parallelism for Scalable Million-Token Inference article for added particulars on context parallelism.
NVIDIA TensorRT 11.0 introduces help for the `IDistCollectiveLayer` primitives required by the varied parallelization methods. The rest of this put up focuses on context parallelism, which straight addresses the dominant price in fashionable generative media pipelines: long-sequence consideration.
Context parallelism for generative media
Diffusion-based picture and video era pipelines spend a big fraction of their compute and reminiscence price range inside consideration blocks working over lengthy token sequences. A high-resolution picture latent or a multi-frame video clip can produce sequences of tens of hundreds of tokens per block, and a focus scales quadratically with sequence size.
AllGather KV
Context parallelism partitions the sequence throughout GPUs. Every rank processes a slice of the queries (Q) comparable to its sequence partition. A simple method to implement context parallelism is the AllGather KV method, the place ranks alternate their key (Ok) and worth (V) shards via an AllGather collective earlier than computing native consideration, enabling every rank to attend over the total sequence. The result’s a per-rank consideration output protecting the total sequence at the price of one further collective per consideration block, whereas the native Q × Kᵀ matrix multiplication shrinks proportionally to the variety of ranks.
For video and high-resolution picture diffusion, this trade-off compounds favorably throughout denoising steps. Communication overhead per step stays bounded by the sequence-dimension AllGather, whereas compute and reminiscence financial savings apply to each consideration layer in each step.


Ring Consideration
Context parallelism might be applied in numerous methods, every presenting distinct trade-offs.
One potential enchancment over the AllGather KV methodology is Ring Consideration, the place communication and computation are overlapped. This allows every GPU to course of its native Q concurrently because the Ok and V repeatedly stream previous in a hoop topology. Ring Consideration additionally reduces the reminiscence footprint: utilizing a web-based softmax, the full-size Ok and V tensors don’t have to be materialized on any GPU. Learn the Ring Consideration with Blockwise Transformers for Close to-Infinite Context article to be taught extra about Ring Consideration.


DeepSpeed Ulysses
For lengthy context (tens of hundreds of tokens), an alternate context parallelism implementation method is DeepSpeed Ulysses. It initially partitions particular person samples alongside the sequence dimension throughout collaborating GPUs. Earlier than the eye computation, it employs an all-to-all communication collective on the partitioned Q, Ok, and V.
This ensures that every GPU receives the total sequence size, however just for a non-overlapping subset of the eye heads, enabling them to compute consideration in parallel. Lastly, a second all-to-all communication gathers the outcomes throughout the eye heads whereas repartitioning them alongside the sequence dimension. Learn extra about context parallelism for lengthy context within the article DeepSpeed Ulysses: System Optimizations for Enabling Coaching of Excessive Lengthy Sequence Transformer Fashions.


Benchmarks: Media era with context parallelism in C++
The next benchmarks consider multi-device TensorRT inference for media era workloads meant for C++ manufacturing deployment. Two consultant generative AI pipelines are used: a video era pipeline primarily based on NVIDIA Cosmos 3 and a picture era pipeline primarily based on FLUX.1.
These pipelines had been first authored in PyTorch, then transformed out of the framework utilizing Torch-TensorRT to supply NVIDIA TensorRT engines appropriate for deployment in C++ inference functions. This workflow permits builders to retain PyTorch because the mannequin growth atmosphere whereas deploying optimized TensorRT engines in manufacturing techniques.
The benchmarks evaluate end-to-end latency throughout totally different context parallelism methods: AllGather KV, Ring Consideration, and Ulysses. All outcomes had been collected on a single node with 8 GPUs.
Video era with NVIDIA Cosmos 3
The NVIDIA Cosmos mannequin platform is a world basis mannequin platform, and the Cosmos3-Nano mannequin can generate photographs, video, audio, and different codecs primarily based on multimodal inputs, together with textual content, photographs, and video. We used the instance immediate file for our benchmarks. Based mostly on these benchmarks, Ulysses is the clear winner when a diffusion mannequin has excessively lengthy context lengths (within the order of tens of hundreds of enter tokens).






Picture era with Flux.1
The FLUX.1-dev mannequin from Black Forest Labs can generate photographs from textual content descriptions. We used the immediate: “a wonderful {photograph} of Mt. Fuji throughout cherry blossom” for our benchmarks. Based mostly on the benchmarks, the Ulysses technique is the winner within the case of picture era as properly, however it’s value noting that Ring Consideration additionally scaled properly to 4 GPUs.






Getting began utilizing TensorRT with the multi-device function
TensorRT helps multi-device inference, enabling a single community to execute throughout a number of GPUs via built-in distributed communication primitives. The core workflow is much like that of single-device TensorRT. The distinction is that the community can now embody distributed communication layers.
On this information, it’s assumed that the identical community is deployed on all GPU ranks, however this isn’t a strict requirement, and, in concept, every rank can run a distinct mannequin.
A working pattern is supplied within the TensorRT repository. The next information gives a step-by-step description of use the brand new multi-device function.
Conditions
Create a community for multi-device inference
On the community stage, multi-device inference is enabled via IDistCollectiveLayer for cross-GPU communication. Collective operations might be added on to a TensorRT community utilizing INetworkDefinition::addDistCollective:
// create empty community
auto community =
std::unique_ptr(builder->createNetworkV2(
1U << static_cast(kSTRONGLY_TYPED)));
auto* enter =
network->addInput(“enter”, DataType::kFLOAT, Dims2{3, 4});
ITensor& inputTensor = *network->getInput(0);
auto* collectiveLayer = network->addDistCollective(
inputTensor,
CollectiveOperation::kALL_REDUCE,
ReduceOperation::kSUM,
-1, // root: -1 for collectives and not using a root rank
nullptr, // teams: nullptr means all ranks take part
0 // groupSize
);
// set the world dimension aka complete variety of GPUs
collectiveLayer->setNbRanks(8);
For discount collectives reminiscent of ALL_REDUCE, REDUCE, and REDUCE_SCATTER, specify a legitimate ReduceOperation, reminiscent of kSUM. For non-reduction collectives reminiscent of ALL_GATHER, BROADCAST, ALL_TO_ALL, GATHER, and SCATTER, use ReduceOperation::kNONE. Root-based operations, together with BROADCAST, REDUCE, GATHER, and SCATTER, require a legitimate root rank.
Construct an engine
auto builderConfig = std::unique_ptr(builder->createBuilderConfig());
// construct engine
auto serializedEngine = std::unique_ptr(builder->buildSerializedNetwork
(*community, *builderConfig));
Create execution context
pattern::gLogger.getTRTLogger()));
Bind IO tensors
char const* outputName = engine->getIOTensorName(1);
std::vector const& inputChunk = (rank == 0) ? config.rank0Input : config.rank1Input;
std::vector outputChunk(config.outputElementCount, 0.0F);
size_t const inputBytes = inputChunk.dimension() * sizeof(float);
size_t const outputBytes = outputChunk.dimension() * sizeof(float);
void* dInput = nullptr;
void* dOutput = nullptr;
CHECK_CUDA(cudaMalloc(&dInput, inputBytes));
CHECK_CUDA(cudaMalloc(&dOutput, outputBytes));
// Copy enter information to GPU asynchronously
CHECK_CUDA(cudaMemcpyAsync(dInput, inputChunk.information(), inputBytes, cudaMemcpyHostToDevice, stream));
// Set enter/output tensor addresses within the execution context
context->setInputTensorAddress(inputName, dInput);
context->setTensorAddress(outputName, dOutput);
context->setInputShape(inputName, Dims2{kINPUT_ROWS, kINPUT_COLS});
Set communicator and enqueue inference
context->enqueueV3(stream);
Be aware: the NCCL communicator should additionally stay legitimate for the lifetime of the execution context that makes use of it.
Kick off inference
Run the appliance with OpenMPI on 8 GPUs. Every rank selects its native CUDA system, initializes NCCL, creates its personal TensorRT engine, creates its personal execution context, and attaches the NCCL communicator.
export TRT_WORLD_SIZE=$OMPI_COMM_WORLD_SIZE;
export TRT_NCCL_ID_FILE=/tmp/nccl_id.txt;
./sample_dist_collective –op all_reduce’
Be taught extra
If you wish to be taught extra concerning the matters launched on this article, we included some helpful hyperlinks for additional studying.
NCCL: NVIDIA Collective Communications Library (NCCL)
Parallelism:
NVIDIA TensorRT:
NVIDIA Torch-TensorRT: Torch-TensorRT Documentation

