導入
前回Fortran + OpenACCの組み合わせでcodeをGPU化しました。NVIDIAのGPUならそれで問題ないのですが、例えばAMDなど他の会社のGPUではOpenACCが使えません。そういう場合に備えてOpenMPの使用を検討してみたのですが、今のところのコンパイラだと難しそうという話をします。今回の実行環境ではOpenMPを使用しますがGPUの種類はNVIDIAのGPU A100です。
問題の概要
前回同様、メモリを全部device側に載せてなるべくホストと通信しないという方向性にしたいのですが、コンパイラが変数を更新する度に毎回ホストと通信してしまう問題が起きました。コンパイラのバージョンはnvhpc/25.7です。将来的にあるいは他のコンパイラであれば解決する可能性もあります（そう願っています）が、今はこういう問題があるということを報告しておきます。ちなみにFortran + OpenACCなら大丈夫ですし、C + OpenMPでも大丈夫でした。
問題の詳細
Fortran + OpenMPの場合
この問題はすごく簡単なコードで検証できます。n-loopの中ではホストと通信せずにデバイスだけで処理して欲しいです。
program main
implicit none
integer,parameter::imax=10000
real(8),dimension(:),allocatable::a
integer,parameter::nmax=10
integer::i,n
allocate(a(imax))
do i=1,imax
a(i) = 1.0d0
enddo
!$omp target enter data map(to:a)
do n=1,nmax
!$omp target teams distribute parallel do
do i=1,imax
a(i) = a(i) + 1.0d0
enddo
enddo
!$omp target exit data map(from:a)
print *, a(1)
end program main
このコードを以下のようにコンパイルしてみます。
nvfortran -O2 -mp=gpu -Minfo=mp,accel,opt -gpu=cc80,mem:separate
-Minfo=mp,accelがついてるのでコンパイルの情報を出してくれます。
main:
10, Recognized memory set idiom
13, Generating target enter data map(to: a(:))
15, !$omp target teams distribute parallel do
15, Generating "nvkernel_MAIN__F1L15_2" GPU kernel
15, Generating implicit map(tofrom:a(:))
21, Generating target exit data map(from: a(:))
上記の15, Generating implicit map(tofrom:a(:)) はhostからdeviceへaのデータを送信して、i-loopが終わったらdeviceからhostへaのデータを送信することを意味しています。これはする必要のない無駄な通信で本来はdevice上で処理が完結しているはずです。gpu=cc80,mem:separateはデータの転送の処理をなるべくしないというオプションなのでこれをつけても通信してしまうと、これに対処するのはなかなか難しいです。
下記のOpenACCの場合、presentという節を使ったりして、転送する必要がない旨をより明確化できるのですが、OpenMPでは今のところそれはできませんでした。
Fortran + OpenACCの場合
OpenACCで同じことをしてみます。
program main
implicit none
integer,parameter:: imax=10000
real(8),dimension(:),allocatable:: a
integer,parameter:: nmax=10
integer::i,n
do i=1,imax
a(i) = 1.0d0
enddo
!$acc data create(a)
do n=1,nmax
!$acc kernels
!$acc loop independent
do i=1,imax
a(i) = a(i) + 1.0d0
enddo
!$acc end kernels
enddo
!$acc update host (a)
!$acc end data
print *, a(1)
end program main
このコードを以下のようにコンパイルしてみます。
nvfortran -O2 -acc -Minfo=accel -gpu=cc80,mem:separate
メッセージは以下です。
main:
11, Generating create(a(:)) [if not already present]
15, Loop is parallelizable
Generating NVIDIA GPU code
15, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
20, Generating update self(a(:))
OpenMPであったGenerating implicit map(tofrom:a(:)) というのがなくなっています。
C + OpenMPの場合
実はC言語のコンパイラならOpenMPでもちゃんと最適化されます。
#include <stdio.h>
#include <stdlib.h>
int main(void){
#define imax 10000
double *a;
a = (double*) malloc(sizeof(double)*imax);
int nmax;
nmax =100;
// initialize
for (int i=0;i<imax;i++){
a[i] = 1.0e0;
}// end of i loop
#pragma omp target enter data map(to:a[imax])
for(int n=0;n<nmax;n++){ // iteration
#pragma omp target teams distribute
for (int i=0;i<imax;i++){
a[i] = a[i] + 1.0e0;
}// end of i loop
}// end of n loop
#pragma omp target exit data map(from:a[imax])
printf("a=%e",a[0]);
}
コンパイルオプションは以下のものです。
nvc -O3 -mp=gpu -Minfo=mp,accel,opt -gpu=cc80,mem:separate -g -traceback
最適化のメッセージは以下です。
11, Recognized memory set idiom
13, Generating target enter data map(to: a[10000])
15, #omp target teams distribute
15, Generating "nvkernel_main_F1L15_2" GPU kernel
17, Loop parallelized across teams, schedule(static)
23, Generating target exit data map(from: a[10000])
ただし、実はこの問題は結構難してくて、以下のようにOpenMP 5の文法であるloop order(concurrent)を使うとGenerating implicit map(tofrom:a[:10000])がでてしまいます。C言語でもOpenMPの使用は少し難しさがあるようです。
#pragma omp target enter data map(to:a[imax])
for(int n=0;n<nmax;n++){ // iteration
#pragma omp target loop order(concurrent)
for (int i=0;i<imax;i++){
a[i] = a[i] + 1.0e0;
}// end of i loop
}// end of n loop
まとめ
上記のことからFortran とOpenMPの組み合わせでは指示文でGPUコードを書くのは現状難しそうということがわかるかと思います。将来、改善されることを期待しています。