PGIアクセラレータコンパイラ > OpenACCサイト・トップ > インデックス

OpenACC ディレクティブによるプログラミング

14章　C 構造体、C++ クラス、Fortran 派生型のデータ管理

1. 構造体、派生型変数の deep copy データ転送

　PGI 14.4 以降の OpenACC 機能では、構造体、クラス、派生型と称される「集合型データ」変数のデータ転送は、構成しているメンバ変数も含めてのデータ転送(data clause)指示が必要とされる。しかしながら、こうした「集合型データ」のメンバを一つひとつ管理してデータ転送指示を与えることは、プログラムのバグの遠因となるし煩雑な行為である。しかし、現在の OpenACC の使用では、こうしたデータ移動を手動で指示する必要がある。「集合型データ」変数のメンバ変数も含めてのデータ転送を一括でできることを一般に deep copy することができると表現する。その第一歩として、「集合型データ」変数名（派生型名、構造体名）だけのコピー操作によって、その配下の階層的に配置されたメンバも含めての full deep copy が行うことができれば、より簡易なディレクティブ指示でプログラムを書くことができる。OpenACC.org においては、Complex Data Management Tech Report (TR-14-1) Nov 2014や、Deep Copy Attach and Detach (TR-16-1) で Deep Copy 機能の仕様について詰めているところであるが、順次、より簡単に deep copy ができるようになることを期待している。

　2017年 8 月、Fortran 言語に関しては、PGI 17.7 において、派生型変数の full deep copy（Betaサポート開始）が可能となった(C/C++ 言語に関しては、今後の OpenACC の正式な規格が定義された後の実装となると思われる）。但し、メンバ要素の部分的なデータ転送についてはまだ実装されていないが、Full deep copy の利用により、Fortran の派生型変数の取扱方が容易くなったと言える。以下に、PGI 17.7 deep copy の利用例を説明することにする。なお、以下の例は、PGI 17.7 バージョンを使用しての結果である。使用しているバージョンを確かめたい場合は -V コマンド・オプションを指定する。

　「集合型データ」変数のデータ転送は、OpenACC 仕様におけるデフォルト動作では、上述の通り、そのメンバ要素も明示的にデータ移動を指示する必要がある。例えば、「集合型データ」変数名だけを acc data copy(変数名）としただけでは、構造体、派生型変数自体のデバイス側の実態（this pointer) は生成されるが、そのメンバを構成している配列要素 (array of pointer) については、ホスト側の配列要素のポインタをコピーするだけなので、デバイス側の実体となる配列要素のポインタが生成されない（この状態を一般に、shallow copy の状態という）。これを実現するにはデバイス側にメンバを構成している当該配列要素のコピーを acc enter data copyin ディレクティブを使って「明示的」に行う必要がある（手動 deep copy)。 OpenACC 2.5 仕様の段階では、C/C++、Fortran 上では、こういった操作が必要となる。ただ、PGI Fortran に限っては、上述の通り、full deep copy 機能が提供されたため、-ta=tesla,deepcopy オプションを付することで、メンバ配列要素の明示的なデータ移動指示なして、deep copy が行われる。なお、構造体、派生型変数のメンバのうち、スカラ変数については、shallow copy の段階で、デバイス側にも自動的に attach されるため、明示的な deep copy は必要とされない。「attach」 とは、デバイスのアドレスへデータメンバのポインタが実際にセットされるということを指す。

　構造体、派生型変数データを管理する方法は以下のものがある。

明示的にメンバ要素のデータ移動をディレクティブで指示する方法（C/C++/Fortran）
PGI Fortran に限り、コンパイラで自動 deep copy を行う方法 (Fortran only)
CUDA Unified Memory を利用して、implicit にデータ管理が行われる方法（C/C++/Cortran）

ページの先頭へ

2. OpenACC における Fortran 派生型の扱い

手動による一般的なディレクティブ指示の方法

　PGI 14.4 以降の OpenACC において、Fortran 派生型の中に、さらに派生型を有するアロケータブルな配列や静的配列を含む配列メンバを取扱えるようになった。いずれの場合でも、適切な data 節ディレクティブの中に派生型それ自身を記述することによって、派生型全体をデバイスメモリに配置する必要がある（一般に言う「shallow copy」の状態である）。また、「集合型データ」の各メンバの実態のデータもデバイスに移動（一般に言う「deep copy」）を行う必要があるが、これを行うディレクティブをユーザが記述する必要がある。

　派生型変数を直接使用している任意の compute 構文に対して、派生型の変数自身は必ず data clause(節)の中に表れることが必要であり、少なくとも present 節で指定する必要がある。これが指定されていない場合、実行時エラーとなる。もう一つの注意点として、派生型変数内のメンバはディレクティブの data 節を利用して明示的に copy あるいは create してデバイス側の割付初期化を行う必要がある。但し、A1%n のようなスカラ変数に関しては、明示的な指示は必要ない（コンパイラ側で行う）が、指定しても良い。

program Test
    implicit none

    type point                ! 派生型の定義
        integer :: n
        real, dimension(:), allocatable :: x  ! Allocatable な配列メンバ
    end type point

    type(point) :: A1         ! 派生型変数 A1 の定義
    integer :: i

    A1%n = 10                 ! A1 のメンバ n の定義
    allocate(A1%x(A1%n))      ! 配列メンバの割付（ホストメモリ側）

!$acc enter data copyin(A1)   ! 派生型変数名 A1 host pointer をデバイス側へ shallow copy
!$acc enter data create(A1%x) ! A1 のメンバの実態配列 x をデバイス側にアロケート指示(attache)
                              ! (メンバの要素の実割付= 手動 deep copy）

!$acc kernels present(A1)    ! compute構文内で派生型変数を使用する場合、
   do i = 1, A1%n             ! A1を 必ずdata clauseで指示必要
      A1%x(i) = 2 * i
   enddo
!$acc end kernels

!$acc update host(A1%x)       ! デバイス上のA1%x内容をホスト側にコピー
   print*,A1%x
!$acc exit data delete(A1%x, A1)  ! メンバ要素、派生型変数名の順番で delete
   deallocate(A1%x)
end program Test

[kato@photon32 17.7]$  pgf90 -acc no-deepcopy.f90 -Minfo=accel -ta=tesla:cc60 -V17.7
test:
     15, Generating enter data copyin(a1)
     16, Generating enter data create(a1%x(:))
     18, Generating present(a1)
     19, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         19, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
     24, Generating update self(a1%x(:))
     26, Generating exit data delete(a1%x(:),a1)

[kato@photon32 17.7]$ a.out
    2.000000        4.000000        6.000000        8.000000
    10.00000        12.00000        14.00000        16.00000
    18.00000        20.00000

full deepcopy 機能を使った場合のディレクティブ指示の方法

　PGI 17.7 の OpenACC Fortran において、派生型の自動ディープコピー(deep copy)の Beta サポートを開始した。この機能を使用すると、OpenACC を使用して深くネストされたデータ構造を持つアプリケーションを Tesla GPU に移植することができる。PGI 17.7 コンパイラでは、OpenACC の　copy、copyin、copyout、update ディレクティブに、集合型の Fortran データオブジェクトをリストして集合型データオブジェクト内のポインタベースのオブジェクトのトラバーサルならびに管理を含めて、ホストとデバイスのメモリ間で移動させることができる。フルディープコピー(full deep copy) を有効にすると、派生型の Fortran 変数をホストからデバイスまたはデバイスからホストに移動する時に、ポインタと割り当て可能な配列を含むデータ構造全体が、ホストとデバイス間、またはデバイスとホストのメモリ間でコピーされる。 deep copyを有効にするには、オプション -ta=tesla に deepcopy サブオプション(-ta=tesla,deepcopy) を指定する。注意すべき点として、多型データ型はサポートされておらず、重複するポインタが存在すると実行時エラーが発生する可能性がある。

　以下のプログラムを見ていただこう。上記のプログラムと較べて、!$acc enter data create(A1%x) の記述がないことが分かる。派生型変数の中の各メンバのデータ移動は記述しなくても、コンパイラがフルに deep copy を行うコードを生成してくれる

program Test
    implicit none
    integer :: i

    type point
        integer :: n
        real, dimension(:), allocatable :: x
    end type point

    type(point) :: A1

    A1%n = 10
    allocate(A1%x(A1%n))

!$acc enter data copyin(A1)  ! 派生型変数名 A1 host pointer をデバイス側へshallow/deep copy
                             ! 派生型変数の各メンバのデータ移動の記述は必要なし
!$acc kernels present(A1)    ! compute構文内で派生型変数を使用する場合、
   do i = 1, A1%n            ! A1を 必ずdata clauseで指示必要
      A1%x(i) = 2 * i
   enddo
!$acc end kernels

!$acc update self(A1%x)
   print*,A1%x
!$acc exit data delete(A1)
   deallocate(A1%x)
end program Test

　PGI 17.7 から -ta=tesla オプションに deepcopy サブオプションがプロダクションサポートされた。これは Fortran 言語における派生型変数に対する full deep copy を実現するためのオプションである。派生型変数を使用している compute 構文に対して、その派生型変数内に置かれているメンバ変数の implicit なデータ転送実現する。但し、派生型変数名は、明示的に data 節で指定し、デバイスメモリ上での割付初期化が必要となる。

[kato@photon32 17.7]$ pgf90 -acc deepcopy.f90 -Minfo=accel -ta=tesla:cc60,deepcopy -V17.7
test:
     15, Generating enter data copyin(a1)
     16, Generating present(a1)
     17, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         17, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
     22, Generating update self(a1%x(:))
     24, Generating exit data delete(a1)
[kato@photon32 17.7]$ a.out
    2.000000        4.000000        6.000000        8.000000
    10.00000        12.00000        14.00000        16.00000
    18.00000        20.00000

CUDA Unified Memory を利用した場合のディレクティブ指示の方法

　PGI 17.7 において、CUDA Unified Memory 上の利用が正式サポートされた。CUDA Fortran または OpenACC において、Allocatable なデータが CUDA Unified Memory に配置されている場合、明示的なデータ移動またはデータ指示は必要とせず、compute 構文ディレクティブの指定だけで OpenACC プログラミングが可能となる。（注意　静的に配列宣言されている配列を含む処理にはこの機能を利用できません。静的な配列データが含まれている場合は、当該データに関して、ディレクティブで明示的にデータ移動する必要があります。）これにより、Allocatable なデータを大量に使用するアプリケーションの GPU アクセラレーションが簡素化され、アルゴリズムの並列化とスケーラビリティに集中できる。特に、Fortran 派生型変数、C/C++構造体、クラスを使用したデータ構造が複雑な場合において便利である。なお、注意が必要なのは、静的データ（C静的および外部変数、Fortran モジュール、共通ブロックおよび保存変数）および関数ローカルデータは、OpenACC ランタイムによって処理されるため、従来通り、ユーザによるデータ移動管理が必要であり、この managed オプションは使用できない。

program Test
    implicit none
    integer :: i

    type point
        integer :: n
        real, dimension(:), allocatable :: x    ! 動的配列のため CUDA Unified Memoryを利用できる
    end type point

    type(point) :: A1

    A1%n = 10
    allocate(A1%x(A1%n))

!$acc kernels
   do i = 1, A1%n
      A1%x(i) = 2 * i
   enddo
!$acc end kernels

   print*,A1%x
end program Test

　OpenACC において、この機能を有効にするには、 -ta=tesla,managed コンパイラオプションを使用する。

[kato@photon32 17.7]$ pgf90 -acc managed.f90 -Minfo=accel -ta=tesla,cc60,managed -V17.7
test:
     15, Generating implicit copyout(a1%x(1:a1%n))
     16, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         16, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
[kato@photon32 17.7]$ a.out
    2.000000        4.000000        6.000000        8.000000
    10.00000        12.00000        14.00000        16.00000
    18.00000        20.00000

ページの先頭へ

3. Fortran 派生型変数の使用例２

手動による一般的なディレクティブ指示の方法

　もう少し、複雑な派生型変数を含む場合の例として、PGI user forum に参考となる例が掲載されていたので、これを引用する。派生型変数 Matrix 自身を !$acc enter data create(Matrix) でデバイス側メモリ上に attach しただけでは、shallow copy となってしまうので、Matrix のメンバ実態の割付ならびにコピーは行われない。したがって、明示的に !$ACC ENTER DATA COPYIN(Matrix(i)%entry, Matrix(i)%col_idx, Matrix(i)%row_ptr) を使ってデータ移動を行う必要がある（手動による deep copy)。こうした概念を理解すれば、必要とする場所に必要とされるディレクティブを記述できるはずだ。

