TSUBAME2.0におけるGPUの活用方法

Size: px

Start display at page:

Download "TSUBAME2.0におけるGPUの活用方法"

なつきいざわ
2 years ago
Views:

1 GPU プログラミング応用編東京工業大学学術国際情報センター

2 はじめに基礎編では以下を説明 GPU プログラミングの基本 GPU 上のスレッドを使った並列プログラミングしかし GPU の特徴を考慮することによって更に高速化が可能同じ計算を行うプログラムでもメモリやスレッドなどの使い方の最適化によって数倍 ~ 数十倍実行速度が違う場合も! 応用編では基礎編に入れられなかったテクニックや高速化のポイントを説明

3 NVIDIA 社の資料についてこのスライドは説明を簡略化しているため正確な情報が必要な場合については下記の NVIDIA 社の公式資料 ( 英語 ) を参照してください CUDA C Programming Guide CUDA API Reference Manual コンテスト内部 Wikiに置いてあります

4 目次 1. はじめに 2. 多次元配列についてのテクニック 3. CUDA プログラムの時間計測に関する注意 4. divergent 分岐の削減による効率化 5. コアレスドアクセスによるメモリアクセス効率化 6. 共有メモリの有効活用 7. おわりに

5 2. 多次元配列についてのテクニック

6 多次元配列についてのテクニック CPU 上では大域変数として多次元配列を使える short int h_array[max1][max2][max3][max4]; のように定義できるサンプルプログラムでも多用している GPUではハードルが高い

7 GPU 上で利用する基本方法は? 多次元配列を GPU に確保コピーして計算に用いるには? 以下が基本の方法 cudamalloc で領域を確保しポインタ型変数を得るたとえば short int *d_array; cudamemcpy で h_array から d_array にコピーカーネル関数内で d_array の内容にアクセス可能しかし一次元配列として使わないといけない d_array[i][j][k][l] コンパイル時エラーになってしまう d_array[i*max2*max3*max4+j*max3*max4+k*max4+l] 動くけどプログラムが結構大変次スライドで少し楽にするテクニックを紹介します

8 GPU 上でも多次元配列! /home/sc12/gpulec/devarray.cu を参考にしてください (1) 変数の定義に device をつけると GPU メモリ上に確保される例 device short int d_array[max1][max2][max3][max4]; (2) cudamemcpytosymbol (CPU GPU の場合 ) や cudamemcpyfromsymbol (GPU CPU) を使う例 cudamemcpytosymbol("d_array", h_array, sizeof(short int)*max1*max2*max3*max4, 0); cudamemcpyfromsymbol(h_array, "d_array", sizeof(short int)*max1*max2*max3*max4, 0); (3) これなら GPU カーネル関数内でも d_array[i][j][k][l] のように使えてちょっと便利! 前スライドの基本の方法を使うか新しい方法を使うかはおまかせします

9 特殊なコピー関数の詳細 cudamemcpytosymbol(char *symbol, const void *src, size_t count, size_t offset) CPU 側のデータを GPU 上の device 変数にコピー symbol 転送先の device 変数の名前を文字列で指定 src 転送元 CPU メモリ count 転送サイズ ( バイト単位 ) offset symbol が表すアドレス +offset を転送先とできる 0 でもよい cudamemcpyfromsymbol(void *dst, char *symbol, size_t count, size_t offset) GPU 上の device 変数のデータを CPU 側にコピー src 転送先 CPU メモリ symbol 転送元の device 変数の名前を文字列で指定 count 転送サイズ ( バイト単位 ) offset symbol が表すアドレス +offset を転送元とできる 0 でもよい

10 3. CUDA プログラムの時間計測に関する注意

11 時間計測に関する注意プログラム中の各部分にかかる時間を測るために clock(), gettimeofday() 関数を使うことはよくある CUDA プログラムで以下を測るとき注意が必要 (a) cudamemcpy( ホストデバイス方向 ) (b) カーネル関数呼び出し本当の時間よりもはるかに短く見えてしまう実際には上記 (a)(b) を実行すると仕事を依頼しただけの状態で実行が帰ってきてしまう ( 非同期呼び出し ) 時刻測定前に cudadevicesynchronize() を行っておくこと cudadevicesynchronize() の意味現在までに GPU に依頼した仕事が全部終了するまで待つ

12 各部分ごとの時間計測を行うには clock_t t1, t2, t3, t4 cudadevicesynchronize(); t1 = clock(); cudamemcpy(, cudamemcpyhosttodevice); cudadevicesynchronize(); t2 = clock(); my_kernel<<<, >>>( ); cudadevicesynchronize(); t3 = clock(); cudamemcpy(, cudamemcpydevicetohost); cudadevicesynchronize(); t4 = clock(); t1 と t2 の差分が cudamemcpy ( ホストからデバイス ) の時間 t2 と t3 の差分がカーネル関数実行にかかった時間 t3 と t4 の差分が cudamemcpy ( デバイスからホスト ) の時間

13 4. DIVERGENT 分岐の削減による効率化

14 GPU でのスレッドの実行のされ方スレッドブロック内のブロック達は ( プログラマからは見えないが )32スレッドごとの塊 (warp) 単位で動作している Warpの中の32スレッドは常に足並みをそろえて動いている If 文などの分岐があるとどうなる? Warp 内のスレッド達の意見がそろうかそろわないかで動作が異なる

15 GPU 上の if 文の実行のされ方 (a) スレッド達の意見がそろう場合全員 xxx>100だとする if (xxx > 100) { } else { } Else 部分は実行せずに飛ばす (b) スレッド達の意見が違う場合あるスレッドでは yyy>100 だが別スレッドは違う場合 if (yyy > 100) { } else { } 一部スレッドを眠らせて Then も else も両方実行これを divergent 分岐と呼ぶ

