OpenACCのCUDA Graphを使ってみた
目次
はじめに
こんにちは、プロメテック・ソフトウェアの阿部です。
エンジニアとして、アプリケーションのGPU高速化に関連する業務や、NVIDIAコンパイラのサポートに携わっております。
NVIDIA HPC SDK 24.7*1から、OpenACCやCUDA FortranでCUDA Graphが利用できるようになりました。
(リリースノートの和訳はこちら*2)
このコラムでは、OpenACCのCUDA Graphの利用方法をサンプルコードを例に紹介いたします。
CUDA Graphとは何か
CUDA Graphとは、コンパイラがキャプチャ領域(後述)内の一連のカーネルとメモリ操作をあらかじめ定義しグラフとして構築する機能です。
CUDA Graphの機能を使うと、グラフ全体が一括して実行され、カーネルローンチやメモリ操作のオーバーヘッドを削減することができます。
CUDA Graph自体は以前からCUDAで使える機能で、グラフ解析の機能ではありません。
従来のOpenACCを使ったGPU上での実行方法では、繰り返し実行されるGPU上での一連の計算タスクについて、コンパイラは個々のカーネルやメモリ操作(データ転送)を一つずつスケジューリングしていました。
OpenACCのCUDA Graphの機能により、繰り返し実行されるキャプチャ領域内の一連の計算タスクについて、GPU上での計算タスク順序や並列性のスケジューリングによる最適化が行われ、ユーザは実行効率の向上が期待できます。
実アプリケーションに対してOpenACCでCUDA Graphを適用する際は、制限事項*3をよく読み、計算結果の妥当性を確認する必要があります。
キャプチャ領域内に指定する処理によっては、処理を一部変更する等、試行錯誤が必要かもしれません。
はじめからプログラムの広い範囲に対してCUDA Graphを適用しようとするのではなく、まずは、プログラム内の狭い範囲をキャプチャ領域として指定し、計算結果が正しいことを確認した後に、プログラムのより広い範囲をキャプチャ領域として指定する、といった段階的な実装方針をオススメします。
OpenACCのCUDA Graphの機能により性能が向上した実アプリケーションとして、例えばGROMACS 2023*4があるようです。
OpenACCのCUDA Graph適用例(C言語)
ここでは、OpenACCによりGPU化されたプログラムに対して、CUDA Graphを適用する方法をソースコードベースで紹介します。
OpenACCによりGPU化されたC言語のプログラムに対して、CUDA Graphを適用したサンプルコードを以下に示します。
1 #include <stdio.h> 2 3 #define CUDAGRAPH 4 5 #ifdef CUDAGRAPH 6 #include <openacc.h> 7 #endif 8 9 int main(){ 10 11 int i, j, k; 12 int n = 5000; 13 double a[n], b[n], c[n]; 14 int async1 = 1; 15 #ifdef CUDAGRAPH 16 acc_graph_t pgraph; 17 #endif 18 19 for(i = 0; i < n; i++) a[i] = 2.0; 20 for(i = 0; i < n; i++) b[i] = 0.5; 21 22 #ifdef _OPENACC 23 #pragma acc data copy(a[0:n]) copyin(b[0:n]) create(c[0:n]) 24 { 25 #endif 26 i = 0; 27 while(i < 1000){ 28 29 #ifdef CUDAGRAPH 30 if(i == 0){ 31 accx_async_begin_capture(async1); // begin capture region 32 #endif 33 j = 0; 34 while(j < 1000){ 35 #pragma acc kernels async(async1) 36 for(k = 0; k < n; k++){ 37 c[k] = a[k] * b[k] * (2.0 + i); 38 } 39 40 #pragma acc kernels async(async1) 41 for(k = 0; k < n; k++){ 42 a[k] = a[k] - c[k] + 2.0 + a[k] * b[k] * i; 43 } 44 j = j + 1; 45 } 46 47 #ifdef CUDAGRAPH 48 accx_async_end_capture(async1, &pgraph); // end capture region 49 } 50 accx_graph_launch(&pgraph, async1); // always launch, then wait 51 #pragma acc wait(async1) 52 #endif 53 i = i + 1; 54 } 55 56 #ifdef CUDAGRAPH 57 accx_graph_delete(&pgraph); 58 #endif 59 60 #ifdef _OPENACC 61 } 62 #endif 63 64 for(i = 0; i < 10; i++) 65 printf("%f\n", a[i]); 66 67 return 0; 68 } |
サンプルコードは、CUDAカーネルが何度も呼ばれるような処理にしました。
_OPENACC
は、コンパイル時にオプション-acc
を付与すると定義されるマクロで、OpenACCを使う際に処理が実行されます。
赤い太字で示した処理は、CUDA Graphを使うために追加しました。
一行目でマクロCUDAGRAPH
が定義され、CUDA Graphに関連する処理が実行されます。
27行目のwhile
文によって、33行目から45行目の処理が何度も繰り返されます。
i=0の場合、ユーザーは33行目から45行目の処理をOpenACCの組み込み関数accx_async_begin_capture
とaccx_async_end_capture
で挟むことによって、それらの処理をキャプチャ領域として指定できます。
キャプチャ領域内のCPUやGPUのタスクはコンパイル時にグラフとして定義され、i=1以降はそのグラフ全体が一括して実行されます。
OpenACCの組み込み関数accx_graph_launch
の処理後は、OpenACCのwait
ディレクティブによりタスクの同期をする必要があります。
27行目のdo while
文を抜けた後、OpenACCの組み込み関数accx_graph_delete
によってグラフが破棄されます。
※上記jループ内のキャプチャ領域の処理がループ変数iに依存しないようなコードの場合は、キャプチャ領域をiループの外に出すことができ、wait
指示文は不要です。
コンパイル・実行した結果は以下です。
$ module load nvhpc/24.9
$ nvc -acc=gpu -Minfo=accel -o test-c test.c
main:
24, Generating copy(a[:n]) [if not already present]
Generating create(c[:n]) [if not already present]
Generating copyin(b[:n]) [if not already present]
36, Loop is parallelizable
Generating NVIDIA GPU code
36, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
41, Loop is parallelizable
Generating NVIDIA GPU code
41, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
$ export NVCOMPILER_ACC_USE_GRAPH=1
$ ./test-c
2.000000
2.000000
2.000000
2.000000
2.000000
2.000000
2.000000
2.000000
2.000000
2.000000
環境の都合上、NVIDIA HPC SDK 24.9を使いました。
以下の2つの設定により、OpenACCのCUDA Graphを使うことができます。
- コンパイル時に
-acc=gpu
のオプションを付ける - 実行時に環境変数
NVCOMPILER_ACC_USE_GRAPH
を1に設定する
OpenACCのCUDA Graph適用例(Fortran言語)
OpenACCによりGPU化されたFortran言語のプログラムに対して、CUDA Graphを適用したサンプルコードを以下に示します。
1 #define CUDAGRAPH 2 3 #ifdef CUDAGRAPH 4 module variables 5 use iso_c_binding 6 implicit none 7 8 type, bind(c) :: acc_graph_t 9 type(c_ptr) :: graph 10 type(c_ptr) :: graph_exec ! This is necessary. 11 end type acc_graph_t 12 13 end module variables 14 #endif 15 16 program main 17 18 #ifdef CUDAGRAPH 19 use openacc 20 use variables 21 #endif 22 23 implicit none 24 integer :: i, j, k 25 integer,parameter :: n = 5000 26 real(8) :: a(n), b(n), c(n) 27 integer,parameter :: async1 = 1 28 #ifdef CUDAGRAPH 29 type(acc_graph_t) :: pGraph ! CUDA Graph 30 #endif 31 32 a(:) = 2.0d0 33 b(:) = 0.5d0 34 35 !$acc data copy(a) copyin(b) create(c) 36 37 i = 0 38 do while(i < 1000) 39 40 #ifdef CUDAGRAPH 41 if(i == 0) then 42 call accx_async_begin_capture(async1) ! begin capture region 43 #endif 44 j = 0 45 do while(j < 1000) 46 !$acc kernels async(async1) 47 c(:) = a(:) * b(:) * (2.0d0 + i) 48 !$acc end kernels 49 50 !$acc kernels async(async1) 51 a(:) = a(:) - c(:) + 2.0d0 + a(:) * b(:) * i 52 !$acc end kernels 53 j = j + 1 54 end do 55 56 #ifdef CUDAGRAPH 57 call accx_async_end_capture(async1, pGraph) ! end capture region 58 endif 59 call accx_graph_launch(pGraph, async1) ! always launch, then wait 60 !$acc wait async(async1) 61 #endif 62 i = i + 1 63 end do 64 65 #ifdef CUDAGRAPH 66 call accx_graph_delete(pGraph) 67 #endif 68 69 !$acc end data 70 71 do i = 1, 10 72 print*, a(i) 73 end do 74 75 end program main |
OpenACCのCUDA GraphをFortran言語で使う場合は、FortranとC言語の相互運用性を提供するiso_c_binding
モジュールを使って、派生型acc_graph_t
を定義する必要があるようです。
それ以外は、C言語の場合とほぼ同様です。
コンパイル・実行した結果は以下です。
$ module load nvhpc/24.9
$ nvfortran -acc=gpu -gpu=cc80 -Minfo=accel -Mpreprocess -o test-f test.f90
35, Generating copy(a(:)) [if not already present]
Generating create(c(:)) [if not already present]
Generating copyin(b(:)) [if not already present]
47, Loop is parallelizable
Generating NVIDIA GPU code
47, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
50, Accelerator serial kernel generated
Generating NVIDIA GPU code
$ export NVCOMPILER_ACC_USE_GRAPH=1
$ ./test-f
2.000000000000000
2.000000000000000
2.000000000000000
2.000000000000000
2.000000000000000
2.000000000000000
2.000000000000000
2.000000000000000
2.000000000000000
2.000000000000000
OpenACCのCUDA Graphを使う際の注意点も、C言語の場合と同様です。
Fortranでマクロを使用するため、コンパイル時に-Mpreprocess
を付与しています。
50行目のOpenACC指示文について、コンパイラは半自動的に並列化せずGPU上で逐次実行していますが、今回は特に気にしないことにします。
性能評価
今回評価したCPUとGPUは以下になります。
製品名 | 理論演算性能 (倍精度浮動小数点演算) |
Clock rate [GHz] |
Memory [GiB] |
core | CUDA version | |
CPU | Intel(R) Xeon(R) Gold 5317 CPU @ 3.00GHz | 1.38 TFLOPS | 3.60 | 314 | 24 | |
GPU | NVIDIA A100 | 9.7 TFLOPS | 1.41 | 80 | CUDA Cores: 6,912 Tensor Cores: 432 |
12.2 |
timeコマンドで実行したところ、実行時間は以下のようになりました。
言語 | CUDA Graph なし(秒) | CUDA Graph あり(秒) | 性能比(倍) |
C | 5.6 | 3.3 | 1.7 |
Fortran | 4.6 | 2.6 | 1.8 |
OpenACCのCUDA Graphを使うことによって、C言語では約1.7倍、Fortran言語では約1.8倍性能が向上しました。
Nsight Systems*5を使って、プロファイリングを実施しました。
C言語の場合の実行方法は以下です。(Fortran言語も同様です。)
$ nsys profile -f true -t cuda,openacc -o report_test-c ./test-c
$ nsys-ui report_test-c.nsys-rep
Nsight SystemsのStats System View内にある、CUDA Summary(API/Kernels/MemOps)の抜粋は以下です。
- C言語
- CUDA Graphなし
Time(%) Total Time(秒) Instances Operation 56.7 5.08 2000000 cuLaunchKernel 21.8 1.95 1000000 _6test_c_main_41_gpu 21.5 1.93 1000000 _6test_c_main_36_gpu - CUDA Graphあり
Time(%) Total Time(秒) Instances Operation 87.2 2.7 1002 cuStreamSynchronize 12.4 0.38 1000 cuGraphLaunch 0.3 0.008 1 cuGraphInstantiate 0.1 0.002 1 cuGraphExecDestroy 0.1 0.002 2000 cuLaunchKernel 0.0 0.001 1 cuGraphDestroy
プロファイリングの結果について(C言語):
_6test_c_main_36_gpuと_6test_c_main_41_gpuは、36行目と41行目のループがそれぞれGPU上で並行実行されたことを示しています。
Instancesは必ずしもループ回数と一致するとは限らないようで深入りはしていませんが、「CUDA Graphなし」のInstancesはそれぞれ1000000であり、(iループ1000回)×(jループ1000回)と解釈できそうです。
一方、「CUDA Graphあり」では、_6test_c_main_36_gpu等の、GPU化されたループに関する情報がなくなり、cuGraphLaunch、cuGraphInstantiate、cuGraphExecDestroyとcuGraphDestroyの情報が出力されました。「CUDA Graphなし」ではcuLaunchKernelは2000000回呼ばれ、おそらく(iループ1000回)×(jループ1000回)×(GPU並列実行ループ2つ)と解釈できそうです。
OpenACCのCUDA Graphを使ったことにより、cuLaunchKernelのInstancesは1/1000になり、cuLaunchKernelの実行時間が約1/2628になりました。 - CUDA Graphなし
- Fortran言語
- CUDA Graphなし
Time(%) Total Time(秒) Instances Operation 63.0 5.5 2000000 cuLaunchKernel 21.9 1.9 1000000 main_47_gpu 15.1 1.3 1000000 main_50_gpu - CUDA Graphあり
Time(%) Total Time(秒) Instances Operation 99.3 2.3 1000 cuGraphLaunch 0.3 0.008 1 cuGraphInstantiate 0.1 0.003 1 cuGraphExecDestroy 0.1 0.002 2000 cuLaunchKernel 0.1 0.001 1 cuGraphDestroy 0.0 5×10-6 2 cuStreamSynchronize
プロファイリングの結果について(Fortran言語):
OpenACCのCUDA Graphを使ったことにより、cuLaunchKernelのInstancesが1/1000になり、cuLaunchKernelの実行時間が約1/3046になりました。 - CUDA Graphなし
C言語とFortran言語の比較:
「CUDA Graphなし」では、C言語とFortran言語いずれも、主要なコストはcuLaunchKernelで約6割のコストです。
一方、「CUDA Graphあり」では、C言語ではcuStreamSynchronizeの約2.7秒が主要なコストであるのに対し、Fortran言語ではcuGraphLaunchの約2.3秒が主要なコストで、C言語とFortran言語でコスト分布が異なります。
「CUDA Graphあり」のcuStreamSynchronizeのInstancesについて、C言語に対する値(1002)はFortran言語に対する値(2)と比べて1000大きく、これがC言語でcuStreamSynchronizeが主要なコストになっている原因であると考えられます。
まとめ
このコラムでは、OpenACCのCUDA Graphの利用方法を紹介いたしました。
今回は簡単なサンプルコードを作成し、CUDA Graphを使うことによって約1.7-1.8倍の性能向上を実現しました。
繰り返し実行されるカーネルローンチやメモリ操作のオーバーヘッドを削減し、さらなる性能向上を目指したい方は、
CUDA Graphの利用を検討してみてください。
参考文献
- https://docs.nvidia.com/hpc-sdk/archive/24.7/hpc-sdk-release-notes/index.html
- https://hpcworld.jp/nvsdk_releasenotes/version-24-7/
- https://docs.nvidia.com/hpc-sdk/compilers/hpc-compilers-user-guide/index.html#openacc-cuda-graphs
- https://developer.nvidia.com/ja-jp/blog/a-guide-to-cuda-graphs-in-gromacs-2023/
- https://developer.nvidia.com/nsight-systems