導入

前回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コードを書くのは現状難しそうということがわかるかと思います。将来、改善されることを期待しています。