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

Transcription

1 GPU のメモリ階層 長岡技術科学大学電気電子情報工学専攻出川智啓

2 今回の内容 GPU のメモリ階層 グローバルメモリ 共有メモリ モザイク処理への適用 コンスタントメモリ 空間フィルタへの適用 577

3 GPU の主要部品 基盤 GPU( チップ )+ 冷却部品 画面出力端子 電源入力端子 メモリ 特性の把握が重要 電源入力端子 画面出力端子 メモリ チップ PCI Ex 端子 new/articles /introducing the geforce gtx 780 に公開されている写真を基に作成 578

4 CPU のメモリ階層 オフチップ (off chip) メモリ CPUのチップ外部に置かれたメモリ メインメモリ ( 主記憶 ) 利用可能なメモリの中で最低速, 最大容量 オンチップ (on chip) メモリ CPU のチップ内部に置かれたメモリ レベル 1(L1) キャッシュ レベル 2(L2) キャッシュ レベル 2(L3) キャッシュ 高速, 容量小 低速, 容量大 579

5 GPU のメモリ階層 オフチップメモリ PCI Exカードの基板上に実装 ビデオメモリ 利用可能なメモリの中で最低速, 最大容量 オンチップメモリ GPU のチップ内部に置かれたメモリ レベル1(L1) キャッシュ レベル2(L2) キャッシュ 高速, 容量小低速, 容量大 CPU の構造に類似 580

6 GPU メモリの独自の特徴 CPU とは異なるメモリを複数搭載 各メモリの特徴を知り, 適材適所で利用する事により高速化 GPU から読み書き可能か 処理を行うスレッドから読み書き可能か, 読み込みのみか 複数のスレッドでデータを共有できるか CPU から読み書き可能か C 言語の機能のみで直接読み書きは不可能 の専用関数 (API) を利用して読み書き 581

7 メモリの階層 CPU のメモリ階層 コアごとに L2 キャッシュ, 全体で L3 キャッシュを持つこともある チップ コア コア コア 演算器 演算器 演算器 演算器 演算器 演算器 L1 キャッシュ L1 キャッシュ L1 キャッシュ L2 キャッシュ メインメモリ 582

8 メモリの階層 GPU のメモリ階層 CPU にはない独自のメモリを複数持つ チップ L1 キャッシュ 共有メモリ L2 キャッシュ テクスチャキャッシュ コンスタントキャッシュ グローバルメモリ ローカルメモリ テクスチャメモリ コンスタントメモリ 583

9 メモリの種類 オンチップメモリ (GPU のチップ内部に置かれたメモリ ) 高速アクセス, 小容量 CPUからはアクセス不可 L1キャッシュと共有メモリは一定サイズを共用 L1 キャッシュ / 共有 ( シェアード ) メモリ 容量小小 速度高速高速 GPU からの読み書き CPU からのアクセス 読み書き可ブロック内の全スレッドが同じアドレスにアクセス ( データを共有 ) することが可能 * 読み書き不可 読み書き可各スレッドが異なるアドレスにアクセス 読み書き不可 * スレッドごとに異なるアドレスにアクセスすることも可能 584

10 メモリの種類 オフチップメモリ (GPU のチップ外部に置かれたメモリ ) 低速アクセス, 大容量 CPUから直接アクセス可能 ローカルメモリだけはアクセス不可 グローバルメモリ ローカルメモリ テクスチャメモリ コンスタントメモリ 容量 大 小 大 小 速度 低速 低速 高速 * 高速 * GPU からの読み書き 読み書き可全てのスレッドが同じアドレスにアクセス可能 ** 読み書き可各スレッドが異なるアドレスにアクセス 読み込み可全てのスレッドが同じアドレスにアクセス可能 ** 読み込み可全てのスレッドが同じアドレスにアクセス可能 ** CPU からのアクセス 読み書き可読み書き不可読み書き可読み書き可 * キャッシュが効く場合 ** スレッドごとに異なるアドレスにアクセスすることも可能 585

