GPU のプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓
今回の内容 GPU プログラミング環境 (CUDA) GPU プログラムの実行の流れ CUDA によるプログラムの記述 カーネル (GPU で処理する関数 ) の構造 記述方法とその理由 GPU 固有のパラメータの確認 405
GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ チップ単体では販売されていない PCI Ex カードで販売 ( チップ単体と区別せずに GPU と呼ぶことも多い ) マザーボードやノート PC に搭載 PCI Ex カードとして販売される GPU には, ビデオメモリと呼ばれる RAM が搭載 406
GPU のハードウェア構造 CUDA Core( 旧 Streaming Processor, SP) と呼ばれる演算器を多数搭載 Streaming Multiprocessor(SM, SMX) が複数の CUDA Core と SFU, メモリをまとめて管理 SFU(Special Function Unit) 数学関数を計算するユニット 複数の SM が集まって GPU を構成 407
Fermi アーキテクチャの構造 Tesla M2050 の仕様 SM 数 14 CUDA Core 数 448(=32 Core/SM 14 SM) 動作周波数 1,150 MHz 単精度演算ピーク性能 1.03 TFLOPS 408
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 ローカルメモリ コンスタントメモリ テクスチャメモリ グローバルメモリ 409
CUDA Compute Unified Device Architecture NVIDIA 社製 GPU 向け開発環境 (Windows,Linux,Mac OS X) 2007 年頃発表 C/C++ 言語 + 独自のGPU 向け拡張 専用コンパイラ (nvcc) とランタイムライブラリ いくつかの数値計算ライブラリ ( 線形代数計算,FFTなど) CUDA 登場以前 グラフィクスプログラミングを利用 足し算を行うために, 色を混ぜる処理を実行 汎用計算のためには多大な労力が必要 410
プログラマブルシェーダを用いた汎用計算 グラフィックス API(DirectX, OpenGL) による描画処理 + シェーダ言語 (HLSL, GLSL) による演算 void gpumain(){ vec4 ColorA = vec4(0.0, 0.0, 0.0, 0.0); vec4 ColorB = vec4(0.0, 0.0, 0.0, 0.0); vec2 TexA = vec2(0.0, 0.0); vec2 TexB = vec2(0.0, 0.0); TexA.x = gl_fragcoord.x; TexA.y = gl_fragcoord.y; TexB.x = gl_fragcoord.x; TexB.y = gl_fragcoord.y; } ColorA = texrect( texunit0, TexA ); ColorB = texrect( texunit1, TexB ); gl_fragcolor = F_ALPHA*ColorA + F_BETA*ColorB; シェーダ言語を用いた配列加算 (c= *a + *b) の例 void main(){ glutinit( &argc, argv ); glutinitwindowsize(64,64);glutcreatewindow("gpgpuhelloworld"); glgenframebuffersext(1, &g_fb); glbindframebufferext(gl_framebuffer_ext, g_fb); glgentextures(4, g_ntexid); // create (reference to) a new texture glbindtexture(opt1, texid); gltexparameteri(opt1, GL_TEXTURE_MIN_FILTER, GL_NEAREST); gltexparameteri(...); glteximage2d(opt1, 0, opt2, width, height, 0, GL_RGBA, GL_FLOAT, 0); ( 以下省略 ) GPU の処理 (GLSL) 各ピクセルに対して実行 CPU の処理 (OpenGL) 411
CUDA によるプログラミング CPU をホスト (Host),GPU をデバイス (Device) と表現 ホスト (CPU) 処理の流れや GPU を利用するための手続きを記述 プログラムの書き方は従来の C 言語と同じ 利用する GPU の決定,GPU へのデータ転送,GPU で実行する関数の呼び出し等 412
CUDA によるプログラミング CPU をホスト (Host),GPU をデバイス (Device) と表現 デバイス (GPU) 処理する内容を関数として記述 引数は利用可能, 返値は利用不可 ( 常にvoid) 関数はkernelと呼ばれる 関数呼び出しはlaunch, invokeなどと呼ばれる 413
Hello World 何を確認するか 最小構成のプログラムの作り方 ファイル命名規則 ( 拡張子は.c/.cpp) コンパイルの方法 (gcc, cl 等を使用 ) #include<stdio.h> int main(void){ printf("hello world n"); } return 0; helloworld.c 414
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 415
CUDA プログラムのコンパイル ソースファイルの拡張子は.cu nvcc を用いてコンパイル CPU が処理する箇所は gcc 等がコンパイル GPU で処理する箇所を nvcc がコンパイル helloworld.cu には CPU で処理する箇所しかない 416
CUDA で Hello World CUDA 専用の処理を追加 #include<stdio.h> global void kernel(){} int main(void){ GPU で実行される関数 ( カーネル ) global が追加されている kernel<<<1,1>>>(); printf("hello world n"); 通常の関数呼出とは異なり, <<<>>> が追加されている } return 0; hellokernel.cu 417
CUDA プログラムの実行 実行時の流れ (CPU 視点 ) 利用するGPUの初期化やデータの転送などを実行 GPUで実行する関数を呼び出し GPUから結果を取得 time CPU 初期化の指示必要なデータのコピーカーネルの実行指示 CPU と GPU は非同期 CPU は別の処理を実行可能 結果の取得 GPU 初期化メモリに書込カーネルを実行実行結果をコピー 418
Hello Thread(Fermi 世代以降 ) printf を GPU から呼び出し, 並列に実行 #include<stdio.h> hellothread.c int hello(){ printf("hello Thread n"); return 0; } int main(void){ hello(); 画面表示 関数呼び出し } return 0; 419
Hello Thread(Fermi 世代以降 ) GPU の各スレッドが画面表示 #include<stdio.h> global void hello(){ printf("hello Thread n"); } int main(void){ hellothread.cu 画面表示 (Fermi 世代以降で可能 ) コンパイル時にオプションが必要 arch=sm_20 以降 } hello<<<1,1>>>(); cudadevicesynchronize(); return 0; カーネル実行 ホストとデバイスの同期をとる CPUとGPUは原則同期しないので, 同期しないとカーネルを呼び出した直後にプログラムが終了 420
CUDA でカーネルを作成するときの制限 printf による画面出力 Fermi 世代以降の GPU で, コンパイルオプションを付与 arch={sm_20 sm_21 sm_30 sm_32 sm_35 sm_50 sm_52} エミュレーションモード GPU の動作 ( 並列実行 ) を CPU で模擬 CUDA4.0 以降では消滅 オプション付きのコンパイル nvcc arch=sm_20 hellothread.cu 421
GPU プログラムへの変更 変更点 関数の前に修飾子 global をつけた 422
変更の理由 変更点 関数の前に修飾子 global をつけた 変更によって実現されること GPU で実行する関数という目印になる 変更が必要な理由 ホスト (CPU) からGPUで実行する関数 ( カーネル ) を呼び出し CPUが処理する箇所とGPUが処理する箇所は別のコンパイラがコンパイル コンパイルの時点でどれがカーネルかを明記 423
GPU プログラムへの変更 変更点 関数 hello の返値を void にした 424
変更の理由 変更点 関数 hello の返値を void にした 変更によって実現されること GPU のハードウェア構造に適したプログラムを作成できる 変更が必要な理由 GPUはホストと別に独立したメモリを持つ GPUは描画情報を受け取り, 画面に出力 GPU CPU の頻繁なデータ転送は苦手 画面出力 描画情報 プログラマがメモリ管理を行い, 無駄なデータ転送による実行速度低下を回避 425
GPU プログラムへの変更 変更点 関数呼出の際に関数名と引数の間に <<<1,1>>> を付けた 426
変更の理由 変更点 関数呼出の際に関数名と引数の間に <<<1,1>>> を付けた 変更によって実現されること GPU のハードウェア構造に適したプログラムを作成できる 変更が必要な理由 GPU には数百から数千の CUDA コアが搭載されており, それらが協調して並列処理を実行 1 スレッドが実行する処理を書くことでカーネルの作成を簡略化 並列処理の度合いはカーネル呼出の際に指定 427
GPU プログラムへの変更 変更点 カーネルを呼び出した後に同期を取る関数を呼んだ 428
変更の理由 変更点 カーネルを呼び出した後に同期を取る関数を呼んだ 変更によって実現されること GPU で実行した結果が正しく得られる 変更が必要な理由 CPU と GPU は非同期に処理を実行 関数を呼んで CPU 側に制御が戻った直後に return 0 でプログラムが終了 ( 画面表示が行われない ) 正しい結果を得るためにカーネルの終了を待つ 429
Hello Thread(Fermi 世代以降 ) <<<, >>> 内の数字を変えると表示される内容が変化 #include<stdio.h> hellothread.cu global void hello(){ printf("hello Thread n"); } int main(void){ } hello<<<?,?>>>(); cudathreadsynchronize(); return 0; <<<>>> 内の数字を変えると画面表示される行数が変わる <<<1,8>>>, <<<8,1>>>, <<<4,2>>> 等 430
<<<,>>> 内の 2 個の数字の意味は? GPUのハードウェアの構成に対応させて並列性を管理 各階層における並列実行の度合を指定 <<<,>>> 内に 2 個の数字を記述して, 各階層の並列度を指定 ハードウェア構成 並列化の階層 CUDA GPU 並列に実行する処理 Grid Streaming Multiprocessor スレッドの集まり Thread Block CUDA Core スレッド Thread 431
GPU の並列化の階層 グリッド-ブロック-スレッドの3 階層 グリッド (Grid) 並列に実行する処理 GPUが処理を担当する領域全体 スレッド (Thread) GPUの処理の基本単位 CPUのスレッドと同じ ブロック (Block) もしくはスレッドブロック (Thread Block)* スレッドの集まり * スレッドブロックだと長いのでブロックで統一 432
GPU の並列化の階層 各階層の情報を参照できる変数 x,y,z をメンバにもつ dim3 型構造体 グリッド (Grid) griddim グリッド内にあるブロックの数 ブロック (Block) blockidx blockdim ブロックに割り当てられた番号ブロック内にあるスレッドの数 スレッド (Thread) threadidx スレッドに割り当てられた番号 433
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<<<?,?>>>(); cudadevicesynchronize(); return 0; <<<>>> 内の数字を変えると画面表示される内容が変わる <<<>>> 内の数字とどのパラメータが対応しているかを確認 hellothreads.cu 434
GPU の構造とカーネルの書き方 GPU はマルチスレッド ( メニースレッド ) で並列処理 数百から数千の CUDA コアが搭載されており, それらが協調して並列処理を実行 カーネルには 1 スレッドが実行する処理を書く カーネルの作成を簡略化 カーネルを呼び出す際に並列処理の度合いを指定 カーネルと引数の間に追加した <<<,>>> で並列処理の度合を指定 435
各階層の値の設定 設定の条件 GPU の世代によって設定できる上限値が変化 確認の方法 pgaccelinfo devicequery GPU Computing SDK に含まれているサンプル CUDA Programming Guide https://docs.nvidia.com/cuda/cuda c programmingguide/#compute capabilities 436
pgaccelinfo の実行結果 Device Number: 0 Device Name: Tesla M2050 Device Revision Number: 2.0 Global Memory Size: 2817982464 Number of Multiprocessors: 14 Number of Cores: 448 Concurrent Copy and Execution: Yes Total Constant Memory: 65536 Total Shared Memory per Block: 49152 Registers per Block: 32768 Warp Size: 32 Maximum Threads per Block: 1024 Maximum Block Dimensions: 1024, 1024, 64 Maximum Grid Dimensions: 65535 x 65535 x 65535 Maximum Memory Pitch: 2147483647B Texture Alignment: 512B Clock Rate: 1147 MHz Initialization time: 4222411 microseconds Current free memory: 2746736640 Upload time (4MB): 2175 microseconds ( 829 ms pinned) Download time: 2062 microseconds ( 774 ms pinned) Upload bandwidth: 1928 MB/sec (5059 MB/sec pinned) Download bandwidth: 2034 MB/sec (5418 MB/sec pinned) 437
選択の際に重要行時のパラメータpgaccelinfo 実行結果 Revision Number: 2.0 GPU の世代 ( どのような機能を有しているか ) 実 Global Memory Size: 2817982464 Warp Size: 32 Maximum Threads per Block: 1024 Maximum Block Dimensions: 1024, 1024, 64 Maximum Grid Dimensions: 65535 x 65535 x 65535 各方向の最大値 1 ブロックあたりのスレッド数は最大 1024 (1024, 1, 1), (1, 1024, 1) (32, 32, 1), (4, 4, 64) など 438
レポート課題 2( 提出期限は 2 学期末 ) hellothreads.cu を実行し, 下記について考察せよ <<<,>>> 内の数字はどの情報を指定しているか 変数名 (griddim.x, etc.) とそれが表す情報の両方について考察 全スレッド数を 2 16, 1 ブロックあたりのスレッド数を 2 6 として実行するには,<<<,>>> 内にどのように記述すればよいか #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<<<?,? >>>(); cudadevicesynchronize(); } return 0; hellothreads.cu 439
レポートの書式 必ず表紙を付けること 授業名, 課題番号, 学籍番号, 氏名, 提出日に加えて課題に要した時間を書く 課題内容, プログラム, 実行結果, 考察で構成 プログラムを実行した tesla?? および GPU の番号も明記すること pdf 形式に変換してメールで提出 宛先 degawa at vos.nagaokaut.ac.jp メール題目 課題 2( 氏名 ) 440