CUDA
カーネルには 1 スレッドが実行する処理を書く
カーネル名と引数の間に <<<1,1>>> を付ける
2015/04/22 GPGPU実践プログラミング
92
Hello Thread （ Fermi 世代以降）
GPU の各スレッドが画面表示
#include<stdio.h>
__global__ void hello(){
printf("Hello Thread¥n");
}
int main(void){
hello<<<1,1>>>();
cudaThreadSynchronize();
return 0;
}
画面表示（Fermi世代以降で可能）
コンパイル時にオプションが必要
‐arch=sm_20以降
・・・
カーネル実行
・・・
ホストとデバイスの同期をとる
CPUとGPUは原則同期しないので，
同期しないとカーネルを実行した 直後にプログラムが終了
・・・
hellothread.cu
Hello Thread （ Fermi 世代以降）
<<< >>> 内の数字で並列度が変わることの確認
2015/04/22 GPGPU実践プログラミング
94
#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
CPU と GPU のやりとり
GPU の想定される使い方
ホスト
(CPU)
からデータを送り，デバイス(GPU)
で計算し，結果を受け取る
CPU
とGPU
のデータのやり取りが必要 GPU は原則データを返さない
PCI‐Ex
経由で描画情報を受け取り，画面に出力 カーネルの返値が
void
の理由CPU と GPU のやりとり
2015/04/22 GPGPU実践プログラミング
96
CUDA 独自の命令と C 言語のポインタを利用
GPU
のメモリ上に計算に必要なサイズを確保 確保したメモリのアドレスを
C
言語のポインタで格納 ポインタの情報を基にデータを送受信
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
CPU と GPU のやり取り（単純な加算）
2015/04/22 GPGPU実践プログラミング
98
関数の返値を void に変更し，メモリの動的確保を使用
#include<stdio.h>
#include<stdlib.h>
void add(int a, int b, int *c){
*c = a + b;
}
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;
}
引数で渡された変数の和を，cが指す アドレスに書き込み
・・・
引数にアドレスを追加
・・・
アドレスを基に結果を参照
・・・
add.c
CPU プログラム（メモリの動的確保）
malloc
指定したバイト数分のメモリを確保
stdlib.h
をインクルードする必要がある sizeof
データ型
1
個のサイズ（バイト数）を求める#include<stdlib.h>
int *a;
a = (int *)malloc( sizeof(int)*100 );
printf("%d, %d¥n", sizeof(float), sizeof(double));
実行すると4,8と表示される
CPU と GPU のやり取り（単純な加算）
2015/04/22 GPGPU実践プログラミング
100
add.c の処理の一部を GPU の処理に置き換え
#include<stdio.h>
__global__ void add(int a, int b, int *c){
*c = a + b;
}
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;
}
__global__を追加
・・・
GPU上のメモリに確保される変数のアドレス
・・・
GPU上にint型変数 一個分のメモリを確保
・・・
↑GPUから結果をコピー メモリを解放
・・・
add.cu
CUDA でカーネルを作成するときの制限
カーネルの引数
値を渡すことができる
GPU
のメモリを指すアドレス
CPU
のメモリを指すアドレスも渡すことは可能 そのアドレスを基にホスト側のメモリを参照することは不可能
printf などの画面出力
Fermi
世代以降のGPU
で，コンパイルオプションを付与
‐arch={sm_20|sm_21|sm_30|sm_32|sm_35|sm_50|sm_52}
エミュレーションモード
新しい
CUDA
（4.0
以降）では消滅CPU プログラムの超簡単移植法
とりあえず GPU で実行すればいいのなら･･･
拡張子を .cu に変更
GPU の都合を反映
関数の返値を
void
にし，__global__
を付ける 関数名と引数の間に
<<<1,1>>>
を付ける
GPU
で使うメモリをcudaMalloc
で確保
malloc
でメモリを確保していればそれをcudaMalloc
に置き換え
GPU
からデータを受け取るためにcudaMemcpy
を追加 最適化は追々考えればいい
2015/04/22 GPGPU実践プログラミング
102
カーネルの完成
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
GPU の並列化の階層
GPU のハードウェアの構成に対応させて並列性を管理
並列化の各階層における情報を利用
GPU
Streaming Multiprocessor
CUDA Core
ハードウェア構成
並列に実行する 処理
スレッドの集 まり
スレッド
並列化の階層
Grid
Block
Thread
CUDA
2015/04/22 GPGPU実践プログラミング
104
GPU の並列化の階層
グリッド－ブロック－スレッドの 3 階層
各階層の情報を参照できる変数
x,y,z
をメンバにもつdim3
型構造体 グリッド（ Grid ）
gridDim
グリッド内にあるブロックの数 ブロック（ Block ）
blockIdx
ブロックに割り当てられた番号
blockDim
ブロック内にあるスレッドの数 スレッド（ Thread ）
threadIdx
スレッドに割り当てられた番号Hello Threads （ Fermi 世代以降）
<<< >>> 内の数字で表示される内容が変化
2015/04/22 GPGPU実践プログラミング
106
#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
GPU の構造とカーネルの書き方
GPU はマルチスレッド
（メニースレッド）で並列処理
関数には 1 スレッドが実行する処理を書く
関数を実行する際に並列処理の度合いを指定
カーネルと引数の間に追加した <<<,>>> で並列処理の 度合を指定
<<<
グリッド内にあるブロックの数,1
ブロックあたりのスレッド の数>>>
プログラム実習
以下のプログラムをコンパイルし，正しく実行できること を確認せよ
helloworld.c
helloworld.cu
hellothread.cu
hellothreads.cu
hellothreads.cu については， <<<>>> 内の数字を変 更し，実行結果がどのように変わるか確認せよ
GPGPU実践プログラミング
108 2015/04/22