11 メモリの種類 共有メモリと L1 キャッシュは一定サイズを共用 グローバルメモリへのアクセスは L2 キャッシュ経由 Fermi 世代以前の GPU はキャッシュ無し * GPU Chip SM L1 キャッシュ 共有メモリ SM L1 キャッシュ 共有メモリ オフチップメモリ オンチップメモリ ホストメモリ ローカルメモリ L2キャッシュコンスタントメモリテクスチャメモリ *Tesla 世代でもテクスチャキャッシュは存在. Read/Write 可能なキャッシュは Fermi 世代から. ローカルメモリ グローバルメモリ 586

12 メモリの種類と並列化階層の対応 オンチップメモリ ブロックまたはスレッドごとに異なる値を持てる が不足するとローカルメモリが使われる オフチップメモリ GPU 全体で共通の値を持てる Grid Block(0,0,0) L1 キャッシュ Thre ad 0 Thre ad 1 共有メモリ Thre ad 2 Thre ad 3 Block(1,0,0) L1 キャッシュ Thre ad 0 Thre ad 1 共有メモリ Thre ad 2 Thre ad 3 各 GPU(Grid) 内でデータを共有各ブロック内でデータを共有各スレッドが個別のデータを保有 ホストメモリ ローカルメモリ ローカルメモリ L2キャッシュコンスタントメモリテクスチャメモリグローバルメモリ 587

13 各スレッドが個別に利用 カーネル内で変数を宣言するとを利用 非常に高速 キャッシュとしても利用可能 少容量 本 * 32bit 利用可能分を超えるとローカルメモリへ追い出される スピル *Kepler からは 本 ホストメモリ GPU Chip SM L1 キャッシュ ローカルメモリ ローカルメモリ 共有メモリ SM L2 キャッシュ L1 キャッシュ コンスタントメモリ テクスチャメモリ グローバルメモリ 共有メモリ 588

14 グローバルメモリ ビデオメモリ ( 数 GB) CPU のメインメモリに相当 読み込みがある一定サイズでまとめて行われる レイテンシが大きい 読み出し命令を発効してからデータが得られるまでの時間 効率よくアクセスするための条件がある コアレスアクセス ( コアレッシング ) 589 ホストメモリ GPU Chip SM L1 キャッシュ ローカルメモリ ローカルメモリ 共有メモリ SM L2 キャッシュ L1 キャッシュ コンスタントメモリ テクスチャメモリ グローバルメモリ 共有メモリ

15 グローバルメモリへのアクセス * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. 16 スレッドが協調して同時にアクセス では 32 スレッドを Warp という単位で管理 Warp の半分 (Half Warp) が協調して読み書き コアレスアクセスか否かで読込速度が大幅に変化 新しい世代のGPUでは速度の落ち込みが緩和 コアレスアクセスはGPUプログラムで最重要 GPUの処理能力と比較するとデータ供給不足が発生 効率よくデータを供給するためにコアレスアクセスは必須 590

16 コアレスアクセスになる条件 * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. データのサイズ 32bit, 64bit, 128bit(4 バイト, 8 バイト, 16 バイト ) アドレスの隣接 16 スレッドがアクセスするアドレスがスレッド番号順に隣接 アクセスする最初のアドレス 16 スレッドがアクセスするアドレスの先頭 ( スレッド 0 がアクセスするアドレス ) が,64 バイトまたは 128 バイト境界 アドレスが 64 の倍数で始まる 64 バイトの領域か, アドレスが 128 の倍数で始まる 128 バイトの領域 591

17 コアレスアクセスの例 * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. データ型が 32bit=4 バイト T0 T1 A128 A132 データ型が 32bit=4 バイト T0 T1 A128 A132 各スレッドが連続して隣接アドレスにアクセス 先頭アドレスが 128 バイト境界 16 スレッド T2 T3 T4 T5 T6 T7 T8 T9 T10 T11 T12 A136 A140 A144 A148 A152 A156 A160 A164 A168 A172 A176 各スレッドが連続して隣接アドレスにアクセス 実際にデータを取得するかは無関係 先頭アドレスが 128 バイト境界 T2 T3 T4 T5 T6 T7 T8 T9 T10 T11 T12 A136 A140 A144 A148 A152 A156 A160 A164 A168 A172 A176 T13 A180 T13 A180 T14 A184 T14 A184 T15 A188 T15 A

