Chelsea John

1.  [Environment
    Setup](#environment-setup){#markdown-toc-environment-setup}
    1.  [Enabling UCC in
        OpenMPI](#enabling-ucc-in-openmpi){#markdown-toc-enabling-ucc-in-openmpi}
    2.  [Enabling NCCL in UCC (Team Layer
        Selection)](#enabling-nccl-in-ucc-team-layer-selection){#markdown-toc-enabling-nccl-in-ucc-team-layer-selection}
    3.  [All The
        Variables](#all-the-variables){#markdown-toc-all-the-variables}
2.  [Results](#results){#markdown-toc-results}
    1.  [1. Plain
        OpenMPI](#1-plain-openmpi){#markdown-toc-1-plain-openmpi}
    2.  [2. OpenMPI with
        UCC](#2-openmpi-with-ucc){#markdown-toc-2-openmpi-with-ucc}
    3.  [3. OpenMPI with
        UCC+NCCL](#3-openmpi-with-uccnccl){#markdown-toc-3-openmpi-with-uccnccl}
3.  [Scaling Plots](#scaling-plots){#markdown-toc-scaling-plots}
    1.  [Average
        Latency](#average-latency){#markdown-toc-average-latency}
    2.  [Bus Bandwidth](#bus-bandwidth){#markdown-toc-bus-bandwidth}
        1.  [Comparing MPI, UCC,
            UCC+NCCL](#comparing-mpi-ucc-uccnccl){#markdown-toc-comparing-mpi-ucc-uccnccl}
        2.  [Comparing UCC+NCCL,
            NCCL](#comparing-uccnccl-nccl){#markdown-toc-comparing-uccnccl-nccl}
4.  [Summary](#summary){#markdown-toc-summary}
5.  [Technical
    Details](#technical-details){#markdown-toc-technical-details}

This post showcases how to use the highly optimized, GPU-based
collectives of Nvidia's ([NCCL library](https:github.com/nvidia/nccl))
by using it as a [UCC](https://github.com/openucx/ucc) backend through
MPI. Further down, we show performance comparisons in form of graphs.

UCC (*Unified Collective Communication*) is a collective communication
operations API and library developed under the umbrella of the [UCX
collaboration](https://openucx.org/). UCC provides various
highly-scalable and performant collectives for HPC, AI/ML, and IO
workloads; together with non-blocking collective operations for MPI and
OpenSHMEM models.

In the following, we take a traditional *allreduce* MPI example code and
accelerate it by using UCC with NCCL through the correct environment
variables. The best thing: The original MPI program stays untouched; no
single NCCL or UCC function is used directly.

## Environment Setup

To setup the MPI → UCC → NCCL stack, the right environment variables
need to be set. Let's go through them using the OpenMPI implementation.

### Enabling UCC in OpenMPI

First, to enable UCC in OpenMPI, two important environment variables
need to be set:

1.  `OMPI_MCA_coll_ucc_enable=1`{.language-plaintext
    .highlighter-rouge}: Sets the MCA component for MPI to use
    UCC-accelerated collectives. Can be disabled by setting
    `OMPI_MCA_coll_ucc_enable=0`{.language-plaintext
    .highlighter-rouge}.
2.  `OMPI_MCA_coll_ucc_priority=100`{.language-plaintext
    .highlighter-rouge}: Is used to give priority for certain
    collective-implementing component. Currently, we must give UCC the
    highest priority (100), so that UCC is selected as the collective
    library over the other default collective
    libraries.^[1](#fn:1){.footnote rel="footnote"}^

Now, on to selecting our desired NCCL backend within UCC!

### Enabling NCCL in UCC (Team Layer Selection)

Different backends are supported in UCC as *Team Layers* (TLs). Beyond
NCCL, also [RCCL, SHARP, and basic UCX transports are
supported](https://github.com/openucx/ucc#supported-transports). The TLs
are thin, composable abstractions that are independent from each other
and are further built up of *Collective Layers* (CLs).

Generally, the internal heuristics of UCC aim to always select the
best-performing algorithm, given available Team and Collective Layers.
However, in some cases these heuristics are not making the optimal
choice, or one wants to select a certain TL for debugging or
experimentation. Then in order to select a specific TL, an according
environment variable has to be set, like:

::: {.language-plaintext .highlighter-rouge}
::: {.highlight}
``` highlight
UCC___TUNE=token1#token2...#tokenN
```
:::
:::

with `#`{.language-plaintext .highlighter-rouge} delimiting a list of
*tokens*. A token has the following structure:

::: {.language-plaintext .highlighter-rouge}
::: {.highlight}
``` highlight
token=coll_type:msg_range:mem_type:team_size:score:alg
```
:::
:::

where `:`{.language-plaintext .highlighter-rouge} is the delimiter for a
list of *qualifiers*. Each token must contain either the
`score`{.language-plaintext .highlighter-rouge} or
`alg`{.language-plaintext .highlighter-rouge} qualifier; the other
qualifiers are optional. The score determines which algorithm of a
certain backend is picked. See also the [UCC Wiki
FAQ](https://github.com/openucx/ucc/wiki/FAQ#6-what-is-tl-scoring-and-how-to-select-a-certain-tl)
for a little more insight.

In order to use the NCCL backend for `allreduce`{.language-plaintext
.highlighter-rouge} operations on device-side CUDA buffers
(`cuda`{.language-plaintext .highlighter-rouge}) for all message sizes,
and give them the highest priority (`inf`{.language-plaintext
.highlighter-rouge}), we need to set:

::: {.language-plaintext .highlighter-rouge}
::: {.highlight}
``` highlight
UCC_TL_NCCL_TUNE=allreduce:cuda:inf
```
:::
:::

Using the NCCL TL with UCC has some restrictions; to be specific
[Multi-Process Service
(MPS)](https://docs.nvidia.com/deploy/mps/index.html) is not supported
i.e. only 1 MPI rank per GPU is supported.

### All The Variables

To sum up, the following variables need to be set to enable NCCL
collectives used by UCC from OpenMPI:

::: {.language-plaintext .highlighter-rouge}
::: {.highlight}
``` highlight
OMPI_MCA_coll_ucc_enable=1
OMPI_MCA_coll_ucc_priority=100
UCC_TL_NCCL_TUNE=allreduce:cuda:inf
```
:::
:::

For more insights and detailed explanations, check out the [UCC user
guide](https://github.com/openucx/ucc/blob/master/docs/user_guide.md).

## Results

To showcase the performance of the different approaches, we use the
*allreduce* latency test of the [*OSU
Micro-Benchmarks*](http://mvapich.cse.ohio-state.edu/benchmarks/) (OSU).
We give `-d cuda`{.language-plaintext .highlighter-rouge} as a runtime
option in order to use device-side buffers.

We use NCCL v2.15.1 and UCC v1.1.0 on our [JUWELS Booster
system](https://apps.fz-juelich.de/jsc/hps/juwels/booster-overview.html)
with Nvidia A100 GPUs (40 GB). 8 tasks are used, with one per GPU --
i.e. two nodes in total

In order to see the potential of UCC+NCCL (UCC with the NCCL backend
enabled) over sole UCC, we use message sizes of the order of 1 GB;
otherwise, the startup overheads of UCC+NCCL will impact the results.

### 1. Plain OpenMPI

::: {.language-plaintext .highlighter-rouge}
::: {.highlight}
``` highlight
OMPI_MCA_coll_ucc_enable=0 \
srun ./osu_allreduce -d cuda -m 1000000000:1000000000 --mem-limit 1048576000
# OSU MPI-CUDA Allreduce Latency Test v7.0
# Size       Avg Latency(us)
1000000000        1183590.64
```
:::
:::

### 2. OpenMPI with UCC

⚠️ Make sure that the NCCL is not available in the environment
(`module unload`{.language-plaintext .highlighter-rouge})! Otherwise,
the UCC heuristics will select the NCCL TL.

::: {.language-plaintext .highlighter-rouge}
::: {.highlight}
``` highlight
OMPI_MCA_coll_ucc_enable=1 OMPI_MCA_coll_ucc_priority=100 \
srun ./osu_allreduce -d cuda -m 1000000000:1000000000 --mem-limit 1048576000
# OSU MPI-CUDA Allreduce Latency Test v7.0
# Size       Avg Latency(us)
1000000000          23284.48
```
:::
:::

### 3. OpenMPI with UCC+NCCL

::: {.language-plaintext .highlighter-rouge}
::: {.highlight}
``` highlight
OMPI_MCA_coll_ucc_enable=1 OMPI_MCA_coll_ucc_priority=100 \
UCC_TL_NCCL_TUNE=allreduce:cuda:inf \
srun ./osu_allreduce -d cuda -m 1000000000:1000000000 --mem-limit 1048576000
# OSU MPI-CUDA Allreduce Latency Test v7.0
# Size       Avg Latency(us)
1000000000          18230.13
```
:::
:::

For this latency, lower is better. Using NCCL with UCC (third variant)
gives us the best result and is 65 times faster than the plain OpenMPI
version.

## Scaling Plots

To investigate the scaling behavior, we run the three variants on a
number of GPUs: 2, 4, 8, 12, 16 GPUs.  
⚠️ Remember: Each JUWELS Booster node has 4 GPUs.

All runs in the scope of this investigation were confined to one
DragonFly group, to limit influences of network traffic. Nevertheless,
similar performance behaviour can be seen when using scattered nodes as
well.

### Average Latency

[![](/assets/img/posts/2023-07-18-mpi-ucc-nccl/OSU_ScalingLatencyPlot_cuda_BCELL09.svg)](/assets/img/posts/2023-07-18-mpi-ucc-nccl/OSU_ScalingLatencyPlot_cuda_BCELL09.svg)

The graph shows average latency plotted over the number of GPUs,
annotated with the number of nodes, number of tasks launched per node,
and average latency values. Again, lower is better.

The order-of-magnitude improvement of any UCC-enabled execution over
plain OpenMPI is consistent over the scope of tested number of GPUs.
UCC+NCCL improves significantly over sole UCC for multi-node setups.
Note that two configurations to use four GPUs are plotted, once using
four GPUs in one node (lower points) and once using two GPUs each in two
nodes (upper points).

The best latency for the largest setup of 16 GPUs is achieved from
UCC+NCCL with only 20 ms GPUs.

### Bus Bandwidth

It is also possible to calculate bandwidths associated to the operation.
The naïve algorithmic bandwidth, \\B\_\text{alg} = S/t\\ (with data size
\\S\\ and execution time \\t\\), can be extended to the *bus bandwidth*
\\B\_\text{bus}\\, which is independent of the number of ranks \\n\\ and
is defined [per NCCL
documentation](https://github.com/NVIDIA/nccl-tests/blob/master/doc/PERFORMANCE.md#allreduce)
to the following:

\begin{equation} B\_\text{bus} = B\_\text{alg} \times \frac{2(n − 1)}{n}
\end{equation}

It is an attempt to show the impact of the hardware bottleneck, i.e.
NVLink, PCIe, and InfiniBand, in our case.

#### Comparing MPI, UCC, UCC+NCCL

When all four GPUs in a node are used (4 GPUs \[nodes=1,
taskspernode=4\]), the theoretical peak unidirectional bandwidth for
[A100's NVLink3](https://de.wikipedia.org/wiki/NVLink) is 300 GB/s, and
when only two GPUs in a node are used (2 GPUs \[nodes=1,
taskspernode=2\]) the peak unidirectional bandwidth is 100 GB/s.

In the case of inter-node communications, [InfiniBand
HDR200](https://network.nvidia.com/files/doc-2020/wp-introducing-200g-hdr-infiniband-solutions.pdf)
offers 200 Gbit/s per adapter. When using all four adapters for
communication (GPUs \> 4 and nodes \> 1), the peak bandwidth is 4×200
Gbit/s = 100 GB/s; and half of it (50 GB/s) when using only two adapters
are used (4 GPUs \[nodes=2, taskspernode=2\]).

All of this aids from the capacity of NCCL to utilize indirect
connections between devices, leading to the use of more resources.
Realistically achievable bandwidths are lower, of course.

[![](/assets/img/posts/2023-07-18-mpi-ucc-nccl/OSU_ScalingBusBWPlot_cuda_BCELL09.svg)](/assets/img/posts/2023-07-18-mpi-ucc-nccl/OSU_ScalingBusBWPlot_cuda_BCELL09.svg)

The graph shows the bandwidth plotted over number of GPUs, annotated
with number of nodes, number of tasks launched per node, and the
bandwidth value. Again, with the NCCL backend of UCC we are able to come
close to theoretical peak bandwidth, while the plain OpenMPI runs are
more than an order of maximum lower.

#### Comparing UCC+NCCL, NCCL

[![](/assets/img/posts/2023-07-18-mpi-ucc-nccl/nccl2_osu_comparison_plusratio.svg)](/assets/img/posts/2023-07-18-mpi-ucc-nccl/nccl2_osu_comparison_plusratio.svg)

Finally, we can also compare the allreduce bandwidth between UCC+NCCL
and sole NCCL, using the [*NCCL
Tests*](https://github.com/NVIDIA/nccl-tests/) in addition to the OSU
benchmarks. As before, the graph plots average bus bandwidth over number
of GPUs for a 1 GB message size. The right y axis shows the relative
difference between OpenMPI with UCC+NCCL and direct NCCL. One can see
that the overhead introduced by UCC is minimal; for multi-node runs it
is about 0.5%.

Furthermore, OSU Micro-Benchmarks have also added support to benchmark
NCCL directly which is also plotted in the
graph.^[2](#fn:confusion){.footnote rel="footnote"}^

## Summary

Enabling NCCL as a backend to UCC, and using UCC from OpenMPI is very
easy and only a matter of setting two (if UCC heuristics select the
desired TL) or three environment variables -- no code changes necessary.
Performance gains are significant, like 65× in our example
configurations shown above. Of course, a proper setup is
needed.^[3](#fn:setuphpc){.footnote rel="footnote"}^

***Acknowledgements**: Thanks go out to [Jiri Kraus
(NVIDIA)](https://www.linkedin.com/in/jirikraus/) for being available
for a multitude of questions arising during writing of this article.
Thanks, Jiri!*

## Technical Details

- [Build UCC](https://github.com/openucx/ucc)
- Modules used for the results, all built according to our [published
  EasyBuild infrastructure](https://github.com/easybuilders/JSC):
  - NCCL/2.15.1
  - Cuda/11.7
  - OpenMPI/4.1.4
  - GCC/11.3.0
  - UCX/1.13.1
  - UCC/1.1.0
- Building OSU:

  ::: {.language-plaintext .highlighter-rouge}
  ::: {.highlight}
  ``` highlight
  >> cd osu-micro-benchmarks-7.0.1
  >> ./configure CC= CXX= --enable-cuda –with-cuda-include= \--with-cuda-libpath= --prefix=$(pwd)
  >> make -j 24; make install
  ```
  :::
  :::
- Sample Script to run on JUWELS Booster

  ::: {.language-plaintext .highlighter-rouge}
  ::: {.highlight}
  ``` highlight
    #!/usr/bin/env bash

    ml purge
    ml load Stages/2023 
    ml load GCCcore/.11.3.0 UCX UCC GCC/11.3.0 OpenMPI/4.1.4 
    ml load NCCL/default-CUDA-11.7
    ml list

    export OMPI_MCA_coll_ucc_enable=1 
    export OMPI_MCA_coll_ucc_priority=100
    export UCC_TL_NCCL_TUNE=allreduce:cuda:inf 
    export OMPI_MCA_coll_ucc_verbose=3
    export UCC_LOG_LEVEL=INFO

    NODES=2
    TASKSPERNODE=4
    GPU=8 

    srun --jobid $SLURM_JOB_ID --nodes=${NODES} --ntasks-per-node=${TASKSPERNODE} --gres=gpu:4 \
    ./osu-micro-benchmarks-7.0.1/c/mpi/collective/osu_allreduce -d cuda -m 1000000000:1000000000 --mem-limit 1048576000 > "OSU_NODE${NODES}_GPU${GPU}_UCCNCCL_${SLURM_JOB_ID}.txt" 2>&1
  ```
  :::
  :::

::: {.footnotes role="doc-endnotes"}
1.  ::: {#fn:1}
    Read more in the according GitHub issue on [why you should set
    `OMPI_MCA_coll_ucc_priority=100`{.language-plaintext
    .highlighter-rouge}?](https://github.com/open-mpi/ompi/issues/9885) [↩](#fnref:1){.reversefootnote
    role="doc-backlink"}
    :::

2.  ::: {#fn:confusion}
    Just to avoid confusion: The *NCCL Tests* directly call NCCL
    functions, no OSU micro-benchmarks involved; *UCC+NCCL* utilize the
    plain old OSU micro-benchmarks calling MPI functions, but enabling
    NCCL through UCC in the backend, unbeknown to the OSU
    micro-benchmarks; *NCCL (OSU)* take the new NCCL variant of the OSU
    micro-benchmarks, in which NCCL functions (like
    `ncclAllReduce()`{.language-plaintext .highlighter-rouge}) are
    called directly! [↩](#fnref:confusion){.reversefootnote
    role="doc-backlink"}
    :::

3.  ::: {#fn:setuphpc}
    A proper setup in HPC is half the battle... Check out [JSC's
    EasyBuild
    repository](https://github.com/easybuilders/JSC)! [↩](#fnref:setuphpc){.reversefootnote
    role="doc-backlink"}
    :::
:::


Environment Setup Enabling UCC in OpenMPI Enabling NCCL in UCC (Team Layer Selection) All The Variables Results 1. Plain OpenMPI 2. OpenMPI with UCC 3. OpenMPI with UCC+NCCL Scaling Plots Average Latency Bus Bandwidth Comparing MPI, UCC, UCC+NCCL Comparing UCC+NCCL, NCCL Summary Technical Details  This post showcases how to use the highly optimized, GPU-based collectives of Nvidia’s (NCCL library) by using it as a UCC backend through MPI.

Rogue Scholar Beiträge

MPI as API: Using UCC’s NCCL Backend for MPI’s Allreduce