Microsoft PowerPoint - GPGPU実践基礎工学（web）.pptx

(1)

GPUのメモリ階層

(2)

今回の内容



_{GPUのメモリ階層}



グローバルメモリ



共有メモリ



モザイク処理への適用



コンスタントメモリ



空間フィルタへの適用

(3)

GPUの主要部品



基盤



画面出力端子



電源入力端子



_{GPU（チップ）+冷却部品}



メモリ



特性の把握が重要

画面出力端子

PCI‐Ex端子

電源入力端子

メモリ

チップ

(4)

CPUのメモリ階層



オフチップ(off‐chip)メモリ



_{CPUのチップ外部に置かれたメモリ}



メインメモリ（主記憶）



利用可能なメモリの中で最低速，最大容量



オンチップ(on‐chip)メモリ



_{CPUのチップ内部に置かれたメモリ}



レジスタ



レベル1(L1)キャッシュ



レベル2(L2)キャッシュ



レベル2(L3)キャッシュ

高速，容量小

低速，容量大

(5)

GPUのメモリ階層



オフチップメモリ



_{PCI‐Exカードの基板上に実装}



ビデオメモリ



利用可能なメモリの中で最低速，最大容量



オンチップメモリ



_{GPUのチップ内部に置かれたメモリ}



レジスタ



レベル1(L1)キャッシュ



レベル2(L2)キャッシュ

CPUの構造に類似

高速，容量小

低速，容量大

(6)

GPUメモリの独自の特徴



_{CPUとは異なるメモリを複数搭載}



各メモリの特徴を知り，適材適所で利用する事により高速化



_{GPUから読み書き可能か}



処理を行うスレッドから読み書き可能か，読み込みのみか



複数のスレッドでデータを共有できるか



_{CPUから読み書き可能か}



_{C言語の機能のみで直接読み書きは不可能}



_{CUDAの専用関数(API)を利用して読み書き}

(7)

メモリの階層



_{CPUのメモリ階層}



コアごとにL2キャッシュ，全体でL3キャッシュを持つこともある

メインメモリ

L2キャッシュ

・・・

チップ

_コア

L1キャッシュ

演算器レジスタレジスタ演算器

L1キャッシュ

演算器レジスタレジスタ演算器

L1キャッシュ

演算器レジスタレジスタ演算器コアコア

(8)

メモリの階層



_{GPUのメモリ階層}



_{CPUにはない独自のメモリを複数持つ}

グローバルメモリ

L2キャッシュ

L1キャッシュ

共有メモリ

CUDA Core レジスタ

チップ

テクスチャ

コンスタント

ローカル

テクスチャ

キャッシュ

コンスタント

キャッシュ

(9)

メモリの種類



オンチップメモリ（GPUのチップ内部に置かれたメモリ）



高速アクセス，小容量



_{CPUからはアクセス不可}



_{L1キャッシュと共有メモリは一定サイズを共用}

L1キャッシュ/共有（シェアー

ド）メモリ

レジスタ

容量

小

速度

高速

GPUからの

読み書き

読み書き可

ブロック内の全スレッドが同じアドレスにアクセス（データを共有）することが可能*

読み書き可

各スレッドが異なるアドレスにアクセス

CPUからの

アクセス

読み書き不可

*スレッドごとに異なるアドレスにアクセスすることも可能

(10)

メモリの種類



オフチップメモリ（GPUのチップ外部に置かれたメモリ）



低速アクセス，大容量



_{CPUから直接アクセス可能}



ローカルメモリだけはアクセス不可

グローバルメモリ

ローカルメモリ

テクスチャメモリ

コンスタントメモリ

容量

大

小

大

小

速度

低速

_高速*

GPUからの

読み書き

読み書き可

全てのスレッドが同じアドレスにアクセス可能**

読み書き可

読み込み可

CPUからの

アクセス

読み書き可

読み書き不可

読み書き可

*キャッシュが効く場合 **スレッドごとに異なるアドレス

(11)



共有メモリとL1キャッシュは

一定サイズを共用



グローバルメモリへのアクセ

スはL2キャッシュ経由



_{Fermi世代以前のGPUは}

