GPU acceleration ofOpenMX with OpenACC
and CUDA
Hiroyuki Kawai
Department of Physics, Niigata University

What is a GPU?
•GPU (Graphics Processing Units) areprocessors that originally specialized inimage processing.
•Recently, the application of GPU
computing resources to general-purpose
computing purposes other than image
processing has become popular.
NVIDIA A100, the GPU used in thisstudy

Comparison of CPU and GPU
Pros and cons of CPU

Pros and cons of GPU

CPUs and GPUs have their own advantages anddisadvantages, so it is best to use a combination of both.
Overall efficiency can be improved by assigning
processing to the processor most suited for the application.
[1] https://www.cc.u-tokyo.ac.jp/events/lectures/207/20230621-1.pdf

CPU-GPU heterogeneous programming

Heterogeneous programming is programming on acomputer system built by combining different types ofprocessors.

GPU supercomputers currently in operation•Japanese domestic•ISSP, System C kugui (AMD EPYC 7763 +NVIDIA A100)•Information Technology Center, The
University of Tokyo, Wisteria/BDEC-01
Aquarius (Intel Xeon Platinum 8360Y +NVIDIA A100)•Overseas•OLCF, Frontier (AMD EPYC 7453s + AMD
Radeon Instinct MI250X)•CSC, LUMI (AMD EPYC 7A53 + AMD
Radeon Instinct MI250X)•CINECA, Leonardo (Intel Xeon 8358 +NVIDIA custom Ampere GPU)Research Objectives
 Currently, many first-principles electronic structure calculation
software such as ABINIT, Quantum ESPRESSO, SIESTA, and VASP
are GPU accelerated.
 As for OpenMX, there has been only one report of an attempt toaccelerate with GPUs[2], and this report was made in the old days
when GPUs did not have sufficient performance, so it is hard to saythat it showed the superiority of GPUs over conventional CPU-basedparallelization (MPI parallelization).
 In this study, we aim to accelerate the bottleneck in OpenMX 3.9,
which is the diagonalization of matrices, using GPU. The acceleration
through GPU involves the use of CUDA math libraries (cuBLAS andcuSOLVER) and OpenACC. Subsequently, we assess the performancethrough benchmark calculations.

 For GPU acceleration, Wisteria/BDEC-01 Aquarius frominformation technology center, the University of Tokyo, wasused. For benchmark calculations, Wisteria/BDEC-01
Aquarius and ISSP's system B ohtaka were used.
[2] J. H. Parq, E. Sevre, and S. M. Lee, Int. j. comput. appl. 98, 20 (2014).

What is CUDA?
 CUDA (Compute Unified Device Architecture) is a GPU
program development environment developed byNVIDIA.
 CUDA enables high-speed parallel processing using theGPU's multiple arithmetic units, using C-like program
descriptions.
 Programming with CUDA can take GPU performance tothe extreme. On the other hand, programming with
CUDA is very time-consuming and difficult.
 Also, CUDA only works on NVIDIA GPUs (CUDA does
not work on AMD GPUs or Intel GPUs).
 GPU versions of BLAS and LAPACK are also includedin CUDA in the form of cuBLAS and cuSOLVER.

What is OpenACC?
 OpenACC is a standard for parallel programming on accelerator
devices, not limited to GPUs, using a directive-based programming
method (i.e., inserting directive lines in the source) like OpenMP.
 OpenACC can easily program GPUs with #pragma acc directive.
 If you aim to maximize performance to the utmost on GPU, CUDA ispreferable. However, if achieving reasonably good performance on
GPU is sufficient, then using OpenACC would be a good choice.
 The code written in OpenACC is expected to run on GPUs from AMD
and Intel as well, but unfortunately, it seems that AMD and Intel do notintend to support OpenACC. OpenACC is supported only on NVIDIA
GPUs, similar to CUDA.
 GPU versions of BLAS and LAPACK are not included in OpenACC.
If you need a GPU version of BLAS or LAPACK, you can use
cuBLAS or cuSOLVER. Alternatively, you can use an external librarycalled MAGMA[3] instead of cuBLAS or cuSOLVER.
 [3] https://icl.utk.edu/magma/

