GPU のアーキテクチャとプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓

Transcription

1 GPU のアーキテクチャとプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓

2 今回の内容 GPU のアーキテクチャ CUDA CUDA によるプログラミング 58

3 GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ チップ単体では販売されていない PCI Ex カードで販売 ( チップ単体と区別せずに GPU と呼ぶことも多い ) ノート PC に搭載 PCI Ex カードとして販売される GPU には, ビデオメモリと呼ばれる RAM が搭載 59

4 GPU(Graphics Processing Unit) の役割 グラフィックスを表示するために様々な処理を行い, 処理の結果をディスプレイに出力 3 次元グラフィックスの発展に伴って役割が大きく変化 描画情報 CPU 3 次元座標変換 ポリゴンとピクセルの対応付け ピクセル色計算テクスチャ参照 フレームバッファ ( ビデオメモリ ) への書き込み ディスプレイ出力 過去 CPU が 3D 描画の演算を実行 GPUが出力 描画情報 3 次元座標変換 ポリゴンとピクセルの対応付け ピクセル色計算テクスチャ参照 フレームバッファ ( ビデオメモリ ) への書き込み ディスプレイ出力 現在 GPUが演算から出力までの全てを担当 CPUは描画情報の生成やGPUへの情報の引き渡し,GPU の制御を行う ディスプレイコントローラ 画面出力 GPU 画面出力 60

5 GPU の描画の流れ 1. CPUからGPUへ描画情報を送信 2. 頂点処理 ( 頂点シェーダ ) 座標変換 画面上での頂点やポリゴンの位置 大きさの決定 頂点単位での照明の計算 3. 頂点やポリゴンからピクセルを生成 ( ラスタライザ ) 4. ピクセル処理 ( ピクセルシェーダ ) 画面上のピクセルの色 テクスチャの模様 5. 画面出力 ピクセルの色情報をフレームバッファに書き込み

6 ビデオカードから GPU へ CG の多様化と共に固定機能の実装が困難に 頂点処理とピクセル処理をユーザが書き換えられるプログラマブルシェーダの実装 グラフィックスカード GPU 頂点処理用回路 頂点シェーダユニット ピクセル処理用回路 ピクセルシェーダユニット 62

7 レンダリングパイプライン処理 頂点情報 光源情報 形状データの画面への投影像 実数演算 投影変換 視野変換 陰影計算 材質情報 クリッピング 投影像を画素へ変換 整数演算とメモリアクセス ビューポート変換 走査変換 合成 テクスチャ 出力画像 63

8 レンダリングパイプライン処理 頂点情報 光源情報 形状データの画面への投影像 実数演算 実数演算を行うハードウェアは高価だった 視野変換陰影計算投影変換クリッピング 材質情報 ハードウェアで処理 ( 固定機能 ) ビューポート変換 走査変換 合成 テクスチャ 出力画像 64

9 レンダリングパイプライン処理 頂点情報 光源情報 ハードウェアで処理 ( 固定機能 ) 投影変換 視野変換 陰影計算 材質情報 クリッピング ハードウェアで処理 ( 固定機能 ) ビューポート変換 走査変換 合成 テクスチャ 出力画像 65

10 レンダリングパイプライン処理 頂点情報 頂点シェーダ 投影変換 視野変換 陰影計算 光源情報 材質情報 クリッピング ピクセルシェーダ ビューポート変換 走査変換 合成 テクスチャ 出力画像 66

11 ビデオカードから GPU へ 頂点処理とピクセル処理をユーザが書き換えられるプログラマブルシェーダの実装 処理によっては利用効率に差が生じる 頂点処理重視の処理 GPU 頂点シェーダユニット ピクセル処理重視の処理 GPU 頂点シェーダユニット 空きユニット ピクセルシェーダユニット ピクセルシェーダユニット 空きユニット 67

12 ビデオカードから GPU へ 頂点シェーダとピクセルシェーダを統合したユニファイドシェーダへの進化 頂点処理とピクセル処理を切り替えることで利用率を高める 頂点処理重視の処理 GPU ユニファイドシェーダユニット ピクセル処理重視の処理 GPU ユニファイドシェーダユニット 68

13 ビデオカードから GPU へ 各ピクセルに対して処理を並列に実行 単純な処理を行う演算器 (Streaming Processor, SP) を大量に搭載 演算器は現在 CUDA Core という名称に変更 高い並列度で処理を行う 69

14 Tesla アーキテクチャの構造 Tesla C1060 の仕様 SM 数 30 CUDA Core 数 240(=8 Core/SM 30 SM) キャッシュを搭載せず 70

