CUDA ---- Kernel性能調節

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

在做各項數據比較的時候需要有個基准，這里使用四個block配置的時間消耗作為基准觀察，分別為（32，32）（32,16）（16,32）和（16,16），本文開始時有提到，第一個參數是x象限維度，第二個參數是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

比較這幾個結果，不難發現，最慢的是第一個（32,32），最快的是第二個（32,16），這里可以猜測的到的是，擁有更多的block並行性更好。這個猜測可以使用nvprof 的achieved_occupancy這個metric參數來驗證。該參數的定義公式在上一篇博文有介紹，實際上就是指每個SM在每個cycle能夠達到的最大active warp數目占總warp的比例。下面是使用該參數后得到的結果：

$ nvprof --metrics achieved_occupancy ./sumMatrix 32 32
sumMatrixOnGPU2D <<<(512,512), (32,32)>>> Achieved Occupancy 0.501071
$ nvprof --metrics achieved_occupancy ./sumMatrix 32 16
sumMatrixOnGPU2D <<<(512,1024), (32,16)>>> Achieved Occupancy 0.736900
$ nvprof --metrics achieved_occupancy ./sumMatrix 16 32
sumMatrixOnGPU2D <<<(1024,512), (16,32)>>> Achieved Occupancy 0.766037
$ nvprof --metrics achieved_occupancy ./sumMatrix 16 16
sumMatrixOnGPU2D <<<(1024,1024),(16,16)>>> Achieved Occupancy 0.810691

從上面的輸出可以得知兩件事兒：

1. 由於第二個配置比第一個有更多的block，device就會達到更多active warp（跟雞蛋放在多個籃子的道理差不多）。也就是第二個性能優於第一個的原因。
2. 第四個的achieved Occupancy最高，但是卻不是最快的，因此，較高的achieved Occupancy並不一定就意味着更好的性能，也就是說還有更多的因素影響着GPU的性能。

checking memory operations with nvprof

對於C[idx] = A[idx] + B[idx]來說共有三個memory操作：兩個memory load和一個memory store。要查看這些操作的效率可以使用nvprof的兩個metric參數，如果想要查看memory的throughput，則可使用gld_throughput：

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

不難看到，第四個擁有最高的load throughput，但是卻比第二個慢（第二個也就是第四個的一半），所以，較高的load throughput也不一定就有較高的性能。之后講到memory transaction時會具體分析這種現象的原因，簡單說，就是高load throughput有可能是一種假象，如果需要的數據在memory中存儲格式未對齊不連續，會導致許多額外的不必要的load操作，所以本文中的efficiency會這么低。

然后，我們可以使用nvprof的gld_efficiency來度量load efficiency，該metric參數是指我們確切需要的global load throughput與實際得到global load memory的比值。這個metric參數可以讓我們知道，APP的load操作利用device memory bandwidth的程度：

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

從上述結果可知，最后兩個的load efficiency只是前兩個的一半。這也可以解釋，為什么較高的throughput和較高的Occupancy卻沒有產生較好的性能。盡管最后兩個的load操作數目要多不少（因為二者throughput較高），但是他們的load effecitiveness卻低不少（由於efficiency較低）。

觀察最后兩個可以發現，他們block的x象限配置是warp的一半，由前文推測知，該象限應該保持為warp大小的整數倍。關於其具體原因將在后續博文詳細解釋。

Exposing More Parallelism

我們現在可以得出一個結論就是blockDim.x應該是warp大小的整數倍。這樣做是很容易就提升了load efficiency。現在，我們可能還有其他疑惑，比如：

• 繼續調整blockDim.x是否會繼續增加load throughput？
• 還有其他方法能增大並行性嗎？

現在，我們重新整一個基准數據出來，這兩個問題可以從這個基准分析個大概：

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

從上面數據，我們能夠分析出來的是：

• 最后一個配置（256,8）不可行，block中總共的thread數目超過了1024，這是GPU的硬件限制。
• 最好的結果是第四個（128，2）。
• 第一個啟動了最多的block，但不是最快的。
• 因為第二個與第四個在一個block中擁有相同數目的thread，本應猜測二者有相同的表現，但是實際卻是第二個略遜色，所以blockDim.x的大小是很關鍵的。
• 剩下的相對第四個都有較少的block數目，所以並行規模也是影響性能的關鍵因素。

現在，我們又得猜測了，擁有block最少的應該會有一個最低的achieved Occupancy吧？而擁有最多block的應該會達到最高的achieved Occupancy吧？為了驗證這些想法，我們再看一組數據：

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

看到了吧，（64,2）的achieved Occupancy竟然是最低的，盡管他有最多的block（高中做物理題也是這感覺），它達到了硬件對block數量的限制。

第四個（128,2）和第七個（256,2）擁有擁有差不多的achieved Occupancy。我們對這兩個再做一個試驗，再次增大，將blockDim.y設置為1，這也減少了block的大小：

$ ./sumMatrix 128 1
sumMatrixOnGPU2D <<<(128,16384),(128,1)>>> elapsed 0.032602 sec
$ ./sumMatrix 256 1
sumMatrixOnGPU2D <<<(64,16384), (256,1)>>> 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（啟發） 為調節性能提供了不錯的切入點。