program foo

 TYPE CSR_MATRIX                        ! 派生型の定義
 SEQUENCE
 INTEGER :: entry_num, row_num, col_num
 REAL(8), ALLOCATABLE :: entry(:)       ! allocatable 配列を含む
 INTEGER, ALLOCATABLE :: col_idx(:)
 INTEGER, ALLOCATABLE :: row_ptr(:)
 END TYPE

 integer :: i,j
 integer,parameter :: N = 100
 TYPE(CSR_MATRIX) :: Matrix(N)          ! 派生型変数 matrix の定義(N要素数）

!$acc enter data create(Matrix)         ! ホスト側 Matrixのデバイスへの shallow copy 
 do i=1,N                               ! このループはホスト側での処理
   allocate(Matrix(i)%entry(N))         ! Matrix の配列要素のメンバの割付（実態）(attach)
   allocate(Matrix(i)%col_idx(N))
   allocate(Matrix(i)%row_ptr(N))
   Matrix(i)%entry_num = i
   Matrix(i)%row_num = i
   Matrix(i)%col_num = 1
                        ! デバイスへ「個数」に係るメンバ要素実態を update(更新, deep copy）
!$ACC UPDATE device(Matrix(i)%entry_num,Matrix(i)%row_num,Matrix(i)%col_num) 
                        ! デバイスへ配列要素のメンバの実態をコピー（deep copy)
!$ACC ENTER DATA COPYIN(Matrix(i)%entry, Matrix(i)%col_idx, Matrix(i)%row_ptr)
 end do

!$acc parallel present(Matrix)         ! compute構文内で派生型変数を使用する場合、
 do j=1,N                              ! present(Matrix)必須、デバイス内で計算処理
 do i=1,N
   Matrix(j)%entry(i) = real(i+j) / real(N+N)
   Matrix(j)%col_idx(i) = Matrix(j)%col_num + i
   Matrix(j)%row_ptr(i) = Matrix(j)%row_num + i
 end do
 end do
!$acc end parallel

#ifdef _OPENACC
 do i=1,N                              ! デバイス側データをホスト側へコピー
!$ACC update host (Matrix(i)%entry(1:N), Matrix(i)%col_idx(1:N), Matrix(i)%row_ptr(1:N))
 end do
#endif

 print *, Matrix(21)%entry(99)
 print *, Matrix(15)%col_idx(3)
 print *, Matrix(67)%row_ptr(97)
 do i=1,N
!$ACC EXIT DATA delete(Matrix(i)%entry, Matrix(i)%col_idx, Matrix(i)%row_ptr)  ! デバイス側配列削除
    deallocate(Matrix(i)%entry)         ! ホスト側配列削除
    deallocate(Matrix(i)%col_idx)
    deallocate(Matrix(i)%row_ptr)
 enddo
!$acc exit data delete(Matrix)

end program foo

　コンパイルした結果は以下のとおりである。

[kato@photon32 Example]$ pgf90 -acc -Minfo=accel -O2 -Mpreprocess -ta=tesla,cc60 test1.f90
foo:
     15, Generating enter data create(matrix(:))
     23, Generating update device(matrix%row_num,matrix%entry_num,matrix%col_num)
     24, Generating enter data copyin(matrix%entry(:),matrix%row_ptr(:),matrix%col_idx(:))
     27, Generating present(matrix(:))
         Accelerator kernel generated
         Generating Tesla code
         28, !$acc loop gang ! blockidx%x
         29, !$acc loop vector(128) ! threadidx%x
     29, Loop is parallelizable
     38, Generating update self(matrix%entry(1:100),matrix%row_ptr(1:100),matrix%col_idx(1:100))
     45, Generating exit data delete(matrix%entry(:),matrix%col_idx(:),matrix%row_ptr(:))
     50, Generating exit data delete(matrix(:))
[kato@photon32 Example]$ a.out
   0.6000000238418579
            4
          164

full deepcopy 機能を使った場合のディレクティブ指示の方法

　上記の例をコンパイラが自動で full deep copy を行えるようにプログラムを変更する。PGI 17.7 においてサポートされた -ta=tesla,deepcopy 機能を使うと以下のような（従来と変わらない）簡素なディレクティブの指定だけで実現する。すなわち、派生型変数の各メンバ要素の明示的なデータ移動は必要ない。

program foo

 TYPE CSR_MATRIX                        ! 派生型の定義
 SEQUENCE
 INTEGER :: entry_num, row_num, col_num
 REAL(8), ALLOCATABLE :: entry(:)       ! allocatable 配列を含む
 INTEGER, ALLOCATABLE :: col_idx(:)
 INTEGER, ALLOCATABLE :: row_ptr(:)
 END TYPE

 integer :: i,j
 integer,parameter :: N = 100
 TYPE(CSR_MATRIX) :: Matrix(N)          ! 派生型変数 matrix の定義(N要素数）
 
 do i=1,N                               ! このループはホスト側での処理
   allocate(Matrix(i)%entry(N))         ! Matrix の配列要素のメンバの割付（実態）
   allocate(Matrix(i)%col_idx(N))
   allocate(Matrix(i)%row_ptr(N))
   Matrix(i)%entry_num = i
   Matrix(i)%row_num = i
   Matrix(i)%col_num = 1
                                 ! デバイスへ「個数」に係るメンバ要素をコピーする必要がない
 end do
     
!$acc parallel loop copy(Matrix)   ! ホスト側 Matrix の デバイスへの full deep copy
 do i=1,N
   Matrix(j)%entry(i) = real(i+j) / real(N+N)
   Matrix(j)%col_idx(i) = Matrix(j)%col_num + i
   Matrix(j)%row_ptr(i) = Matrix(j)%row_num + i
 end do
 end do
!$acc end parallel

 print *, Matrix(21)%entry(99)
 print *, Matrix(15)%col_idx(3)
 print *, Matrix(67)%row_ptr(97)

 do i=1,N
    deallocate(Matrix(i)%entry)         ! ホスト側配列削除
    deallocate(Matrix(i)%col_idx)
    deallocate(Matrix(i)%row_ptr)
 enddo

end program foo

　PGI 17.7 以降にサポートされた派生型変数に対する full deep copy を実現するためのオプション deepcopy を -ta=tesla のサブオプションとして指定する必要がある。これを行うことにより、派生型配列 Matrix のメンバを含めた full deep copy が自動的に行われる。

$ pgf90 -acc -Minfo=accel -O2 -Mpreprocess -ta=tesla,cc60,deepcopy test2.F90
foo:
     24, Generating copy(matrix(:))
         Accelerator kernel generated
         Generating Tesla code
         25, !$acc loop gang ! blockidx%x
         26, !$acc loop vector(128) ! threadidx%x
     26, Loop is parallelizable
$ a.out
   0.6000000238418579
            4
          164

CUDA Unified Memory を利用した場合のディレクティブ指示の方法

　 OpenACC において allocatable なデータが CUDA Unified Memory に配置されている場合、明示的なデータ移動やデータ指示は必要としない。以下の例のように、!$acc parallel loop のみで compute 構文のみの指定で、データの移動は CUDA Unified memoryが管理して実行している。（注意　静的に配列宣言されている配列を含む処理にはこの機能を利用できません。静的な配列データが含まれている場合は、当該データに関して、ディレクティブで明示的にデータ移動する必要があります。）

program foo

 TYPE CSR_MATRIX                        ! 派生型の定義
 SEQUENCE
 INTEGER :: entry_num, row_num, col_num
 REAL(8), ALLOCATABLE :: entry(:)       ! allocatable 配列を含む、動的配列のため CUDA Unified Memoryを利用できる
 INTEGER, ALLOCATABLE :: col_idx(:)
 INTEGER, ALLOCATABLE :: row_ptr(:)
 END TYPE

 integer :: i,j
 integer,parameter :: N = 100
 TYPE(CSR_MATRIX) :: Matrix(N)          ! 派生型変数 matrix の定義(N要素数）
 
 do i=1,N                               ! このループはホスト側での処理
   allocate(Matrix(i)%entry(N))         ! Matrix の配列要素のメンバの割付（実態）
   allocate(Matrix(i)%col_idx(N))
   allocate(Matrix(i)%row_ptr(N))
   Matrix(i)%entry_num = i
   Matrix(i)%row_num = i
   Matrix(i)%col_num = 1
                                 ! デバイスへ「個数」に係るメンバ要素をコピーする必要がない
 end do
     
!$acc parallel loop   ! compute 構文の指定のみ
 do i=1,N
   Matrix(j)%entry(i) = real(i+j) / real(N+N)
   Matrix(j)%col_idx(i) = Matrix(j)%col_num + i
   Matrix(j)%row_ptr(i) = Matrix(j)%row_num + i
 end do
 end do
!$acc end parallel

 print *, Matrix(21)%entry(99)
 print *, Matrix(15)%col_idx(3)
 print *, Matrix(67)%row_ptr(97)

 do i=1,N
    deallocate(Matrix(i)%entry)         ! ホスト側配列削除
    deallocate(Matrix(i)%col_idx)
    deallocate(Matrix(i)%row_ptr)
 enddo

end program foo

　OpenACC において、この機能を有効にするには、 -ta=tesla,managed コンパイラオプションを使用する必要がある。

$ pgf90 -acc -Minfo=accel -O2 -Mpreprocess -ta=tesla,cc60,managed managed.F90
foo:
     24, Accelerator kernel generated
         Generating Tesla code
         25, !$acc loop gang ! blockidx%x
         26, !$acc loop vector(128) ! threadidx%x
     24, Generating implicit copyin(matrix(:))
         Generating implicit copyout(matrix(0))
     26, Loop is parallelizable
[kato@photon32 Example]$ a.out
   0.6000000238418579
            4
          164

ページの先頭へ

4. OpenACC における C 構造体データ

手動による一般的なディレクティブ指示の方法

　以下の C プログラム例は、動的に割り付けられるデータをメンバとした C 構造体（vector変数）を使用する例である。明示的に構造体メンバ要素のデータコピーを行う例は、vector * allocData(size_t size1) 関数を見てもらいたい。まず、vector * tmp 配列をホスト側に割り付けた後、#pragma acc enter data create(tmp[0:size1])によって、デバイス側にも tmp[0:size1] を生成（attach) する。これは、shallow copy の状態であり、その配下のスカラ変数を除く配列形式のメンバ要素は、個別にデバイス側に割付する必要がある。ホスト側で配列形式のメンバ要素を割り付けたと同時に、#pragma acc enter data create(tmp[i].data[0:tmp[i].size])　によって、デバイス側にも対応するメンバ要素の割付を行う。基本的には、C 構造体データは、deepcopy 機能がサポートされるまではこのような形で明示的なデータ移動が必要とされる。なお、C 構造体を用いたプログラムのポーティング・チュートリアルは、以下の「6. C 構造体データを使用したプログラムのポーティング・チュートリアル」で詳しく解説しているので参照して欲しい。

array_of_structs.c

! Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
!
! Redistribution and use in source and binary forms, with or without modification,
! are permitted provided that the following conditions are met:
!
! *Redistributions of source code must retain the above copyright notice,
!  this list of conditions and the following disclaimer.
! *Redistributions in binary form must reproduce the above copyright notice,
!  this list of conditions and the following disclaimer in the documentation
!  and/or other materials provided with the distribution.
! *Neither the name of NVIDIA CORPORATION nor the names of its contributors
!  may be used to endorse or promote products derived from this software
!  without specific prior written permission.

! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
! EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
! WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
! DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE
! FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
! DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
! SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
! CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
! OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE
! USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

#include <stdlib.h>
#include <stdio.h>

#ifndef N
#define N 32 
#endif

typedef struct {             // vector構造体の配列の定義 (array of struct)
   int size;                 // スカラ変数
   double * data;            // メンバは動的割付データ
} vector; 

vector * allocData(size_t size1);
int deleteData(vector* A, size_t size1);
int initData(vector *A, size_t size1, double val);
int printData(vector *A, size_t size1);

int main() {

    vector *A, *B;          // vector構造体の配列 A, B
    size_t size1, i, j;
    size1 = N;
   
    A=allocData(size1);
    B=allocData(size1);
    initData(B,size1,2.5);

/* Perform the computation on the device */
#pragma acc parallel loop gang present(A,B)   // present(A,B) 節の指定必要
    for (j=0; j < size1; ++j) {
       int size2 = A[j].size;
#pragma acc loop vector
       for (i=0; i < size2; ++i) {
          A[j].data[i]= B[j].data[i] + (double) ((j*size2)+i);
       }
    } 
#ifdef _OPENACC
/* Copy back the results 各vector配列のメンバは個々にコピーする必要あり */ 
    for (j=0; j < size1; ++j) {
#pragma acc update self (A[j].data[0:A[j].size])  
    }
#endif

    printData(A,size1);
    deleteData(A,size1);
    deleteData(B,size1);
    exit(0);
}
/* vector構造体 配列の割付ルーチン */
vector * allocData(size_t size1) {
    vector * tmp;      
    int i;
    tmp = (vector*) malloc(size1*sizeof(vector));　 // vector構造体自体の配列領域割付

/* Create an array of pointers デバイス側にポインタの配列作成 (Shallow copy) */
#pragma acc enter data create(tmp[0:size1])　
    for (i=0; i < size1; ++i) {   // vector構造体の dataメンバの配列領域割付(サイズが異なる）
       tmp[i].size = i+10;
       tmp[i].data = (double *) malloc(tmp[i].size*sizeof(double));
/* Create the vector and attach it to the pointer array  デバイス側にdataメンバの配列領域割付(Deep copy) */ 
#pragma acc enter data create(tmp[i].data[0:tmp[i].size])
/* Update the device's size */   // スカラのメンバsizeはデバイス側にコピー
#pragma acc update device(tmp[i].size)                    
    }
    return tmp;
}

int deleteData(vector * A, size_t size1) {
    int i;
    for (i=0; i < size1; ++i) {
       free(A[i].data);
#pragma acc exit data delete(A[i].data) // 個々のvector構造体の dataメンバを先に削除
  } 
#pragma acc exit data delete(A)         // vector構造体 配列自体を削除
    free(A);                            // ホスト側の構造体を削除
} 

int initData(vector *A, size_t size1, double val) {
    size_t i,j;
    for (j=0; j < size1; ++j) {
       int size2=A[j].size;
       for (i=0; i < size2; ++i) {
	  A[j].data[i] = val;
       }
/* Update the device with the initial values */    // 個々の dataメンバをデバイス側へコピー
#pragma acc update device(A[j].data[0:size2])
    }
}

int printData(vector *A, size_t size1) {
    size_t i,j;
    printf("Values:\n");
    for (i=0; i < 5; ++i) {   
        int last = A[i].size-1;
	printf("A[%d].data[0]=%f A[%d].data[%d]=%f\n",i,A[i].data[0],i,last,A[i].data[last]);
    } 
    printf("....\n");
    for (i=size1-5; i < size1; ++i) {   
        int last = A[i].size-1;
	printf("A[%d][0]=%f A[%d][%d]=%f\n",i,A[i].data[0],i,last,A[i].data[last]);
    } 
}

[kato@photon32]$ pgcc -acc -O2 -Minfo=accel array_of_structs.c  -ta=tesla,cc60,cuda8.0
main:
     29, Generating present(B[:],A[:])
         Accelerator kernel generated
         Generating Tesla code
         30, #pragma acc loop gang /* blockIdx.x */
         33, #pragma acc loop vector(128) /* threadIdx.x */
     33, Loop is parallelizable
     39, Generating update self(A->data[:A->size])
allocData:
     56, Generating enter data create(tmp[:size1])
     61, Generating enter data create(tmp->data[:tmp->size])
     64, Generating update device(tmp->size)
deleteData:
     72, Generating exit data delete(A->data[:1])
     74, Generating exit data delete(A[:1])
initData:
     87, Generating update device(A->data[:size2])

[kato@photon32 Chap5]$ a.out
Values:
A[0].data[0]=2.500000 A[0].data[9]=11.500000
A[1].data[0]=13.500000 A[1].data[10]=23.500000
A[2].data[0]=26.500000 A[2].data[11]=37.500000
A[3].data[0]=41.500000 A[3].data[12]=53.500000
A[4].data[0]=58.500000 A[4].data[13]=71.500000
....
A[27][0]=1001.500000 A[27][36]=1037.500000
A[28][0]=1066.500000 A[28][37]=1103.500000
A[29][0]=1133.500000 A[29][38]=1171.500000
A[30][0]=1202.500000 A[30][39]=1241.500000
A[31][0]=1273.500000 A[31][40]=1313.500000

ページの先頭へ

CUDA Unified Memory を利用した場合のディレクティブ指示の方法

　上記の C プログラムを Unified Memory を使用すること前提にすると、明示的な OpenACC ディレクティブによるデータ移動やデータ指示は必要ありません。OpenACC compute 構文による並列処理の指示だけで済みます。ただし、GPU 上の処理において使用される配列は、動的に割り付ける allocatable データで構成されているものに限ります。これには構造体データも含みます。一方、静的に割り付けられた配列が一部で使用される場合は、当該データに関しては OpenACC data ディレクティブで指示する必要があります。

array_of_structs.c

! Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
!
! Redistribution and use in source and binary forms, with or without modification,
! are permitted provided that the following conditions are met:
!
! *Redistributions of source code must retain the above copyright notice,
!  this list of conditions and the following disclaimer.
! *Redistributions in binary form must reproduce the above copyright notice,
!  this list of conditions and the following disclaimer in the documentation
!  and/or other materials provided with the distribution.
! *Neither the name of NVIDIA CORPORATION nor the names of its contributors
!  may be used to endorse or promote products derived from this software
!  without specific prior written permission.

! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
! EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
! WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
! DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE
! FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
! DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
! SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
! CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
! OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE
! USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

#include <stdlib.h>
#include <stdio.h>

#ifndef N
#define N 32 
#endif

typedef struct {             // vector構造体の配列の定義 (array of struct)
   int size;                 // スカラ変数
   double * data;            // メンバは動的割付データ
} vector; 

vector * allocData(size_t size1);
int deleteData(vector* A, size_t size1);
int initData(vector *A, size_t size1, double val);
int printData(vector *A, size_t size1);

int main() {

    vector *A, *B;          // vector構造体の配列 A, B
    size_t size1, i, j;
    size1 = N;
   
    A=allocData(size1);
    B=allocData(size1);
    initData(B,size1,2.5);

/* Perform the computation on the device */
#pragma acc parallel loop gang present(A,B)   // present(A,B) 節の指定必要
    for (j=0; j < size1; ++j) {
       int size2 = A[j].size;
#pragma acc loop vector
       for (i=0; i < size2; ++i) {
          A[j].data[i]= B[j].data[i] + (double) ((j*size2)+i);
       }
    } 
#ifdef _OPENACC
/* Copy back the results 必要なし */ 
//    for (j=0; j < size1; ++j) {
// #pragma acc update self (A[j].data[0:A[j].size])
//    }
#endif

    printData(A,size1);
    deleteData(A,size1);
    deleteData(B,size1);
    exit(0);
}
/* vector構造体 配列の割付ルーチン */
vector * allocData(size_t size1) {
    vector * tmp;      
    int i;
    tmp = (vector*) malloc(size1*sizeof(vector));   // vector構造体自体の配列領域割付

/* Create an array of pointers */
// #pragma acc enter data create(tmp[0:size1])     // デバイス側へのコピー指示は必要なし
    for (i=0; i < size1; ++i) {   
       tmp[i].size = i+10;
       tmp[i].data = (double *) malloc(tmp[i].size*sizeof(double));
/* Create the vector and attach it to the pointer array  */ 
// #pragma acc enter data create(tmp[i].data[0:tmp[i].size]) // 必要なし
/* Update the device's size */  
// #pragma acc update device(tmp[i].size)          // 必要なし      
    }
    return tmp;
}

int deleteData(vector * A, size_t size1) {
    int i;
    for (i=0; i < size1; ++i) {
       free(A[i].data);
// #pragma acc exit data delete(A[i].data)    // 必要なし
  } 
// #pragma acc exit data delete(A)      
    free(A);                          
} 

int initData(vector *A, size_t size1, double val) {
    size_t i,j;
    for (j=0; j < size1; ++j) {
       int size2=A[j].size;
       for (i=0; i < size2; ++i) {
	  A[j].data[i] = val;
       }
/* Update the device with the initial values */    
// #pragma acc update device(A[j].data[0:size2])    // 必要なし
    }
}

int printData(vector *A, size_t size1) {
    size_t i,j;
    printf("Values:\n");
    for (i=0; i < 5; ++i) {   
        int last = A[i].size-1;
	printf("A[%d].data[0]=%f A[%d].data[%d]=%f\n",i,A[i].data[0],i,last,A[i].data[last]);
    } 
    printf("....\n");
    for (i=size1-5; i < size1; ++i) {   
        int last = A[i].size-1;
	printf("A[%d][0]=%f A[%d][%d]=%f\n",i,A[i].data[0],i,last,A[i].data[last]);
    } 
}

　OpenACC において、この機能を有効にするには、 -ta=tesla,managed コンパイラオプションを使用する必要がある。

[kato@photon32 Managed]$ pgcc -acc -ta=tesla,managed,cc60,cuda8.0 -O2 -Minfo=accel array_of_structs.c
main:
     29, Generating present(B[:],A[:])
         Accelerator kernel generated
         Generating Tesla code
         30, #pragma acc loop gang /* blockIdx.x */
         33, #pragma acc loop vector(128) /* threadIdx.x */
     33, Loop is parallelizable
[kato@photon32 Managed]$ a.out
Values:
A[0].data[0]=2.500000 A[0].data[9]=11.500000
A[1].data[0]=13.500000 A[1].data[10]=23.500000
A[2].data[0]=26.500000 A[2].data[11]=37.500000
A[3].data[0]=41.500000 A[3].data[12]=53.500000
A[4].data[0]=58.500000 A[4].data[13]=71.500000
....
A[27][0]=1001.500000 A[27][36]=1037.500000
A[28][0]=1066.500000 A[28][37]=1103.500000
A[29][0]=1133.500000 A[29][38]=1171.500000
A[30][0]=1202.500000 A[30][39]=1241.500000
A[31][0]=1273.500000 A[31][40]=1313.500000

ページの先頭へ

5. OpenACC における C++クラスデータ

手動による一般的なディレクティブ指示の方法

　上記 C プログラム中の vector 構造体を使用した同じような例を C++ の List クラスに拡張してみよう。デバイス側のデータ生成とデータの削除は、クラスのコンストラクタとデストラクタの中に記述する。これによって、クラス内のオブジェクトが生成あるいは削除される際に、インプリシットに実行される。また、以下の例では、異なるデータ型に対して使用可能なクラス定義を行うためにテンプレートを使用する。クラスの記述の中で、 enter data create の引数リストにクラスの this ポインタを使用している。この this ポインタは、あるオブジェクトのデータメンバの実インスタンスを意味する。プログラマからは見えないのだが、オブジェクトのデータメンバへの全てのアクセスは、this ポインタを通して実行される。"_size=0" を使ってメンバをセットすると、プログラムは実際、"this->_size = 0" にアクセスする。this ポインタがデバイス上に作成される際に、領域は、オブジェクトのデータメンバ毎に生成される。例えば、this ポインタが copyin の中で指定された場合、これは、データメンバの shallow copy が行われる。もし、その中のデータメンバがポインタ（配列）である場合、対応するホスト側のポインタが、単にデバイス側にコピーされることに注意する必要がある（この時点で、デバイス側でのデータメンバのポインタ（配列）の attach は行われていない。

　次に、動的なデータメンバは acc enter data create(_A[0:_size]) によってデバイス上にアロケートされ、this ポインタに attach される。「attach」 とは、デバイスデータのアドレスへデータメンバのポインタが実際にセットされるということを指す。 shallow copy した後に deep copy するという順番が大事である。もし、動的なデータメンバが、this ポインタの生成より前に、attach 行為を行った場合、attach する場所がないため、この行為はスキップされる。一方、デストラクタ上では、この順番とは逆に、データメンバを先に削除し、最後に this ポインタを削除する順番となる。

　myList.cpp メインプログラムでは、データ移動の同期ポイント（accUpdateDevice、accUpdateSelf）を考慮するだけで、データ移動に係る他の操作（create, delete, update) は行う必要がない。処理の並列化の指示 acc parallel loop だけで良い。なお、ここでは例示しないが、CUDA Unified Memory を使用する際は、データ移動に係るディレクティブの挿入は一切必要なくなり、処理の並列化の指示のみでよい。

myList.h

! Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
!
! Redistribution and use in source and binary forms, with or without modification,
! are permitted provided that the following conditions are met:
!
! *Redistributions of source code must retain the above copyright notice,
!  this list of conditions and the following disclaimer.
! *Redistributions in binary form must reproduce the above copyright notice,
!  this list of conditions and the following disclaimer in the documentation
!  and/or other materials provided with the distribution.
! *Neither the name of NVIDIA CORPORATION nor the names of its contributors
!  may be used to endorse or promote products derived from this software
!  without specific prior written permission.

! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
! EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
! WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
! DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE
! FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
! DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
! SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
! CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
! OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE
! USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.


#ifdef _OPENACC
#include <openacc.h>
#endif

using namespace std;


template<typename T>
class myList {

   private:
      T* _A{nullptr};   
      size_t _size{0};  
  
   public:

    #pragma acc routine seq
    T& operator[](size_t idx) { return _A[idx]; };

    #pragma acc routine seq
    const T& operator[](size_t idx) const { return _A[idx]; };

    size_t size() const {
	return _size;
    }

    explicit myList() { }
    explicit myList(size_t size) {
	_size = size;
        _A = new T[_size];
	#pragma acc enter data copyin(this)         // this pointer のコピー(shallow copy)
        #pragma acc enter data create(_A[0:_size])  // その後、データメンバの attach (deep copy)
    }

    ~myList() {
        #pragma acc exit data delete(_A[0:_size])   // データメンバから削除
	#pragma acc exit data delete(this)          // 最後に this pointer の削除
	delete [] _A;
	_A=NULL;
        _size=0;
    }

    inline void accUpdateSelf() {
        #pragma acc update self(_A[0:_size])    // ホスト側へデータコピー
    } 
    inline void accUpdateDevice() {
        #pragma acc update device(_A[0:_size])  // デバイス側へデータコピー
    } 
};

myList.cpp

#ifndef N
#define N 1024
#endif
#include <iostream>
#include <string.h>
#include "myList.h" 

int main() {

    myList <double> A(N), B(N);
    for (int i=0; i < B.size(); ++i) {
        B[i]=2.5;
    } 
    B.accUpdateDevice();
    #pragma acc parallel loop present(A,B)  // present(A,B) は必須
    for (int i=0; i < A.size(); ++i) {
	   A[i]=B[i]+i;
    } 
    A.accUpdateSelf();
    for(int i=0; i<10; ++i) {
	cout << "A[" << i << "]: " << A[i] << endl;
    }
    exit(0);
}

[kato@photon32 Chap5]$ pgc++ -acc -Minfo=accel -O2 myList.cpp -std=c++11 -ta=tesla,cc60,cuda8.0
main:
     14, Generating present(B,A)
         Accelerator kernel generated
         Generating Tesla code
         16, #pragma acc loop gang /* blockIdx.x */
     16, Scalar last value needed after loop for A._A-> at line 21
myList::operator [](unsigned long):
      6, include "myList.h"
          19, Generating acc routine seq
              Generating Tesla code
myList::size() const:
      6, include "myList.h"
          24, Generating implicit acc routine seq
              Generating acc routine seq
              Generating Tesla code
myList::myList(unsigned long):
      6, include "myList.h"
          34, Generating enter data copyin(this[:1])
              Generating enter data create(_A[:_size])
myList::~myList():
      6, include "myList.h"
          39, Generating exit data delete(this[:1],_A[:_size])
myList::accUpdateSelf():
      6, include "myList.h"
          46, Generating update self(_A[:_size])
myList::accUpdateDevice():
      6, include "myList.h"
          49, Generating update device(_A[:_size])
std::basic_ostream<T1, T2> & std::endl<char, std::char_traits<char>>(std::basic_ostream &):
      6, include "myList.h"

[kato@photon32 Chap5]$ pgprof a.out
==12294== PGPROF is profiling process 12227, command: a.out
Values:
A[0].data[0]=2.500000 A[0].data[9]=11.500000
A[1].data[0]=13.500000 A[1].data[10]=23.500000
A[2].data[0]=26.500000 A[2].data[11]=37.500000
A[3].data[0]=41.500000 A[3].data[12]=53.500000
A[4].data[0]=58.500000 A[4].data[13]=71.500000
....
A[27][0]=1001.500000 A[27][36]=1037.500000
A[28][0]=1066.500000 A[28][37]=1103.500000
A[29][0]=1133.500000 A[29][38]=1171.500000
A[30][0]=1202.500000 A[30][39]=1241.500000
A[31][0]=1273.500000 A[31][40]=1313.500000
==12294== Profiling application: a.out
==12294== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 96.34%  307.37us         1  307.37us  307.37us  307.37us  main_14_gpu
  2.40%  7.6480us         6  1.2740us  1.2160us  1.3760us  [CUDA memcpy HtoD]
  1.26%  4.0320us         2  2.0160us  1.7280us  2.3040us  [CUDA memcpy DtoH]

==12294== API calls:
Time(%)      Time     Calls       Avg       Min       Max  Name
 58.56%  141.71ms         1  141.71ms  141.71ms  141.71ms  cuDevicePrimaryCtxRetain
 32.84%  79.466ms         1  79.466ms  79.466ms  79.466ms  cuDevicePrimaryCtxRelease
  4.96%  12.005ms         1  12.005ms  12.005ms  12.005ms  cuMemHostAlloc
  2.77%  6.7027ms         1  6.7027ms  6.7027ms  6.7027ms  cuMemFreeHost
  0.37%  888.05us         6  148.01us  3.8970us  309.99us  cuMemAlloc
  0.29%  697.99us         1  697.99us  697.99us  697.99us  cuMemAllocHost
  0.13%  319.37us         2  159.69us  4.5040us  314.87us  cuMemcpyDtoHAsync
  0.03%  70.598us         1  70.598us  70.598us  70.598us  cuModuleLoadData
  0.02%  46.747us         6  7.7910us  3.4430us  19.489us  cuMemcpyHtoDAsync
  0.01%  25.949us         6  4.3240us     658ns  7.2380us  cuStreamSynchronize
  0.01%  24.738us         1  24.738us  24.738us  24.738us  cuLaunchKernel
  0.01%  23.730us         1  23.730us  23.730us  23.730us  cuStreamCreate
  0.00%  8.1400us        10     814ns     257ns  2.4640us  cuPointerGetAttributes
  0.00%  2.6650us         3     888ns     221ns  2.2100us  cuDeviceGetCount
  0.00%  2.3080us         3     769ns     521ns  1.0500us  cuCtxSetCurrent
  0.00%  2.2030us         2  1.1010us     440ns  1.7630us  cuEventCreate
  0.00%  1.5710us         1  1.5710us  1.5710us  1.5710us  cuModuleGetFunction
  0.00%  1.5470us         1  1.5470us  1.5470us  1.5470us  cuEventRecord
  0.00%  1.2940us         1  1.2940us  1.2940us  1.2940us  cuMemFree
  0.00%     945ns         4     236ns     157ns     391ns  cuDeviceGetAttribute
  0.00%     855ns         3     285ns     176ns     410ns  cuDeviceGet
  0.00%     804ns         1     804ns     804ns     804ns  cuEventSynchronize
  0.00%     217ns         1     217ns     217ns     217ns  cuCtxGetCurrent
  0.00%     206ns         1     206ns     206ns     206ns  cuDeviceComputeCapability

==12294== OpenACC (excl):
Time(%)      Time     Calls       Avg       Min       Max  Name
 95.65%  12.680ms         4  3.1701ms  7.8520us  12.335ms  acc_enter_data@myList.cpp:34
  2.38%  316.08us         1  316.08us  316.08us  316.08us  acc_enqueue_download@myList.cpp:14
  0.63%  83.908us         1  83.908us  83.908us  83.908us  acc_device_init@myList.cpp:34
  0.35%  46.577us         4  11.644us  4.1260us  24.707us  acc_enqueue_upload@myList.cpp:34
  0.21%  27.843us         1  27.843us  27.843us  27.843us  acc_enqueue_launch@myList.cpp:14 (main_14_gpu)
  0.15%  19.755us         1  19.755us  19.755us  19.755us  acc_compute_construct@myList.cpp:14
  0.12%  16.457us         2  8.2280us  8.0640us  8.3930us  acc_wait@myList.cpp:34
  0.09%  11.658us         1  11.658us  11.658us  11.658us  acc_enqueue_download@myList.cpp:46
  0.08%  11.243us         1  11.243us  11.243us  11.243us  acc_wait@myList.cpp:46
  0.06%  7.5190us         1  7.5190us  7.5190us  7.5190us  acc_update@myList.cpp:49
  0.05%  6.4190us         1  6.4190us  6.4190us  6.4190us  acc_exit_data@myList.cpp:14
  0.05%  6.3820us         1  6.3820us  6.3820us  6.3820us  acc_wait@myList.cpp:49
  0.05%  6.0200us         1  6.0200us  6.0200us  6.0200us  acc_update@myList.cpp:46
  0.04%  5.2980us         1  5.2980us  5.2980us  5.2980us  acc_enqueue_upload@myList.cpp:14
  0.03%  4.5980us         1  4.5980us  4.5980us  4.5980us  acc_enqueue_upload@myList.cpp:49
  0.03%  3.5480us         1  3.5480us  3.5480us  3.5480us  acc_enter_data@myList.cpp:14
  0.02%  3.1260us         2  1.5630us  1.1740us  1.9520us  acc_wait@myList.cpp:14
  0.00%       0ns         4       0ns       0ns       0ns  acc_alloc@myList.cpp:34
  0.00%       0ns         4       0ns       0ns       0ns  acc_create@myList.cpp:34

======== CPU profiling result (bottom up):
Time(%)      Time  Name
 73.12%  687.99ms  ???
 73.12%  687.99ms  | start_thread
 73.12%  687.99ms  |   clone
 15.05%  141.64ms  cuDevicePrimaryCtxRetain
 15.05%  141.64ms  | __pgi_uacc_cuda_init_device
 15.05%  141.64ms  |   __pgi_uacc_cuda_select_valid
 15.05%  141.64ms  |     __pgi_uacc_select_devid
 15.05%  141.64ms  |       __pgi_uacc_dataenterstart
 15.05%  141.64ms  |         myList::__complete_object_constructor__(unsigned long)
  8.60%   80.94ms  cuDevicePrimaryCtxRelease
  8.60%   80.94ms  | __pgi_uacc_cuda_release_buffer
  8.60%   80.94ms  |   __run_exit_handlers
  8.60%   80.94ms  |     ???
  8.60%   80.94ms  |       main
  8.60%   80.94ms  |         ???
  1.08%  10.117ms  cuMemHostAlloc
  1.08%  10.117ms  | __pgi_uacc_cuda_get_buffer
  1.08%  10.117ms  |   __pgi_uacc_cuda_dataup1
  1.08%  10.117ms  |     __pgi_uacc_dataup1
  1.08%  10.117ms  |       __pgi_uacc_dataupx
  1.08%  10.117ms  |         __pgi_uacc_dataonb
  1.08%  10.117ms  |           myList::__complete_object_constructor__(unsigned long)
  1.08%  10.117ms  cuMemFreeHost
  1.08%  10.117ms  | __pgi_uacc_cuda_free_device_buffers
  1.08%  10.117ms  |   __pgi_uacc_cuda_release_buffer
  1.08%  10.117ms  |     __run_exit_handlers
  1.08%  10.117ms  |       ???
  1.08%  10.117ms  |         main
  1.08%  10.117ms  |           ???
  1.08%  10.117ms  cuInit
  1.08%  10.117ms    __pgi_uacc_cuda_init
  1.08%  10.117ms      __pgi_uacc_enumerate
  1.08%  10.117ms        __pgi_uacc_initialize
  1.08%  10.117ms          __pgi_uacc_dataenterstart
  1.08%  10.117ms            myList::__complete_object_constructor__(unsigned long)

======== Data collected at 100Hz frequency

ページの先頭へ

6.　C 構造体データを使用したプログラムのポーティング・チュートリアル

　以下に示す ufbase.c プログラムは、複数のポインタメンバを含んだ、一つだけの構造体変数（スカラ構造体）を持つプログラムである。このプログラムは、数値計算上、機能的な意味を持つのではなく、C 構造体を持つプログラムを OpenACC 化する際に注意すべきことを解説するためのものである。最初に、このプログラムのポーティングから注意すべきポイント（誤った使い方により、OpenACC 実行がエラーとなる原因）を示しておく。

配列構成ではない「一つ」の構造体（スカラ構造体）変数でかつ、「ポインタ変数」として宣言された場合の OpenACC directive 上での当該構造体の表記方法。表記の方法を誤ってしまう例を解説する。「一つ」と言う意味は、構造体が配列（ベクトル）構成ではないという意味で、ここでは「スカラ構造体」と称することにする。
OpenACC 環境で、メンバを含めて構造体をデバイス側に attach する方法。「attach」とは、デバイスデータのアドレスへデータのポインタが実際にセットされるということを指す。OpenACC の data 構文 cluase を使って、デバイス側に構造体のデータ領域確保、メンバのデータ転送を行う場合、必ず、構造体「本体」変数の copyin or create を行なった後に、各メンバの data 構文 clause によるディレクティブを記述する必要がある（手動 deep copy）。

　OpenACC において構造体に対する扱い方の留意点とは別に、以下のプログラムの主要なループは、デフォルトのコンパイルオプションの指定だけでは、ベクトル化あるいは並列化できない。これは、ループ内で構造体のメンバ同士の演算が行われており、当該データ間の依存性について、コンパイラは判断できないことに因る。OpenACC 用のオプションをつけても同様にループは、並列化出来ない。こうしたポインタ変数を使用したベクトル・並列化を阻害する問題は、C/C++ 特有の問題である。

ループ内でポインタ配列や構造体のメンバ配列を使用する場合、そのデータ間に「データ依存性（領域の重なり）」がない場合は、コンパイルオプション -Msafeptr を付けると、このような問題は回避され最適化される。

① ベースとなるプログラム（構造体をポインタ変数ではない実体変数として宣言）

ufbase.c

#include <stdlib.h>
#include <stdio.h>

typedef struct {
  double *a;
  double *b;
  double **c;
} structdata;

int main() {
// 構造体 ポインタ変数ではない実体変数 "data" として宣言
  structdata data;
  int i,j;
  int n = 10;
  int m = 10;

  data.a = (double*)calloc(n, sizeof(double));
  data.b = (double*)calloc(n, sizeof(double));
  data.c = (double**)calloc(n, sizeof(double*));
  data.c[0] = (double*)calloc(m, sizeof(double*));
  for (int i = 1; i < n; i++) {
        data.c[i] = data.c[i-1] + m;
        printf("%d %d %lf\n", n, m, data.c[i]);
  }

// Initilize data members
  for (j=0; j<n; j++) {
    data.a[j] = 0.0;
    data.b[j] = 0.0;
    for (i = 1; i < m; i++) {
      data.c[j][i] = 0.0;
    }
  }
// main loop
  for (i = 1; i < n; i++) {
          data.a[i] = i + 2;
          data.b[i] = i;
  }
  for (j = 0; j < n; j++) {
          data.c[j][0] = j*2 ;
    for (i = 1; i < m; i++) {
          data.c[j][i] += data.a[i] + data.b[i];
    }
  }

  printf("j i a b c\n");
  for (j = 0; j < n; j++) {
        for (i = 0; i < m; i++) {
          printf("%d %d %lf %lf %lf\n", j, i, data.a[j], data.b[i], data.c[j][i]);
        }
  }

  exit(0);
}

　上記のプログラムを最適化レベル -O2 でコンパイルすると、以下のメッセージが出力される。左端の番号はソースの行番号である。「data dependency」、「Loop not vectorized/parallelized」等の文字列が並び、ループは並列化もベクトル化も出来ない状態である。

$ pgcc -O2 -Minfo ufbase.c
main:
     21, Loop not vectorized/parallelized: contains call
     30, Memory zero idiom, loop replaced by call to __c_mzero8
     35, Loop not vectorized: data dependency
         Loop unrolled 16 times
         Generated 2 prefetches in scalar loop
     39, Generated 1 prefetches in scalar loop
     41, Loop not vectorized: data dependency
         Loop unrolled 4 times
     48, Loop not vectorized/parallelized: contains call

　コンパイルオプション -Msafeptr を付けるとこの問題は回避され、ベクトル化(vector simd)されたコードが生成される。OpenACC 化に当っては、まず、ループ自体がベクトル化あるいは並列化できる条件を満たす必要がある。（ただし、コンパイラはベクトル化、並列化できないループでも、デバイス側でスカラ実行するコードは作成するが性能的に意味はない。）

$  pgcc -O2 -Minfo ufbase.c -Msafeptr
main:
     21, Loop not vectorized/parallelized: contains call
     30, Memory zero idiom, loop replaced by call to __c_mzero8
     35, Generated vector simd code for the loop
     41, Generated vector simd code for the loop
         Generated 3 prefetch instructions for the loop
     48, Loop not vectorized/parallelized: contains call

② ベースを OpenACC でデバイスコード化（構造体をポインタ変数ではない実体変数として宣言）

　以下に示す uf0.c プログラムは、main loop に対して OpenACC のデータ領域指示と並列化指示を行なった例である。ここでの説明は、構造体が「ポインタ変数」ではない「実体変数」として宣言された場合の #pragma acc data 構文を用いて、構造体の shallow copy と deep copy を行う部分についてである。

　以下の main loop の部分を見て欲しい。デバイス側で処理したいブロックを #pragma acc data 構文で囲み、create, copy, copyin, copyout 等の clause を使って、デバイス側で使用するデータを attach する。ループ内に対しては、#pragma acc parallel あるいは kernels の構文を挿入して、並列化・ベクトル化を指示する。単純にこれだけの挿入により、コンパイラは GPU デバイス用の並列コードを生成できる。

　このプログラムで注目すべき点は、#pragma acc data 構文を用いて、デバイス側にデータを attach する際の data cluase の順番と構造体本体変数(dataという変数）の表記法についてである。

uf0.c

#include <stdlib.h>
#include <stdio.h>

typedef struct {
  double *a;
  double *b;
  double **c;
} structdata;

int main() {
// 構造体 ポインタ変数ではない実体変数 "data" として宣言
  structdata data;
  int i,j;
  int n = 10;
  int m = 10;

  data.a = (double*)calloc(n, sizeof(double));
  data.b = (double*)calloc(n, sizeof(double));
  data.c = (double**)calloc(n, sizeof(double*));
  data.c[0] = (double*)calloc(m, sizeof(double*));
  for (int i = 1; i < n; i++) {
    data.c[i] = data.c[i-1] + m;
    printf("%d %d %lf\n", n, m, data.c[i]);
  }

// Initilize data members
  for (j=0; j<n; j++) {
     data.a[j] = 0.0;
     data.b[j] = 0.0;
    for (i = 1; i < m; i++) {
     data.c[j][i] = 0.0;
    }
  }

// main loop
// 構造体　ポインタではない実体変数
// Allocate a struct variable, then allocate pointer members in the device
  #pragma acc data create(data) copyout(data.a[1:n],data.b[1:n]) copy(data.c[0:n][1:m])
  {
    #pragma acc parallel loop present(data.a[1:n], data.b[1:n])
    for (i = 1; i < n; i++) {
          data.a[i] = i + 2;
          data.b[i] = i;
    }
    #pragma acc kernels loop present(data.a[1:n], data.b[1:n], data.c[0:n][1:m])
    for (j = 0; j < n; j++) {
          data.c[j][0] = j*2 ;
      for (i = 1; i < m; i++) {
          data.c[j][i] += data.a[i] + data.b[i];
      }
    }
  }

  printf("j i a b c\n");
  for (j = 0; j < n; j++) {
    for (i = 0; i < m; i++) {
      printf("%d %d %lf %lf %lf\n", j, i, data.a[j], data.b[i], data.c[j][i]);
    }
  }

  exit(0);
}

　コンパイルした結果は以下のとおりである。

$  pgcc -O2 -Minfo uf0.c -Msafeptr -acc -ta=tesla,cc60
main:
     21, Loop not vectorized/parallelized: contains call
     30, Generated vector simd code for the loop
         Residual loop unrolled 1 times (completely unrolled)
     38, Generating create(data)
         Generating copyout(data.a[1:n],data.b[1:n])
         Generating copy(data.c[:n][1:m])
     40, Generating present(data.b[1:n],data.a[1:n])
         Accelerator kernel generated
         Generating Tesla code
         41, #pragma acc loop gang, vector(9) /* blockIdx.x threadIdx.x */
     45, Generating present(data.a[1:n],data.c[:n][1:m],data.b[1:n])
     46, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         46, #pragma acc loop gang /* blockIdx.x */
         48, #pragma acc loop vector(32) /* threadIdx.x */
     48, Loop is parallelizable
     56, Loop not vectorized/parallelized: contains call

　OpenACC 2.5 規約(執筆段階ではPGI 17.9）では、C/C++ の構造体への自動 deep copy の指示はできない。規約自体がまだ固まっていない。したがって、プログラム上に deep copy と同じことを directive を使って記述する必要がある。C/C++ 構造体は、内部的に「構造体本体（枠）」とその配下の「構造体メンバの変数実体」の二つからなり、これらをデバイス側にコピーしなければならない。「構造体本体」とは、構造体変数自体のことであり、最初に、これをデバイス側にコピーあるいは create しなければならない。これによってデバイス側の構造体本体のポインタ自身の領域を確保する。このコピーを一般に shallow copy と称する。このコピーだけでは、自動的にその配下の「構造体メンバの変数実体」のコピーは行われない。但し、スカラ変数メンバに限っては shallow copy の操作でコピー（領域確保）される。したがって、その後に、構造体メンバ（静的配列、ポインタメンバ）のコピー(copyin) を明示的にサイズを指定して行う必要がある。これがいわゆる deep copy 操作である。copyin clause 以外でも、create、copy あるいは、copyout の clause を使用しても、初期段階の領域の割付(attach) を行うための指示として機能する。

　ここで重要な点は、shallow copy を行なってから、その後、deep copy の指示をするという順番である。この指示の順番を逆にしてしまうと、デバイス側に構造体自身の枠が出来ていないため、構造体メンバ領域の attach も出来ない状態となり、以降の処理でエラーが生じる。

#pragma acc data create(data) copyout(data.a[1:n],data.b[1:n]) copy(data.c[0:n][1:m])

 1) create(data) : 構造体変数自身dataのshallow copyの指示
 2) copyout(data.a[1:n],data.b[1:n]) copy(data.c[0:n][1:m]) 
　　　　　   　　: 構造体メンバのdeep copyを指示、明示的なサイズを指示する必要あり。

　次に、構造体本体変数(dataという変数）の表記法について説明しよう。このプログラム上では、スカラ構造体 data は、「ポインタ変数」としてではなく「実体変数」として宣言されている。C 言語上、実体変数のメンバの表記は、例えば data.a[0] といったドット演算子表記となる。OpenACC directive 上での当該構造体の表記方法も同様な方法をとって良い。

structdata data;　実体変数として宣言

実体変数として扱う場合のメンバの表記方法
　data.a[0]
　data.b[0]
OpenACC directive 上での当該構造体の表記方法も同じ
　data.a[0]
　data.b[0]

　スカラ構造体 data を「ポインタ変数」として宣言されている場合はどうか？。C 言語上、ポインタ変数のメンバの表記は、例えば data->a[0] といったアロー演算子表記となる。OpenACC directive 上での当該構造体の表記方法も同様な方法をとって良い。一方、構造体の実体が配列構成の場合、当該ポインタ型構造体をアロー演算子表記ではなく、実体変数表記（配列の場合は実体配列表記）でコーディングする場合も少なくない。実体が配列構成の場合は問題は発生しないが、「スカラ構造体」を扱っている場合でコーディング中に混乱し、ポインタ宣言された「スカラ構造体」のメンバの表記を data.a[0] （正しくは、data->a[0]）と書いてしまう場合もある。コンパイラは、C プログラム自体にこのような構文を記述した場合は構文エラーとするので直ぐ分かるが、OpenACC の directive clause の中に、data.a[0] といった記述をしてしまうと、エラーとしないため問題の発見が遅れる。OpenACCの記述ルールでは、ポインタ宣言されたスカラ構造体は、要素数 1 の配列とみなす。すなわち、OpenACC directive 内の記述では、 data[0].a[0:n] といった記述をしなければならない。

structdata *data;　ポインタ変数として宣言

ポインタ変数として扱う場合のメンバの表記方法
　data->a[0]
　data->b[0]
OpenACC directive 上での当該構造体の表記方法も同じ
　data->a[0]
　data->b[0]

*data宣言変数をプログラムの中で、スカラ構造体のポインタ変数ではなく、
実体変数として扱いたい場合の表記法（OpenACC)
　data[0].a[0]
　data[0].b[0]

以下のような表記を OpenACC directive で行うと実行時エラーとなる。スカラ構造体だと誤り易い。
　data.a[0]

　もう一点、「スカラ構造体」自体の shallow copy を行う場合の OpenACC directive 内の表記は、要素数 1 の配列とみなすということから以下のようにする。create(data[0:1]) or copy(datain[0:1]) 等の記述となる。配列要素の notiation ルールで「0 を始点として1個」という記述を行う。create(data[0])といった記述は誤りである。

#pragma acc data create(data[0:1]) copyout(data[0].a[1:n],data[0].b[1:n]) copy(data[0].c[0:n][1:m])

　スカラ構造体 data を「ポインタ変数」として宣言されている場合のプログラムは、後述するが上記の点について十分留意する必要がある。

③ ②のプログラムをプロシジャに分離（構造体をポインタ変数ではない実体変数として宣言）

　以下に示す uf2-func2.c プログラムは、②のプログラムをベースとして、構造体のアロケーション部分と main loop 部分をプロシジャに分離したものである。

　main プログラムの中で、プロシジャ mainloop を呼び、その動作領域を #pragma acc data 構文を使ってデータ領域として定義する。 data 構文の中では、スカラ構造体 data の shallow copy とそれに引き続き、各メンバの attach を行っている。

　main プログラム内のスカラ構造体は、ポインタではなく「実体変数」として宣言されているため、OpenACC directive 内の構造体本体の表記法は、acc data create(data) という風にスカラ変数 data 自身を指定すれば良い。また、各メンバの指定方法は、data.a[1:n]と言うような C言語表記と同じ、ドット演算子表記で良い。

一方、プロシジャ mainloop 内では、スカラ構造体 data がアドレスで渡されているため、ポインタとしてのハンドリングとなる。OpenACC directive の表記としては、構造体メンバは、#pragma acc parallel loop present(data->a[1:n], data->b[1:n])　という風にアロー演算子表記となる。なお、present clause を使って、すでにデバイス側に構造体データは attach されていることをコンパイラに知らせることは必須である。

uf2-func2.c

#include <stdlib.h>
#include <stdio.h>

typedef struct {
  double *a;
  double *b;
  double **c;
} structdata;

void allocData(structdata *data, int n, int m);
void mainloop(structdata *data,int n, int m);

int main() {
// 構造体 ポインタ変数ではない実体変数 "data" として宣言
  structdata data;
  int i,j;
  int n = 10;
  int m = 10;

  allocData(&data,n,m);

//Initilize data members

  for (j=0; j<n; j++) {
     data.a[j] = 0.0;
     data.b[j] = 0.0;
    for (i = 1; i < m; i++) {
     data.c[j][i] = 0.0;
    }
  }

  //Allocate a struct variable, then allocate pointer members in the device
  #pragma acc data create(data) copyout(data.a[1:n],data.b[1:n]) copy(data.c[0:n][1:m])
  {
  mainloop(&data, n,  m);
  }

  printf("j i a b c\n");
  for (j = 0; j < n; j++) {
    for (i = 0; i < m; i++) {
      printf("%d %d %lf %lf %lf\n", j, i, data.a[j], data.b[i], data.c[j][i]);
    }
  }

  exit(0);
}

void allocData(structdata *data, int n, int m) {

  data->a = (double*)calloc(n, sizeof(double));
  data->b = (double*)calloc(n, sizeof(double));
  data->c = (double**)calloc(n, sizeof(double*));
  data->c[0] = (double*)calloc(m, sizeof(double*));

  for (int i = 1; i < n; i++) {
        data->c[i] = data->c[i-1] + m;
  }

}

void
mainloop(structdata *restrict data, int n,int m) {
  int i,j;
    #pragma acc parallel loop present(data->a[1:n], data->b[1:n])
    for (i = 1; i < n; i++) {
          data->a[i] = i + 2;
          data->b[i] = i;
    }
    #pragma acc kernels loop present(data->a[1:n], data->b[1:n], data->c[0:n][1:m])
    for (j = 0; j < n; j++) {
          data->c[j][0] = j*2 ;
      for (i = 1; i < m; i++) {
          data->c[j][i] += data->a[i] + data->b[i];
      }
    }
}

　コンパイルした結果は以下のとおりである。なお、ここではコンパイルオプションに、-Msafeptr を使用していない。mainloop(structdata *restrict data, int n,int m) プロシジャの引数に、restrict 修飾子を使って、ポインタ構造体 *data のメンバにデータの重なり合いがないことをコンパイラに知らせることで、手続内ループの並列化、ベクトル化が可能となる。　

$ pgcc -Minfo -O2 -acc -ta=tesla,cc60 uf2-func2.c
main:
     27, Generated vector simd code for the loop
         Residual loop unrolled 1 times (completely unrolled)
     33, Generating create(data)
         Generating copyout(data.a[1:n])
         Generating copy(data.c[:n][1:m])
         Generating copyout(data.b[1:n])
     40, Loop not vectorized/parallelized: contains call
allocData:
     55, Loop not vectorized: data dependency
         Loop unrolled 8 times
         Generated 1 prefetches in scalar loop
mainloop:
     64, Generating present(data->a[1:n])
         Accelerator kernel generated
         Generating Tesla code
         65, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     64, Generating implicit copyout(data[:1])
         Generating present(data->b[1:n])
     69, Generating present(data->a[1:n],data->b[1:n])
         Generating implicit copyout(data[:1])
         Generating present(data->c[:n][1:m])
     70, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         70, #pragma acc loop gang /* blockIdx.x */
         72, #pragma acc loop vector(128) /* threadIdx.x */
     72, Loop is parallelizable

④ スカラ構造体をポインタ変数として宣言

　以下に示す uf4.c プログラムは、③のプログラムをベースとして、スカラ構造体変数 *data をポインタ宣言したものに変更したものである。なお、ここでは便宜上、data_body という構造体実体も定義した。

　③プログラムとの構成上の違いは、void allocDataプロシジャ内で、デバイス上でのデータ使用開始を指示する「非構造化データ領域ディレクティブ」を使用した点である。実行フローの中で、この時点からデバイス側のデータ利用の開始が行われ、Exit Data ディレクティブが現れるまでデバイス上のデータが保持される。

　void allocDataプロシジャ内で、#pragma acc enter data 構文で、create(data[0:1]) を指示し、スカラ構造体のポインタ data の shallow copy を行う。その次に、各メンバーを data->a[1:n] のようにアロー演算子表記で copyin する。この表記法とは別に、スカラ構造体のポインタ変数であれば、OpenACC では１個の要素を持つ配列と解釈するので、その構造体実体を a[0] といった形で配列表記ができる。この表記を使用すると data[0].a[1:n] のようなドット演算子表記が可能となる（以下のプログラムでは青字で示した directive）。

mainloop プロシジャ内では、#pragma acc parallel loop present(data) というポインタ構造体変数名の指定だけで動作する。このように、当該ループ内で使用する構造体のメンバーのデータ要素を指定しなくても動作する場合がある。上述した②と③のプログラムでは、mainloop 内で使用する構造体のメンバーのデータ要素を明示的に指定しなければ実行時エラーとなった。しかし、本プログラムでは、ポインタ構造体変数名の指定だけで動作する。本来、C/C++言語の場合は、Fortran 言語のように内部的に配列の descriptor を持っていないため、配列の組成状態を有していない。したがって、data clause の中のメンバ配列要素の個数の記述は明示的に行う必要がある（この辺りが Fortran と違って面倒くさい部分である）。本プログラムが動作する理由は、OpenACC 仕様書にも明記されていないため筆者も分からない。ただ、②と③のプログラムと異なる点は、当該プログラムが実行フローの中で、acc data 構文を使用しないで、acc enter data 構文と acc exit data 構文を使って「非構造化データ領域」の指示行を使用していることである。この点に関しては、プログラム開発時の留意事項として欲しい。

　最終的には mainloop 終了後、#pragma acc exit data copyout 指示で、ホスト側に必要なデータをコピーすることで終了する。

uf4.c

#include <stdlib.h>
#include <stdio.h>

typedef struct {
  double *a;
  double *b;
  double **c;
} structdata;

void allocData(structdata *data, int n, int m);
double *allocateDouble(int size);
double **allocateDouble2D(int n, int m);
void mainloop(structdata *restrict data,int n, int m);

int main() {

  structdata *data, data_body ;             // 構造体ポインタ *data 宣言, data_body は実体
  int i,j;
  int n = 10;
  int m = 10;

  data = &data_body;
  allocData(data,n,m);                      // 実引数 ポインタ渡し

  mainloop(data,n,m);

  // copyout a, b, c members　非構造化データ領域終了
  #pragma acc exit data copyout(data->a[1:n],data->b[1:n],data->c[0:n][1:m])
  // #pragma acc exit data copyout(data[0].a[1:n],data[0].b[1:n],data[0].c[0:n][1:m])

  printf("j i a b c\n");
  for (j = 0; j < n; j++) {
    for (i = 0; i < m; i++) {
      printf("%d %d %lf %lf %lf\n", j, i, data->a[j], data->b[i], data->c[j][i]);
    }
  }

  exit(0);
}

// *data ポインタ渡し
void allocData(structdata *data, int n, int m) {

  data->a = allocateDouble(n);
  data->b = allocateDouble(m);
  data->c = allocateDouble2D(n,m);

  // Initilize data members
  for (int j=0; j<n; j++) {
     data->a[j] = 0.0;
     data->b[j] = 0.0;
    for (int i = 1; i < m; i++) {
       data->c[j][i] = 0.0;
    }
  }

  // Allocate a struct variable, then allocate(attach) pointer members in the device
  // ポインタの構造体として指示、非構造化データ領域開始
  #pragma acc enter data create(data[0:1]) copyin(data->a[1:n],data->b[1:n],data->c[0:n][1:m])
  // #pragma acc enter data create(data[0:1]) copyin(data[0].a[1:n],data[0].b[1:n],data[0].c[0:n][1:m])
}


double *allocateDouble(int n) {

  double *array;
  array = (double*)calloc(n, sizeof(double));
  return array;
}

double **allocateDouble2D(int n, int m) {

  double **array;
  array = (double**)calloc(n, sizeof(double*));
  array[0] = (double*)malloc(n*m*sizeof(double));

  for (int i = 1; i < n; i++) {
        array[i] = array[i - 1] + m;
        printf("%d %d %lf\n", n, m, array[i]);
  }
  return array;
}

void
mainloop(structdata *restrict data, int n,int m) {
  int i,j;

    #pragma acc parallel loop present(data)
    for (i = 1; i < n; i++) {
          data->a[i] = i + 2;
          data->b[i] = i;
    }
    #pragma acc kernels loop present(data)
    for (j = 0; j < n; j++) {
          data->c[j][0] = j*2 ;
      for (i = 1; i < m; i++) {
          data->c[j][i] += data->a[i] + data->b[i];
      }
    }

}

　コンパイルした結果は以下のとおりである。

$ pgcc -Minfo -O2 -acc -ta=tesla,cc60  uf4.c
     28, Generating exit data copyout(data->b[1:n],data->c[:n][1:m],data->a[1:n])
     33, Loop not vectorized/parallelized: contains call
allocData:
     52, Memory zero idiom, loop replaced by call to __c_mzero8
     59, Generating enter data copyin(data->b[1:n],data->a[1:n])
         Generating enter data create(data[:1])
         Generating enter data copyin(data->c[:n][1:m])
allocateDouble2D:
     77, Loop not vectorized/parallelized: contains call
mainloop:
     88, Generating present(data[:])
         Accelerator kernel generated
         Generating Tesla code
         89, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     93, Generating present(data[:])
     94, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         94, #pragma acc loop gang /* blockIdx.x */
         96, #pragma acc loop vector(128) /* threadIdx.x */
     96, Loop is parallelizable

⑤【最終形】構造体を利用する場合の OpenACC プログラム例

　上述のプログラムをもう少し、スマートに書いたものが以下ののものである。ここではスカラ構造体ではなく、配列型の構造体を利用できるように変更している。また、#pragma acc update device(DD->a[1:n],DD->b[1:n],DD->c[0:n][1:m]) を使用してデータの更新を行う方法も見て取れるかと思う。

uf5-advance.c

#include <stdlib.h>
#include <stdio.h>

typedef struct {
  double *a;
  double *b;
  double **c;
} structdata;

//prototype
structdata * allocData(size_t size1, int n, int m);
double *allocateDouble(int size);
double **allocateDouble2D(int n, int m);
int deleteData(structdata* DD, int n, int m);
int initData  (structdata * DD, int n, int m, double val);
void mainloop (structdata *restrict data,int n, int m);

int main() {

  structdata *data;              // 構造体ポインタ *data 宣言
  int i,j;
  int n = 10;
  int m = 10;
  size_t size1;
  double val;

  size1 = 1;                     // data[0:1]  個数

  data = allocData(size1, n, m); // Allocation data

  val =5.0;
  initData(data, n, m, val);     // Initilize data
  mainloop(data, n, m);          // kernel loopp in the device

  // copyout a, b, c members to host
  #pragma acc exit data copyout(data->a[1:n],data->b[1:n],data->c[0:n][1:m])
  //#pragma acc exit data copyout(data[0].a[1:n],data[0].b[1:n],data[0].c[0:n][1:m])

  printf("j i a b c\n");
  for (j = 0; j < n; j++) {
    for (i = 0; i < m; i++) {
      printf("%d %d %lf %lf %lf\n", j, i, data->a[j], data->b[i], data->c[j][i]);
    }
  }

  deleteData(data, n, m);        // Delete an allocation area both device and host

  exit(0);
}

structdata * allocData(size_t size1, int n, int m) {
  structdata * data;

  data = (structdata*) malloc(size1*sizeof(structdata));
  #pragma acc enter data create(data[0:size1])              // shallow

  data->a = allocateDouble(n);
  data->b = allocateDouble(m);
  data->c = allocateDouble2D(n,m);

  // Allocate a struct variable, then allocate(attach) pointer members in the device
  // ポインタの構造体として指示
  #pragma acc enter data create(data->a[1:n],data->b[1:n],data->c[0:n][1:m])
  //#pragma acc enter data create(data[0].a[1:n],data[0].b[1:n],data[0].c[0:n][1:m])

  return data;
}

int deleteData(structdata * DD, int n, int m) {
    int i;
    free(DD->a);
    free(DD->b);
    free(DD->c[0]);
    free(DD->c);
    #pragma acc exit data delete(DD->a[0:n], DD->b[0:n], DD->c[:n][:m])
    #pragma acc exit data delete(DD)
    free(DD);
}

int initData(structdata * DD, int n, int m, double val) {
    int i,j;

  // Initilize data members
     for (int j = 0; j<n; j++) {
       DD->a[j] = 0.0;
       DD->b[j] = 0.0;
       for (int i = 1; i < m; i++) {
         DD->c[j][i] = val;
       }
     }
  // Update the device with the initial values
     #pragma acc update device(DD->a[1:n],DD->b[1:n],DD->c[0:n][1:m])
}


double *allocateDouble(int n) {
  double *array;
  array = (double*)calloc(n, sizeof(double));
  return array;
}

double **allocateDouble2D(int n, int m) {
  double **array;
  array = (double**)calloc(n, sizeof(double*));
  array[0] = (double*)malloc(n*m*sizeof(double));

  for (int i = 1; i < n; i++) {
        array[i] = array[i - 1] + m;
        printf("%d %d %lf\n", n, m, array[i]);
  }
  return array;
}

void
mainloop(structdata *restrict data, int n,int m) {
  int i,j;

    //#pragma acc parallel loop present(data[0].a[1:n], data[0].b[1:n])
    #pragma acc parallel loop present(data)
    for (i = 1; i < n; i++) {
          data->a[i] = i + 2;
          data->b[i] = i;
    }
    //#pragma acc kernels loop present(data[0].a[1:n], data->b[1:n], data->c[0:n][1:m])
    #pragma acc kernels loop present(data)
    for (j = 0; j < n; j++) {
          data->c[j][0] = j*2 ;
      for (i = 1; i < m; i++) {
          data->c[j][i] += data->a[i] + data->b[i];
      }
    }

}

　コンパイルした結果は以下のとおりである。

$ pgcc -Minfo -O2 -acc -ta=tesla,cc60 uf5-advance.c
main:
     36, Generating exit data copyout(data->b[1:n],data->c[:n][1:m],data->a[1:n])
     41, Loop not vectorized/parallelized: contains call
allocData:
     55, Generating enter data create(data[:size1])
     63, Generating enter data create(data->c[:n][1:m],data->b[1:n],data->a[1:n])
deleteData:
     75, Generating exit data delete(DD->b[:n],DD->a[:n],DD->c[:n][:m])
     76, Generating exit data delete(DD[:1])
initData:
     87, Memory set idiom, loop replaced by call to __c_mset8
     93, Generating update device(DD->c[:n][1:m],DD->b[1:n],DD->a[1:n])
allocateDouble2D:
    107, Loop not vectorized/parallelized: contains call
mainloop:
    119, Generating present(data[:])
         Accelerator kernel generated
         Generating Tesla code
        120, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
    125, Generating present(data[:])
    126, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
        126, #pragma acc loop gang /* blockIdx.x */
        128, #pragma acc loop vector(128) /* threadIdx.x */

　実行時に、GPU 上で動作しているかどうか確認する場合は、環境変数 PGI_ACC_NOTIFY をセットして実行する。デバイスへのイベント、アクションがあったときに以下のようなメッセージを出力する。

$ export PGI_ACC_NOTIFY=3
$ a.out
（略）
upload CUDA data  file=/home/kato/GPGPU/OpenACC/StructClass/IHI-sasou/STEPS/uf5-advance.c function=allocData
 line=63 device=0 threadid=1 variable=.pointer. bytes=8
launch CUDA kernel  file=/home/kato/GPGPU/OpenACC/StructClass/IHI-sasou/STEPS/uf5-advance.c 
 function=mainloop line=119 device=0 threadid=1 num_gangs=1 num_workers=1 vector_length=128 
 grid=1 block=128
（略）

ページの先頭へ

[Reference]

Editor by Rob Farber, Parallel Programming with OPENACC Chapter 5

前章へ

次章へ

OpenACCプログラミングのインデックスへ

ページの先頭へ

OpenACC ディレクティブによるプログラミング

14章 C 構造体、C++ クラス、Fortran 派生型のデータ管理

1. 構造体、派生型変数の deep copy データ転送

2. OpenACC における Fortran 派生型の扱い

手動による一般的なディレクティブ指示の方法

full deepcopy 機能を使った場合のディレクティブ指示の方法

CUDA Unified Memory を利用した場合のディレクティブ指示の方法

3. Fortran 派生型変数の使用例２

手動による一般的なディレクティブ指示の方法

full deepcopy 機能を使った場合のディレクティブ指示の方法

CUDA Unified Memory を利用した場合のディレクティブ指示の方法

4. OpenACC における C 構造体データ

手動による一般的なディレクティブ指示の方法

array_of_structs.c

CUDA Unified Memory を利用した場合のディレクティブ指示の方法

array_of_structs.c

5. OpenACC における C++クラスデータ

手動による一般的なディレクティブ指示の方法

myList.h

myList.cpp

6. C 構造体データを使用したプログラムのポーティング・チュートリアル

① ベースとなるプログラム（構造体をポインタ変数ではない実体変数として宣言）

ufbase.c

② ベースを OpenACC でデバイスコード化（構造体をポインタ変数ではない実体変数として宣言）

uf0.c

③ ②のプログラムをプロシジャに分離（構造体をポインタ変数ではない実体変数として宣言）

uf2-func2.c

④ スカラ構造体をポインタ変数として宣言

uf4.c

⑤【最終形】構造体を利用する場合の OpenACC プログラム例

uf5-advance.c

14章　C 構造体、C++ クラス、Fortran 派生型のデータ管理

6.　C 構造体データを使用したプログラムのポーティング・チュートリアル