15 Tesla アーキテクチャの構造 Tesla C1060 の仕様 CUDA コア数 ( 単精度 ) CUDA コアクロック周波数 240 Cores 1,296 MHz 単精度演算ピーク性能 622* 1 (933* 2 ) GFLOPS 倍精度演算ユニット数 30* 3 Units 倍精度演算ピーク性能 メモリクロック周波数 メモリバス幅 最大メモリバンド幅 * 4 78 GFLOPS 800 MHz 512 bit 102 GB/s * 1 単精度演算ピーク性能 = コアクロック周波数 コア数 命令の同時発行数 (2) * 2 CUDA Core と SFU が同時に命令を発行できれば 1296 MHz * 3 一つの SM に倍精度演算器が一つ搭載 * 4 最大メモリバンド幅 = メモリクロック周波数 メモリバス幅 /8 2(Double Data Rate) 71

16 Fermi アーキテクチャの構造 Tesla M2050 の仕様 SM 数 14 CUDA Core 数 448(=32 Core/SM 14 SM) L1/L2 キャッシュを搭載 ECC( 誤り訂正機能 ) を搭載 72

17 Fermi アーキテクチャの構造 Tesla M2050 の仕様 CUDAコア数 ( 単精度 ) CUDAコアクロック周波数単精度演算ピーク性能倍精度演算ユニット数倍精度演算ピーク性能メモリクロック周波数メモリバス幅最大メモリバンド幅 448 Cores 1,150 MHz 1.03 TFLOPS 0* 1 Unit 515 GFLOPS 1.55 GHz 384 bit 148 GB/s * 1 単精度 CUDA Core を 2 基使って倍精度演算を実行 73

18 Kepler アーキテクチャの構造 Tesla K20c/m の仕様 SMX 数 13 Streaming Multiprocessor extreme (?) CUDA Core 数 2,496(=192 Core/SM 13 SMX) 74

19 Kepler アーキテクチャの構造 Tesla K20c/m の仕様 CUDAコア数 ( 単精度 ) 2,496 Cores CUDAコアクロック周波数 706 MHz 単精度演算ピーク性能 3.52 TFLOPS 倍精度演算ユニット数 832* 1 Units 倍精度演算ピーク性能 1.17 TFLOPS メモリクロック周波数 2.6 GHz メモリバス幅 320 bit 最大メモリバンド幅 208 GB/s * 1 64 基 /SMX 13 基 75

20 Maxwell アーキテクチャ GeForce GTX TITAN X の仕様 SM 数 24 CUDA Core 数 3,072(=128 Core/SM 24 SM) 76

21 Maxwell アーキテクチャ GeForce GTX TITAN X の仕様 * CUDA コア数 ( 単精度 ) CUDA コアクロック周波数 単精度演算ピーク性能 3,072 Cores 1,002 MHz 6.14 TFLOPS 倍精度演算ユニット数 0* 1 Units 倍精度演算ピーク性能 192 GFLOPS* 2 メモリクロック周波数 3.5 GHz* 3 メモリバス幅 最大メモリバンド幅 * top gpus/geforce gtx titan x/specifications 384 bit GB/s * 1 * 2 倍精度演算は単精度演算の性能の 1/32 (1/16 Flop/Core/clock) * 3 DDR(Double Data Rate) 7GHz 相当と書かれている場合もある 77

22 Pascal アーキテクチャ 2016 年にリリース予定 倍精度演算器を搭載予定 NVLink GPU 同士や GPU と CPU を接続する独自の方式 通信 (CPU メモリ PCI Express メモリ GPU) のボトルネックを解消 (PCI Express3.0 の 5~12 倍 ) 複数の GPU を使って大規模な計算が可能 3D メモリ (High Bandwidth Memory, HBM)* 3 次元積層技術を利用し, メモリの容量と帯域を大幅に増加 最大 32GB, メモリ帯域 1TB/s * 78

23 Volta アーキテクチャ Pascal の後継 詳しい情報は不明 アメリカの次世代スーパーコンピュータへ採用予定 オークリッジ国立研究所 SUMMIT 150~300PFLOPS ローレンス リバモア研究所 SIERRA 100PFLOPS 以上 地球シミュレータと同等の演算性能を 1 ノードで実現 現在 Top500 2 位のスーパーコンピュータと同じ電力で 5~10 倍高速, サイズは 1/5 * 79

