導入
前回Fortran + OpenACCの組み合わせでcodeをGPU化しました。NVIDIAのGPUならそれで問題ないのですが、例えばAMDなど他の会社のGPUではOpenACCが使えません。そういう場合に備えてOpenMPの使用を検討してみたのですが、今のところのコンパイラだと難しそうという話をします。今回の実行環境ではOpenMPを使用しますがGPUの種類はNVIDIAのGPU A100です。
問題の概要
前回同様、メモリを全部device側に載せてなるべくhostと通信しないという方向性にしたいのですが、コンパイラが変数を更新する度に毎回hostと通信してしまう問題が起きました。コンパイラのバージョンは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))
!$omp target enter data map(alloc:a)
do i=1,imax
a(i) = 1.0d0
enddo
!$omp target update 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 update from(a)
!$omp target exit data map(delete:a)
print *, a(1)
end program main
このコードを以下のようにコンパイルしてみます。
nvfortran -O2 -mp=gpu -Minfo=mp,accel,opt -gpu=cc80,mem:separate
-Minfo=mp,accelがついてるのでコンパイルの情報を出してくれます。
9, Generating target enter data map(create: a(:))
10, Recognized memory set idiom
13, Generating update 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 update from(a(:))
22, Generating target exit data map(delete: a(:)) finalize
上記の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
allocate(a(imax))
!$acc data create(a)
do i=1,imax
a(i) = 1.0d0
enddo
!$acc update device (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:
9, Generating create(a(:)) [if not already present]
13, Generating update device(a(:))
17, Loop is parallelizable
Generating NVIDIA GPU code
17, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
22, 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);
#pragma omp target enter data map(alloc:a[0:imax])
int nmax;
nmax =100;
// initialize
for (int i=0;i<imax;i++){
a[i] = 1.0e0;
}// end of i loop
#pragma omp target update to (a[0: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[0:imax])
printf("a=%e",a[0]);
}
コンパイルオプションは以下のものです。
nvc -O3 -mp=gpu -Minfo=mp,accel,opt -gpu=cc80,mem:separate -g -traceback
最適化のメッセージは以下です。
10, Generating target enter data map(create: a[:10000])
12, Recognized memory set idiom
14, Generating update to(a[:10000])
16, #omp target teams distribute
16, Generating "nvkernel_main_F1L16_2" GPU kernel
18, Loop parallelized across teams, schedule(static)
24, 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コードを書くのは現状難しそうということがわかるかと思います。将来、改善されることを期待しています。