キャッシュ無し*

メモリの種類

オフチップメモリ

オンチップメモリ

*Tesla世代でもテクスチャキャッシュは存在．

ホスト

メモリ

L2キャッシュ

ローカル

メモリ

コンスタントメモリ

テクスチャメモリ

GPU

Chip

レジスタレジスタレジスタレジスタ CUDA Core CUDA Core CUDA Core CUDA Core L1キャッシュ共有メモリ

SM

グローバルメモリ

ローカル

メモリ

(12)

メモリの種類と並列化階層の対応



オンチップメモリ



ブロックまたはスレッドごとに

異なる値を持てる



レジスタが不足するとローカ

ルメモリが使われる



オフチップメモリ



_{GPU全体で共通の値を持てる}

各GPU(Grid)内

でデータを共有

各ブロック内で

データを共有

各スレッドが個別の

データを保有

ホスト

メモリ

L2キャッシュ

ローカル

メモリ

コンスタントメモリ

テクスチャメモリ

Grid

レジスタレジスタレジスタレジスタ Thre ad 0 Thre ad 1 Thre ad 2 Thre ad 3 L1キャッシュ共有メモリ

Block(0,0,0)

レジスタレジスタレジスタレジスタ Thre ad 0 Thre ad 1 Thre ad 2 Thre ad 3 L1キャッシュ共有メモリ

Block(1,0,0)

グローバルメモリ

ローカル

メモリ

(13)

レジスタ



各スレッドが個別に利用



カーネル内で変数を宣言する

とレジスタを利用



非常に高速



キャッシュとしても利用可能



少容量



_{32768本*×32bit}



利用可能分を超え

るとローカルメモリ

へ追い出される



レジスタスピル

ホスト

メモリ

L2キャッシュ

ローカル

メモリ

コンスタントメモリ

テクスチャメモリ

GPU

Chip

SM

グローバルメモリ

ローカル

メモリ

(14)

グローバルメモリ



ビデオメモリ（数GB）



_{CPUのメインメモリに相当}



読み込みがある一定サイズ

でまとめて行われる



レイテンシが大きい



読み出し命令を発効してから

データが得られるまでの時間



効率よくアクセス

するための条件

がある



コアレスアクセス

（コアレッシング）

ホスト

メモリ

L2キャッシュ

ローカル

メモリ

コンスタントメモリ

テクスチャメモリ

GPU

Chip

SM

グローバルメモリ

ローカル

メモリ

(15)

グローバルメモリへのアクセス

*



_{16スレッドが協調して同時にアクセス}



_{CUDAでは32スレッドをWarpという単位で管理}



_{Warpの半分（Half Warp）が協調して読み書き}



コアレスアクセスか否かで読込速度が大幅に変化



新しい世代のGPUでは速度の落ち込みが緩和



コアレスアクセスはGPUプログラムで最重要



_{GPUの処理能力と比較するとデータ供給不足が発生}



効率よくデータを供給するためにコアレスアクセスは必須

*かなり古い情報のため要注意． Fermi世代以降では状況が大きく異なる．

(16)

コアレスアクセスになる条件

*



データのサイズ



_{32bit, 64bit, 128bit（4バイト, 8バイト, 16バイト）}



アドレスの隣接



_{16スレッドがアクセスするアドレスがスレッド番号順に隣接}



アクセスする最初のアドレス



_{16スレッドがアクセスするアドレスの先頭（スレッド0がアクセ}

スするアドレス）が，64バイトまたは128バイト境界



アドレスが64の倍数で始まる64バイトの領域か，アドレスが128の倍

数で始まる128バイトの領域

(17)

コアレスアクセスの例

*



データ型が

32bit=4バイト



各スレッドが連

続して隣接アド

レスにアクセス



先頭アドレスが

128バイト境界

₁₆

スレッド



データ型が

32bit=4バイト



各スレッドが連

続して隣接アド

レスにアクセス



実際にデータ

を取得するか

は無関係



先頭アドレス

が128バイト

境界