24 GPU の模式図 GPU Chip Streaming Multiprocessor SM SM SM SM SM SM SM SM レジ ローカルメモリ SM SM SM SM L2 キャッシュ L1 キャッシュ スタ CUDA Core レジスタ CUDA Core 共有メモリ レジスタ CUDA Core レジスタ CUDA Core GPU Streaming Multiprocessor Streaming Multiprocessor CUDA Core CUDA Core ローカルメモリ コンスタントメモリ テクスチャメモリ グローバルメモリ 80

25 GPU の並列化の階層 グリッド-ブロック-スレッドの3 階層 グリッド (Grid) 並列に実行する処理 GPUが処理を担当する領域全体 スレッド (Thread) GPUの処理の基本単位 CPUのスレッドと同じ ブロック (Block) もしくはスレッドブロック スレッドの集まり 81

26 GPU の並列化の階層 GPU のハードウェアの構成に対応させて並列性を管理 ハードウェア構成 GPU 並列化の階層 並列に実行する処理 CUDA Grid Streaming Multiprocessor スレッドの集まり Block CUDA Core スレッド Thread 82

27 CUDA Compute Unified Device Architecture NVIDIA 社製 GPU 向け開発環境 (Windows,Linux,Mac OS X) 2007 年頃発表 C/C++ 言語 + 独自のGPU 向け拡張 専用コンパイラ (nvcc) とランタイムライブラリ いくつかの数値計算ライブラリ ( 線形代数計算,FFTなど) CUDA 登場以前 グラフィックスプログラミングを利用 足し算を行うために, 色を混ぜる処理を実行 汎用計算のためには多大な労力が必要 83

28 CUDA によるプログラミング CPU をホスト (Host),GPU をデバイス (Device) と表現 ホスト (CPU) 処理の流れや GPU を利用するための手続きを記述 プログラムの書き方は従来の C 言語と同じ 利用する GPU の決定,GPU へのデータ転送,GPU で実行する関数の呼び出し等 84

29 CUDA によるプログラミング CPU をホスト (Host),GPU をデバイス (Device) と表現 デバイス (GPU) 処理する内容を関数として記述 引数は利用可能, 返値は利用不可 ( 常にvoid) 関数はkernelと呼ばれる 関数呼び出しはlaunch, invokeなどと呼ばれる 85

30 Hello World 何を確認するか 最小構成のプログラムの作り方 ファイル命名規則 ( 拡張子は.c/.cpp) コンパイルの方法 (gcc, cl 等を使用 ) #include<stdio.h> int main(void){ printf("hello world n"); } return 0; helloworld.c 86

31 CUDA で Hello World 何を確認するか 最小構成のプログラムの作り方 ファイル命名規則 ( 拡張子は.cu) コンパイルの方法 (nvcc を使用 ) #include<stdio.h> int main(void){ #include<stdio.h> int main(void){ } printf("hello world n"); printf("hello world n"); return 0; helloworld.cu 違いは拡張子だけ? } return 0; helloworld.c 87

32 CUDA プログラムのコンパイル ソースファイルの拡張子は.cu nvcc を用いてコンパイル CPU が処理する箇所は gcc 等がコンパイル GPU で処理する箇所を nvcc がコンパイル helloworld.cu には CPU で処理する箇所しかない 88

33 CUDA で Hello World CUDA 専用の処理を追加 #include<stdio.h> global void kernel(){} int main(void){ GPU で実行される関数 ( カーネル ) global が追加されている } kernel<<<1,1>>>(); printf("hello world n"); return 0; 通常の関数呼出とは異なり, <<<>>> が追加されている helloworld_kernel.cu 89

34 CUDA プログラムの実行 実行時の流れ (CPU 視点 ) 利用するGPUの初期化やデータの転送などを実行 GPUで実行する関数を呼び出し GPUから結果を取得 time CPU 初期化の指示必要なデータのコピーカーネルの実行指示 CPU と GPU は非同期 CPU は別の処理を実行可能 結果の取得 GPU 初期化メモリに書込カーネルを実行実行結果をコピー 90

35 GPU の構造とカーネルの書き方 プログラムから GPU で実行する関数を呼出 GPU で実行する関数という目印が必要 GPU は PCI Ex バスを経由してホストと接続 GPU はホストと別に独立したメモリを持つ 関数の実行に必要なデータは GPU のメモリに置く GPU はマルチスレッド ( メニースレッド ) で並列処理 関数には 1 スレッドが実行する処理を書く 関数を実行する際に並列処理の度合いを指定 91

