2.1 NVCC IDENTIFICATION MACRO

最近nvccのドキュメント読んでるので一応記録。
もうそろそろCUDA 6がリリースされそうなので、リリースされたら少し変わるかもしれない。

今日は2.1節のみ。

nvccでは__NVCC__マクロと__CUDACC__マクロが、あらかじめ定義されている。

__NVCC__

このマクロは、ソースファイルがnvccでコンパイルされているかどうかを
確認するために使うらしい。

thrustライブラリから例を持って来る。コピーライト等のコメントと一部のコードは
省略しています。
次のコードはthrust/detail/config/global_workarounds.hから。

...
#if defined(THRUST_GCC_VERSION) && (THRUST_GCC_VERSION >= 40800)
#  if defined(__NVCC__) && (CUDA_VERSION >= 6000)
#    pragma GCC diagnostic ignore "-Wunused-local-typedefs"
#  endif // nvcc & cuda 6+
#endif // gcc 4.8

__CUDACC__

このマクロは、ソースファイルがCUDAファイルとして扱われているかどうかを
確認するために使うらしい。
これはヘッダを書くときに便利みたい。

再びthrustライブラリから例を提示。先ほどと同様にコピーライト等は省略している。
次のコードはthrust/detail/config/compiler.hから。

...
#ifdef __CUDACC__

#include <cuda.h>

// Thrust supports CUDA >= 3.0
#if CUDA_VERSION < 3000
#error "CUDA v3.0 or newer is required"
#endif // CUDA_VERSION

#endif // __CUDACC__

最後に

thrustへのリンクを張っておきます。
https://github.com/thrust/thrust

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

説明が下手ですみません。

6.1

テキストに載ってるコードには無駄があるので改良しなさい、という問題。

 まず図6.2のコードを改良したコードを次に示す。

配列のサイズのスレッドでは半分のスレッドが動作をしないので、実行コンフィギュレーションで配列のサイズの半分のスレッドを起動させる。

そうすると元のカーネルでは上手く動かないのでいくつか訂正した。

元のカーネルでは最初に偶数番目のスレッドが加算を行っていたが、半数のスレッドしかない改良版の場合、全てのスレッド番号を２倍して元のカーネルと同様の状況を作る。

またループの条件も変更した。元のカーネルのままだと最後の答えを求める前に終わってしまうので < から <= に変更した。

加えて、元のカーネルは１つのスレッドが自分の担当する配列の要素をshared memoryに格納していたがスレッド数を元の半分にしたことにより１つのスレッドが２つの要素をshared memoryに格納する必要が出て来ることに注意しなければならない。

元のコードに比べて乗算のコードを増やし、for文の条件に使われる演算子も変更した。加えてロード演算も増えたため、演算コストは増えたと思われる（不等号の演算子の変更によってコストが増えるのかはわからない）。

リソース制限については、元々のコードは配列のサイズと同数のスレッドを起動していたのでSMの最大スレッド数のサイズの配列までしか還元出来なかったが、改良したことによりその倍のサイズの配列まで還元出来るようになった。

次に図6.4のコードの改良版を示す。

先ほどのコードと同様に起動させるスレッド数は配列のサイズの半分にしている。

カーネル自体の変更はfor文の初期化の所のみ。
blockIdx.x>>1 を blockIdx.x と変更した。

先ほどの改良版と同様にロード演算が増えていることに注意。

先ほどの改良版とは異なり、演算は減ったが、ロード演算が増えた。

またリソース制限については先ほどの改良版と同様の対処が可能になった。

6.2

先ほどの２つのコードを比べて追加の演算が少ないのはどちらかという問題。

図6.2の改良版は乗算が１つ増え、ロード演算が１つ増えた。

図6.4の改良版は右シフト演算が減り、ロード演算が１つ増えた。

以上により図6.4の改良版の方が追加の演算が少ないと言える。

6.3

後日。

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

CUDAプログラミング実践講座読んだ

一通り読んだ。でも理解した所はほんの一部だと思う。なので今また読んでいる。７章以降の理解が曖昧な気がするのでそこをより重点的に読んで行きたいと思う。

今の時点でまだ４章以降の演習問題を行っていないので読むのに合わせて行って行きたい。

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

次に読む本

もう読んでいるのだけれどCode Readingを読んでいる。一回読んだことがあるけれど理解が浅かったかなと思い、もう一度読むことにした。なるべく練習問題をやっていきたいと思う。ただその練習問題が全体的に難しかった気がする。

また練習問題をしたらブログにアップするつもりだ。

Code Reading―オープンソースから学ぶソフトウェア開発技法

atomicAdd

悩んだこと

今日少し詰まったことがあった。それは次のようなもの。 

訪れたことがない顧客を候補としてcandidates配列に格納し、candidatesのサイズを増やすというもの。候補者を探すのをスレッドで並列に行う。

しかし、こうすると上手くいかない。

本当はあるスレッドがcandidates配列にcustomerを代入したら直後にcandidateSizeを増やしたかったのだが、これだとそうはいかない。あるスレッドがcandidates配列にcustomerを代入した直後に他のスレッドがcandidates配列の同じ場所にcustomerを格納する可能性がある。

どうしたか

自分がやりたいことをどう表現するか１時間ほど悩んだ後で先生にやり方を聞いてみた。それは次のようなもの。

