CUDA学習ノート6


Exposing Parallelism
この部分は主に並列分析を紹介し,nvprofのいくつかのmetricパラメータを把握することに関連し,具体的なこれらの調節がなぜ性能に影響を及ぼすのかを後続の博文で説明する.
コードの準備
次はkernel関数sumMatrixOnGPUDです.

__global__ void sumMatrixOnGPU2D(float *A, float *B, float *C, int NX, int NY) {
    unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
    unsigned int idx = iy * NX + ix;
    if (ix < NX && iy < NY) {
        C[idx] = A[idx] + B[idx];
    }
}            

16384要素を含む比較的大きなデータ行列を指定します.

int nx = 1<<14;
int ny = 1<<14;

次のコードはmain関数のパラメータ、すなわちblockの次元構成を構成するために使用されます.

if (argc > 2) {
    dimx = atoi(argv[1]);
    dimy = atoi(argv[2]);
}
dim3 block(dimx, dimy);
dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

コンパイル:

$ nvcc -O3 -arch=sm_20 sumMatrix.cu -o sumMatrix

Checking Active Warps with nvprof
各データの比較には基準が必要であり,ここでは4つのblock構成の時間消費を基準として,(32,32)(32,16)(16,32)と(16,16)を用い,本稿の開始時に述べたように,第1パラメータはx象限次元,第2パラメータはy象限次元である.
以下に、いくつかの構成の時間消費出力結果を示します.

$ ./sumMatrix 32 32
sumMatrixOnGPU2D <<< (512,512), (32,32) >>> elapsed 60 ms
$ ./sumMatrix 32 16
sumMatrixOnGPU2D <<< (512,1024), (32,16) >>> elapsed 38 ms
$ ./sumMatrix 16 32
sumMatrixOnGPU2D <<< (1024,512), (16,32) >>> elapsed 51 ms
$ ./sumMatrix 16 16
sumMatrixOnGPU2D <<< (1024,1024),(16,16) >>> elapsed 46 ms

これらの結果を比較すると、最も遅いのは1番目(32,32)、最も速いのは2番目(32,16)であり、ここで推測できるのは、より多くのblock並列性を持つほうが良いことである.この推測はnvprofの
achieved_occupancyというmetricパラメータを検証します.このパラメータの定義式は、前のブログで紹介されていますが、実際には、cycleごとにSMが達成できる最大active warp数がwarp全体に占める割合を指します.次に、このパラメータを使用した結果を示します.

$ nvprof --metrics achieved_occupancy ./sumMatrix 32 32
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.501071
$ nvprof --metrics achieved_occupancy ./sumMatrix 32 16
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.736900
$ nvprof --metrics achieved_occupancy ./sumMatrix 16 32
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.766037
$ nvprof --metrics achieved_occupancy ./sumMatrix 16 16
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.810691

上の出力から2つのことがわかります.
2番目の構成が1番目より多くのblockを持つため、deviceはより多くのactive warpに達する(卵を複数のかごに置く道理とは差が少ない).つまり2番目の性能が1番目より優れている理由である.4番目のachieved Occupancyは最も高いが、最も速いわけではないため、より高いachieved Occupancyは必ずしもより良い性能を意味するわけではない.つまり、GPUの性能に影響を与える多くの要素がある.
checking memory operations with nvprof
C[idx]=A[idx]+B[idx]には、2つのmemory loadと1つのmemory storeの3つのmemory操作があります.これらの操作の効率を表示するにはnvprofの2つのmetricパラメータを使用し、memoryのthroughputを表示するにはgld_を使用します.throughput:

$ nvprof --metrics gld_throughput./sumMatrix 32 32
sumMatrixOnGPU2D <<>> Global Load Throughput 35.908GB/s
$ nvprof --metrics gld_throughput./sumMatrix 32 16
sumMatrixOnGPU2D <<>> Global Load Throughput 56.478GB/s
$ nvprof --metrics gld_throughput./sumMatrix 16 32
sumMatrixOnGPU2D <<>> Global Load Throughput 85.195GB/s
$ nvprof --metrics gld_throughput./sumMatrix 16 16
sumMatrixOnGPU2D <<>> Global Load Throughput 94.708GB/s

4番目に高いload throughputを持っていますでも2番目より遅い(2番目、つまり4番目の半分)なので、高いload throughputでも必ずしも高いパフォーマンスが得られるとは限らない.後にmemory transactionについて述べると、この現象の原因を具体的に分析する.簡単に言えば、高load throughputは偽のイメージであり、必要なデータがmemoryに格納されているフォーマットが整列していないと、多くの追加の不要を招く可能性があるのload操作なので、本稿のefficiencyはこんなに低いです.
そして、nvprofのgld_を使用することができます.efficiencyは、global load throughputと実際に得られたglobal load memoryとの比を正確に必要とするload efficiencyを測定するために使用される.このmetricパラメータは、アプリのload操作がdevice memory bandwidthを利用する程度を知ることができます.

$ nvprof --metrics gld_efficiency ./sumMatrix 32 32
sumMatrixOnGPU2D <<>> Global Memory Load Efficiency 100.00%
$ nvprof --metrics gld_efficiency ./sumMatrix 32 16
sumMatrixOnGPU2D <<>> Global Memory Load Efficiency 100.00%
$ nvprof --metrics gld_efficiency ./sumMatrix 16 32
sumMatrixOnGPU2D <<>> Global Memory Load Efficiency 49.96%
$ nvprof --metrics gld_efficiency ./sumMatrix 16 16
sumMatrixOnGPU2D <<>> Global Memory Load Efficiency 49.80%