36 GPU の構造とカーネルの書き方 GPU で実行する関数 ( カーネル ) という目印 修飾子 global を付ける GPU は PCI Ex バスを経由してホストと接続 GPU はホストと別に独立したメモリを持つ カーネルの返値を void にする GPU はマルチスレッド ( メニースレッド ) で並列処理 カーネルには 1 スレッドが実行する処理を書く カーネル名と引数の間に <<<1,1>>> を付ける 92

37 Hello Thread(Fermi 世代以降 ) GPU の各スレッドが画面表示 #include<stdio.h> global void hello(){ printf("hello Thread n"); } int main(void){ 画面表示 (Fermi 世代以降で可能 ) コンパイル時にオプションが必要 arch=sm_20 以降 } hello<<<1,1>>>(); cudathreadsynchronize(); return 0; カーネル実行 ホストとデバイスの同期をとる CPUとGPUは原則同期しないので, 同期しないとカーネルを実行した直後にプログラムが終了 hellothread.cu 93

38 Hello Thread(Fermi 世代以降 ) <<< >>> 内の数字で並列度が変わることの確認 #include<stdio.h> global void hello(){ printf("hello Thread n"); } int main(void){ } hello<<<?,?>>>(); cudathreadsynchronize(); return 0; <<<>>> 内の数字を変えると画面表示される行数が変わる <<<1,8>>>, <<<8,1>>>, <<<4,2>>> 等 hellothread.cu 94

39 CPU と GPU のやりとり GPU の想定される使い方 ホスト (CPU) からデータを送り, デバイス (GPU) で計算し, 結果を受け取る CPU と GPU のデータのやり取りが必要 GPU は原則データを返さない PCI Ex 経由で描画情報を受け取り, 画面に出力 カーネルの返値が void の理由 NVIDIA 社ホームページより引用 95

40 CPU と GPU のやりとり CUDA 独自の命令と C 言語のポインタを利用 GPU のメモリ上に計算に必要なサイズを確保 確保したメモリのアドレスを C 言語のポインタで格納 ポインタの情報を基にデータを送受信 96

41 CPU と GPU のやり取り ( 単純な加算 ) int 型の変数 2 個を引数として受け取り,2 個の和を返す C 言語らしい書き方 #include<stdio.h> int add(int a, int b){ return a + b; } 引数で渡された変数の和を返す int main(void){ int c; c = add(6, 7); printf("6 + 7 = %d n", c); 関数呼び出し } return 0; add_naive.c 97

42 CPU と GPU のやり取り ( 単純な加算 ) 関数の返値を void に変更し, メモリの動的確保を使用 #include<stdio.h> #include<stdlib.h> void add(int a, int b, int *c){ *c = a + b; } 引数で渡された変数の和を,c が指すアドレスに書き込み int main(void){ int c; int *addr_c; addr_c = (int *)malloc(sizeof(int)); add(6, 7, addr_c); c = *addr_c; printf("6 + 7 = %d n", c); 引数にアドレスを追加アドレスを基に結果を参照 } return 0; add.c 98

43 CPU プログラム ( メモリの動的確保 ) malloc 指定したバイト数分のメモリを確保 stdlib.hをインクルードする必要がある #include<stdlib.h> int *a; a = (int *)malloc( sizeof(int)*100 ); sizeof データ型 1 個のサイズ ( バイト数 ) を求める printf("%d, %d n", sizeof(float), sizeof(double)); 実行すると 4,8 と表示される 99

44 CPU と GPU のやり取り ( 単純な加算 ) add.c の処理の一部を GPU の処理に置き換え #include<stdio.h> global void add(int a, int b, int *c){ *c = a + b; } global を追加 int main(void){ int c; int *dev_c; cudamalloc( (void **)&dev_c, sizeof(int) ); } add<<<1, 1>>>(6, 7, dev_c); cudamemcpy(&c, dev_c, sizeof(int), cudamemcpydevicetohost); printf("6 + 7 = %d n", c); cudafree(dev_c); return 0; GPU 上のメモリに確保される変数のアドレス GPU 上にint 型変数一個分のメモリを確保 GPU から結果をコピーメモリを解放 add.cu 100

45 CUDA でカーネルを作成するときの制限 カーネルの引数 値を渡すことができる GPUのメモリを指すアドレス CPU のメモリを指すアドレスも渡すことは可能 そのアドレスを基にホスト側のメモリを参照することは不可能 printf などの画面出力 Fermi 世代以降の GPU で, コンパイルオプションを付与 arch={sm_20 sm_21 sm_30 sm_32 sm_35 sm_50 sm_52} エミュレーションモード 新しい CUDA(4.0 以降 ) では消滅 101

