PGIトップ PGI技術情報・TIPS › PGI テクニカル情報・コラム › Multi-GPU プログラミング (UVA)

CUDA Fortran による CUDA 4.0 Multi-GPU プログラミング (2)

(Unified Virtual Addressing の使用法)

キーワード GPGPU、CUDA Fortran、UVA, GPUdirect

 マルチGPUを備えたシステム上で、CUDA 4.0 からの機能である Unified Virtual Addressing(UVA) や GPUdirect v2.0 の機能の使用方法を解説する。今回は、ホスト側のメモリ領域も含めて Unified Virtual 空間として捉えた場合のプログラミングの方法を解説する。
2011年10月11日 Copyright © 株式会社ソフテック 加藤

CUDA 4.0 Unified Virtual Addresing 機能

 前回のコラム、CUDA Fortran による CUDA 4.0 Multi-GPU プログラミング (1)(GPUdirect™ V2.0 / peer-to-peer Access) では、マルチデバイス環境での devie to device 間のデータ通信の方法を CUDA 4.0 の GPUdirect 機能で行う方法を解説しました。今回は、Unified Virtual Addresing の機能に焦点を当てます。その例として、ホスト側のメモリ領域を Unified Virtual Addresing で使用する方法を具体的に説明します。Unified Virtual Addresing(UVA)機能とは、ホスト・メモリや複数のデバイス・メモリの空間を一つのメモリ空間として扱う技術です。これは CUDA 4.0 でリリースされた技術です。この UVA の機能は、64ビットモードの Linux あるいは、Windows 上で動作する NVIDIA Fermi GPU(compute capability 2.0以上)以降で利用できます。Windows の場合は、さらに TCC ドライバーモードでしか動作しませんので、Tesla GPU 以上のボードの動作だけに制約されます。

使用する CUDA Runtime API

 CUDA Fortran から NVIDIA CUDA C Library に含まれるRuntime APIを使用する際の注意点は、必ず、PGI CUDA Fortran Programming Guide の API の説明を確認することです。CUDA Fortran が提供している NVIDIA CUDA C Library へのインタフェースは、Fortran スタイルで使えるように、引数の単位や型等を本来の CUDA C の関数のものとは異なったインタフェースで提供しているものがあります。例えば、CudaMemcpy... 関係の API の場合、本来の C 言語型の API では、src/dst を C 型ポインタ(c_ptr)とバイト数を引数としていますが、CUDA Fortran では、この仕様の他に、Fortran型のポインタ(アドレス渡し)とデータ個数を引数とするインタフェースも使えるようにしています。こうした Fortran ライクな引数も使用できるものもありますが、ほとんどの CUDA API では、CUDA C の API 関数の引数の属性(C型ポインタ)をそのまま使えるようになっています。CUDA Fortran 上の CUDA API ライブラリのデフォルトは、NVIDIA CUDA C言語で構築された API をそのままインタフェースする形でしょうから、その引数が C 言語型のポインタの場合、これを想定した Fortran 上のハンドリングを行う必要があります。

 NVIDIA CUDA C 言語の世界で作成された CUDA ライブラリは、基本的に C 言語の世界のプログラム仕様となっています。こうした API 関数を CUDA Fortran 上で使用する場合、Fortran の世界と C の世界の「橋渡し」が必要です。この「橋渡し」とは、例えば、変数の「型」のマッチングや「C ポインタ世界と Fortran ポインタの世界」の変換と言ったことです。これをきちんと体系立てて行うための構文(C と Fortran の相互互換性維持のための構文)が、Fortran 2003(F2003) から提供されています。以下のプログラム例では、F2003 から導入された iso_c_binding モジュールを使って、これを実現するための方法が記述されていますので参考にして下さい。大事な点は、ライブラリ、API の生い立ちが C 言語体系の中で作成されたルーチンの場合、これを直接 Fortran ベースで使用するには、必ず、その引数を C 言語体系の世界へ接続するための(変換)処理が必要です。これを実現するのが、F2003 の iso_c_binding モジュール並びに、その補助ルーチン(c_f_pointer関数等)です。

 以下に示した CUDA Fortran 構文は、Fortran 2003 を使って、C と Fortran の相互互換を行うために使用する構文の例です。プログラム中で使う CUDA Fortran runtime API の引数の中で、「ポインタ」変数に関しては、それが CUDA C 言語の世界のポインタ属性なのか、Fortran の世界のポインタ属性なのかを確認しながらコーディングすると間違いがありません。なお、CUDA Fortran でコーディングしている場所では、その引数の引き渡しの際に、必ず、Fortran の世界のポインタ変数(従来のアドレス渡しと言う意味)を使うことが必要です。なお、C 型ポインタと Fortran 型ポインタを変換するために、F2003 のライブラリ関数 c_f_pointer() をが用意されています。これで余計なことを考えずに変換できます。