18 コアレスアクセスにならない例 * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. T0 A128 T0 A128 T0 A128 T0 A128 T1 A132 T1 A132 T1 A132 T1 T2 A136 T2 A136 T2 A136 T2 T3 A140 T3 A140 T3 A140 T3 A140 T4 A144 T4 A144 T4 A144 T4 T5 A148 T5 A148 T5 A148 T5 T6 A152 T6 A152 T6 A152 T6 A152 T7 T10 A156 各スレッドが番 T8 A160 号順にアクセス T9 A164 していない A168 T7 T10 A156 先頭が128バイト T8 A160 境界ではない T9 A164 ( 現在は緩和 ) A168 T7 A156 アドレスが連続し T8 A160 ていない T9 A164 T10 A168 T7 データが 32bit, 64bit, T8 128bit T9 ではないA164 ( 構造体など ) T10 T11 A172 T11 A172 T11 A172 T11 T12 A176 T12 A176 T12 A176 T12 A176 T13 A180 T13 A180 T13 A180 T13 T14 A184 T14 A184 T14 A184 T14 T15 A188 T15 A188 T15 A188 T15 A

19 コアレスアクセスにならない例 * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. 128 バイト境界からわずかにずれている場合 T0 T1 T2 A128 A132 A136 Tesla 世代以降は 64 バイトブロックと 32 バイトブロックに分けて読込 メモリアクセス要求は 2 回 コアレスアクセスの半分程度の性能は得られる T3 T4 T5 T6 T7 T8 T9 T10 T11 A140 A144 A148 A152 A156 A160 A164 A168 A バイトブロックでデータ読込 T12 A176 T13 A180 T14 A184 T15 A188 A バイトブロックでデータ読込 594

20 モザイク処理 前回授業で取り上げたモザイク処理 画像を小さな領域に分け, その領域を全て同じ色にする 領域内の全画素を, 領域内の画素の平均値に置き換える 原画像処理画像 595

21 モザイク処理 前回授業で取り上げたモザイク処理 高速化が今ひとつ 他の処理より一桁遅い 処理 処理時間 [ms] CPU GPU ネガティブ処理 水平反転 垂直反転 空間フィルタ モザイク処理

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 597

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[] 598

24 モザイク処理が高速化されない原因 各ブロックの 1 スレッドのみが処理を実行 グローバルメモリから画素情報を読込 ブロック内の画素の平均値を計算 グローバルメモリに画素の平均値を書込 グローバルメモリへコアレスアクセスできていない 複数のスレッドが協調し, アドレスが隣接したメモリを読むと高速 ある 1 スレッドが 1 画素ずつメモリアドレスを変えながら読むのは最悪の処理 599

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; 600

26 メモリの種類 オンチップメモリ (GPU のチップ内部に置かれたメモリ ) 高速アクセス, 小容量 CPUからはアクセス不可 L1キャッシュと共有メモリは一定サイズを共用 L1 キャッシュ / 共有 ( シェアード ) メモリ 容量小小 速度高速高速 GPU からの読み書き CPU からのアクセス 読み書き可ブロック内の全スレッドが同じアドレスにアクセス ( データを共有する ) ことが可能 * 読み書き不可 読み書き可各スレッドが異なるにアクセス 読み書き不可 * スレッドごとに異なるアドレスにアクセスすることも可能 601

27 共有 ( シェアード ) メモリ ブロック内のスレッドが共通のデータにアクセス可能 1 回目のアクセスに時間がかかるが, それ以降は非常に高速にアクセス可能 Fermi 世代以前の GPU で管理可能なキャッシュとして利用 1 ブロックあたり 16kB~48kB ホストメモリ GPU Chip SM L1 キャッシュ ローカルメモリ ローカルメモリ 共有メモリ SM L2 キャッシュ L1 キャッシュ コンスタントメモリ テクスチャメモリ グローバルメモリ 共有メモリ 602

28 共有 ( シェアード ) メモリの宣言 カーネル内で修飾子 shared を付けて宣言 配列として宣言 配列サイズを静的 ( コンパイル時 ) に決定する場合 shared 型変数名 [ 要素数 ] 多次元配列も宣言可能 配列サイズを動的 ( カーネル実行時 ) に決定する場合 extern shared 型変数名 [] サイズはカーネル呼出時のパラメータで指定 <<< ブロック数, スレッド数, 共有メモリのバイト数 >>> 603

