OpenACC プログラミング入門
エヌビディア合同会社 CUDA エンジニア 村上 真奈
GPU コンピューティングとは?
アプリケーションコード
計算負荷が高い関数
GPU
その他の逐次処理
CPUで実行
5% of Code
+
CPU
GPU アクセラレーションの実現方法
アプリケーション
GPU
ライブラリ
ライブラリを呼び出すだけ
簡単に高速化を実現
OpenACC
ディレクティブ
CUDA
既存コードにディレクティブ
を挿入して高速化
重要なコードをCUDAで記述
最も自由度が高い
main()
{
<serial code>
#pragma acc kernels
イリノイ大学
MRI画像再構成
//自動的にGPUで実行
{
OpenACC
<parallel code>
}
}
70倍 高速化
シンプル | 強力 | 柔軟
理研
NICAM- 大気現象シミュレーション
OpenACC利用者
8000人+
5% のコード変更で
7~8倍 高速化
OpenACC ディレクティブ
CPU
GPU
コンパイラへシンプルなヒント
Program myscience
... serial code ...
!$acc kernels
do k = 1,n1
do i = 1,n2
... parallel code ...
enddo
enddo
!$acc end kernels
...
End Program myscience
Fortran または C言語の
オリジナルコード
コンパイラがコードを並列化
コンパイラへの
OpenACC
ヒント
並列部はGPUで
逐次処理はCPUで動作
OpenMP と OpenACC の比較
OpenMP
CPU
main() {
double pi = 0.0; long i;
#pragma omp parallel for reduction(+:pi)
for (i=0; i<N; i++)
{
double t = (double)((i+0.05)/N);
pi += 4.0/(1.0+t*t);
}
CPUコアに
printf(“pi = %f¥n”, pi/N);
計算処理を分散
}
OpenACC
CPU
GPU
main() {
double pi = 0.0; long i;
#pragma acc kernels
for (i=0; i<N; i++)
{
double t = (double)((i+0.05)/N);
pi += 4.0/(1.0+t*t);
}
printf(“pi = %f¥n”, pi/N);
}
GPUコアに
計算処理を分散
new OpenACC ツールキット(2015/07~)
大学関係者の方は無償で使用可能
• Linux x64/ノードロック・ライセンス
• 大学関係者以外でも90日間無料トライアル可能
下記のサイトからOpenACC ツールキットをダウンロード
https://developer.nvidia.com/openacc
OpenACCを始める為のコンパイラ・解析
ツールなど一式が簡単にインストール
OpenACC ツールキット
PGIコンパイラ
Free OpenACC compiler for academia
NVProfプロファイラ
Easily find where to add compiler directives
GPUウィザード
Identify which GPU libraries can jumpstart code
サンプルコード
http://developer.nvidia.com/openacc
Learn from examples of real-world algorithms
ドキュメント
Quick start guide, Best practices, Forums
OpenACC を始めましょう
OpenACC ディレクティブ構文
C/C++
#pragma acc 指示行 [節[,]節] …]
{ structured block }
Fortran
!$acc 指示行 [節[,]節] …]
{ structured block }
!$acc end directive
例: SAXPY (Y=A*X+Y)
OpenACC
OpenMP
void saxpy(int n,
float a,
float *x,
float *restrict y)
{
#pragma omp parallel for
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}
void saxpy(int n,
float a,
float *x,
float *restrict y)
{
#pragma acc parallel
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}
...
saxpy(N, 3.0, x, y);
...
...
saxpy(N, 3.0, x, y);
...
 omp  acc
 データの移動