use, intrinsic :: iso_c_binding
real, dimension(:), pointer :: ha                    ! Fortran 型のポインタ配列
type(c_ptr) :: local_ha                              ! C 型のポインタ
real, dimension(:), device, allocatable :: da        ! Fortran 型のデバイス配列
type(c_devptr) :: local_da                           ! C 型のデバイスポインタ
call c_f_pointer( local_da, da, (/ n /) )            ! C 型ポインタから Fortran 型ポインタへ変換

 以下のプログラム例で使用する、UVA 環境の記述で必要な CUDA API関数を列挙しました。これらの関数の中で使用されている引数が、C 型ポインタ(c_ptr)であるかどうかは、CUDA Fortran Programming Guide and Reference で確認することが必要です。引数が、C型ポインタの場合は、その型が、c_ptr(ホスト側のC型ポインタ) あるいは、c_devptr(デバイス側のC型ポインタ)として定義されています。

istat = cudasetdeviceflags(cudaDeviceMapHost)
        cudaSetDeviceFlags records how the CUDA runtime interacts with this host thread
        ホスト側スレッドのメモリに mapped(pinned) ホストメモリを割り付けることを指示する
        
istat = cudahostalloc( local_ha, sizeof(1.0)*n, cudaHostAllocMapped )
        type(C_PTR) :: hostptr = local_ha
        cudaHostAlloc allocates pinned memory on the host. It returns in hostptr the address o
        the pagelocked allocation, or returns an error if the memory is unavailable. Size is
        in bytes. The flags argument enables different options to be specified that affect 
        the allocation. The normal iso_c_binding subroutine c_f_pointer can be used to move 
        the type(c_ptr) to a Fortran pointer.

istat = cudahostgetdevicepointer( local_da, local_ha, 0 )
        type(C_DEVPTR) :: devptr = local_da
        type(C_PTR) :: hostptr   = local_ha
        integer :: flags
        cudaHostGetDevicePointer returns a pointer to a device memory address corresponding
        to the pinned memory on the host. hostptr is a pinned memory buffer that was allocated
        via cudaHostAlloc(). It returns in devptr an address that can be passed to, and read
        and written by, a kernel which runs on the device. The flags argument is provided for
        future releases. The normal iso_c_binding subroutine c_f_pointer can be used to move 
        the type(c_devptr) to a device array.

GPUデバイスからホスト側のメモリ領域を Unified Virtual Addressing(UVA) の管理下で使用する

 ホスト側メモリ領域に配列を確保し、これを UVA 配下のアドレス空間の一つとして設定します。 GPU 側の kernel 実行する際に、このホスト上の配列を使用する簡単なプログラム例を説明します。以下に、コメント付きのサンプルプログラムを示しますので、これを参考にしていただきたいと思います。プログラムの流れは、以下の通りですが、CUDA Fortran からこうした UVA 関係の CUDA API を使用する際は、必ず、PGI CUDA Fortran Programming Guide の API の説明を参照して下さい。特に、CUDA API関数の引数が、「C型ポインタ」として扱われているのかどうかの確認が必要です。

  1. ホスト側のメモリが pinned モードで割付可能な場合、cudasetdeviceflags(cudaDeviceMapHost)関数を使用して、pinned メモリ割付を行うことを指示する。
  2. cudahostalloc() を使用し、ホスト側に pinned memory 領域にアロケートする。さらに、これに対して Zero-copy memory として mapped 属性を指示する。その際の CUDA API 管理下のホスト側領域のポインタ(C 型ポインタ)を local_ha とする。
  3. NVIDIA CUDA API の世界のポインタ local_ha(C 型ポインタ)を Fortran の世界のポインタ ha に変換する。 ホスト側の配列 ha にデータを入れる。即ち、これは、local_ha の場所にデータが入ると言う意味となる。
  4. ホスト側の local_ha ポインタ領域を UVA 下の「一つのデバイス」領域として認識し、UVA 管理下のポインタ(local_da) に付け替える。これが、Unified Virtual Addressing の核となるところである。これは、cudahostgetdevicepointer( local_da, local_ha, 0 ) と c_f_pointer( local_da, da, (/ n /) ) で行う。
  5. カーネルの実行を行う。カーネルへの引数には、CUDA Fortran の世界のデバイス側の配列 da を指定する。
  6. カーネルの実行後、その計算結果は、ホスト側の ha 配列にストアされている。デバイス側の da から ha への明示的なコピーは必要ない。
module square_kernel
 implicit none
contains
 ! 簡単な kernel プログラム
 ! Square the entries of 'a'
 attributes(global) subroutine s_kernel( a )
  real, dimension(*) :: a
  integer i
  i = (blockidx%x-1)*blockdim%x + threadidx%x
  a(i) = a(i) * a(i)
 end subroutine
end module