46 CPU プログラムの超簡単移植法 とりあえず GPU で実行すればいいのなら 拡張子を.cu に変更 GPU の都合を反映 関数の返値を void にし, global を付ける 関数名と引数の間に <<<1,1>>> を付ける GPU で使うメモリを cudamalloc で確保 malloc でメモリを確保していればそれを cudamalloc に置き換え GPU からデータを受け取るために cudamemcpy を追加 最適化は追々考えればいい カーネルの完成 102

47 Hello Thread(Fermi 世代以降 ) <<< >>> 内の数字で並列度が変わる この情報を利用すれば並列処理が可能 #include<stdio.h> global void hello(){ printf("hello Thread n"); } int main(void){ } hello<<<?,?>>>(); cudathreadsynchronize(); return 0; <<<>>> 内の数字を変えると画面表示される行数が変わる <<<1,8>>>, <<<8,1>>>, <<<4,2>>> 等 hellothread.cu 103

48 GPU の並列化の階層 GPU のハードウェアの構成に対応させて並列性を管理 並列化の各階層における情報を利用 ハードウェア構成 並列化の階層 CUDA GPU 並列に実行する処理 Grid Streaming Multiprocessor スレッドの集まり Block CUDA Core スレッド Thread 104

49 GPU の並列化の階層 グリッド - ブロック - スレッドの 3 階層 各階層の情報を参照できる変数 x,y,z をメンバにもつ dim3 型構造体 グリッド (Grid) griddim グリッド内にあるブロックの数 ブロック (Block) blockidx blockdim ブロックに割り当てられた番号ブロック内にあるスレッドの数 スレッド (Thread) threadidx スレッドに割り当てられた番号 105

50 Hello Threads(Fermi 世代以降 ) <<< >>> 内の数字で表示される内容が変化 #include<stdio.h> global void hello(){ printf("griddim.x=%d, blockidx.x=%d, blockdim.x=%d, threadidx.x=%d n", griddim.x, blockidx.x, blockdim.x, threadidx.x); } int main(void){ } hello<<<?,?>>>(); cudathreadsynchronize(); return 0; <<<>>> 内の数字を変えると画面表示される内容が変わる <<<>>> 内の数字とどのパラメータが対応しているかを確認 hellothreads.cu 106

51 GPU の構造とカーネルの書き方 GPUはマルチスレッド ( メニースレッド ) で並列処理 関数には1スレッドが実行する処理を書く 関数を実行する際に並列処理の度合いを指定 カーネルと引数の間に追加した <<<,>>> で並列処理の度合を指定 <<< グリッド内にあるブロックの数,1 ブロックあたりのスレッドの数 >>> 107

52 プログラム実習 以下のプログラムをコンパイルし, 正しく実行できることを確認せよ helloworld.c helloworld.cu hellothread.cu hellothreads.cu hellothreads.cu については,<<<>>> 内の数字を変更し, 実行結果がどのように変わるか確認せよ 108

53 レポート課題 1( 提出期限は 1 学期末 ) 二つの値を交換する関数 swap を GPU に移植せよ 並列化する必要はなく,1 スレッドで実行すればよい #include<stdio.h> void swap(int *addr_a, int *addr_b){ int c; //c は値を一時的に保持するための変数 c = *addr_a; //* は間接参照演算子 *addr_a = *addr_b; // メモリアドレス (=addr_a,addr_b の値 ) にある変数の値を参照 *addr_b = c; } int main(void){ int a=1,b=2; printf("a = %d, b = %d n", a, b); swap(&a, &b); // 変数 a, b のメモリアドレスを渡す.& はアドレス演算子 printf("a = %d, b = %d n", a, b); } return 0; swap.c 109

54 レポート課題 1( 提出期限は 1 学期末 ) #include<stdio.h> カーネルという目印 void swap(int *a, int *b){ } int main(void){ int a=1,b=2; GPU で使う変数を宣言? printf("a = %d, b = %d n", a, b); GPU 上のメモリを確保 (a の分 ) GPU 上のメモリを確保 (b の分 ) CPU から GPU にメモリの内容をコピー (a の分 ) CPU から GPU にメモリの内容をコピー (b の分 ) swap 実行時の並列度の指定 (GPU で使う変数, GPU で使う変数 ); GPU のメモリの内容を CPU にコピー (a の分 ) GPU のメモリの内容を CPU にコピー (b の分 ) printf("a = %d, b = %d n", a, b); } return 0; 110