A certain engineer "COMPLEX"

.NETでGPUPUを試してみる CUDA編 第3回

前回はCUDAの性能を改善しましたが、今一歩及ばずというところです。

Introduction


ところで、これまで.NETが一度も出てきていないのにこのタイトルは詐欺なんじゃないですかねぇ(震え声)

Explanation


Tread!thread!!Thread!!!

前回はスレッドを使って大幅に性能を改善しました。
が、それでもCPUが純粋にシーケンシャルで処理する方が速かったです。

CUDAに対する理解が全く足りていないのでもう少し調べてみます。
前回登場した、CUDA_C_Programming_Guide.pdf で、スレッドに関する記述を見てみます。

2.2. Thread Hierarchy
For convenience, threadIdx is a 3-component vector, so that threads can be identified
using a one-dimensional, two-dimensional, or three-dimensional thread index, forming
a one-dimensional, two-dimensional, or three-dimensional block of threads, called a
thread block. This provides a natural way to invoke computation across the elements in a
domain such as a vector, matrix, or volume.


訳:
2.2. スレッドの階層
便利なことに、threadIdxは3成分のベクトルであり、つまりスレッドは1次元、2次元、3次元のスレッドインデックス、で識別することができ、スレッドブロックと呼ばれる1次元、2次元、3次元のスレッドのブロックを形作る。これは、ベクトル、行列、ボリューム (訳注:ここでいうベクトルは、力学のベクトルではなく、配列のような1次元データで、行列は2次元、ボリュームは3次元という意味) のようなドメイン内の要素を横断する計算を実行するための自然な方法を提供する。

といっています。
ブロックという単語が出てきました、これは前回でも出てきましたが、流しました。
要するに、スレッドはブロックという単位になっていることです。

The index of a thread and its thread ID relate to each other in a straightforward way:
For a one-dimensional block, they are the same; for a two-dimensional block of size (D_x, D_y)
,the thread ID of a thread of index (x, y) is (x + y, D_x); for a three-dimensional block of
size (D_x, D_y, D_x), the thread ID of a thread of index (x, y, x) is (x + y D_x + z D_x D_y).


訳:
スレッドのインデックスとそのスレッドIDは素直なやり方で互いに関連している。1次元のブロックではそれらは同一であり、2次元のブロックのサイズ (Dx, Dy)、スレッドのインデックス (x, y) のスレッドIDは (x + y, Dx); 3次元のブロックのサイズ (Dx, Dy, Dx)、スレッドのインデックス (x, y, z) のスレッドIDは (x + y Dx + z Dx Dy)である。

つまるところ、各ブロックにスレッドが存在するっていってます。

There is a limit to the number of threads per block, since all threads of a block are
expected to reside on the same processor core and must share the limited memory
resources of that core. On current GPUs, a thread block may contain up to 1024 threads.
However, a kernel can be executed by multiple equally-shaped thread blocks, so that the
total number of threads is equal to the number of threads per block times the number of
blocks.
Blocks are organized into a one-dimensional, two-dimensional, or three-dimensional
grid of thread blocks as illustrated by Figure 6. The number of thread blocks in a grid is
usually dictated by the size of the data being processed or the number of processors in
the system, which it can greatly exceed.

Figure6


訳:
1ブロック毎のスレッド数には制限があり、一つのブロックの全スレッドは、プロセッサーコアから予測でき、コアの制限されたメモリリソースを共有しなくてはならない。現在のGPUにて、スレッドブロックは1024スレッドまで内包するかもしれない。しかし、カーネルは複数の同一形状のスレッドブロックで実行され、ゆえにスレッドの合計数はブロックの個数と1ブロックあたりのスレッドの個数の積になる。
ブロックは、Figure 6で図示された、スレッドブロックの1次元、2次元、3次元のグリッドを形作る。1つのグリッド内におけるスレッドブロックの数は通常、処理されるデータのサイズまたはシステムのプロセッサー数から規定される。(訳注:最後のwhichはどこにかかっているのかわかりませんでしたので、最後のwhich以降は無視します)

スレッドの合計数はブロックの個数と1ブロックあたりのスレッドの個数の積

なんか凄いこと書いてあります。
ってか、前回書いた記事で1ブロックあたり最大512スレッド、って書きましたけど、あれは間違い?
当時のCUDA搭載GPUの最大値なのか?

とりあえず、自分のGPU性能を詳しく知りたい。
第3回 CUDA4.0のインストールによれば、deviceQuery.exe というプログラムが性能を詳しく教えてくれるとのこと。
これはCUDA Toolkitインストール時にコピーされており、デフォルトは。C:\ProgramData\NVIDIA Corporation\CUDA Samples\v7.5\1_Utilities\deviceQuery に入っていると。
とりあえず、自分でビルドして実行

deviceQuery

結果として、1プロセッサーあたり1536スレッド、1ブロックあたり最大1024スレッド起動できることがわかりました。最大1536スレッドでしょぼくない?と思ったら、上に (4) Multiprocessors ってあるので、全部で1536*4=6144スレッドってこと?

