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

}