OpenACC 構文: parallel 指示行
• parallel : 並列に実行される領域を指示行で指定
#pragma acc parallel
for(int i=0;i<n;i++){
a[i] = 0.0;
b[i] = 1.0;
c[i] = 2.0;
}
kernel 1
Kernel(カーネル):
GPU上で実行される
関数
OpenACC 構文: kernels 指示行
• kernels : 複数のカーネルを作成
#pragma acc kernels
for(int i=0;i<n;i++){
a[i] = 0.0;
b[i] = 1.0;
c[i] = 2.0;
}
#pragma acc kernels
for(int i=0;i<n;i++){
a[i] = b[i] + c[i];
}
kernel 1
Kernel(カーネル):
GPU上で実行される
関数
kernel 2
簡単にコンパイル
OpenMP / OpenACC
void saxpy(int n, float a,
float *x,
float *restrict y)
$ pgcc –acc {–ta=nvidia –Minfo=accel saxpy.c
#pragma acc parallel copy(y[:n]) copyin(x[:n])
saxpy:
#pragma
omp parallel for
16, Generating
present_or_copy(y[:n])
for present_or_copyin(x[:n])
(int i = 0; i < n; ++i)
Generating
+= code
a*x[i];
Generatingy[i]
Tesla
} parallelizable
19, Loop is
Accelerator kernel generated
...
19, #pragma
acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
saxpy(N, 3.0, x, y);
...
例:ヤコビ反復法(アルゴリズム)
A(i,j+1)
while ( error > tol ) {
error = 0.0;
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]));
}
}
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
A(i-1,j)
A(i,j) A(i+1,j)
A(i,j-1)
並列領域 (OpenACC)
while ( error > tol ) {
error = 0.0;
#pragma acc kernels
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma acc kernels
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
 Parallels と Kernels
— 並列領域を指示
 Parallels
— 並列実行スタート
 Kernels
— 複数のカーネル
簡単に解析(nvprof)
OpenMP / OpenACC
void saxpy(int n, float a,
float *x,
float *restrict y)
{
$ pgcc -Minfo -acc
saxpy.c
$ nvprof ./a.out
#pragma acc kernels copy(y[:n]) copyin(x[:n])
saxpy:
==10302==
NVPROF
ispresent_or_copy(y[:n])
profiling
process
#pragma
omp parallel
for 10302, command: ./a.out
16, Generating
==10302==
Profiling
application:
for present_or_copyin(x[:n])
(int
i = 0; i ./a.out
< n; ++i)
Generating
==10302==
Profiling
result:
y[i]
+= code
a*x[i];
Generating
Tesla
Time(%)
Calls
Avg
Min
Max Name
} parallelizable
19, Loop Time
is
62.95% Accelerator
3.0358ms
1.5179ms 1.5172ms 1.5186ms [CUDA memcpy HtoD]
kernel2generated
31.48% 19,
1.5181ms
1 1.5181ms
1.5181ms/*1.5181ms
[CUDA
memcpy DtoH]
...
#pragma
acc loop
gang, vector(128)
blockIdx.x
threadIdx.x
*/
5.56% 268.31us
268.31us 268.31us saxpy_19_gpu
saxpy(N, 3.0,1 x,268.31us
y);
...
ボトルネックの可視化(nvvp)
1 サイクル
GPU
カーネル
GPU
カーネル
過剰なメモリ転送
while ( error > tol ) {
error = 0.0;
#pragma acc kernels
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma acc kernels
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
配列Aへメモリ転送
(CPU->GPU)
配列Anewへメモリ転送
(GPU->CPU)
配列Anewへメモリ転送
(CPU->GPU)
配列Aへメモリ転送
(GPU->CPU)
メモリ転送のタイミングを制御
#pragma acc data pcopy(A, Anew)
while ( error > tol ) {
error = 0.0;
#pragma acc kernels pcopy(Anew[:][:]) pcopyin(A[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma acc kernels pcopy(A[:][:]) pcopyin(Anew[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
データ制御ディレクティブ
copyin (CPUGPU)
copyout (CPUGPU)
copy
create
present
pcopyin
pcopyout
pcopy
pcreate
性能比較
Time(sec)
Speedup
CPU1 OpenMP thread
26.186
--
OpenACC(最適化前)
13.638
1.92x
OpenACC(最適化後)
0.773
33.88x
Ubuntu 14.04LTS 64bit
CPU:Intel Xeon CPU E5-1603 [email protected] x 4
GPU:Tesla K40c
Thank you