CUDA プログラミング実践講座 5.6演習

４章に引き続き５章の演習。

5.1

行列の和で共有メモリを使って帯域幅を減らせるかどうかという問題。

人にどう説明していいのかわからないけど、結論としては帯域幅は減らせない。

行列の和の計算なので、２つの入力行列の各要素を足してその結果を１つの出力行列に書き込む。

ということは、グローバルメモリから２つの値（各入力行列から１つずつ）を読み込む。それらを足し合わせて、グローバルメモリに書き込む、という動作を行う。

この動作の中でスレッド間で共有すべきことはないのでグローバルメモリの帯域幅は減らせない。

ということでいいのかな。というよりこの説明では人を納得させれるのかわからん。もっといい説明が思いつけば、また編集しよう。

5.2

表を書く問題。下付き文字の使い方がわからないため画像で提供しようとしたが非常にひどい方法だと思ったため、省略。

5.3

__syncthreads()を使うのを忘れた場合どういう誤動作が起こるか、という問題。

11行目の__syncthreads()を使うのを忘れた場合、全てのスレッドが自分の担当する行列の成分をshared memoryに格納する前に内積を計算してしまう可能性がある。そのため内積の計算を正確に行えない。

また14行目の__syncthreads()を使うのを忘れた場合、全てのスレッドが内積を計算する前にshared memoryに値を代入する。そのためまだ内積の計算を終わっていないスレッドがその新たに代入された値を使用してしまう可能性があり、これもまた内積の計算を正確に行えない。

同時に忘れた場合は、上記のような理由からshared memoryに値を代入する前に内積の計算を行ってしまったり、内積の計算をする前にshared memoryに次の値を代入してしまったりするため、結局内積の計算を正確に行えない。

このような書き方で伝わるのだろうか。

5.4

グローバルメモリから取り出された値を保持するのにレジスタよりshared memoryの方が有効である場合を挙げなさい、という問題。

同一ブロック内のスレッドでデータを共有する場合に有効であるのではないか。

なぜならshared memoryは同一ブロックの全てのスレッドから可視であり、レジスタはスレッドローカルな記憶領域であるため同一ブロッく内のスレッドでデータを使うことは出来ないからである。

他にもレジスタより共有メモリを使うのが有効な場合はあるのだろうか。

CUDAプログラミング実践講座 ? 超並列プロセッサにおけるプログラミング手法