program main
 ! C言語関数との interface が必須となるため iso_c_binding 使用
 use, intrinsic :: iso_c_binding
 use cudafor
 use square_kernel
 implicit none

 ! CUDA Runtime API の中のC言語ベースのライブラリを使用するため、C ポインタ体系の世界
 ! そこで、Fortran上では、C ポインタ と Fortran ポインタの変換が必要となる
 ! 二つのポインタを用意しておく
 real, dimension(:), pointer :: ha      ! host pointer (Fortran pointer)
 type(c_ptr) :: local_ha                ! C pointer which we convert to ha
 real, dimension(:), device, allocatable :: da   ! device pointer (Fortran pointer)
 type(c_devptr) :: local_da             ! C pointer which we convert to da

 integer :: i, n, err, istat
 real :: e

 ! Properties check
 type(cudaDeviceProp) :: prop
 integer :: idevice, ilen

 ! Device properties
 idevice = 0
 istat = cudaSetDevice(idevice)
 istat = cudaGetDeviceProperties(prop,idevice)
 ilen = verify(prop%name, ' ', .true.)
 write (*,900) prop%name(1:ilen), &
               real(prop%clockRate)/1000.0, &
               real(prop%totalGlobalMem)/1024.0/1024.0, &
               prop%canMapHostMemory
900 format(/,' Device name: ',a,/,1x, f6.1,' MHz clock, ',f6.1,' MB memory.'/, &
           ' canMapHostMemory Flag : ', i2/)

     ! ホスト側で pinned メモリの使用が可能かどうか確認する
     ! UVA 下では、必ず mapped (Zero-copy memory) ホストメモリを割り付けることを指示する
     ! (cudaDeviceMapHost)
     ! This flag must be set in order to allocate pinned host memory
     ! that is accessible to the device.

 if (prop%canMapHostMemory == 1) then
  istat = cudasetdeviceflags(cudaDeviceMapHost)
 !istat = cudasetdeviceflags(cudaDeviceScheduleBlockingSync)
 end if

 n = 2048

     ! cudahostalloc を使用し、ホスト側に pinned and mapped memory 領域をアロケートする
     ! その際の CUDA API 管理下の(世界の)ホスト側領域のポインタを local_ha とする
     ! Allocate (mapped) pinned host memory.
     ! Maps the allocation into the CUDA address space. The device pointer to the memory
     ! may be obtained by calling cudaHostGetDevicePointer().
     ! local_ha : type(c_ptr) = C pointer

 istat = cudahostalloc( local_ha, sizeof(1.0)*n, cudaHostAllocMapped )

     ! CUDA API の世界のポインタ local_ha を Fortran の世界のポインタに変換する
     ! Convert C(local_ha) pointer to Fortran pointer

 call c_f_pointer( local_ha, ha, (/ n /) )

     ! ホスト側の配列 ha にデータを入れる= local_ha の場所にデータが入る
 do i = 1, n
  ha(i) = i+2
 enddo

     ! ホスト側の local_ha ポインタ領域を UVA 下の「一つのデバイス」領域として認識し、
     ! UVA上のポインタ(local_da) に付け替える
     ! convert the host pointer to a device address
     ! Passes back the device pointer corresponding to the mapped,
     ! pinned host buffer allocated by cudaHostAlloc() .
     ! local_da : Returned device pointer for mapped memory : type(C_DEVPTR)
     ! local_ha : Requested host pointer mapping : type(C_PTR)

 istat = cudahostgetdevicepointer( local_da, local_ha, 0 )

     ! Convert C pointer to Fortran pointer

 call c_f_pointer( local_da, da, (/ n /) )
 
      ! GPUデバイスには、UVA 管理下のポインタ(local_da)を渡すことが必要
      ! プログラムは Fortran の世界で coding しているため、s_kernel への引数には、
      ! Fortran の世界の da ポインタを渡す必要あり
      
 call s_kernel <<< 16,128 >>> ( da )
     
     ! Wait for the kernel to finish before testing the results
 istat = cudaThreadSynchronize();

     ! UVA 下では、デバイス上の演算結果は、 da ポインタ= ha ポインタが示すホスト側領域に
     ! ストアされているため、デバイス側からホストへのデータ転送は必要ない
     ! At thist point, da(i) ( = ha(i) ) should be (i+2)**2
 err = 0
 do i = 1, n
  e = float(i+2)
  e = e * e
  if( ha(i) - e .ne. 0 )then
   if( err <= 10 )then
    write(*,10) i, ha(i), e
10  format( 'a(',i4,') = ', e14.8, ', expecting ', e14.8 )
   endif
  endif
 enddo
 print *, err, ' errors found'

 ! shutdown ... host pinned memory
 istat = cudaFreehost( local_ha )
!print *, 'error code = ', cudaGetErrorString( cudaGetLastError() )

end program

 現時点での最新バージョン PGI 11.9 を使用してコンパイルを行います。PGI 11.6~11.9 では、そのデフォルトが CUDA 3.2 の toolkit を利用しますので、明示的に -Mcuda:4.0 として CUDA 4.0 の toolkit を利用してバイナリを作成する必要があります。このプログラムでは、演算の結果に問題があれば、エラーを発行します。エラーが無ければ、問題なくホスト側のメモリを使用してGPU上で計算が行われています。但し、性能は遅いでしょうが・・・

[kato@photon29]$ pgf90 -Mcuda:4.0 -fast -Minfo uva.cuf
main:
     64, Loop unrolled 8 times
     87, Loop not vectorized/parallelized: contains call
[kato@photon29]$ ./a.out

 Device name: GeForce GTX 580
 1544.0 MHz clock, 1535.2 MB memory.
 canMapHostMemory Flag :  1

            0  errors found