More about GPU acceleration
3•The typical O(N ) band calculations (collinear, non-collinear), divide conquer (DC) method, divideconquer method with localized natural orbitals (DCLNO) method and NEGF calculations were performedon GPU.
•We cannot say for sure because we have not done
enough benchmark calculations, but it appears that thecalculation speed of the GPU exceeds that of the CPU
on large systems of 5000 basis or more.
•Similarly, in the NEGF calculation, the GPU does notseem to perform well unless the system is large,
although benchmark calculations are not sufficient.
•The most successful GPU speedups were achieved with
the DC and DC-LNO methods.

More about GPU acceleration•For band calculations, the CPU versions of ELPA[4] and ScaLAPACK
were accelerated on the GPU by replacing the GPU version of ELPAand ScaLAPACK (cuScaLAPACK[5]), respectively.
•For the DC-LNO method, the real matrix product (DGEMM) and realsymmetric eigenvalue problem solver (dsyevd) of BLAS and LAPACK
were replaced by the corresponding routines of cuBLAS andcuSOLVER, respectively, for GPU acceleration.
•The DC method was accelerated on the GPU by offloading some forloops to the GPU with OpenACC and by replacing the LAPACK realsymmetric eigenvalue problem solver (dsyevx) with the corresponding
cuSOLVER routines.
•The same procedure was used for the NEGF calculations, which werealso performed by GPU.
•In all cases, no modifications to the source code were made at thealgorithm level.
• [4] https://github.com/marekandreas/elpa• [5] https://github.com/smorita/cuscalapack

Comparison of computation time between GPU and CPU for bandcalculations (collinear)

The diagonalization time and total computation time, except for the case with 32GPU/CPU numbers, consistently show the GPUs to be just under 2 times faster thanthe CPU. The scaling of the GPUs is better than that of the CPUs, and it scales nearlylinearly. This represents an ideal scaling.
Measured by calculation of 384 atoms of silicon supercell (total basis number is 9984),
8×8×8 k-mesh, 1SCF only.
GPU: NVIDIA A100 (8 GPUs per node in Wisteria/BDEC-01 Aquarius) , 4 MPI processes are allocated per GPU.
CPU: AMD EPYC 7763 (2 CPUs per node in System B ohtaka)， 16 MPI process allocated per CPU.

Comparison of computation time between GPU and CPU for bandcalculations (non-collinear)

The diagonalization time and total computation time are about 20% to 30% faster onGPUs than on CPUs, regardless of the number of GPUs/CPUs, but scaling on GPUs isworse than on CPUs.
Measured by calculation of 192 atoms of silicon supercell (total basis number is 4992),
9×9×9 k-mesh, 1SCF only.
GPU: NVIDIA A100 (8 GPUs per node in Wisteria/BDEC-01 Aquarius) , 3 MPI processes are allocated per GPU.
CPU: AMD EPYC 7763 (2 CPUs per node in System B ohtaka)， 8 MPI process allocated per CPU.

Comparison of computation time between GPU and CPU for DC-LNO
methodIn terms of diagonalization time, GPUs are 4 to 5 times faster than CPUs for both
GPU/CPU counts, and GPU scaling is about the same as that of CPUs. In termsof total computation time, GPUs are about 2 times faster than CPUs for all
GPU/CPU numbers, and GPU scaling is about the same as that of CPUs.
Measured by calculation of 650 atoms of DNA (total basis number is 12980)
GPU: NVIDIA A100 (8 GPUs per node in Wisteria/BDEC-01 Aquarius) , 9 MPI processes are allocated per GPU.
CPU: AMD EPYC 7763 (2 CPUs per node in System B ohtaka)， 64 MPI process allocated per CPU.

Summary•The typical O(N3) band calculations (collinear, noncollinear), DC method, DC-LNO method, and NEGF
calculations in OpenMX 3.9 were computed on GPUs.
•These GPU calculations were somewhat faster, except for
the NEGF calculation.
•In particular, for the DC-LNO method, the GPU
calculations achieved a 4 to 5 times speedup in thecomparison of diagonalization times compared to CPU
calculations.

Future works•Increase the processes offloaded to the GPU.
•Currently, only a portion of the matrix diagonalization
calculations are done on the GPU, but the goal is toperform all matrix diagonalization calculations on theGPU.
•In the future, efforts will be made to utilize the GPU fortasks beyond matrix diagonalization as much as possible.
•Although related to the previous section, rewrite the code
so that the OpenMP parallelization part is performed by
the GPU as much as possible.
•Only processes that the GPU is not good at should bedone on the CPU.