上記の結果から,最後の2つのload efficiencyは前の2つの半分にすぎないことが分かった.これはまた、より高いthroughputとより高いOccupancyがより良い性能を生じていない理由を説明することができる.最後の2つのload操作の数は多くなるが(両者はthroughputが高いため)、彼らのload effecitivenessはかなり低い(efficiencyが低いため).
最後の2つを観察すると、彼らのblockのx象限配置はwarpの半分であり、前述の推測から、この象限はwarpサイズの整数倍に維持されるべきであることが分かった.その具体的な原因については、後続のブログで詳しく説明します.
Exposing More Parallelism
私たちが今結論を出すことができるのはblockDimです.xはwarpサイズの整数倍であるべきである.これでload efficiencyを簡単にアップグレードできます.他にも疑問があるかもしれません

blockDimの調整を継続する.xはload throughputを増加し続けますか?

並列性を増大させる他の方法はありますか?

次に、基準データを再構築します.この2つの問題は、この基準から大まかに分析することができます.

$ ./sumMatrix 64 2
sumMatrixOnGPU2D <<>> elapsed 0.033567 sec
$ ./sumMatrix 64 4
sumMatrixOnGPU2D <<>> elapsed 0.034908 sec
$ ./sumMatrix 64 8
sumMatrixOnGPU2D <<>> elapsed 0.036651 sec
$ ./sumMatrix 128 2
sumMatrixOnGPU2D <<>> elapsed 0.032688 sec
$ ./sumMatrix 128 4
sumMatrixOnGPU2D <<>> elapsed 0.034786 sec
$ ./sumMatrix 128 8
sumMatrixOnGPU2D <<>> elapsed 0.046157 sec
$ ./sumMatrix 256 2
sumMatrixOnGPU2D <<>> elapsed 0.032793 sec
$ ./sumMatrix 256 4
sumMatrixOnGPU2D <<>> elapsed 0.038092 sec
$ ./sumMatrix 256 8
sumMatrixOnGPU2D <<>> elapsed 0.000173 sec
Error: sumMatrix.cu:163, code:9, reason: invalid configuration argument

上記のデータから分析できるのは、

の最後の構成(256,8)は実行できません.blockの合計thread数は1024を超えています.これはGPUのハードウェア制限です.

の最良の結果は4番目(128,2)であった.

は最初に最も多くのblockを起動したが、最も速くはなかった.

2 2 2番目と4番目が1つのblockで同じ数のthreadを持っているため、両者が同じ表現をしていると推測すべきだったが、実際には2番目に少し劣っていたのでblockDim.xの大きさが肝心です.

の残りの比較的4番目はブロック数が少ないため、並列規模も性能に影響を与える重要な要素である.

今、私たちはまた推測しなければなりません.blockを持っている人が一番少ないのは、最も低いachieved Occupancyがあるのではないでしょうか.最も多くのblockを持つ人は最高のachieved Occupancyに達するだろうか.これらのアイデアを検証するために、データのセットを見てみましょう.

$ nvprof --metrics achieved_occupancy ./sumMatrix 64 2
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.554556
$ nvprof --metrics achieved_occupancy ./sumMatrix 64 4
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.798622
$ nvprof --metrics achieved_occupancy ./sumMatrix 64 8
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.753532
$ nvprof --metrics achieved_occupancy ./sumMatrix 128 2
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.802598
$ nvprof --metrics achieved_occupancy ./sumMatrix 128 4
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.746367
$ nvprof --metrics achieved_occupancy ./sumMatrix 128 8
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.573449
$ nvprof --metrics achieved_occupancy ./sumMatrix 256 2
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.760901
$ nvprof --metrics achieved_occupancy ./sumMatrix 256 4
sumMatrixOnGPU2D <<>> Achieved Occupancy 0.595197

見たでしょう、(64,2)のachieved Occupancyは意外にも最も低く、彼が最も多くのblockを持っているにもかかわらず(高校で物理の問題をするのもこのような感じ)、それはハードウェアのblockの数に対する制限に達しています.
4番目(128,2)と7番目(256,2)は、差の少ないachieved Occupancyを持っています.この2つについてもう1つの実験を行い、再び増大し、blockDim.yを1に設定します.これにより、blockのサイズも減少します.

$ ./sumMatrix 128 1
sumMatrixOnGPU2D <<>> elapsed 0.032602 sec
$ ./sumMatrix 256 1
sumMatrixOnGPU2D <<>> elapsed 0.030959 sec

今回の構成では、特に(256,1)が(128,1)よりも優れていることを確認し、achieved Occupancy、load throughput、load efficiencyを再確認します.

$ nvprof --metrics achieved_occupancy ./sumMatrix 256 1
$ nvprof --metrics gld_throughput ./sumMatrix 256 1
$ nvprof --metrics gld_efficiency ./sumMatrix 256 1

出力:

Achieved Occupancy 0.808622
Global Load Throughput 69.762GB/s
Global Memory Load Efficiency 100.00%

最適な構成は、最高achieved Occupancyでも最高load throughputでもないことがわかります.計算性能を最適化するために唯一のmetricは存在しません.多くのmetricからバランスを求める必要はありません.
まとめ

ほとんどの場合、性能を正確に最適化できる唯一のmetricは存在しない.

のどのmetricまたはeventが性能に与える影響が大きいかはkernelの具体的なコードによって決定される.

は、多くの関連するmetricおよびeventにおいてバランスを求めている.

Grid/blcok heuristics(ヒント)パフォーマンスの調整に良い切り込みポイントを提供します.