Transcript 資料 - GPU コンピューティング研究会

PGIコンパイラーによる
ディレクティブベースのGPGPU
株式会社 ベストシステムズ
石川 直太
[email protected]
PGIによるGPGPUの概要



アクセラレータ対応ライセンスが必要
Fortran 95 または C99 のコードにディレクテ
ィブ

!$acc ..

!$acc& … 継続行

#pragma acc …
オプションを付けてコンパイル

-ta=nvidia

-ta=nvidia,time 実行時に統計情報を表示

-Minfo,accel コンパイル時に詳細な情報を表示
主要なディレクティブ

acc region --- GPUによる処理の開始

acc end region --- GPUによる処理の終了

copyin, copy, copyout --- まとめてコピー

local --- GPU側だけで使う変数

acc do vector --- GPUによる並列処理

private --- スレッドごとにインスタンスを持つ
変数
多重ループのベクトル化
!$acc to vector(4)
do k = 1, km
!$acc do vector(4)
do j = 1, jm
!$acc do vector(16)
do i = 1, im


多重ループをベクト
ル化できる
それぞれのループの
並列度を定数で指定
する

並列度の積は256以下

Cut and try

明示的に指定しなく
ても、自動ベクトル
化
PGI 9.0 から 10.0 への改良

総和演算が可能

姫野ベンチマークのコードより

GOSA = GOSA + SS * SS

すべてのスレッドに渡る総和演算

9.0 ではコンパイルエラー

10.0 ではコンパイル、計算可能
PGI 9.0から10.0への改良




コンパイラーまかせでの、性能向上
姫野ベンチマークに、「!$acc region」と「
!$acc end region」だけを追加したコード

PGI 9.0-4 では 0.229 GFLOPS (CPUより遅い)

PGI 10.0 では 15.3 GFLOPS
チューニングすると、9.0でも10.0でも、20.5
GFLOPS
10.1、10.2 では性能向上なし
GPGPU化可能なコード

並列化可能が大前提

リストベクトルでなく多重ループ

OpenMP から GPGPU への移行は容易

倍精度演算よりも単精度演算が高速
ディレクティブベースGPGPUの特
徴

CUDAプログラミングよりも容易

ディレクティブを無視すれば、通常のC/Fortran


GPUによる計算とCPUによる計算の比較が容易

オリジナルコードの変更への対処が比較的容易
ハードウェアに依存しない

理論的には、AMDのGPUや将来のCPUに対応可能

ヘテロジニアスマルチコアCPUの可能性？
事例１：姫野ベンチマーク

連立一次方程式をヤコビ法で解く

メモリ性能が現われる

http://accc.riken.jp/HPC/HimenoBMT.html
GPU版姫野ベンチのコード主要部
!$acc region &
!$acc copyin(aa(1:mimax,1:12,1:mjmax,1:mkmax)) &
!$acc copyin( p(1:mimax,1:mjmax,1:mkmax)) &
!$acc copyout(p(1:mimax,1:mjmax,1:mkmax)) &
#ifdef USE_WW
!$acc copyout(ww(1:mimax,1:2,1:mjmax,1:mkmax)) &
#else
!$acc copyout(wrk2(1:mimax,1:mjmax,1:mkmax)) &
!$acc copyout(gosatmp(1:mimax,1:mjmax,1:mkmax)) &
#endif ! USE_WW
!$acc local(k,j,i,s0,ss) &
!$acc copyin(kmax,jmax,imax,omega)
!$acc do host
do loop=1,nn
!$acc do parallel, vector(4) ! Cut and try.
do k=2,kmax-1
!$acc do parallel !!!, vector(2) ! Cut and try.
do j=2,jmax-1
!$acc do vector(64) ! Cut and try.
do i=2,imax-1
姫野ベンチの配列のパディング

Portland Groupによる改良
–
配列の第一次元の大きさを調節する。
–
#ifdef PAD
–
mimax = 272 ! GPUに適するマジックナンバー?
–
#else
–
mimax = 257 ! 姫野ベンチオリジナルの値
–
#endif

CPUによる計算：850 MFLOPS

ＧＰＵによる計算：20292 MFLOPS --- 23.8倍
マクロによる添え字のすり替え
!!!! BEGIN BSI HACK !!!!
#define a1(i,j,k) aa(i,1,j,k)
#define a2(i,j,k) aa(i,2,j,k)
#define a3(i,j,k) aa(i,3,j,k)
#define b1(i,j,k) aa(i,4,j,k)
#define b2(i,j,k) aa(i,5,j,k)
#define b3(i,j,k) aa(i,6,j,k)
#define c1(i,j,k) aa(i,7,j,k)
#define c2(i,j,k) aa(i,8,j,k)
#define c3(i,j,k) aa(i,9,j,k)
#define a4(i,j,k) aa(i,10,j,k)
#define bnd(i,j,k) aa(i,11,j,k)
#define wrk1(i,j,k) aa(i,12,j,k)
#ifdef USE_WW
#define gosatmp(i,j,k) ww(i,1,j,k)
#define wrk2(i,j,k) ww(i,2,j,k)
#endif ! USE_WW
!!!! END BSI HACK !!!!
Fortranのマクロ

C/C++のマクロと同様

配列構造の試行錯誤に便利

大文字小文字の区別に注意



implicit none との併用をお勧め
PGIではソースの拡張子が「F」または「F90」
(大文字)の場合と、オプション「-Mpreprocess
」で有効
Intel コンパイラーではオプション「-fpp」で有
効
姫野ベンチで解ったこと