16 Divergent 分岐はなぜ非効率? CPUの常識では if 文はthen 部分とelse 部分の片方しか実行しないので片方だけの実行時間がかかる Divergent 分岐があると then 部分とelse 部分の両方の時間がかかってしまう

17 5. コアレスドアクセスによるメモリアクセス効率化

18 グローバルメモリのアクセスの効率化コアレスドアクセスメモリの性質上近い ( たとえば番号が隣りの ) スレッドが近いアドレスを同時にアクセスするのが効率的コアレスドアクセス (coalesced access) と呼ぶ隣り合ったスレッドが配列の隣の要素をアクセスコアレスドアクセスになっており高速各スレッドがばらばらの要素をアクセスコアレスドアクセスではなく低速基礎編の inc_par プログラムはコアレスドアクセスになっていた

19 6. 共有メモリの有効活用

20 CUDA メモリモデルスレッドが階層化されているのと同様メモリも階層化されているスレッド固有レジスタ局所変数を格納高速だが容量小ブロック内共有共有メモリ本スライドで登場高速だが容量小 (L1 キャッシュ ) グリッド内 ( 全スレッド ) 共有グローバルメモリ global 変数や cudamalloc で利用容量大きいが低速 (L2 キャッシュ ) それぞれ速度と容量にトレードオフ有 ( 高速 & 小容量 vs. 低速 & 大容量 ) メモリアクセスの局所性が重要 Host (Device) Grid Block (0, 0) レジスタ Thread (0, 0) 共有メモリレジスタ Thread (1, 0) Block (1, 0) レジスタ Thread (0, 0) 共有メモリレジスタ Thread (1, 0) グローバルメモリ ( 本スライドではデバイスメモリとも呼ぶ ) Source Kirk and Hwu, ECE 498AL, UIUC

共有メモリの利用によるプログラム効率化基礎編のようにプログラムを書くと通常はレジスタとグローバルメモリのみを利用共有メモリとはブロック内のスレッド達で共有されるメモリ領域高速容量は小さい ( ブロックあたり 16KB 以下 ) shared int a[16]; のように書くと共有メモリ上に置かれる

21 共有メモリの利用によるプログラム効率化基礎編のようにプログラムを書くと通常はレジスタとグローバルメモリのみを利用共有メモリとはブロック内のスレッド達で共有されるメモリ領域高速容量は小さい ( ブロックあたり 16KB 以下 ) shared int a[16]; のように書くと共有メモリ上に置かれる Host (Device) Grid Block (0, 0) レジスタ Thread (0, 0) 共有メモリレジスタ Thread (1, 0) Block (1, 0) レジスタ Thread (0, 0) 共有メモリレジスタ Thread (1, 0) グローバルメモリ ( 本スライドではデバイスメモリとも呼ぶ )

22 共有メモリをどういう時に使うと効果的? 一般的にはグローバルメモリの同じ場所をブロック内の別スレッドが使いまわす場合に効率的たとえば matmul_par プログラムでは A, B の要素は複数スレッドによって読み込まれる一度グローバルメモリから共有メモリに明示的にコピーしてから使いまわすと有利カーネル関数の書き換えが必要ただし GPUにはキャッシュもあるため共有メモリで本当に高速化するか? は場合による

23 共有メモリを使った行列積プログラム最適化前 (matmul_par) スレッド t i, t i+1 はそれぞれ同一行をロード matmul_shared 最適化後 (matmul_shared) 各行列を 16x16 要素のタイルに分けて考える各スレッドブロックは 16x16 のスレッドを持つとするスレッド ti, ti+1 はそれぞれ 1 要素のみをロード計算は共有メモリ上の値を利用 t i t i+1 タイルタイル共有メモリグローバルメモリ

24 matmul_shared の流れこのプログラムでは 1 スレッドブロックが C の 1 タイル分を計算 1 スレッドが C の 1 要素を計算 1. 行列 A B 共にその一部のタイルをグローバルメモリから共有メモリにコピー 2. syncthreads() により同期 3. 共有メモリを用いてタイルとタイルのかけ算 4. 次のタイルのために 1へ戻る 5. 各スレッドは自分が計算したC i,j をグローバルメモリに書き込む 2. の syncthreads() とは? スレッドブロック内の全スレッドの足並みをそろえる ( 同期 ) この命令を呼ぶまでは共有メモリに書いた値が必ずしも他のスレッドへ反映されない

25 共有メモリを使った高速化の結果サイズ 1024x1024 の行列 A, B, C があるとき C=A B を計算する matmul_cpu.c CPU で計算約 8.3 秒 (gcc O2 でコンパイルした場合 ) matmul_seq.cu GPU の 1 スレッドで計算約 200 秒 matmul_par.cu GPU の複数スレッドで計算約秒 matmul_shared.cu GPU の複数スレッドで計算し共有メモリも利用約秒 (!)

26 おわりに GPUプログラムにおいて基礎編の知識だけではちょっと不便な点の改良何が起こってしまうと非効率的になってしまうか何をできるだけ避けるべきかについていくつかポイントを説明した

1. GPU コンピューティング GPU コンピューティング GPUによる汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境 Lin

Windows で始める CUDA 入門 GTC 2013 チュートリアルエヌビディアジャパン CUDA エンジニア森野慎也 1. GPU コンピューティング GPU コンピューティング GPUによる汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境

TSUBAME2.0におけるGPUの 活用方法

TSUBAME2.0におけるGPUの活用方法