CUDA Fortran による CUDA 4.0 Multi-GPU プログラミング (1)
（GPUdirect™ V2.0 / peer-to-peer Access)
GPGPU、CUDA Fortran、UVA, GPUdirect, peer to peer
マルチGPUを備えたシステム上で、CUDA 4.0 からの機能である Unified Virtual Addressing や GPUdirect v2.0 の機能の使用方法を解説する。今回は、デバイス間の peer-to-peer メモリアクセスを利用する際のプログラムの方法を説明する。
2011年10月6日 Copyright © 株式会社ソフテック 加藤
CUDA 4.0 以降のマルチ GPU プログラミング
CUDA 3.2 以前においては、一つのホスト・アプリケーションから複数の GPU (CUDA devices) を使用する場合は、必ず、1スレッド（プロセス）に特定の 1GPU device （以下、「デバイス」と言う。）を割り当てると言う構成でソフトウェアを組む必要がありました。すなわち、一つの「ホスト・スレッド」から「複数」の GPU device（以下、「デバイス」と言う。）を使用することができませんでした。従って、MultiGPU を有する筐体マシンでは、ホスト側において、MPI 等でマルチ・プロセスに分割し、その個々のプロセスに個別のデバイスを割当て、並列 GPU デバイス実行を行うと言う方法がとられてきました。もちろん、こうした構成は現状でも問題はなく、素直なデバイスの使用法と思います。
CUDA 4.0 になると、大きな変化がありました。一番大きな改善点は、ホスト側の一つのスレッド上で、複数のデバイスを切り替えて使用できると言う点です。即ち、複数のデバイスを自由に選択してハンドリングできるため、並列実行プログラミングの自由度が増しました。ホスト側のプログラム・シーケンスの中で、cudaSetDevice() API を使用して特定デバイスを指定すると、これが使用すべき「カレント・デバイス」となります。それ以降のシーケンスでそのデバイスへの命令を発行できます。もし、別のデバイスを使用したいならば、再度 cudaSetDevice() で切り替えると言う方法をとります。こうした方法をとることにより、ホスト側のシングルスレッドのプログラムから、複数のGPUへデータの移動、カーネルプログラムの投入等が自由にできます。これは、マルチ GPU プログラミングが容易になったと言うことに他なりません。
このコラムでは、まず最初に、(1) 一つのホスト側プログラムから2つのGPUデバイス上のデータを ping-pong する CUDA Fortran プログラムを例示します。次に、(2) 片側の GPU 上のカーネル実行において、もう一つのデバイス側のデータを使用する方法を例示します。これらは、CUDA 4.0 (Fermi以上のGPU) 以降でサポートした、Unified Virtual Addressing(UVA) と GPUdirect™ v2.0 peer-to-peer 通信機能を使用することで実現できます。なお、この UVA 関連の機能は、64ビットモードの Linux あるいは、Windows 上で動作する NVIDIA Fermi GPU（compute capability 2.0以上）以降で利用できます。Windows の場合は、さらに TCC ドライバーモードでしか動作しませんので、Tesla GPU 以上のボードの動作だけに制約されます。Linux の場合は、以下のように GeForce 系の GPU でも動作しています。
次回のコラムでは、ホスト側のメモリと Unified Virtual Addressing を行う方法について解説します。
（1）2つの GPU デバイス間でデータを交換する GPUdirect v2.0 peer-to-peer アクセス
一つのホスト側プログラムから2つのGPUデバイス上のデータを ping-pong する CUDA Fortran プログラムを用いて、UVA 配下でのデバイス間 peer to peer データ移動（GPUdirect v2.0）の方法を説明します。以下に、コメント付きのサンプルプログラムを示します。プログラムの流れは、以下の通りです。
- システム内の GPU デバイス数を確認し、その二つの GPU を使用する
- 各デバイスが、peer-to-peer アクセスが可能なデバイスかをチェックする（cudaDeviceCanAccessPeerを使用）。また、UVA 機能を使用できるかを確認する。
- cudaDeviceEnablePeerAccess の APIで、device 0 から device 1 への peer access を有効にする（デフォルトは disable となっている）。また、device 1 から device 0 への peer access も有効にする。これで、peer to peer アクセスの準備が完了する。
- 各デバイス上に、配列 A と B のバッファ・エリアをアロケートする。これは、cudaSetDevice()関数で、デバイスを切り替えてアロケートする。ホスト側の A(1,N) データを device 0 の Ad0(1:N) にコピーする。また、同時にホスト側の B(1,N) を device 1 の Bd1(1:N) にコピーする。これらのデータを各相手方のデバイスへコピーする（ping-pongする）。
- デバイス間のデータ通信に、UVA 環境下での cudaMemcpyPeer 関数を使用する。device 0 から device 1 へのコピーは、cudaMemcpyPeer(Ad1,dev1,Ad0,dev0,N)、逆に、 device 1 から device 0 へのコピーは、cudaMemcpyPeer(bd0,dev0,bd1,dev1,N) として、実行する。これらの関数は、各デバイス上での処理は同期型となっているので、命令処理が終了するまで、次の実行シーケンスへは移らない。これを双方向（ping-pong) で同時実行する。この ping-pong を50回繰り返して、平均の通信バンド幅を計算する。また、4byte 送信の片方向通信時間を計測して、レイテンシーを計測する。
- 最後に、データが正常にコピーしているかを検証するために、device 0 にコピーされたBd0(1:N)の内容をホスト側 C 配列にコピーして、その値の検証を行う。エラーが無ければ、出力の最後に no errors と表示される。
module ping_pong
use cudafor
! integer(c_int), parameter :: cudaMemcpyDefault=4
! Peer-to-Peer Memory Access under Unified virtual addressing
! Ping-Pong Test
! Kato, SofTek
contains
subroutine pingpong( A, B, C )
implicit none
real(4), dimension(:) :: A, B, C
real(4), device, allocatable, dimension(:):: Ad0, Bd0
real(4), device, allocatable, dimension(:):: Ad1, Bd1
integer :: N, istat, dev0, dev1, canaccess0_1, canaccess1_0
integer :: ndevices, uva0, uva1
integer :: nreps = 100, i
logical :: has_uva
type(cudaEvent) :: startEvent, stopEvent
real (4) :: time, gbytes
! Properties check
type(cudaDeviceProp) :: prop
integer :: ilen
N = size( A, 1 )
! アクティブな GPU デバイス数をチェック
print *,"\nChecking for multiple GPUs..."
istat = cudaGetDeviceCount(ndevices)
print '(2x,i2,a)', ndevices, " GPU devices are found."
if ( ndevices < 2 ) then
print '(2x,a/)', " This system don't have two GPUs and above."
stop
end if
! 各デバイスのプロパティ（名前）をプリント
do i = 0, ndevices - 1
istat = cudaSetDevice(i)
istat = cudaGetDeviceProperties(prop,i)
ilen = verify(prop%name, ' ', .true.)
print '(1x, a, a, i1,a)', prop%name(1:ilen), " (GPU",i, ") is available"
end do
! 二つのデバイスに論理デバイス番号を定義
print *, "\n Now, using 2 devices (GPU0 and GPU1) "
dev0=0
dev1=1
print *, "\nChecking GPU(s) for support of peer to peer memory access...\n"
! peer-to-peer アクセスが可能なデバイスかをチェック
! Determine if Peer to Peer communication is supported
istat = cudaDeviceCanAccessPeer(canaccess0_1,dev0,dev1) ! check dev0 to dev1
istat = cudaDeviceCanAccessPeer(canaccess1_0,dev1,dev0) ! check dev1 to dev0
! 双方向でpeer-to-peerアクセスが可能であれば、以下実行
if (canaccess0_1 == 1 .and. canaccess1_0 == 1 ) then
print '(a, i4/,a,i4/)'," canaccess device 0 to 1 =", canaccess0_1, &
" canaccess device 1 to 0 =", canaccess1_0
! dev0 に対して UVA 機能が使用可能か？尋ねる
istat = cudaSetDevice(dev0)
istat = cudaGetDeviceProperties(prop,dev0)
print '(1x,a,i2)', "device0 supports UVA ? : ",prop%unifiedAddressing
uva0 = prop%unifiedAddressing
! dev1 への peer access を enable にする
istat = cudaDeviceEnablePeerAccess(dev1,0)
! dev1 に対して UVA 機能が使用可能か？尋ねる
istat = cudaSetDevice(dev1)
istat = cudaGetDeviceProperties(prop,dev1)
print '(1x,a,i2)', "device1 supports UVA ? : ",prop%unifiedAddressing
uva1 = prop%unifiedAddressing
! dev0 への peer access を enable にする
istat = cudaDeviceEnablePeerAccess(dev0,0)
! 両デバイスともUVA をサポートしているかを確認する
! supprt UVA?
has_uva = ( uva0 == 1 .and. uva1 == 1 )
if (has_uva) then
print *, "\nBoth GPUs can support UVA, enabling\n"
else
print *,"\nAt least one of the two GPUs does NOT support UVA\n"
stop 999
end if
! current device = dev1、デバイス上に配列割付
! allocate device 1's arrays
allocate( Ad1(N), Bd1(N) )
! ホスト側 B(1,N) を Bd1(1:N) にコピーする
! copy host data to device 1 (Bd1)
Bd1 = B(1:N)
!デバイス切替 current device = dev0、デバイス上に配列割付
! allocate device 0's arrays
istat = cudaSetDevice(dev0)
allocate( Ad0(N), Bd0(N) )
! ホスト側 A(1,N) を Ad0(1:N) にコピーする
! copy host data to device 0 (Ad0)
Ad0 = A(1:N)
! 時間計測用の event 作成
istat = cudaEventCreate(startEvent)
istat = cudaEventCreate(stopEvent)
! timing experiment
time = 0.0e0
istat = cudaEventRecord(startEvent, 0)
! ping-pong 往復 50回繰り返し、時間測定
do i = 0, nreps-1
! ======= ! ping-pong between dev0 and dev1
if (mod(i,2) == 0) then
! copy Ad0 data from device 0 to device 1 i が偶数の場合
istat = cudaMemcpyPeer(Ad1,dev1,Ad0,dev0,N)
else
! copy bd1 data from device 1 to device 0 i が奇数の場合
istat = cudaMemcpyPeer(bd0,dev0,bd1,dev1,N)
end if
! 以下の API でも通信が可能。PGI 11.8 では cudaMemcpydefault が未サポート
! Direct access to peer memory on "UVA"
! It is necessary that cudaDeviceEnablePeerAccess() is enable.
! istat = cudaMemcpy(Ad1,Ad0,N,cudaMemcpyDeviceToDevice)
! istat = cudaMemcpy(Bd0,Bd1,N,cudaMemcpyDeviceToDevice)
! cannot use 'cudaMemcpydefault' on PGI 11.8. It may not support yet.
! istat = cudaMemcpy(Ad1,Ad0,N,cudaMemcpydefault)
! istat = cudaMemcpy(Bd0,Bd1,N,cudaMemcpydefault)
! =======
end do
! 時間測定終了
istat = cudaEventRecord(stopEvent, 0)
istat = cudaThreadSynchronize()
istat = cudaEventElapsedTime(time, startEvent, stopEvent)
time = time / (nreps*1.0d3)
gbytes = N * sizeof(a(1)) / ( 1024.*1024.*1024. )
! 通信バンド幅性能を印字
print '(a,F7.2,a)', "PingPong performance : ", gbytes*1024. ," Mbytes Xfer Test"
print '(1x,a,F6.2,a)', "cudaMemcpyPeer / cudaMemcpy between device 0 and device 1 :", &
gbytes/time , " (GB/sec)"
! データ転送の結果を検証、一例で、dev0 にコピーされたBd0(1:N)の値を検証、ホスト側のC配列にコピー
! verify Bd0() as a result : store Bd0 to C array
istat = cudaSetDevice(dev0)
C(1:N) = Bd0
! 通信レイテンシーを測定、ここでは 4-byte 送出時のレイテンシー
! Latency (4-byte transfer time)
istat = cudaEventCreate(startEvent)
istat = cudaEventCreate(stopEvent)
! timing experiment
time = 0.0e0
istat = cudaEventRecord(startEvent, 0)
! 4byte dev0 to dev1 通信時間測定
do i = 0, nreps-1
! copy Ad0 4bytes data from device 0 to device 1
istat = cudaMemcpyPeer(Ad1,dev1,Ad0,dev0,1)
end do
istat = cudaEventRecord(stopEvent, 0)
istat = cudaThreadSynchronize()
istat = cudaEventElapsedTime(time, startEvent, stopEvent)
time = time / (nreps) / 1.d3 * 1.d6
print '(a,f5.2,a)', " 4-byte Ping Latency (one-way) : ", time ," usec"
! Shutdown devices
print *, "\nShutting down...\n"
! 各デバイスのクリーンアップ
! deallocate the arrays on device 1 and cleanup
istat = cudaSetDevice(dev1)
deallocate( Ad1, Bd1 )
istat = cudaDeviceDisablePeerAccess(dev1)
! deallocate the arrays on device 0 and cleanup
istat = cudaSetDevice(dev0)
deallocate( Ad0, Bd0 )
istat = cudaDeviceDisablePeerAccess(dev0)
istat = cudaEventDestroy(startEvent)
istat = cudaEventDestroy(stopEvent)
endif
!print *, 'error code = ', cudaGetErrorString( cudaGetLastError() )
end subroutine
end module
program p
use ping_pong
use cudafor
implicit none
integer :: nargs, n, i, ierr
real, allocatable, dimension(:) :: a, b, c
external iargc
integer iargc
character*10 arg
real(8) :: Bytes, Mbytes
! for timing
real(8) :: time
integer(4) :: hz, clock0, clock1
nargs = iargc()
if( nargs == 0 )then
print *, 'Instruction : ./a.out xfer-size (Unit: Mbyte)'
print *, 'For example, ./a.out 1.0 (this means 1.0MB)'
return
endif
if( nargs >= 1 )then
call getarg( 1, arg )
read(arg,'(f)') Mbytes
endif
if ( Mbytes < 1.0d-8) then
print *, "Mbyte-size(argument) should be defined as a floating-point value"
stop
endif
!print *, MBytes
Bytes = MBytes * (1024.*1024.)
n = Bytes / 4.0
!print *, n , sizeof(a(1))
allocate( a(n), b(n), c(n) )
do i = 1,n
a(i) = i
b(i) = 2*i
c(i) = 0.
enddo
! Initialize time function
call system_clock(count_rate=hz)
! print *, "system_clock resolution: ", real(1.d0/hz)
! Start a clock
call system_clock(count=clock0)
call pingpong( a, b, c ) ! ping-pongの実行
call system_clock(count=clock1)
! time= real((clock1-clock0)) / real (hz)
! print *, "Total Time (sec): ",time
! C 配列の内容を検証する
ierr = 0
do i = 1,n
if( c(i) .ne. b(i) )then
ierr = ierr + 1
if( ierr <= 10 )then
print 30, i, c(i), b(I)
30 format( 'c(',i,') = ', f12.5, ' should be ', f12.5 )
endif
endif
enddo
if( ierr == 0 )then
print *, 'no errors'
else
print *, ierr, ' errors found'
endif
end program
現時点での最新バージョン PGI 11.8 を使用してコンパイルを行います。PGI 11.8 では、そのデフォルトが CUDA 3.2 の toolkit を利用しますので、明示的に -Mcuda:4.0 として CUDA 4.0 の toolkit を利用してバイナリを作成する必要があります。以下は、0.1 MB (100KB) 転送を指定したときの実行結果です。実行時の引数に 0.1 を指定します。なお、この引数は、「整数」の指定ではなく、小数点数値で指定して下さい。なお、私の GPU デバイス環境では、通信バンド幅が 4.9GB/sec （仕様上は 6GB/sec?)で飽和するようです。大事な点は、通信レイテンシーが 6 μ秒程度で小さいと言うことが挙げられます。ここで示した性能は、UVA下の GPUdirect v2.0 の機能を使用したもので、PCIバスに繋がるホスト側インタフェースを経由しない device to device の性能です。
[kato@photon29]$ pgf90 -Mcuda:4.0 pingpong.cuf -V
PGF90/x86-64 Linux 11.8-0
Copyright 1989-2000, The Portland Group, Inc. All Rights Reserved.
Copyright 2000-2011, STMicroelectronics, Inc. All Rights Reserved.
[kato@photon29]$ ./a.out 0.1
Checking for multiple GPUs...
2 GPU devices are found.
GeForce GTX 580 (GPU0) is available
GeForce GTX 480 (GPU1) is available
Now, using 2 devices (GPU0 and GPU1)
Checking GPU(s) for support of peer to peer memory access...
"1" means "Yes", "0" means No"
canaccess device 0 to 1 = 1
canaccess device 1 to 0 = 1
device0 supports UVA ? : 1
device1 supports UVA ? : 1
Both GPUs can support UVA, enabling
PingPong performance : 0.10 Mbytes Xfer Test
cudaMemcpyPeer / cudaMemcpy between device 0 and device 1 : 3.15 (GB/sec)
4-byte Ping Latency (one-way) : 5.78 usec
Shutting down...
no errors
それでは、ホスト側インタフェースを staging として経由すると peer-to-peer 性能はどのようになるでしょうか？上記のプログラムをちょっと変更すると、その性能を測定できます。以下のように、cudaDeviceEnablePeerAccess（）をコメント化して、PeerAccessを disable にすることです。この状態で、cudaMemcpyPeer() APIを使用しても、device to device コピーは正常に動作しますが、その性能は劣化します。小さいデータ転送時のバンド幅が劣化し、さらに、レイテンシーも 20 μ秒弱に増加しています。一般に、バンド幅より、レイテンシーの方がアプリケーション性能に寄与しますので、CUDA 4.0 の inter-Device のインタフェースは、デバイス間の低レイテンシーの通信機能を提供するものとなります。
上記のソースの中の以下の関数をコメント化する。
! istat = cudaDeviceEnablePeerAccess(dev1,0)
! istat = cudaDeviceEnablePeerAccess(dev0,0)
[kato@photon29]$ ./a.out 0.1 (snip) PingPong performance : 0.10 Mbytes Xfer Test cudaMemcpyPeer / cudaMemcpy between device 0 and device 1 : 1.93 (GB/sec) 4-byte Ping Latency (one-way) : 19.88 usec Shutting down...
さて、peer-to-peer メモリコピーを行う API は他にもありますが、これらは内部的には全て cudaMemcpyPeer() 関数に置換され使用されます。例えば、cudaMemcpy() の kind = cudaMemcpyDeviceToDevice と指定した場合、CUDA 3.2 以前は、1 デバイス内のメモリコピー機能でしたが、もし、cudaDeviceEnablePeerAccess() を enable にしておけば、デバイスが異なっていても、コピーできます。また、UVA 環境下では、cudaMemcpy() の kind 引数に、cudaMemcpydefault が使用できるようになりましたが、CUDA Fortran 11.8/11.9 リリースでは、まだ、サポートしておりません。cudaMemcpydefault は、runtime 時にドライバーが src, dst 等のデバイス（ポインタ）等を判断するように指示するための引数なので、あえて使用しなくても良いと思います。プログラム・コーディング上は明示的な Memcpy 等の API を使用した方が、間違いがなく、後々分かり易い記述となります。
! Direct access to peer memory on "UVA"
! It is necessary that cudaDeviceEnablePeerAccess() is enable.
istat = cudaMemcpy(Ad1,Ad0,N,cudaMemcpyDeviceToDevice)
istat = cudaMemcpy(Bd0,Bd1,N,cudaMemcpyDeviceToDevice)
! cannot use 'cudaMemcpydefault' on PGI 11.8. It may not support yet.
istat = cudaMemcpy(Ad1,Ad0,N,cudaMemcpydefault)
istat = cudaMemcpy(Bd0,Bd1,N,cudaMemcpydefault)
(2) デバイス 0 上のデータを使って、デバイス 1 上の kernel で計算処理する
次に、実際の kernel プログラムの中で、異なるデバイス上のデータを使用する方法の一例を以下に記します。プログラムの前半部は、上述(1) のプログラムと同じです。GPUデバイスのチェックを行って、cudaDeviceEnablePeerAccess() を enable にしています。これを行うことが、GPUdirect 機能を使うために必須条件となります。以下のプログラム例では、vadd_kernel と言うカーネルプログラムの引数に、使用する配列（アドレス）を渡します。この引数に、異なるデバイスの配列名を記すことで、異なるデバイス・メモリ上のデータをアクセスして演算を進めることになります。もちろん、こうした構成の場合、その演算性能は、デバイス上の自身のメモリデータを使うより、性能が大きく劣化します。これは、演算中、必要なデータを peer to peer でコピーするオーバーヘッドが加算される訳ですから、当然のことです。皆さん、誤解してはならないことは、UVA や peer to peer アクセスが可能なプログラム環境とは、簡単に言えばプログラムの書き易さを提供するものであり、一方、その性能は必ず劣化します。どう言った所に UVA 構成を利用するのか、予め性能設計が必要であり、適材適所で使い分ける必要があります。性能を考えた場合、まあ、MPIコーディングと同じ考えで、並列分割時の袖データの交換にのみ使うような構成で使用すると言うのが、一つの方法なのでしょう。
! using UVA
! launch the kernel on device 1 using Ad0/Bd0 on device0 memory
! device 1 に投入する kernel 引数に device 0 上のデータ Ad0, Bd0 のアドレスを渡す
istat = cudaSetDevice(dev1)
call vadd_kernel <<< (N+63)/64, 64 >>> ( Ad0, Bd0, Cd1, N )
module vector_add
use cudafor
! Peer-to-Peer Memory Access under Unified virtual addressing
! Vector add kernel using another device memory data.
! Kato, SofTek
contains
attributes(global) subroutine vadd_kernel(A,B,C,N)
implicit none
real(4), device :: A(N), B(N), C(N)
integer, value :: N
integer :: i
i = (blockidx%x-1)*blockdim%x + threadidx%x
if( i <= N ) C(i) = A(i) + B(I)
! if( i <= N ) print *, c(i)
end subroutine
subroutine vadd( A, B, C )
implicit none
real(4), dimension(:) :: A, B, C
real(4), device, allocatable, dimension(:):: Ad0, Bd0
real(4), device, allocatable, dimension(:):: Ad1, Bd1, Cd1
integer :: i, N, istat, dev0, dev1, canaccess0_1, canaccess1_0
integer ndevices, uva0, uva1
integer :: nreps = 1000
logical has_uva
type(cudaEvent) :: startEvent, stopEvent
real (4) :: time
! Properties check
type(cudaDeviceProp) :: prop
integer :: ilen
N = size( A, 1 )
print *,"\nChecking for multiple GPUs..."
istat = cudaGetDeviceCount(ndevices)
print '(2x,i2,a)', ndevices, " GPU devices are found."
if ( ndevices < 2 ) then
print '(2x,a/)', " This system don't have two GPUs and above."
stop
end if
do i = 0, ndevices - 1
istat = cudaSetDevice(i)
istat = cudaGetDeviceProperties(prop,i)
ilen = verify(prop%name, ' ', .true.)
print '(1x, a, i1, 5x, a, a, i1,a)',"device", i, prop%name(1:ilen), " (GPU",i, ") is available"
end do
print *, "\n Now, using 2 devices (GPU0 and GPU1) "
dev0 = 0
dev1 = 1
print *, "\nChecking GPU(s) for support of peer to peer memory access...\n"
! Determine if Peer to Peer communication is supported
istat = cudaDeviceCanAccessPeer(canaccess0_1,dev0,dev1)
istat = cudaDeviceCanAccessPeer(canaccess1_0,dev1,dev0)
if (canaccess0_1 == 1 .and. canaccess1_0 == 1 ) then
print '(a, i4/,a,i4/)'," canaccess device 0 to 1 =", canaccess0_1, &
" canaccess device 1 to 0 =", canaccess1_0
istat = cudaSetDevice(dev0)
istat = cudaGetDeviceProperties(prop,dev0)
print '(1x,a,i2)', "device0 supports UVA ? : ",prop%unifiedAddressing
uva0 = prop%unifiedAddressing
istat = cudaDeviceEnablePeerAccess(dev1,0)
istat = cudaSetDevice(dev1)
istat = cudaGetDeviceProperties(prop,dev1)
print '(1x,a,i2)', "device1 supports UVA ? : ",prop%unifiedAddressing
uva1 = prop%unifiedAddressing
istat = cudaDeviceEnablePeerAccess(dev0,0)
! supprt UVA?
has_uva = ( uva0 == 1 .and. uva1 == 1 )
if (has_uva) then
print *, "\nBoth GPUs can support UVA, enabling\n"
else
print *,"\nAt least one of the two GPUs does NOT support UVA\n"
stop 999
end if
! allocate device 1's arrays
allocate( Ad1(N), Bd1(N), Cd1(N) )
! allocate device 0's arrays
istat = cudaSetDevice(dev0)
allocate( Ad0(N), Bd0(N) )
! copy host data to device 0
Ad0 = A(1:N)
Bd0 = B(1:N)
! using UVA
! launch the kernel on device 1 using Ad0/Bd0 on device0 memory
! device 1 に投入する kernel 引数に device 0 上のデータ Ad0, Bd0 のポインタを渡す
print *, "Run kernel on GPU 1, taking GPU 0 data(Ad0,Bd0), the result is on GPU1."
istat = cudaSetDevice(dev1)
istat = cudaEventCreate(startEvent)
istat = cudaEventCreate(stopEvent)
! timing experiment
time = 0.0e0
istat = cudaEventRecord(startEvent, 0)
do i = 1, nreps
call vadd_kernel <<< (N+63)/64, 64 >>> ( Ad0, Bd0, Cd1, N )
enddo
istat = cudaEventRecord(stopEvent, 0)
istat = cudaThreadSynchronize()
istat = cudaEventElapsedTime(time, startEvent, stopEvent)
time = time / (nreps) * 1.e3 !(usec)
print '(1x,a,F8.2)', "Total Time (usec): ",time
! get the results from device 1 for verification
C(1:N) = Cd1
! launch the kernel on device 1 using Ad1/Bd1 on device1 memory
! device 1 に投入する kernel 引数に device 1 上のデータ Ad1, Bd1 のポインタを渡す
print *, "\n Run kernel on GPU 1, using GPU 1 data(Ad1,Bd1), the result is on GPU1."
! timing experiment
time = 0.0e0
istat = cudaEventRecord(startEvent, 0)
do i = 1, nreps
call vadd_kernel <<< (N+63)/64, 64 >>> ( Ad1, Bd1, Cd1, N )
enddo
istat = cudaEventRecord(stopEvent, 0)
istat = cudaThreadSynchronize()
istat = cudaEventElapsedTime(time, startEvent, stopEvent)
time = time / (nreps) * 1.e3 !(usec)
print '(1x,a,F8.2)', "Total Time (usec): ",time
! deallocate the arrays on device 0 and cleanup
deallocate( Ad1, Bd1, Cd1 )
istat = cudaDeviceDisablePeerAccess(dev1)
! deallocate the arrays on device 1 and cleanup
istat = cudaSetDevice(dev0)
deallocate( Ad0, Bd0 )
istat = cudaDeviceDisablePeerAccess(dev0)
endif
end subroutine
end module
program p
use vector_add
use cudafor
implicit none
integer :: nargs, n, i, ierr
real, allocatable, dimension(:) :: a, b, c
external iargc
integer iargc
character*10 arg
! for timing
! real(8) :: time
integer(4) :: hz, clock0, clock1
nargs = iargc()
if( nargs == 0 )then
print *, 'Instruction : ./a.out N-size'
return
endif
if( nargs >= 1 )then
call getarg( 1, arg )
read(arg,'(i)') n
endif
allocate( a(n), b(n), c(n) )
do i = 1,n
a(i) = i
b(i) = 2*i
c(i) = 0.
enddo
! Initialize time function
! call system_clock(count_rate=hz)
! print *, "system_clock resolution: ", real(1.d0/hz)
! Start a clock
! call system_clock(count=clock0)
call vadd( a, b, c )
! call system_clock(count=clock1)
! time= real((clock1-clock0)) / real (hz)
! print *, "Total Time (sec): ",time
ierr = 0
do i = 1,n
if( c(i) .ne. a(i) + b(i) )then
ierr = ierr + 1
if( ierr <= 10 )then
print 30, i, c(i), a(i)+b(I)
30 format( 'c(',i,') = ', f12.5, ' should be ', f12.5 )
endif
endif
enddo
if( ierr == 0 )then
print *, '\n no errors'
else
print *, ierr, '\n errors found'
endif
end program
このプログラムも、現時点での最新バージョン PGI 11.8 を使用してコンパイルを行います。PGI 11.8 では、デフォルトが CUDA 3.2 の toolkit を利用しますので、明示的に -Mcuda:4.0 として CUDA 4.0 の toolkit を利用してバイナリを作成してください。実行時の引数には、何個の vector add 計算を行うかと言う個数を指定します。以下の例では 100000個の加算演算を行うことになります。さて、このプログラムの中には、同じデバイス上のデータを使用して演算した時間計測も行っております。これが、いわゆる GPU の primitive な性能ということになりますが、どの程度の性能差が現れるのかがよく分かります。以下の例では、性能で10倍以上の開きがあることが分かります。
[kato@photon29 2585-cuda40feature]$ pgf90 -Mcuda:4.0 -fast p2p.cuf
[kato@photon29 2585-cuda40feature]$ a.out 100000
Checking for multiple GPUs...
2 GPU devices are found.
device0 GeForce GTX 580 (GPU0) is available
device1 GeForce GTX 480 (GPU1) is available
Now, using 2 devices (GPU0 and GPU1)
Checking GPU(s) for support of peer to peer memory access...
canaccess device 0 to 1 = 1
canaccess device 1 to 0 = 1
device0 supports UVA ? : 1
device1 supports UVA ? : 1
Both GPUs can support UVA, enabling
Run kernel on GPU 1, taking GPU 0 data(Ad0,Bd0), the result is on GPU1.
Total Time (usec): 162.24
Run kernel on GPU 1, using GPU 1 data(Ad1,Bd1), the result is on GPU1.
Total Time (usec): 13.95
no errors
