GPU編程(四): 並行規約優化

目錄

• 前言
• cuda-gdb
• 未優化並行規約
• 優化後並行規約
• 結果分析
• 最後

---

前言

• 以前第三篇也看到了, 並行方面GPU真的是無往不利, 如今再看下第二個例子, 並行規約. 經過此次的例子會發現, 須要瞭解GPU架構, 而後寫出與之對應的算法的, 二者結合才能獲得使人驚歎的結果.
• 此次也會簡要介紹下cuda-gdb的用法, 其實和gdb用法幾乎同樣, 也就是多了個cuda命令.

---

cuda-gdb

若是以前沒有用過gdb, 能夠速學一下, 就幾個指令. 想要用cuda-gdb對程序進行調試, 首先你要確保你的gpu沒有在運行操做系統界面, 比方說, 我用的是ubuntu, 我就須要用sudo service lightdm stop關閉圖形界面, 進入tty1這種字符界面. 固然用ssh遠程訪問也是能夠的. 接下來, 使用第二篇中矩陣加法的例子. 可是注意, 編譯的使用須要改變一下, 加入**-g -G**參數, 其實和gdb是類似的.

nvcc -g -G CUDAAdd.cu -o CUDAAdd.o
複製代碼

而後使用cuda-gdb CUDAAdd.o便可對程序進行調試.

在調試以前, 我把代碼貼出來:

#include <stdio.h>

__global__ void add(float * x, float *y, float * z, int n){
        int index = threadIdx.x + blockIdx.x * blockDim.x;
        int stride = blockDim.x * gridDim.x;

        for (int i = index; i < n; i += stride){
                z[i] = x[i] + y[i];
        }
}

int main()
{
    int N = 1 << 20;
    int nBytes = N * sizeof(float);

    float *x, *y, *z;
    cudaMallocManaged((void**)&x, nBytes);
    cudaMallocManaged((void**)&y, nBytes);
    cudaMallocManaged((void**)&z, nBytes);

    for (int i = 0; i < N; ++i)
    {
        x[i] = 10.0;
        y[i] = 20.0;
    }

    dim3 blockSize(256);
    // 4096
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);

    add << < gridSize, blockSize >> >(x, y, z, N);

    cudaDeviceSynchronize();

    float maxError = 0.0;
    for (int i = 0; i < N; i++){
    	        maxError = fmax(maxError, (float)(fabs(z[i] - 30.0)));
    }
    printf ("max default: %.4f\n", maxError);

    cudaFree(x);
    cudaFree(y);
    cudaFree(z);

    return 0;
}
複製代碼

以後就是常規操做了, 添加斷點, 運行, 下一步, 查看想看的數據. 不一樣點是cuda的指令, 例如cuda block(1,0,0)能夠從一開始block(0,0,0)切換到block(1,0,0).

---

未優化並行規約

若是按照常規的思路, 兩兩進行進行加法運算. 每次步長翻倍便可, 從算法的角度來講, 這是沒啥問題的. 可是沒有依照GPU架構進行設計.

#include <stdio.h>

const int	threadsPerBlock = 512;
const int	N		= 2048;
const int	blocksPerGrid	= (N + threadsPerBlock - 1) / threadsPerBlock; /* 4 */

__global__ void ReductionSum( float * d_a, float * d_partial_sum )
{
	/* 申請共享內存, 存在於每一個block中 */
	__shared__ float partialSum[threadsPerBlock];

	/* 肯定索引 */
	int	i	= threadIdx.x + blockIdx.x * blockDim.x;
	int	tid	= threadIdx.x;

	/* 傳global memory數據到shared memory */
	partialSum[tid] = d_a[i];

	/* 傳輸同步 */
	__syncthreads();

	/* 在共享存儲器中進行規約 */
	for ( int stride = 1; stride < blockDim.x; stride *= 2 )
	{
		if ( tid % (2 * stride) == 0 )
			partialSum[tid] += partialSum[tid + stride];
		__syncthreads();
	}

	/* 將當前block的計算結果寫回輸出數組 */
	if ( tid == 0 )
		d_partial_sum[blockIdx.x] = partialSum[0];
}


int main()
{
	int size = sizeof(float);

	/* 分配顯存空間 */
	float	* d_a;
	float	* d_partial_sum;

	cudaMallocManaged( (void * *) &d_a, N * size );
	cudaMallocManaged( (void * *) &d_partial_sum, blocksPerGrid * size );

	for ( int i = 0; i < N; ++i )
		d_a[i] = i;

	/* 調用內核函數 */
	ReductionSum << < blocksPerGrid, threadsPerBlock >> > (d_a, d_partial_sum);

	cudaDeviceSynchronize();

	/* 將部分和求和 */
	int sum = 0;
	for ( int i = 0; i < blocksPerGrid; ++i )
		sum += d_partial_sum[i];

	printf( "sum = %d\n", sum );

	/* 釋放顯存空間 */
	cudaFree( d_a );
	cudaFree( d_partial_sum );

	return(0);
}
複製代碼

---

優化後並行規約

其實須要改動的地方很是小, 改變步長便可.

__global__ void ReductionSum( float * d_a, float * d_partial_sum )
{
    // 相同, 略去
	/* 在共享存儲器中進行規約 */
	for ( int stride = blockDim.x / 2; stride > 0; stride /= 2 )
	{
		if ( tid < stride )
			partialSum[tid] += partialSum[tid + stride];
		__syncthreads();
	}
    // 相同, 略去
}
複製代碼

---

結果分析

以前的文章裏面也說過warp. warp: GPU執行程序時的調度單位, 目前cuda的warp的大小爲32, 同在一個warp的線程, 以不一樣數據資源執行相同的指令, 這就是所謂SIMT. 說人話就是, 這32個線程必需要幹相同的事情, 若是有線程動做不一致, 就須要等待一波線程完成本身的工做, 而後再去作另一件事情. 因此, 用圖說話就是, 第二種方案能夠更快將warp閒置, 交給GPU調度, 因此, 確定是第二種更快.

圖一在運算依次以後, 沒有warp能夠空閒, 而圖二直接空閒2個warp. 圖一到了第二次能夠空閒2個warp, 而圖二已經空閒3個warp. 我這副圖只是示意圖, 若是是實際的, 差距會更大.

因此來看下運行耗時, 會發現差距仍是很大的, 幾乎是差了一倍. 不過GPU確實算力太猛, 這樣看還不太明顯, 有意放大數據量會更加明顯.

最後

因此GPU又一次展現了強大的算力, 並且, 此次也看到了只是小小變更, 讓算法更貼合架構, 就讓運算耗時減半, 因此在優化方面能夠作的工做真的是太多了, 以後還有更多優化相關的文章, 有意見或者建議, 評論區見哦~