とりあえず可能な限りスレッドを活用するべくコードを修正。
また、スレッド数を活用できるようなカーネルに内容を変更します。
サンプルコードはページの末尾を参照。

変更点は

  • ベクトルの和の計算を行列乗算の計算に変更。1000x1000の行列なので1000^3 の計算
  • スレッド数指定

const int SIZE = 32;
const int MATRIX_SIZE = 1000;

cudaError_t addWithCuda(float *c, const float *a, const float *b);

__global__ void addKernel(float *c, const float *a, const float *b)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

if (row < MATRIX_SIZE && col < MATRIX_SIZE) { auto v = 0.f; for (auto x = 0; x < MATRIX_SIZE; x++) { v += a[row * MATRIX_SIZE + x] * b[x * MATRIX_SIZE + col]; }
c[row * MATRIX_SIZE + col] = v;
}
}

こんな感じ。
ホスト側は省略。

結果は


threadsPerBlock.X =32, threadsPerBlock.y = 32
numBlocks.X =32, numBlocks.y = 32
CUDA is
time = {117}
No CUDA is
time = {7076}

やりました。60倍のパフォーマンスをたたき出しています。

Conclusion


そもそもの原因として、CUDAの得意領域に持ち込んでいなかったのが問題でした。
如何にスレッドを有効活用するか、というところに主眼をおいて実装を進める必要があります。

Source Code


https://github.com/takuya-takeuchi/Demo/tree/master/CUDA3

3 thoughts on “.NETでGPUPUを試してみる CUDA編 第3回

  1. 匿名

    古い記事にコメントってのはアレですし、そもそも既知かと思われますが
    カーネルの内部でループ処理を行うことがそもそものGPGPU処理として不適です。
    この場合の行列乗算であれば、カーネル内部のループ処理も分割して処理を行い、
    スレッド間で求めた要素値をReductionするほうが遥かに高速に処理が行えます。
    (cuBLAS使えばいいじゃんって話はここでは無しの方向で)

    以下は具体例となります。

    GPU内におけるスレッド処理は基本的に32個セットで同時に動作するため、行列の要素数は1024*1024に設定すると最適な動作が望めます。

    #define MATRIX_SIZE 1024

    __global__ void addKernel(float* c, const float* __restrict__ a, const float* __restrict__ b)
    {
    __shared__ float v_sh[32];

    int num = threadIdx.x;
    int col = blockIdx.y;
    int row = blockIdx.z;

    if (num < MATRIX_SIZE && col < MATRIX_SIZE && row < MATRIX_SIZE)
    {
    float v = a[row * MATRIX_SIZE + num] * b[num * MATRIX_SIZE + col];

    v += __shfl_xor(v, 16);
    v += __shfl_xor(v, 8);
    v += __shfl_xor(v, 4);
    v += __shfl_xor(v, 2);
    v += __shfl_xor(v, 1);

    v[num / 32] = v;
    __syncthreads();

    v += __shfl_xor(v, 16);
    v += __shfl_xor(v, 8);
    v += __shfl_xor(v, 4);
    v += __shfl_xor(v, 2);
    v += __shfl_xor(v, 1);

    if (num == 0)
    c[row * MATRIX_SIZE + col] = v;
    }
    }

    スレッドの数の立て方は以下で変更します
    dim3 threadsPerBlock(1024, 1, 1);
    dim3 numBlocks(1, 1024, 1024);

    これで動かせば100倍以上の高速化が見込めるかと思われます。

  2. 匿名

    処理が一部間違っておりましたので修正です。

    __global__ void addKernel(float* c, const float* __restrict__ a, const float* __restrict__ b)
    {
    __shared__ float v_sh[32];

    int num = threadIdx.x;
    int col = blockIdx.y;
    int row = blockIdx.z;

    if (num < MATRIX_SIZE && col < MATRIX_SIZE && row < MATRIX_SIZE)
    {
    float v = a[row * MATRIX_SIZE + num] * b[num * MATRIX_SIZE + col];

    v += __shfl_xor(v, 16);
    v += __shfl_xor(v, 8);
    v += __shfl_xor(v, 4);
    v += __shfl_xor(v, 2);
    v += __shfl_xor(v, 1);

    if (num % 32 == 0)
    v_sh[num / 32] = v;
    __syncthreads();

    if (num < 32)
    {
    v = v_sh[num];

    v += __shfl_xor(v, 16);
    v += __shfl_xor(v, 8);
    v += __shfl_xor(v, 4);
    v += __shfl_xor(v, 2);
    v += __shfl_xor(v, 1);

    if (num == 0)
    c[row * MATRIX_SIZE + col] = v;
    }
    }

    1. Takuya Takeuchi Post author

      古い記事の誤りを訂正していただきありがとうございます。
      ネットの広い海で間違った情報が漂うのは辛いので、あたなの行動に深く感謝いたします。

コメントを残す

メールアドレスが公開されることはありません。

%d人のブロガーが「いいね」をつけました。