長岡技術科学大学 電気電子情報工学専攻 出川智啓
今回の内容
GPUプログラミング環境（CUDA）
GPUプログラムの実行の流れ
CUDAによるプログラムの記述
カーネル（GPUで処理する関数）の構造 記述方法とその理由 GPU固有のパラメータの確認
GPU（Graphics Processing Unit）とは
画像処理専用のハードウェア
具体的には画像処理用のチップ チップ単体では販売されていない PCI‐Exカードで販売（チップ単体と区別せずにGPUと呼ぶことも多い） マザーボードやノートPCに搭載 PCI‐Exカードとして販売されるGPUには，ビデオメモリと呼ばれ るRAMが搭載GPUのハードウェア構造
CUDA Core（旧Streaming Processor, SP）と呼ばれ
る演算器を多数搭載
Streaming Multiprocessor（SM, SMX）が複数の
CUDA CoreとSFU，メモリをまとめて管理
SFU(Special Function Unit) 数学関数を計算するユニット 複数のSMが集まってGPUを構成
Fermiアーキテクチャの構造
Tesla M2050の仕様
SM数 14 CUDA Core数 448(=32 Core/SM×14 SM) 動作周波数 1,150 MHz 単精度演算ピーク性能 1.03 TFLOPSCUDA CoreCUDA CoreCUDA CoreCUDA
CoreCUDA CoreCUDA CoreCUDA CoreCUDA Core
Streaming Multiprocessor
GPUの模式図
GPU Streaming Multiprocessor L2キャッシュ ローカル メモリ コンスタントメモリ テクスチャメモリ GPU Chip グローバルメモリ ローカル SM SM SM ・・・ SM ・・・ SM SM SM ・・・ SM レジ スタ レジ スタ レジ スタ レジ スタ CUDA Core CUDA Core CUDA Core CUDA Core L1キャッ シュ 共有 メモリ Streaming Multiprocessor SM SM SM ・・・ SMCUDA
Compute Unified Device Architecture
NVIDIA社製GPU向け開発環境(Windows,Linux,Mac OS X) 2007年頃発表 C/C++言語＋独自のGPU向け拡張 専用コンパイラ（nvcc）とランタイムライブラリ いくつかの数値計算ライブラリ（線形代数計算，FFTなど） CUDA登場以前
グラフィクスプログラミングを利用 足し算を行うために，色を混ぜる処理を実行 汎用計算のためには多大な労力が必要プログラマブルシェーダを用いた汎用計算
グラフィックス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; } 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) シェーダ言語を用いた配列加算 （c=*a + *b）の例
CUDAによるプログラミング
CPUをホスト（Host），GPUをデバイス（Device）と表現
ホスト（CPU）
処理の流れやGPUを利用するための手続きを記述 プログラムの書き方は従来のC言語と同じ 利用するGPUの決定，GPUへのデータ転送，GPUで実行する関 数の呼び出し等CUDAによるプログラミング
CPUをホスト（Host），GPUをデバイス（Device）と表現
デバイス（GPU）
処理する内容を関数として記述 引数は利用可能，返値は利用不可（常にvoid） 関数はkernelと呼ばれる 関数呼び出しはlaunch, invokeなどと呼ばれるHello World
何を確認するか
最小構成のプログラムの作り方 ファイル命名規則（拡張子は.c/.cpp） コンパイルの方法(gcc, cl等を使用) #include<stdio.h> int main(void){ printf("hello world¥n"); return 0; } helloworld.cCUDAでHello World
何を確認するか
最小構成のプログラムの作り方 ファイル命名規則（拡張子は.cu） コンパイルの方法（nvccを使用） #include<stdio.h> int main(void){ printf("hello world¥n"); return 0; } #include<stdio.h> int main(void){ printf("hello world¥n"); return 0; } 違いは拡張子だけ？CUDAプログラムのコンパイル
ソースファイルの拡張子は
.cu
nvccを用いてコンパイル
CPUが処理する箇所はgcc等がコンパイル GPUで処理する箇所をnvccがコンパイル helloworld.cuにはCPUで処理する箇所しかない
CUDAでHello World
CUDA専用の処理を追加
#include<stdio.h> __global__ void kernel(){} int main(void){ kernel<<<1,1>>>(); printf("hello world¥n"); return 0; } GPUで実行される関数（カーネル） __global__が追加されている ・・・ 通常の関数呼出とは異なり， <<<>>>が追加されている ・・・ hellokernel.cuCUDAプログラムの実行
実行時の流れ
（CPU視点） 利用するGPUの初期化やデータの転送などを実行 GPUで実行する関数を呼び出し GPUから結果を取得 初期化の指示 初期化 カーネルの実行指示 カーネルを実行 実行結果をコピー time CPUとGPUは非同期 CPUは別の処理を実行可能 必要なデータのコピー メモリに書込CPU
GPU
Hello Thread（Fermi世代以降）
printfをGPUから呼び出し，並列に実行
#include<stdio.h> int hello(){ printf("Hello Thread¥n"); return 0; } int main(void){ hello(); return 0; } 画面表示 ・・・ 関数呼び出し ・・・ hellothread.cHello Thread（Fermi世代以降）
GPUの各スレッドが画面表示
#include<stdio.h> __global__ void hello(){ printf("Hello Thread¥n"); } int main(void){ hello<<<1,1>>>(); cudaDeviceSynchronize(); return 0; } 画面表示（Fermi世代以降で可能） コンパイル時にオプションが必要 ‐arch=sm_20以降 ・・・ カーネル実行 ・・・ ホストとデバイスの同期をとる CPUとGPUは原則同期しないので， 同期しないとカーネルを呼び出し た直後にプログラムが終了 ・・・ hellothread.cuCUDAでカーネルを作成するときの制限
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.cuGPUプログラムへの変更
変更点
変更の理由
変更点
関数の前に修飾子__global__をつけた 変更によって実現されること
GPUで実行する関数という目印になる 変更が必要な理由
ホスト（CPU）からGPUで実行する関数（カーネル）を呼び出し CPUが処理する箇所とGPUが処理する箇所は別のコンパイラ がコンパイル コンパイルの時点でどれがカーネルかを明記GPUプログラムへの変更
変更点
変更の理由
変更点
関数helloの返値をvoidにした 変更によって実現されること
GPUのハードウェア構造に適したプログラムを作成できる 変更が必要な理由
GPUはホストと別に独立したメモリを持つ GPUは描画情報を受け取り，画面に出力 GPU→CPUの頻繁なデータ転送は苦手 プログラマがメモリ管理を行い，無駄なデータ転送による実行 描画情報 画面出力GPUプログラムへの変更
変更点
変更の理由
変更点
関数呼出の際に関数名と引数の間に<<<1,1>>>を付けた 変更によって実現されること
GPUのハードウェア構造に適したプログラムを作成できる 変更が必要な理由
GPUには数百から数千のCUDAコアが搭載されており，それらが 協調して並列処理を実行 1スレッドが実行する処理を書くことでカーネルの作成を簡略化 並列処理の度合いはカーネル呼出の際に指定GPUプログラムへの変更
