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

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

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

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

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

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

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

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

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

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

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

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

メモリの種類と並列化階層の対応オンチップメモリブロックまたはスレッドごとに異なる値を持てるが不足するとローカルメモリが使われるオフチップメモリ 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

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

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

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

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

コアレスアクセスの例 * * かなり古い情報のため要注意. 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 A188 592

コアレスアクセスにならない例 * * かなり古い情報のため要注意. 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 A188 593

コアレスアクセスにならない例 * * かなり古い情報のため要注意. 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 A172 64 バイトブロックでデータ読込 T12 A176 T13 A180 T14 A184 T15 A188 A190 32 バイトブロックでデータ読込 594

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

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

モザイク処理 ( 前回授業で作成 ) 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

モザイク処理 ( 前回授業で作成 ) 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

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

モザイク処理が高速化されない原因せめてコアレスアクセスしたい平均を並列に計算するのは中級レベルの処理カーネル内で変数を宣言するとが使われる下の例では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

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

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

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

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

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

共有メモリの宣言 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

共有メモリの宣言 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

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

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

モザイク処理 ( 共有メモリを利用 ) 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

モザイク処理 ( 共有メモリを利用 ) 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

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

モザイク処理 ( 共有メモリを利用 ) 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

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

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

空間フィルタカーネルは 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; 輪郭抽出 0 1 0 1 4 1 0 1 0 float laplacian[9] ={ 0.0f, 1.0f, 0.0f, 1.0f, 4.0f, 1.0f, 0.0f, 1.0f, 0.0f; 616

空間フィルタ ( 前回授業で作成 ) 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

空間フィルタ ( 前回授業で作成 ) 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[]

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

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

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

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

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

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

メイン関数 ( コンスタントメモリの宣言と転送 ) :( 省略 ) : 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

空間フィルタ ( コンスタントメモリ利用 ) 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

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

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

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