T15 T14 T13 T12 T11 T10 T9 T8 T7 T6 T5 T4 T3 T2 A188 A184 A180 A176 A172 A168 A164 A160 A156 A152 A148 A144 A140 A136 T1 A132 T0 A128 T15 T14 T13 T12 T11 T10 T9 T8 T7 T6 T5 T4 T3 T2 A188 A184 A180 A176 A172 A168 A164 A160 A156 A152 A148 A144 A140 A136 T1 A132 T0 A128 *かなり古い情報のため要注意． Fermi世代以降では状況が大きく異なる．

(18)

T15 T14 T13 T12 T11 T10 T9 T8 T7 T6 T5 T4 T3 T2 T1 T0 A188 A176 A164 A152 A140 A128 T15 T14 T13 T12 T11 T10 T9 T8 T7 T6 T5 T4 T3 T2 A188 A184 A180 A176 A172 A168 A164 A160 A156 A152 A148 A144 A140 A136 T1 A132 T0 A128 T15 T14 T13 T12 T11 T10 T9 T8 T7 T6 T5 T4 T3 T2 A188 A184 A180 A176 A172 A168 A164 A160 A156 A152 A148 A144 A140 A136 T1 A132 T0 A128 T15 T14 T13 T12 T11 T10 T9 T8 T7 T6 T5 T4 T3 T2 A188 A184 A180 A176 A172 A168 A164 A160 A156 A152 A148 A144 A140 A136 T1 A132 T0 A128

コアレスアクセスにならない例

*

各スレッドが番

号順にアクセス

していない

先頭が128バイト

境界ではない

（現在は緩和）

アドレスが連続し

ていない

データが32bit,

64bit, 128bit

ではない

（構造体など）

(19)

コアレスアクセスにならない例

*



_{128バイト境界からわずかに}

ずれている場合



_{Tesla世代以降は64バイトブ}

ロックと32バイトブロックに分

けて読込



メモリアクセス要求は2回



コアレスアクセスの半分程度

の性能は得られる

A190

64バイトブロック

でデータ読込

32バイトブロック

でデータ読込

T15 T14 T13 T12 T11 T10 T9 T8 T7 T6 T5 T4 T3 T2 A188 A184 A180 A176 A172 A168 A164 A160 A156 A152 A148 A144 A140 A136 T1 A132 T0 A128 *かなり古い情報のため要注意． Fermi世代以降では状況が大きく異なる．

(20)

モザイク処理



前回授業で取り上げたモザイク処理



画像を小さな領域に分け，その領域を全て同じ色にする



領域内の全画素を，領域内の画素の平均値に置き換える

(21)

モザイク処理



前回授業で取り上げたモザイク処理



高速化が今ひとつ



他の処理より一桁遅い

処理

処理時間[ms]

CPU

GPU

ネガティブ処理

₁₇₅

_1.17

水平反転

₁₈₇

_1.18

垂直反転

₁₈₅

_1.18

空間フィルタ

₅₅₃

_4.13

モザイク処理

₂₆₀

_38.5

(22)

