2024年10月31日 / 最終更新日時 : 2024年10月31日 hpcworldNEWS 技術情報

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になりました。

  • 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になりました。

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の利用を検討してみてください。

参考文献

  1. https://docs.nvidia.com/hpc-sdk/archive/24.7/hpc-sdk-release-notes/index.html
  2. https://hpcworld.jp/nvsdk_releasenotes/version-24-7/
  3. https://docs.nvidia.com/hpc-sdk/compilers/hpc-compilers-user-guide/index.html#openacc-cuda-graphs
  4. https://developer.nvidia.com/ja-jp/blog/a-guide-to-cuda-graphs-in-gromacs-2023/
  5. https://developer.nvidia.com/nsight-systems
カテゴリー
技術情報
タグ
技術情報

前の記事

OpenACCのworker並列を使ってプログラムを高速化する
2024年09月17日