配列の最も左側の添え字(Fortranの場合)を、最
も内側のループで、1づつ増やすとよい。
コンパイル時に「Non-stride-1 accesses」と表
示された場合には、性能が出にくい。
copyin, copyout ディレクティブが重要。
配列の構造を変えると、性能が上がる可能性が
あるが、若干工数を要する。
vector(64) のパラメーターは試行錯誤。
姫野ベンチで効果がなかったこと
•
•
private ディレクティブ
–
スレッドごとに別々のインスタンスを持つと指
定
–
省略しても、コンパイラーが自動的に判断
ストライド 0
–
構造体の配列と等価なデータ構造
–
CPUでは、キャッシュのヒット率向上に効果
–
GPUでは、遅くなる
姫野ベンチの他の研究との比較

富士通研究所 (情報処理学会HPC研究会)




CUDAプログラミングで、69.7 GFLOPS
メモリ転送速度がピーク性能の80%を超えるチュー
ニング
これと比較して本実験は 0.28倍
ソフテック
– PGIコンパイラーを使って、20457.79 MFLOPS

NEC (2009年9月2日 セミナー資料)
– PGIコンパイラーを使って、18477.78 MFLOPS
事例２：行列積


SGEMM

BLASに含まれるサブルーチン

S --- 単精度実数

GE --- 一般的な行列

MM --- 行列×行列
Netlibでソースコード公開

http://www.netlib.org/blas/index.html
オリジナルコード主要部
*
*
Form C := alpha*A*B + beta*C.
*
!$acc region
DO 90, J = 1, N
IF( BETA.EQ.ZERO )THEN
DO 50, I = 1, M
C( I, J ) = ZERO
50
CONTINUE
ELSE IF( BETA.NE.ONE )THEN
DO 60, I = 1, M
C( I, J ) = BETA*C( I, J )
60
CONTINUE
END IF
DO 80, L = 1, K
IF( B( L, J ).NE.ZERO )THEN
TEMP = ALPHA*B( L, J )
DO 70, I = 1, M
C( I, J ) = C( I, J ) + TEMP*A( I, L )
70
CONTINUE
END IF
80
CONTINUE
90
CONTINUE
!$acc end region
1570 MFLOPS (CPU よ り 遅い )
ないほうがよい条件分枝を除去
*
*
Form C := alpha*A*B + beta*C.
*
!$acc region
DO 90, J = 1, N
CBSI2010
IF( BETA.EQ.ZERO )THEN
CBSI2010
DO 50, I = 1, M
CBSI2010
C( I, J ) = ZERO
CBSI2010 50
CONTINUE
CBSI2010
ELSE IF( BETA.NE.ONE )THEN
DO 60, I = 1, M
C( I, J ) = BETA*C( I, J )
60
CONTINUE
CBSI2010
END IF
DO 80, L = 1, K
CBSI2010
IF( B( L, J ).NE.ZERO )THEN
TEMP = ALPHA*B( L, J )
DO 70, I = 1, M
C( I, J ) = C( I, J ) + TEMP*A( I, L )
70
CONTINUE
CBSI2010
END IF
80
CONTINUE
90
CONTINUE
!$acc end region
1583 MFLOPS
ループの回し方を変更
*
*
Form C := alpha*A*B + beta*C.
*
!$acc region local(temp,i,j,l)
!$acc do vector(16)
DO J = 1, N
!$acc do vector(16) private(temp)
DO I = 1, M
! C( I, J ) = BETA*C( I, J )
temp = 0.0
DO L = 1, K
! C( I, J ) = C( I, J ) + ALPHA * B(L,J) *A( I, L )
temp = temp + B(L,J) *A( I, L )
enddo
c(i,j) = c(i,j) * beta + temp * alpha
enddo
enddo
!$acc end region
不可解な現象



vector(16) ディレクティブがあると 1 GFLOPS
vector(16) ディレクティブがないと 23
GFLOPS
診断メッセージによると、どちらも16x16ブロ
ック

試行錯誤が必要

PGI コンパイラーは、まだ発展途上か？
実習
行列積のコードをGPUで計算しましょう
用意してあるファイル

sample.f --- GPU化していないサブルーチンコ
ード

sgemm-4.f --- GPU化の例

Makefile

test-sgemm.f --- 評価用メインプログラム
まずはCPUで実行

make sample

./sample
最初の一歩

高速化したいブロックの最初に !$acc region

終わりに !$acc end region

転置行列の積を計算するブロックもあるが、と
りあえずは、最初のブロックだけ

make sample

./sample
チューニング(1)


オリジナルコードには、0による乗算を避ける
ための条件分枝があります。
なくてもよい条件分枝を削除しましょう。
チューニング(2)

ループの回し方を変えてみましょう。

ヒントは「sgemm-4.f」
試行錯誤の例

明示的な、「local」、「private」

!$acc do vector(並列度)

「copyin」、「copy」、「copyout」

コンパイルオプション「-ta=nvidia,mul24」
終わりに、最新 Bad know how

2010年3月5日に、PGI 10.3 リリース

PGI 10.3 でサンプルコードがコンパイルエラー

アクセラレータコンパイラーは枯れていない


ライセンスファイルとライセンスサーバーが新
しければ、古いコンパイラーも動く
複数のバージョンのコンパイラーをインストー
ルして、パスの設定で選択可能
お問い合わせは

価格表、オンライン見積もりによる割引


ご注文窓口


http://www.bestsystems.co.jp/
[email protected]
技術ご質問、ライセンス発行窓口

[email protected]
Happy hacking!