29 共有メモリを使ったモザイク処理の高速化 ブロック内の全スレッドで共有したい値 各スレッドがグローバルメモリから読んだ画素情報 ブロック内の画素の平均値 処理の流れ 1. ブロック内の全スレッドがグローバルメモリから画素の値を読み出し, 共有メモリに置く ( コアレスアクセス ) 2. ある1スレッドが共有メモリに置かれた画素の値を読み, 画素の平均値を計算し, 共有メモリに置く 3. ブロック内の全スレッドが共有メモリに置かれた画素の平均値を読み, グローバルメモリに書き込む ( コアレスアクセス ) 604

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; 605

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[][] 606

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 607

33 共有メモリの宣言 filtered[i + WIDTH*j] = (unsigned char)average; 各スレッドが average を読み込み グローバルメモリの各位置に画素の平均値を書き込む filtered[] average 608

34 実行結果 正しく処理できていない ブロック内のスレッドの協調が不十分 原画像 処理画像 609

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; スレット 0 スレット 0 以外 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; 他のスレッドが共有メモリに書き込む前に cache[][] にアクセスする可能性がある 610

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; スレット 0 スレット 0 以外 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 以外は平均値を計算しないので, 直ちにこの行に到達し, averageの値が確定する前に値を読んでfilteredに書き込む可能性がある 611

37 ブロック内でのスレッドの同期 syncthreads(); カーネル実行中にスレッドの同期を取る syncthreads() が書かれた行にスレッドが到達すると, 同一ブロック内の他の全てのスレッドがその行に達するまで待機 異なるブロック間での同期は不可能 if の中に記述するとカーネルが終了しないことがある if( 条件 ){ syncthreads(); // 条件が真にならないスレッドは if の中に入らないため, // カーネルが永久に終わらない 612

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; スレット 0 スレット 0 以外 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 文を実行しないが, スレッド 0 が到達する ( 平均値を計算し終わる ) まで syncthreads() で停止 if 文の前でブロック内の全スレッドが同期しているので, cache[][] には画素情報が入っている ブロック内で同期 ブロック内で同期 imageproc_mem.cu 613

39 実行結果 正しく処理できている 実行時間 260 ms(cpu) 38.5 ms 18.7 ms 処理時間が約 1/2 に短縮原画像 処理画像 614

40 空間フィルタ ある画素とその周囲の画素を使って処理 処理の仕方を規定したカーネルを定義 カーネルは 1 次元配列で表現 原画像 フィルタ ( カーネル ) 輪郭抽出 a b c d e f g h i = b+d 4e+f+h 615

41 空間フィルタ カーネルは 1 次元配列で表現 ぼかし ( 平均フィルタ ) 1/9 1/9 1/9 1/9 1/9 1/9 1/9 1/9 1/9 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; 616

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 617

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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 618 p[]

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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 619 j p[] i

45 空間フィルタ ( 前回授業で作成 ) 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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 620 j p[] i

46 空間フィルタ ( 前回授業で作成 ) 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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 621 j p[] i

47 空間フィルタ ( 前回授業で作成 ) 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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 622 j p[] i

48 空間フィルタ ( 前回授業で作成 ) 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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 623 j p[] i

49 空間フィルタ ( 前回授業で作成 ) 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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 624 j p[] i

50 空間フィルタ ( 前回授業で作成 ) 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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 625 j p[] i

51 空間フィルタ ( 前回授業で作成 ) 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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 626 j p[] i

52 空間フィルタ ( 前回授業で作成 ) 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)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 627 j p[] i

53 空間フィルタ処理の高速化 フィルタ ( 配列 filter) へのアクセス コアレスアクセスできていない L2 キャッシュに入る可能性は高いが, 有効活用されているかは不明 配列 filter へのアクセスの高速化 共有メモリを利用すると処理が冗長 1 ブロックから 9 スレッドを選び, グローバルメモリから共有メモリへコピーし, 同期をとる コンスタントメモリが活用できる 628