なるほどなぁ。確実に自分には思いつかなかった。悩んでた時はatomicAddの返り値を知らなかったのだが知ってたとしても思いつかなかったと思う。

このコードのパターンはよくあるのかもしれない。覚えておこう。

やはり聞ける人がいるというのは大きいなと感じた。

cudaMemcpyでハマったことの覚え書き

題名通り、cudaMemcpyの際にハマったことの覚え書き。

構造体のメンバに構造体

まず次のコードを見てください（プログラムの一部のみです）。

今日ハマったコードはこれではないですが、同様の問題を含んでいます。

problem_create()はsizeof(problem)バイトをホストメモリに確保した後、メンバを初期化しproblem *を返します。

17行目でcudaMemcpy()を用いてpをdev_pにコピーしています。
これだとランタイムエラーが生じた。

なぜならdev_p->dist.costはp->dist.costをコピーしているからだ。
すなわちdev_p->dist.costはホストのアドレスを指していることになる。

今思うと何故これでいけると思っていたのか。恥ずかしい。
改めて勉強し直さないといけないなと感じた。

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プログラミング実践講座 ? 超並列プロセッサにおけるプログラミング手法

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

CUDA プログラミング実践講座の紹介

この本の原著は"Programmin Massively Prallel Processors A Hands-on Approarch"。2nd Editionが発売されたということもあるのか、1st EditionはPDFで公開されている。

2010年に出版されたということで本の中で出て来るGPUは古い(G80やGT200)。それでもGPUの内部がどのように設計され、どのように動作するのかということはそこまで変わっていないと思うので有用であると思う。

まだ５章までしか読んでおらず、自分の理解も浅いと思うので上の感想は変化する可能性がある。引き続き精進していこうと思う。

この本は４章から演習があるのでその解答を書いていこうと思う。問題文は省いている。

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

 4.1

まずG80(GeForce 8800 GTX)について調べる必要があったので調べてみた。といってもG付録BにCompute Capability 1.0のpropertyが載っているのであまり調べる必要はなかった。

問題には、1024 x 1024の行列のかけ算をするのに1024個のスレッドブロックを用いると書いているので各スレッドブロックは1024個の行列の積を結果として出力する。

さらにスレッドブロックの各スレッドが行列の積の要素を１つ計算すると書いているので各スレッドブロックのサイズは1024となる。

しかしG80の１ブロック当たりのスレッド数の上限は512なので、この制限に引っかかる。

よってこのタイル分割による行列の積は出来ない。

このような感じだろうか。解答の書き方としては、結論を書いて理由を書くのが正しいのだろうが自分の思考の流れとしてはこんな感じだ。

ちなみにこの問題を少し変更した問題を掲載したスライドを見つけたのでリンクを張っておく。嬉しいことに答えも載っている。

4.2

まず初めにこの問題のプログラムを見ていて気になったことは次の２行(詳しくソースが見たいかたは上の原著のPDFをご覧ください)。

blockA[threadIdx.y][threadIdx.x] = A_elements[baseIdx];

A_elements[baseIdx] = blockA[threadIdx.x][threadIdx.y]; 

 これって間に__syncthreads()入れないといけない気がする。

でないとblockA[threadIdx.x][threadIdx.y]に想定した値が入らない可能性があると思うんだけど。

おそらくこれを聞いてる問題なのだろう。

最大ワープサイズが32と仮定するとBLOCK_SIZEが５までしか正しく動く保証が無い。

BLOCK_SIZEが5だとすると１ブロック当たりのスレッド数が25になる。これは最大ワープサイズ以下の値なので、この25スレッドは同時に動作する。なので上の２行も正しく動く。

なぜなら各スレッドは同時に動くので１行目を終えた時点で
blockA[threadIdx.x][threadIdx.y]に想定している値が入っていることが保証されるため。

BLOCK_SIZEが６以上だと１ブロック当たりのスレッド数が最大ワープサイズの32を超えてしまう。そのため、あるワープが１行目を実行した後に２行目のblockA[threadIdx.x][threadIdx.y]に想定した値が入っていることは保証出来ない。

なのでBLOCK_SIZEが１から５の時に正しく動作すると言える。

ところで、この日本語の訳は間違っている気がする。

問題文にはBLOCK_SIZEの可能な範囲の値以外に...と書いてるんだけど、BLOCK_SIZEの可能な範囲の中で...とかじゃないのかな？原著にはOut of ...と書いてるんだけど、自分はさほど英語出来ないのでわからないのだが。

とりあえずここでは後者の可能な範囲の中で...の意味でこの問題を解いている。

※この4.2の文は汚いのでまた整理して書き直します。

4.3

全てのBLOCK_SIZEの値に対して動作するようにコードを書き直す問題。カーネルのみ示す。

__global__ void

BlockTranspose (float* A_elements, int A_width, int A_height)

{

    __shared__ float blockA[BLOCK_SIZE][BLOCK_SIZE];

    int baseIdx = blockIdx.x * BLOCK_SIZE + threadIdx.x;

    baseIdx += (blockIdx.y * BLOCK_SIZE + threadIdx.y) * A_width;

    blockA[threadIdx.y][threadIdx.x] = A_elements[baseIdx];

    __syncthreads();

    A_elements[baseIdx] = blockA[threadIdx.x][threadIdx.y];

}