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又一次展现了强大的算力, 并且, 此次也看到了只是小小变更, 让算法更贴合架构, 就让运算耗时减半, 因此在优化方面能够作的工做真的是太多了, 以后还有更多优化相关的文章, 有意见或者建议, 评论区见哦~