54 メモリの種類 オフチップメモリ (GPU のチップ外部に置かれたメモリ ) 低速アクセス, 大容量 CPUから直接アクセス可能 ローカルメモリだけはアクセス不可 グローバルメモリ ローカルメモリ テクスチャメモリ コンスタントメモリ 容量 大 小 大 小 速度 低速 低速 高速 * 高速 * GPU からの読み書き 読み書き可全てのスレッドが同じアドレスにアクセス可能 ** 読み書き可各スレッドが異なるアドレスにアクセス 読み込み可全てのスレッドが同じアドレスにアクセス可能 ** 読み込み可全てのスレッドが同じアドレスにアクセス可能 ** CPU からのアクセス 読み書き可読み書き不可読み書き可読み書き可 * キャッシュが効く場合 ** スレッドごとに異なるアドレスにアクセスすることも可能 629

55 コンスタントメモリ GPU 全体で同じメモリにアクセス GPU Chip SM SM メモリを読み取り専用とすることで値をキャッシュし, 一度読んだ値を再利用 GPU 全体で 64kB L1 キャッシュ ローカルメモリ 共有メモリ L2 キャッシュ L1 キャッシュ 共有メモリ ホストメモリ ローカルメモリ コンスタントメモリ テクスチャメモリ グローバルメモリ 630

56 コンスタントメモリの宣言 グローバル領域で修飾子 constant を付けて宣言 配列サイズは静的に決定 constant 型変数名 ; constant 型変数名 [ 要素数 ]; 配列としても宣言可能 サイズはコンパイル時に確定している必要がある cudamalloc() や cudafree() は不要 グローバル変数として宣言し, 複数のカーネルからアクセスすることが多い 読込専用なので許される 書込可能なメモリでは厳禁 631

57 コンスタントメモリの利用 メモリは読込専用 CPUからは変更可能 専用のメモリ転送命令でコピー cudamemcpytosymbol CPU 上のメモリをコンスタントメモリにコピーする cudamemcpytosymbol( 転送先変数名, 転送元アドレス, バイト数, オフセット, 方向 ); オフセット, 方向は省略可 632

58 コンスタントメモリへのアクセス コンスタントメモリへ高速にアクセスできる要因 コンスタントメモリはオフチップメモリ コンスタントメモリへのアクセス自体は高速ではない 1. データの配分 32 スレッド (Warp) 単位でアクセスし,1 回の読込を 32 スレッドで共有できる 2. キャッシュによる値の再利用 他の Warp がキャッシュされたデータへアクセスすることで, コンスタントメモリから直接読むよりも高速化 633

59 コンスタントメモリを使った空間フィルタの高速化 空間フィルタに用いるカーネル 1 次元の配列として宣言,GPU( グローバルメモリ ) へ転送 ( 端を除く ) 全スレッドからアクセス 値は固定値で,GPUから書き換えない コンスタントメモリを利用 1 次元の配列として宣言,GPU( コンスタントメモリ ) へ転送 全スレッドがコンスタントメモリにアクセス コンスタントキャッシュが有効利用される 634

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 635

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 636

62 実行結果 実行時間 553 ms(cpu) 4.13 ms 3.38 ms CPU からかなり高速化されていたが, さらに 2 割短縮 原画像 処理画像 637

63 レポート課題 4( 提出期限は 2 学期末 ) ガウシアンフィルタ ( ガウスぼかし ) を実装せよ 実行の条件 画像の形状は自身で定め, プログラム中で生成せよ 既存の画像を読む機能を実装できる場合は, 既存の画像を用いてよい ガウシアンフィルタのカーネルサイズは 5 5 とする 小さいサイズの画像を用いて 1. 原画像が正しく生成できている事を確認せよ 2. フィルタが正しくかけられている事を確認せよ 大きいサイズの画像を用いて 1. フィルタのカーネルをグローバルメモリから読んだ場合とコンスタントメモリから読んだ場合の実行時間の違いを比較せよ 2.1ブロックあたりのスレッド数を変更し, 実行時間がどのように変化するかを考察せよ 638

64 レポートの書式 必ず表紙を付けること 授業名, 課題番号, 学籍番号, 氏名, 提出日に加えて課題に要した時間を書く 課題内容, プログラム, 実行結果, 考察で構成 プログラムを実行した tesla?? および GPU の番号も明記すること 利用する GPU を cudasetdevice 命令で選択すること pdf 形式に変換してメールで提出 宛先 degawa at vos.nagaokaut.ac.jp メール題目 課題 4( 氏名 ) 639