__global__ void mosaic(unsigned char *p, unsigned char *filtered, int mosaicSize){ int i,j, isub,jsub; int average; i = blockIdx.x*blockDim.x + threadIdx.x; j = blockIdx.y*blockDim.y + threadIdx.y; if(threadIdx.x == 0 && threadIdx.y == 0){//ブロック内の1スレッドのみが処理 //領域内の画素の平均値を計算 average = 0; for(jsub = 0; jsub<mosaicSize; jsub++){ for(isub = 0; isub<mosaicSize; isub++){ average += p[(i+isub) + WIDTH*(j+jsub)]; } } average /= (mosaicSize*mosaicSize); //領域内の画素を計算した平均値で塗りつぶす for(jsub = 0; jsub<mosaicSize; jsub++){ for(isub = 0; isub<mosaicSize; isub++){ filtered[(i+isub) + WIDTH*(j+jsub)] = (unsigned char)average; } } } }

モザイク処理（前回授業で作成）

imageproc.cu

(23)

i = blockIdx.x*blockDim.x + threadIdx.x;

j = blockIdx.y*blockDim.y + threadIdx.y;

if(threadIdx.x == 0 && threadIdx.y == 0){

//1スレッドのみが処理

//領域内の画素の平均値を計算

//計算した平均値をグローバルメモリへ書き込む

}

モザイク処理（前回授業で作成）



ブロックに分割し，ブロック内の1ス

レッドのみが動作



ブロックの数だけ並列に処理



ブロックの中では1スレッドのみが処理



グローバルメモリから画素情報を読込



ブロック内の画素の平均値を計算



グローバルメモリに画素の平均値を書込



残りのスレッドはi,jの計算をするだけ

p[]

(24)

モザイク処理が高速化されない原因



各ブロックの1スレッドのみが処理を実行



グローバルメモリから画素情報を読込



ブロック内の画素の平均値を計算



グローバルメモリに画素の平均値を書込



グローバルメモリへコアレスアクセスできていない



複数のスレッドが協調し，アドレスが隣接したメモリを読むと高速



ある1スレッドが1画素ずつメモリアドレスを変えながら読むのは

最悪の処理

(25)

モザイク処理が高速化されない原因



せめてコアレスアクセスしたい



平均を並列に計算するのは中級レベルの処理



カーネル内で変数を宣言するとレジスタが使われる



下の例ではaverageとcacheはレジスタに確保



ブロック内で値を共有できない

__global__ void mosaic(unsigned char *p, unsigned char *filtered, int mosaicSize){ int i,j, isub,jsub;

int cache, average; //cacheとaverageはレジスタに確保（各スレッドが異なる値を保持）

i = blockIdx.x*blockDim.x + threadIdx.x; j = blockIdx.y*blockDim.y + threadIdx.y; cache = p[(i+isub) + WIDTH*(j+jsub)];//グローバルメモリから読み込み（コアレスアクセス） if(threadIdx.x == 0 && threadIdx.y == 0){ //領域内の画素の平均値を計算したいが，他のスレッドが持つcacheの値を読む事は不可能 } //グローバルメモリに書き出し（コアレスアクセス） filtered[(i+isub) + WIDTH*(j+jsub)] = (unsigned char)average; }

(26)

メモリの種類



オンチップメモリ（GPUのチップ内部に置かれたメモリ）



高速アクセス，小容量



_{CPUからはアクセス不可}



_{L1キャッシュと共有メモリは一定サイズを共用}

L1キャッシュ/共有（シェアー

ド）メモリ

レジスタ

容量

小

速度

高速

GPUからの

読み書き

読み書き可

ブロック内の全スレッドが同じアドレスにアクセス（データを共有する）ことが可能*

読み書き可

各スレッドが異なるレジスタにアクセス

CPUからの

アクセス

読み書き不可

*スレッドごとに異なるアドレスにアクセスすることも可能

(27)

共有（シェアード）メモリ



ブロック内のスレッドが共通

のデータにアクセス可能



_{1回目のアクセスに時間がか}

かるが，それ以降は非常に

高速にアクセス可能



_{Fermi世代以前の}

GPUで管理可能な

キャッシュとして利

用



_{1ブロックあたり}

16kB～48kB

ホスト

メモリ

L2キャッシュ

ローカル

メモリ

コンスタントメモリ

テクスチャメモリ

GPU

Chip

SM

グローバルメモリ

ローカル

メモリ

(28)

共有（シェアード）メモリの宣言



カーネル内で修飾子

_{shared}

を付けて宣言



配列として宣言



配列サイズを静的（コンパイル時）に決定する場合



shared 型変数名[要素数]



多次元配列も宣言可能



配列サイズを動的（カーネル実行時）に決定する場合



_extern

_{shared 型変数名[]}



サイズはカーネル呼出時のパラメータで指定



_{<<<ブロック数，スレッド数，}

共有メモリのバイト数

_>>>

(29)

共有メモリを使ったモザイク処理の高速化



ブロック内の全スレッドで共有したい値



各スレッドがグローバルメモリから読んだ画素情報



ブロック内の画素の平均値



処理の流れ

1. ブロック内の全スレッドがグローバルメモリから画素の値を読

み出し，共有メモリに置く（コアレスアクセス）

2. ある1スレッドが共有メモリに置かれた画素の値を読み，画素

の平均値を計算し，共有メモリに置く

3. ブロック内の全スレッドが共有メモリに置かれた画素の平均

値を読み，グローバルメモリに書き込む（コアレスアクセス）

(30)

__global__ void mosaic_shared(unsigned char *p,unsigned char *filtered, int mosaicSize){ int i,j, isub,jsub;

__shared__ int average;

__shared__ unsigned char cache[THREADX][THREADY]; i = blockIdx.x*blockDim.x + threadIdx.x; j = blockIdx.y*blockDim.y + threadIdx.y; cache[threadIdx.x][threadIdx.y] = p[i + WIDTH*j]; if(threadIdx.x == 0 && threadIdx.y == 0){ average = 0; for(jsub = 0; jsub<mosaicSize; jsub++){ for(isub = 0; isub<mosaicSize; isub++){ average += cache[isub][jsub]; } } average /= (mosaicSize*mosaicSize); } filtered[i + WIDTH*j] = (unsigned char)average; }

モザイク処理（共有メモリを利用）

(31)

shared

int average;

shared unsigned char cache[THREADX][THREADY];

i = blockIdx.x*blockDim.x + threadIdx.x;

j = blockIdx.y*blockDim.y + threadIdx.y;

cache[threadIdx.x][threadIdx.y] = p[i + WIDTH*j];

共有メモリの宣言



ブロック内のスレッド数

分の共有メモリを確保



画素情報を共有するのは1

ブロックの中だけなので，

配列サイズはブロック内の

スレッド数分でよい



各スレッドがグローバル

メモリからデータを読み，

共有メモリに書込む

p[]

cache[][]

(32)

if(threadIdx.x == 0 && threadIdx.y == 0){

average = 0;

for(isub = 0; isub<mosaicSize; isub++){

for(jsub = 0; jsub<mosaicSize; jsub++){

average += cache[isub][jsub];

}

average /= (mosaicSize*mosaicSize);

}

共有メモリの宣言



_{1スレッドが共有メモリに置かれた画素}

の値を読み，画素の平均値を計算



画素の平均値averageも共有メモリに

存在



ブロック内の全スレッドがaverageにアクセ

ス可能

cache[][]

average

(33)

filtered[i + WIDTH*j] = (unsigned char)average;

共有メモリの宣言



各スレッドがaverageを

読み込み



グローバルメモリの各位

置に画素の平均値を書

き込む

filtered[]

average

(34)

実行結果



正しく処理できていない



ブロック内のスレッドの協調が不十分

(35)

int i,j, isub,jsub; __shared__ int average; __shared__ unsigned char cache[THREADX][THREADY]; i = blockIdx.x*blockDim.x + threadIdx.x; j = blockIdx.y*blockDim.y + threadIdx.y; cache[threadIdx.x][threadIdx.y] = p[i + WIDTH*j]; if(threadIdx.x == 0 && threadIdx.y == 0){ average = 0; for(jsub = 0; jsub<mosaicSize; jsub++){ for(isub = 0; isub<mosaicSize; isub++){ average += cache[isub][jsub]; } } average /= (mosaicSize*mosaicSize); } filtered[i + WIDTH*j] = (unsigned char)average;

モザイク処理（共有メモリを利用）

ｽﾚｯﾄﾞ

₀

ｽﾚｯﾄﾞ

_0以外

他のスレッドが共有メモリに書き込む前にcache[][]にアクセスする可能性がある

(36)

int i,j, isub,jsub; __shared__ int average; __shared__ unsigned char cache[THREADX][THREADY]; i = blockIdx.x*blockDim.x + threadIdx.x; j = blockIdx.y*blockDim.y + threadIdx.y; cache[threadIdx.x][threadIdx.y] = p[i + WIDTH*j]; if(threadIdx.x == 0 && threadIdx.y == 0){ average = 0; for(jsub = 0; jsub<mosaicSize; jsub++){ for(isub = 0; isub<mosaicSize; isub++){ average += cache[isub][jsub]; } } average /= (mosaicSize*mosaicSize); } filtered[i + WIDTH*j] = (unsigned char)average;

モザイク処理（共有メモリを利用）

ｽﾚｯﾄﾞ

₀

ｽﾚｯﾄﾞ

_0以外

スレッド0以外は平均値を計算しないので，直ちにこの行に到達し， averageの値が確定する前に値を読んでfilteredに書き込む可能性がある

(37)

ブロック内でのスレッドの同期



_{__syncthreads();}



カーネル実行中にスレッドの同期を取る



_{__syncthreads()が書かれた行にスレッドが到達すると，}

同一ブ

ロック内の他の全てのスレッド

がその行に達するまで待機



異なるブロック間での同期は不可能



_{ifの中に記述するとカーネルが終了しないことがある}

if(条件){

__syncthreads();

//条件が真にならないスレッドはifの中に入らないため，

//カーネルが永久に終わらない

}

(38)

int i,j, isub,jsub; __shared__ int average; __shared__ unsigned char cache[THREADX][THREADY]; i = blockIdx.x*blockDim.x + threadIdx.x; j = blockIdx.y*blockDim.y + threadIdx.y; cache[threadIdx.x][threadIdx.y] = p[i + WIDTH*j]; __syncthreads(); if(threadIdx.x == 0 && threadIdx.y == 0){ average = 0; for(isub = 0; isub<mosaicSize; isub++){ for(jsub = 0; jsub<mosaicSize; jsub++){ average += cache[isub][jsub]; } } average /= (mosaicSize*mosaicSize); } __syncthreads(); filtered[i + WIDTH*j] = (unsigned char)average;

モザイク処理（共有メモリを利用）

ｽﾚｯﾄﾞ

₀

ｽﾚｯﾄﾞ

_0以外

if文の前でブロック内の全スレッドが同期しているので， cache[][]には画素情報が入っている

ブロック内で同期

スレッド0以外はif文を実行しないが，スレッド0が到達する（平均値を計算し終わる）まで__syncthreads()で停止

ブロック内で同期

imageproc_mem.cu

(39)

実行結果



正しく処理できている



実行時間

_{260 ms(CPU) →}

_{38.5 ms →}

_18.7 ms



処理時間が約1/2に短縮

(40)

空間フィルタ



ある画素とその周囲の画素を使って処理



処理の仕方を規定したカーネルを定義



カーネルは1次元配列で表現

原画像

輪郭抽出

0 1 0 1 ‐4 1 0 1 0 = b+d‐4e+f+h フィルタ（カーネル） a b c d e f g h i

(41)

空間フィルタ



カーネルは1次元配列で表現



ぼかし（平均フィルタ）



輪郭抽出

1/9 1/9 1/9 1/9 1/9 1/9 1/9 1/9 1/9 0 1 0 1 ‐4 1 0 1 0

float blur[9] ={1.0f/9.0f,1.0f/9.0f,1.0f/9.0f,

1.0f/9.0f,1.0f/9.0f,1.0f/9.0f,

1.0f/9.0f,1.0f/9.0f,1.0f/9.0f};

float laplacian[9] ={ 0.0f, 1.0f, 0.0f,

1.0f,‐4.0f, 1.0f,

0.0f, 1.0f, 0.0f};

(42)

__global__ void boxfilter(unsigned char *p,unsigned char *filtered, float *filter){ int i,j;

int result = BLACK;

i = blockIdx.x*blockDim.x + threadIdx.x; j = blockIdx.y*blockDim.y + threadIdx.y;

if(0<i && i<WIDTH‐1 && 0<j && j<HEIGHT‐1){ //端の画素は処理をしないため，ifで処理を分岐 result = filter[0]*p[(i‐1) + WIDTH*(j‐1)] +filter[1]*p[(i ) + WIDTH*(j‐1)] +filter[2]*p[(i+1) + WIDTH*(j‐1)] +filter[3]*p[(i‐1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i‐1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; } if(result<BLACK) result = ‐result; //数値が負になっていれば‐1をかける if(result>WHITE) result = WHITE; //数値が255を超えていれば255に収める filtered[i+WIDTH*j] = (unsigned char)result; }

空間フィルタ（前回授業で作成）

imageproc.cu

(43)

if(0<i && i<WIDTH‐1 && 0<j && j<HEIGHT‐1){ //端の画素は処理をしないため，ifで処理を分岐 result = filter[0]*p[(i‐1) + WIDTH*(j‐1)] +filter[1]*p[(i ) + WIDTH*(j‐1)] +filter[2]*p[(i+1) + WIDTH*(j‐1)] +filter[3]*p[(i‐1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i‐1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; }

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

(44)

if(0<i && i<WIDTH‐1 && 0<j && j<HEIGHT‐1){ //端の画素は処理をしないため，ifで処理を分岐 result = filter[0]*p[(i‐1) + WIDTH*(j‐1)] +filter[1]*p[(i ) + WIDTH*(j‐1)]

+filter[2]*p[(i+1) + WIDTH*(j‐1)] +filter[3]*p[(i‐1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i‐1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; }

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

j

(45)

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

j

(46)

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

j

(47)

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

j

(48)

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

j

(49)

+filter[2]*p[(i+1) + WIDTH*(j‐1)] +filter[3]*p[(i‐1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )]

+filter[6]*p[(i‐1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; }

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

j

(50)

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

j

(51)

+filter[2]*p[(i+1) + WIDTH*(j‐1)] +filter[3]*p[(i‐1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i‐1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)]

+filter[8]*p[(i+1) + WIDTH*(j+1)]; }

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

j

(52)

空間フィルタ（前回授業で作成）

p[]

filter[]



メモリアクセス



画像（配列p)へは複数スレッドが隣接し

たメモリアドレスにアクセス



コアレスアクセス可能



フィルタ（配列filter)へは複数スレッ

ドが同じ要素にアクセス



コアレスアクセス不可能



_{L2キャッシュに入る可能性はある}

j

(53)

空間フィルタ処理の高速化



フィルタ（配列filter）へのアクセス



コアレスアクセスできていない



_{L2キャッシュに入る可能性は高いが，有効活用されているか}

は不明



配列filterへのアクセスの高速化



共有メモリを利用すると処理が冗長



_{1ブロックから9スレッドを選び，グローバルメモリから共有メモリへコ}

ピーし，同期をとる



コンスタントメモリが活用できる

(54)

メモリの種類



オフチップメモリ（GPUのチップ外部に置かれたメモリ）



低速アクセス，大容量



_{CPUから直接アクセス可能}



ローカルメモリだけはアクセス不可

グローバルメモリ

ローカルメモリ

テクスチャメモリ

コンスタントメモリ

容量

大

小

大

小

速度

低速

_高速*

GPUからの

読み書き

読み書き可

読み込み可

CPUからの

アクセス

読み書き可

読み書き不可

読み書き可

*キャッシュが効く場合 **スレッドごとに異なるアドレス

(55)

ホスト

メモリ

L2キャッシュ

ローカル

メモリ

コンスタントメモリ

テクスチャメモリ

GPU

Chip

SM

グローバルメモリ

ローカル

メモリ

コンスタントメモリ



_{GPU全体で同じメモリにアク}

セス



メモリを読み取り専用とする

ことで値をキャッシュし，一

度読んだ値を再利用



_{GPU全体で64kB}

(56)

コンスタントメモリの宣言



グローバル領域で修飾子

_{constant}

を付けて宣言



配列サイズは静的に決定



_{constant 型変数名;}



_{constant 型変数名[要素数];}



配列としても宣言可能



サイズはコンパイル時に確定している必要がある



_{cudaMalloc()やcudaFree()は不要}



グローバル変数として宣言し，複数のカーネルからアクセ

スすることが多い



読込専用なので許される



書込可能なメモリでは厳禁

(57)

コンスタントメモリの利用



メモリは読込専用



_{CPUからは変更可能}



専用のメモリ転送命令でコピー



_{cudaMemcpyToSymbol}



_{CPU上のメモリをコンスタントメモリにコピーする}



_{cudaMemcpyToSymbol(}

転送先変数名

_{, 転送元アドレス,}

バイト数, オフセット, 方向);



オフセット，方向は省略可

(58)

コンスタントメモリへのアクセス



コンスタントメモリへ高速にアクセスできる要因



コンスタントメモリはオフチップメモリ



コンスタントメモリへのアクセス自体は高速ではない

1. データの配分



_{32スレッド（Warp）単位でアクセスし，1回の読込を32スレッド}

で共有できる

2. キャッシュによる値の再利用



他のWarpがキャッシュされたデータへアクセスすることで，コ

ンスタントメモリから直接読むよりも高速化

(59)

コンスタントメモリを使った空間フィルタ

の高速化



空間フィルタに用いるカーネル



_{1次元の配列として宣言，GPU（グローバルメモリ）へ転送}



（端を除く）全スレッドからアクセス



値は固定値で，GPUから書き換えない



コンスタントメモリを利用



_{1次元の配列として宣言，GPU（コンスタントメモリ）へ転送}



全スレッドがコンスタントメモリにアクセス



コンスタントキャッシュが有効利用される

(60)

メイン関数

（コンスタントメモリの宣言と転送）

:（省略）

:

constant float cfilter[9];

//コンスタントメモリにフィルタのカーネル分のメモリを確保

//mainの外で宣言しているので，どの関数からもアクセス可能

int main(void){

:（省略）

float laplacian[9] ={ 0.0f, 1.0f, 0.0f,

1.0f,‐4.0f, 1.0f,

0.0f, 1.0f, 0.0f};

//グローバルメモリに確保していたフィルタのカーネルは不要

//float *filter;

//cudaMalloc( (void **)&filter, sizeof(float)*9);

//cudaMemcpy(filter, laplacian, sizeof(float)*9, cudaMemcpyHostToDevice);

//コンスタントメモリにフィルタのカーネルを送る

cudaMemcpyToSymbol(cfilter, laplacian, sizeof(float)*9);

:（省略）

boxfilter_constant<<<block,thread>>>(dev_p,dev_filtered);

:（省略）

}

imageproc_mem.cu

(61)

__global__ void boxfilter_constant(unsigned char *p, unsigned char *filtered){ int i,j;

int result = BLACK;

i = blockIdx.x*blockDim.x + threadIdx.x; j = blockIdx.y*blockDim.y + threadIdx.y;

if(0<i && i<WIDTH‐1 && 0<j && j<HEIGHT‐1){ //端の画素は処理をしないため，ifで処理を分岐 result = cfilter[0]*p[(i‐1) + WIDTH*(j‐1)]

+cfilter[1]*p[(i ) + WIDTH*(j‐1)] +cfilter[2]*p[(i+1) + WIDTH*(j‐1)] +cfilter[3]*p[(i‐1) + WIDTH*(j )] +cfilter[4]*p[(i ) + WIDTH*(j )] +cfilter[5]*p[(i+1) + WIDTH*(j )] +cfilter[6]*p[(i‐1) + WIDTH*(j+1)] +cfilter[7]*p[(i ) + WIDTH*(j+1)] +cfilter[8]*p[(i+1) + WIDTH*(j+1)]; } if(result<BLACK) result = ‐result; //数値が負になっていれば‐1をかける if(result>WHITE) result = WHITE; //数値が255を超えていれば255に収める filtered[i+WIDTH*j] = (unsigned char)result; }

空間フィルタ（コンスタントメモリ利用）

imageproc_mem.cu

(62)

実行結果



実行時間 553 ms(CPU) → 4.13 ms →

_3.38 ms



_{CPUからかなり高速化されていたが，さらに2割短縮}

(63)

レポート課題

4（提出期限は2学期末）



ガウシアンフィルタ（ガウスぼかし）を実装せよ



実行の条件



画像の形状は自身で定め，プログラム中で生成せよ



既存の画像を読む機能を実装できる場合は，既存の画像を用いてよい



ガウシアンフィルタのカーネルサイズは5×5とする



小さいサイズの画像を用いて

1. 原画像が正しく生成できている事を確認せよ

2. フィルタが正しくかけられている事を確認せよ



大きいサイズの画像を用いて

1. フィルタのカーネルをグローバルメモリから読んだ場合とコンスタント

メモリから読んだ場合の実行時間の違いを比較せよ

2. 1ブロックあたりのスレッド数を変更し，実行時間がどのように変化す