Various notes from the Accelerating Devices Lab (X-Dev) of Jülich Supercomputing Centre

Chelsea John

The Supercomputing Conference 2023 took place in Denver, Colorado, from November 12th to 17th. For the

 Women in HPC

workshop, we submitted a paper, which focused on benchmarking different accelerators for AI. The paper was accepted and I was invited to hold a lightning talk to show the work, spun off our OpenGPT-X project.

SC23 WHPC Workshop Paper: OpenGPT-X – Novel Architecture Exploration

Stepan Nassyr

** Poster in institute repository:** http://dx.doi.org/10.34734/FZJ-2023-03437 During the RISC-V Summit Europe 2023


in Barcelona we presented our work generating highly optimized RISC-V and ARM GEMM microkernels for BLIS using a custom software tool.

 1

We presented results on the Fujitsu A64FX processor, the in-development RISC-V

 VEC

processor from the EUPILOT project using an FPGA SDV (RVV 0.7.1, later RVV

Poster: GEMMERATOR – GEMM Kernel Generator

Andreas Herten

Together with our colleagues from Helmholtz-Zentrum Dresden-Rossendorf (HZDR) and in collaboration with HIDA and OpenHackathons, we hosted the

 Helmholtz GPU Hackathon 2023

in Jülich in May.

 I’ve blogged about the event for the
 
 Zweikommazwei
 
 blog of Forschungszentrum;

Helmholtz GPU Hackathon 2023

At the end of May, Forschungszentrum Jülich got a new website. Finally, the group also has a dedicated page there, listing the things we do, the project we are part of, and all the members in the group. And thanks to a photo shooting by Sebastian, we have great portraits and an amazing group picture.

X-Dev Website at fz-juelich.de

Jayesh Badwaik

** Poster in institute repository:** http://dx.doi.org/10.34734/FZJ-2023-04519  At the ISC High Performance Conference 2023 a little while ago, we presented a project poster on the SCALABLE project, embedded at the end of this post. The work was also presented as a paper at the the Computing Frontiers 2023 conference previously.

ISC23 Project Poster: SCALABLE: Scalable Lattice-Boltzmann Leaps to Exascale

Sebastian Achilles

MSA Concept MSA Software Building Blocks Workshop Exercises 1: Hello World! 2. GPU Hello World! 3: CPU-GPU Ping Pong Slides  On May 29, we held a workshop about using the Modular Supercomputing Architecture (MSA) together with project partners from ParTec.

MSA Introduction Workshop

** Poster publication:** http://hdl.handle.net/2128/34532  The ISC High Performance Conference 2023 was held at Hamburg, Germany from 21st May to 25th May. At the conference, we presented a project poster on the OpenGPT-X project, outlining the progress and initial exploration results. The poster was even featured in HPCWire’s May 24 recap of ISC within the AI segment!

ISC23 Project Poster: OpenGPT-X – Training Large Language Models on HPC Systems

For a recent talk at DKRZ in the scope of the natESM project, I created a table summarizing the current state of using a certain programming model on a GPU of a certain vendor, for C++ and Fortran. Since it lead to quite a discussion in the session, I made a standalone version of it with some updates and elaborations here and there.

GPU Vendor/Programming Model Compatibility Table

TL;DR: I held a HPC intro talk. Slides are below.

In MAELSTROM, we connect three areas of science:

 🌍Weather and climate simulation

with

 🤖Machine Learning methods and workflows

using

 📈HPC techniques and resources

. Halfway into the project, we held a boot camp at JSC to teach this Venn diagram to a group of students a few days ago. Some were ML experts, but had never used a HPC system.

Talk: Introduction to HPC

Albert Kahira

A few days ago, OPTIMA announced the release of deliverable 3.5, to which I contributed. This deliverable is part of a set of five deliverables under work package 3. But first, let’s talk about

 OPTIMA

. OPTIMA is an EU-funded project whose goal is to prove that several HPC applications can take advantage of the future highly heterogeneous FPGA-populated HPC systems.

OOPS Version 1 Release

On June 21 and 22, we held a workshop looking back on the last ten years of our lab together with NVIDIA – the NVIDIA Application Lab at Jülich (or

 NVLab

, as I sometimes abbreviate it). The material can be found in the agenda at Indico. We invited a set of application owners with which we worked together during that time to present past developments, recent challenges, and future plans.

10 Year Anniversary Workshop of NVIDIA Application Lab at Jülich

In November of 2022, I created a table comparing GPU programming models and their  support on GPUs of the three vendors (AMD, Intel, NVIDIA) for a talk. The audience liked it, so I beefed it up a little and posted it in this very blog.

Paper: Many Cores, Many Models – GPU Programming Model vs. Vendor Compatibility Overview

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.

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

A few months ago, we extended the JURECA Evaluation Platform

 1

at JSC by two nodes with AMD Instinct MI250 GPUs (four GPUs each). The nodes are Gigabyte G262-ZO0 servers, each with a dual socket AMD EPYC 7443 processor (24 cores per socket, SMT-2) and with four MI250 GPUs (128 GB memory). OSU Bandwidth Micro-Benchmark A100 Comparison GPU STREAM Variant Data Size Scan Threads and Data Sizes Conclusion Technical Details OSU

First Benchmarks with AMD Instinct MI250 GPUs at JSC

Carolin Penke

** Poster publication:** http://hdl.handle.net/2128/32006  The 14th JLESC workshop (JLESC: Joint Laboratory for Extreme-Scale Computing) was hosted by the National Center for Supercomputing Applications (NCSA) in Urbana, Illinois from 28th September to 30th September. We had the opportunity to present the OpenGPT-X project in form of a poster.

Poster: OpenGPT-X - Training Large Language Models on HPC Systems

About    This blog post is based on a presentation I held at the “New Trends in Computational Science in Engineering and Industrial Mathematics” workshop in Magdeburg on 01/07/2022. My goal is to give a brief introduction to the state of current large language models, the OpenGPT-X project, and the transformer neural network architecture for people unfamiliar with the subject. About What is a language model?

A mathematician’s introduction to transformers and large language models

1

This blog is an experiment. We want to share bits and pieces of our work; the reports we write, the presentations we hold, or the little discoveries we make, or even some first, water-testing investigations; and all the rest. It’s a

 documentation

of what we do. Little bits of science, collected in the open, and sometimes even not that little.

DOIng it Right! (DOIs for This Blog)

JSC Accelerating Devices Lab

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"}
    :::
:::
