Can Fortran’s `do concurrent’ replace
directives for accelerated computing?
Miko Stulajter, Ronald M. Caplan, and Jon A. Linker
www.predsci.com

Introduction• Directives (e.g. OpenMP/OpenACC) are popular for parallelization onCPUs and GPUs

• Standard languages have begun to add features that compilers can use
to parallelize code•

C++17’s Standard Parallel Algorithms•

Fortran’s do concurrent• GPU-offload of directives supported by many compilers:•NVIDIA HPC SDK, Intel OpenAPI HPC Toolkit, LLVM Flang (along with AMD AOCC/AOMP
extensions), IBM’s XL, HPE’s Cray Fortran• Here, we want to test the current status of being able to replace
directives with do concurrent for accelerated computingDirectives
Directives are widely used for parallelizing codes
PROS• Performance canbe similar to lowlevel APIs

CONS
• Not always
supported• Portability• APIs can change,
requiring re-writes• Minimal code
interference• Can make code
harder to readfor (i=0; i<N; i++)y[i] = a*x[i] + y[i];#pragma acc kernelsfor (i=0; i<N; i++)y[i] = a*x[i] + y[i];“OpenACC is a user-driven directivebased performance-portable parallel
programming model. It is designed forscientists and engineers interested in
porting their codes to a wide-variety ofheterogeneous HPC hardware platforms
and architectures with significantly lessprogramming effort than required with a
low-level model.” - openacc.org“OpenMP is a specification for a set ofcompiler directives, library routines, andenvironment variables that can be used
to specify high-level parallelism in
Fortran and C/C++ programs.” openmp.org

Code Description• Space weather events can cause interference& damage to electronic infrastructure• New NASA+NSF program called SWQU toimprove models of solar wind and solarstorms• HipFT: A flux-evolution code to generate
observation-based model boundary
conditions• Most expensive computation of HipFT
encapsulated in BC-smoothing tool DIFFUSE• Here we use DIFFUSE as a mini-app ofHipFT for DC tests

Code Description cont. / Test Problem• Integrates the sphericalsurface heat equation

For a test problem, we select a real-world exampleof using DIFFUSE, that of smoothing the ‘Native res
PSI map‘ described in• Logically rectangular
non-uniform grid
• Operator is discretizedwith a second-order
central finite-difference
scheme• Time integration withsecond-order Legendre
polynomial extended
stability Runge-Kutta
scheme (RKL2)

The grid has a resolution of 3974x2013 in theta-phi
and the test requires 40,260 iterations of thediffusion operator

Computational Environment• Utilized the latest compilers attime of testing• Singularity containers used to

streamline testing and provide
reproducibility• Containers provide

performance comparable tobare metal• Computational resources provided
by NSF's XSEDE program and theCSRC at SDSU

Baseline Performance Results••••

OpenACC and OpenMP similar performance for CPU
All compilers similar for CPU
Note DIFFUSE is memory-bound so low speed-up over CPU cores is not unusual
nvfortran faster than gfortran for OpenACC GPU

Implementation: Fortran’s `do concurrent’• Introduced in ISO StandardFortran 2008

• Indicates loop can be run without-of-order execution• Can be seen as hint to thecompiler that loop may
be parallelizable• Current specification has nosupport of reductions oratomics

Implementation: Code Versions• Original: Uses directives and data movement directives• New: Uses do concurrent except for reductions and datamovement directives• Serial: No directives or do concurrent loops• Experimental: All directives removed, and all parallelizable loops
utilize do concurrent including reductions.
This represents "ideal" situation of having no directives

Implementation: Compiler Options• Utilized -03 and -march=<ARCH> for all compilers• gfortran:• CPU: -fopenmp and/or -ftree-parallelize-loops=<N>• GPU: -fopenacc and -foffload=nvptx-none• nvfortran:• CPU: -acc=multicore and/or -stdpar=multicore• GPU: -stdpar=gpu and/or -acc=gpu and -gpu=cc<X><Y>,cuda<X>.<Y>• Note unified memory is enabled by default(can turn off with -gpu=nomanaged)• ifort:• CPU: -fopenmp• GPU: No support for NVIDIA GPUs

Results: nvfortran• GPU performance stayed consistent• CPU performance stayed consistent• Experimental code worked correctlyResults: gfortran• No GPU support of do concurrent

parallelism• CPU do concurrent support relies on
auto parallelization• Small performance loss with autoparallelization

Results: ifort•

do concurrent gave better
performance on CPU•

Currently no support on the GPUfor do concurrent•

The experimental code gave thewrong answer•

Compiler failed to correctly
detect reductions

Discussion• Compatibility on the GPU:
• Currently only nvfortran has do concurrent support• Using do concurrent we lose gfortran GPU support• Planned ifort support of do concurrent on Intel GPUs
• Removing data movement directives and relying on unified memory
could cause performance loss, but doesn’t here

• Portability:• CPU multicore parallelization was not lost (except for Experimentalcode)• nvfortran and ifort have direct support of do concurrent onCPUs• gfortran relies on auto-parallelization detection• Implicit reductions with do concurrent not supported everywhere

Discussion Continued• Performance:• Comparable performance for the original and new code for both
CPUs and GPUs

• Unified memory on GPU gave comparable performance to manual
data management with directives in this case

• Summary:• do concurrent allows cleaner looking code and adds robustness• nvfortran allowed for the elimination of all directives• Using a combination of directives and do concurrent gives bettercross compiler/hardware compatibility at this time

Conclusion

Can Fortran's do concurrent replace directives for accelerated computing?

• With nvfortran and NVIDIA GPUs, for some codes (such as ours) theanswer is YES, and with no (or minimal) loss of performance.

• Upcoming language features and compiler implementations, may allow
more complicated codes to eventually be parallelized without directives

Reproducibility with Singularity

predsci.com/papers/dc
Make sure singularity container does nothave overhead:Thank you for your attention.
Questions?
https://arxiv.org/abs/2110.10151