OpenACC ディレクティブによるプログラミング
12章 OpenACC と CUDA C/Fortran との相互運用性
OpenACCとの相互運用性
OpenACC for NVIDIA device は、CUDA の他の言語環境と相互運用が可能なプログラミングモデルである。例えば、CUDA C で記述されたルーチンと OpenACC で記述されたルーチンは、そこで使用されているデータの受け渡しが可能であり、デバイス上のメモリを同じ土俵で利用可能である。従って、このようなルーチンを組み合わせて一つの実行モジュールを作成できる。同様に、OpenACC Fortran と CUDA Fortran の組み合わせや CUBLAS、CUFFT 等のライブラリも、OpenACC の環境から呼び出すことが可能である。また、CUDA C や CUDA Fortran で記述されたプログラムの中で、まだ CUDA kernel に書き下すことができていない（つまらない）部分をより簡単な OpenACC のディレクティブでプログラミングを行うことにより、cuda kernel 化の促進による並列化率を向上させることが出来る。こうした用途で、より積極的に相互運用するべきであると思う。ここでは、非常に簡単なプログラム例を使って、それぞれの組み合わせにおけるプログラミングの方法を説明する。
相互運用が可能と言うことは、それぞれのプログラミングモデル、言語体系で、共通のデバイス側のメモリ、すなわち「デバイスポインタ」を利用可能であると言うことである。しかも、プログラム上の同じデバイス変数名、配列名を用いて、プログラミングすることが出来ると言うことでもある。このように、OpenACC for NVIDIA Device と CUDA C、CUDA Fortran 環境は、デバイス上のポインタを共通に利用出来るお膳立てが出来ている。
相互運用において、配列、変数のデバイスメモリのマネージメントをどちら側のプログラミングモデルで行うかによって、プログラミングの対応は異なるが、ここでは、OpenACC 側で配列の宣言を行った場合、逆に CUDA C あるいは CUDA Fortran 側でデバイス配列の宣言を行った場合の方法を以下に例示する。
なお、以下の例は、PGI 13.10 バージョンを使用しての結果である。使用しているバージョンを確かめたい場合は、以下のように -V コマンド・オプションを指定する。
$ pgcc -V pgcc 13.10-0 64-bit target on x86-64 Linux -tp nehalem The Portland Group - PGI Compilers and Tools Copyright (c) 2013, NVIDIA CORPORATION. All rights reserved. $ pgfortran -V pgfortran 13.10-0 64-bit target on x86-64 Linux -tp nehalem The Portland Group - PGI Compilers and Tools Copyright (c) 2013, NVIDIA CORPORATION. All rights reserved.
OpenACC と CUDA C の相互運用 (1)
最初に、OpenACC C と CUDA C の相互運用の方法について説明する。以下は、非常に簡単なベクトルの和を求めるためのプログラムである。相互運用を説明するために、メインプログラムとサブルーチンの二つの構成とした。各々のルーチンを別々のプログラミングモデルを使ってコンパイルし、二つのオブジェクトをリンクして executable を作成すると言う形をとる。
以下の openacc_main.c は、手続き vecadd を呼ぶためのメインプログラムである。C 言語のプログラムに OpenACC のディレクティブが挿入されている。#pragma acc data と言うデータ領域の指定がなされている中に、vecadd(n, a, b, c); と言う CUDA C で記述されたルーチンを call している。これをメモリマネージメントの観点で見ると、OpenACC の世界で宣言された変数、配列体系を使用してデバイス側のメモリへのインターフェースを行うと言うことになる。すなわち、OpenACC の内部的に使用している「デバイスポインタ」を CUDA C のルーチンに受け渡すと言うことが必要となる。ここでお気づきの方もいるかと思うが、OpenACC では、ホスト側の配列とデバイス側の配列と言ったような二つのメモリ上の配列を明示的に宣言する必要がない。CUDA C や CUDA Fortran では、Unified Memory 機能を使用しない限り、二つのメモリ上の割付をプログラム上で宣言して使用する。これに対して、OpenACC の利点は、ユーザ視点では配列や変数がホスト側、デバイス側を問わず、一元化して扱えると言うことである。しかし、内部的にはもちろん、ホスト側とデバイス側の双方ポインタをハンドリングしているのであるが、これはユーザには見えない。こうした場合、OpenACC 環境上で当該変数等のデバイス側のポインタを扱うための仕組みが必要となる。OpenACC では、このためのディレクティブが用意されている。host_data ディレクティブである。その clause である use_device の引数は、当該変数あるいは配列のデバイスポインタを低レベル API(CUDA Cあるいは CUDA Fortran) に渡すために使用する。
以下の例では、use_device(a,b,c) としているので、a、b、c 配列のデバイス側のポインタを vecadd ルーチンへ渡すと言うことを指示している。この場合、当然ながら vecadd ルーチンは低レベルの CUDA C 言語で記述されたものでなければならない。 CUDA C 言語プログラムは、直接デバイスポインタを使ってハンドリングされているため、こうした対応となる。
ソースプログラム： openacc_main.c、vecadd.cu
openacc_main.c
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
extern void vecadd(int,float*,float*,float*);
int main()
{
float *a, *b, *c;
int i;
int n;
n = 1024;
a = (float*)malloc(n*sizeof(float));
b = (float*)malloc(n*sizeof(float));
c = (float*)malloc(n*sizeof(float));
acc data 領域の構築を行う
#pragma acc data create(a[0:n],b[0:n]) copyout(c[0:n])
{
#pragma acc kernels acc kernels でデバイス上で a, b 配列初期化
for( i = 0; i < n; i++)
{
a[i] =(float)i;
b[i] =(float)i*2.0f;
} a,b,c はデバイスポインタで渡すことを指示する
#pragma acc host_data use_device(a,b,c)
{
vecadd(n, a, b, c);
}
}
for( i = 0; i < 10; i++)
fprintf(stdout, "c[%d] = %f\n",i,c[i]);
return 0;
}
vecadd.cu
__global__ void vecadd_kernel(int n, float *a, float *b, float *c) { // cuda c カーネル using device pointers int i = blockDim.x * blockIdx.x + threadIdx.x; if ( i < n ) c[i] = a[i] + b[i]; } extern "C" void vecadd(int n ,float *a, float *b, float *c) { dim3 griddim, blockdim; blockdim = dim3(128,1,1); griddim = dim3(n/blockdim.x,1,1); vecadd_kernel<<<griddim,blockdim>>>(n,a,b,c); }
コンパイルの方法は、以下のようになる。openacc_main.c は、-acc オプションを付して pgcc コンパイラでそのオブジェクト(openacc_main.o、Windows上では openacc_main.obj) を作成する。次に、CUDA C プログラムは、NVIDIA CUDA C コンパイラ nvcc でオブジェクト(vecadd.o、Windows 上では vecadd.obj) を作成する。最後の、これらのオブジェクトを pgcc コンパイラでリンク結合する。その際のオプションは、-acc と -Mcuda は必須となる。これらは、OpenACC に必要なシステム・ライブラリと cuda に関するシステムライブラリをリンクする際に必要なオプションである。
[kato@photon29 ACC+cudaC]$ pgcc -acc -Minfo -O2 -c openacc_main.c
main:
18, Generating create(b[0:n])
Generating create(a[0:n])
Generating copyout(c[0:n])
20, Generating present_or_create(b[0:n])
Generating present_or_create(a[0:n])
Generating NVIDIA code
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
Generating compute capability 3.0 binary
21, Complex loop carried dependence of '*(a)' prevents parallelization
Loop carried dependence of '*(b)' prevents parallelization
Loop carried backward dependence of '*(b)' prevents vectorization
Accelerator scalar kernel generated
Loop not vectorized: data dependency
Loop unrolled 8 times
32, Loop not vectorized/parallelized: contains call
[kato@photon29 ACC+cudaC]$ nvcc -c -O2 vecadd.cu
[kato@photon29 ACC+cudaC]$ pgcc -acc -Mcuda openacc_main.o vecadd.o
[kato@photon29 ACC+cudaC]$ ./a.out
c[0] = 0.000000
c[1] = 3.000000
c[2] = 6.000000
c[3] = 9.000000
c[4] = 12.000000
c[5] = 15.000000
c[6] = 18.000000
c[7] = 21.000000
c[8] = 24.000000
c[9] = 27.000000
OpenACC と CUDA C の相互運用 (2)
上述の例とは全く逆の例である。すなわち、メインプログラム cuda_main.cu は CUDA C でデバイス上の変数、配列は宣言、アロケートされている。Low レベル CUDA C の デバイスポインタの世界でマネージメントされているプログラムとなっている。ここに、OpenACC で作成されたサブルーチン(手続き）をコールする構成を考える。CUDA C の cuda_main.cu プログラムの方は変更はない。一方、OpenACC による openacc_vecadd.c は、kernels の中で使用される変数、配列が Low レベルのデバイスポインタであることをコンパイラに指示するために、deviceptr clause を使用している。すなわち、この deviceptr clause を使うと、この OpenACC kernels の領域では、引数で渡されてきたデバイスポインタを使用するコードが作られることになる。
ソースプログラム： cuda_main.cu、openacc_vecadd.c
cuda_main.cu
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <cuda_runtime.h>
extern "C" void vecadd(int,float*,float*,float*);
extern "C" void set_value (int, float*, float*);
int main()
{
float *a, *b, *c;
float *result;
int i;
int n;
n = 1024;
result = (float*)malloc(n*sizeof(float));
cudaMalloc((void**)&a,(size_t)n*sizeof(float));
cudaMalloc((void**)&b,(size_t)n*sizeof(float));
cudaMalloc((void**)&c,(size_t)n*sizeof(float));
set_value(n, a, b); OpenACCで作成されたルーチンを呼ぶ
vecadd(n, a, b, c); OpenACCで作成されたルーチンを呼ぶ
cudaMemcpy (result, c, sizeof(float)*n, cudaMemcpyDeviceToHost);
for( i = 0; i < 10; i++)
fprintf(stdout, "c[%d] = %f\n", i, result[i]);
return 0;
}
openacc_vecadd.c
void vecadd(int n, float* restrict a, float* restrict b, float* restrict c)
{
#pragma acc kernels deviceptr(a,b,c) device pointer として a, b, c を使うことを指示
{
for(int i=0; i<n; i++)
{
c[i] = a[i] + b[i];
}
}
}
void set_value(int n, float * restrict aa, float * restrict bb)
{
#pragma acc kernels deviceptr(aa,bb) device pointer として a1, bb を使うことを指示
{
for(int i=0; i<n; i++)
{
aa[i] = (float)i ;
bb[i] = (float)i*2.0f;
}
}
}
コンパイルの方法は、以下のようになる。cuda_main.cu は nvcc でコンパイルする。次に、openacc_vecadd.c は、pgcc コンパイラ でオブジェクトを作成する。最後に、これらのオブジェクトを pgcc コンパイラでリンク結合する。その際のオプションは、-acc と -Mcuda は必須となる。
[kato@photon29 cudaC+ACC]$ nvcc -c -O2 cuda_main.cu
[kato@photon29 cudaC+ACC]$ pgcc -c -acc -O2 -Minfo openacc_vecadd.c
vecadd:
3, Generating NVIDIA code
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
Generating compute capability 3.0 binary
5, Loop is parallelizable
Accelerator kernel generated
5, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
Generated 4 alternate versions of the loop
Generated vector sse code for the loop
Generated 2 prefetch instructions for the loop
set_value:
13, Generating NVIDIA code
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
Generating compute capability 3.0 binary
15, Loop is parallelizable
Accelerator kernel generated
15, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
Generated 5 alternate versions of the loop
Generated vector sse code for the loop
[kato@photon29 cudaC+ACC]$ pgcc -acc -Mcuda cuda_main.o openacc_vecadd.o
[kato@photon29 cudaC+ACC]$ ./a.out
c[0] = 0.000000
c[1] = 3.000000
c[2] = 6.000000
c[3] = 9.000000
c[4] = 12.000000
c[5] = 15.000000
c[6] = 18.000000
c[7] = 21.000000
c[8] = 24.000000
c[9] = 27.000000
OpenACC と CUDA Fortran の相互運用 (1)
openacc_main.f90 と言うメインプログラムが OpenACC で構築され、OpenACC 配下のメモリマネージメントを採用している。その中から、Low レベル CUDA Fortran のルーチンが call される構図となっている。Fortran の場合も上記で説明した C プログラムの場合と同じような考え方となる。CUDA Fortran で記述されたルーチン vecadd.cuf をコールする際に渡す実引数がデバイスポインタであることをコンパイルに伝えるために、host_data use_device(a,b,c) を使う。これが厳密な使い方となるが、実は Fortran の場合はこのディレクティブを省略しても良い。このプログラムは、PGI CUDA Fortran と PGI OpenACC Fortran を使う。実はこれらは、pgfortran と言う一つのコンパイラであり、-Mcuda あるいは -acc オプションを指定することにより当該機能が有効化される。同じ PGI Fortran コンパイラの配下では、OpenACC コンパイラ機能が vecadd_wrapper(n, a, b, c) の引数が、デバイスポインタを利用するかどうか自動的に判断するため、host_data ディレクティブは省略できる（Fortran Module ブロックを使用しているため、引数の interface がコンパイラが事前に理解できることも奏功している）。ただ、私は Fortran の場合でも host_data ディレクティブをきちんと指定することを推奨する。
ソースプログラム： openacc_main.f90、vecadd.cuf
openacc_main.f90
program main use vecadd_module integer(4), parameter :: n = 2**20 real(4), allocatable, dimension(:) :: a, b, c allocate(a(n)) allocate(b(n)) allocate(c(n)) !$acc data create(a,b) copyout(c) acc data 領域の構築を行う !$acc kernels do i = 1, n a(i) = i b(i) = i * 2.0 end do !$acc end kernels a,b,c はデバイスポインタで渡すことを指示する !$acc host_data use_device(a,b,c) call vecadd_wrapper(n, a, b, c) !$acc end host_data !$acc end data print '(1x, i10, F15.0)', (i, c(i), i=n-10,n) end program
vecadd.cuf
module vecadd_module
contains
subroutine vecadd_wrapper(n, a, b, c)
use cudafor
real, device :: a(:), b(:), c(:)
integer :: n
type(dim3) :: blocks
type(dim3) :: threads
blocks = dim3(n/128,1,1)
threads = dim3(128, 1, 1)
call vecadd_kernel<<<blocks, threads>>>(n, a, b, c)
end subroutine
attributes(global) subroutine vecadd_kernel (n, a, b, c)
real(4), device :: a(:), b(:), c(:)
integer(4),value :: n,i
i = (blockIdx%x-1)*blockDim%x + threadIdx%x
if (i<=n) c(i) = a(i) + b(i)
end subroutine
end module vecadd_module
コンパイルの方法は、以下のようになる。vecadd.cuf は pgfortran の -Mcuda オプションを付してコンパイルする。次に、メインプログラムは、pgfortran コンパイラ -acc オプションを付してオブジェクトを作成する。最後に、これらのオブジェクトを pgfortran コンパイラでリンク結合する。その際のオプションは、-acc と -Mcuda は必須となる。
[kato@photon29 ACC+cudaF]$ pgfortran -Mcuda -O2 -Minfo -c vecadd.cuf
[kato@photon29 ACC+cudaF]$ pgfortran -acc -O2 -Minfo -c openacc_main.f90
main:
10, Generating copyout(c(:))
Generating create(b(:))
Generating create(a(:))
12, Generating present_or_create(b(:))
Generating present_or_create(a(:))
Generating NVIDIA code
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
Generating compute capability 3.0 binary
13, Loop is parallelizable
Accelerator kernel generated
13, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
Generated an alternate version of the loop
Generated vector sse code for the loop
27, Loop not vectorized/parallelized: contains call
[kato@photon29 ACC+cudaF]$ pgfortran -acc -Mcuda openacc_main.o vecadd.o
[kato@photon29 ACC+cudaF]$ ./a.out
1048566 3145698.
1048567 3145701.
1048568 3145704.
1048569 3145707.
1048570 3145710.
1048571 3145713.
1048572 3145716.
1048573 3145719.
1048574 3145722.
1048575 3145725.
1048576 3145728.
OpenACC と CUDA Fortran の相互運用 (2)
以下のプログラムは、CUDA Fortran プログラムが配列のメモリマネージメントを行っている例である。CUDA Fortran ではデバイスの配列は宣言文で指定できる。この「デバイス配列」をそのまま OpenACC のループ内で利用させるには、data、parallel、kernels 構文の data clause である deviceptr を使用する。この引数にデバイス属性を持つ配列、変数を指定することにより、OpenACC カーネル内部で直接、デバイス配列等が使用できる。
ソースプログラム： cuf_main.cuf、vecadd_openacc.f90
cuf_main.cuf
program main use submod integer(4), parameter :: n = 2**20 real(4), device, allocatable, dimension(:) :: a_d, b_d, c_d デバイス配列の宣言 real(4), allocatable, dimension(:) :: c allocate(a_d(n)) デバイス配列の割付を行う allocate(b_d(n)) allocate(c_d(n)) allocate(c (n)) !$cuf kernel do(1) <<< *, * >>> CUDA Fortran 拡張機能、簡易kernelの作成 do i = 1, n a_d(i) = i b_d(i) = i * 2.0 end do call vecadd(n, a_d, b_d, c_d) デバイス配列 a_d, b_d, c_d を直接渡している c = c_d print '(1x, i10, F15.0)', (i, c(i), i=n-10,n) end program
vecadd_openacc.f90
module submod このルーチンは、CUDA Fortran 構文と OpenACC 構文のミックスである contains subroutine vecadd (n, a, b, c) real(4), device :: a(:), b(:), c(:) CUDA Fortranでデバイス配列宣言する integer(4) :: n !$acc kernels deviceptr(a,b,c) OpenACC に、引数がデバイス配列であることを伝える do i = 1, n c(i) = a(i) + b(i) end do !$acc end kernels end subroutine end module submod
vecadd_openacc.f90 は、デバイス配列の宣言を行っているため、CUDA Fortran 構文を使用している。また、ループに対して OpenACC kernels の適用を行っているため OpenACC 構文も使用している。すなわち、このサブルーチンは、CUDA Fortran + OpenACC の両機能を使用するものである。従って、コンパイルする際のオプションは -Mcuda -acc の二つのオプションを指定してオブジェクトを作成する必要がある。cuf_main.cuf は、CUDA Fortran プログラムであるため、-Mcuda オプションだけで良い。作成された二つのオブジェクトは、pgfortran コンパイラでリンク結合される。その際のオプションは、-acc と -Mcuda は必須となる。
[kato@photon29 cudaF+ACC]$ pgfortran -Mcuda -acc -O2 -Minfo -c vecadd_openacc.f90
vecadd:
9, Loop is parallelizable
Accelerator kernel generated
9, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
[kato@photon29 cudaF+ACC]$ pgfortran -Mcuda -O2 -Minfo -c cuf_main.cuf
main:
15, CUDA kernel generated
15, !$cuf kernel do <<< (*), (128) >>>
24, Loop not vectorized/parallelized: contains call
[kato@photon29 cudaF+ACC]$ pgfortran -Mcuda -acc cuf_main.o vecadd_openacc.o
[kato@photon29 cudaF+ACC]$ ./a.out
1048566 3145698.
1048567 3145701.
1048568 3145704.
1048569 3145707.
1048570 3145710.
1048571 3145713.
1048572 3145716.
1048573 3145719.
1048574 3145722.
1048575 3145725.
1048576 3145728.
[Reference]
- MJeff Larkin, NVIDIA, OpenACC Interoperability Tricks, PGInsider Newsletters
