[菜鸟每天来段CUDA_C] GPU上实现直方图计算

本文通过在GPU上计算直方图说明GPU计算中的原子操作。原子操作是计算中不能分解为更小的部分的操作。当有数千个线程在内存访问上发生竞争时，这些操作能够确保在内存上实现安全的操作。即数据能按照实际的顺序进行读写以至于不发生错误。

 CPU上的直方图计算比较简单（以100M随机生成的无符号字符数据为例），一个for循环扫描一遍数据集就能统计各数据出现的频率。主要代码如下：

 

	unsigned char *buffer = (unsigned char*)big_random_block(_SIZE);

	unsigned int hist[256];
	for (int i=0; i<256; i++)
	{
		hist[i] = 0;
	}

	for (int i=0; i<_SIZE; i++)
	{
		hist[buffer[i]]++;
	}

而在GPU上进行直方图计算时，多线程同时对数据集中的数据进行统计。为了保证计算过程中只有一个线程在对计算结果做改动，CUDA C使用原子操作的方式，调用函数atomicAdd(头文件sm_12_atomic_functions.h)。核函数如下：

__global__ void histoKernel(unsigned char* buffer, long size, unsigned int* histo)
{
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	int strides = blockDim.x * gridDim.x;
	while(i<size)
	{
		atomicAdd(&(histo[buffer[i]]), 1);
		i += strides;
	}
}

实验结果发现，GPU进行的直方图计算性能比CPU上还低(GTX260需要2011.1ms)。原因在于核函数中只包含了非常少的计算，反而要在全局内存上加入原子操作，使得入不敷出性能降低。为此使用共享内存对以上核函数作相应改动。主要是想把每个线程上的统计结果先写到共享内存，线程是哪个的工作结束后再写全局内存，从而降低数千个线程在少量内存地址上的竞争产生的等待(运算时间降低到69.2ms，提高了约30倍)。修改后的核函数如下：

__global__ void histoKernel(unsigned char* buffer, long size, unsigned int* histo)
{
	__shared__ unsigned int temp[256];
	temp[threadIdx.x] = 0;
	__syncthreads();

	int i = threadIdx.x + blockIdx.x * blockDim.x;
	int offset = blockDim.x * gridDim.x;
	while(i<size)
	{
		atomicAdd(&(temp[buffer[i]]), 1);
		i += offset;
	}

	__syncthreads();
	atomicAdd(&(histo[threadIdx.x]), temp[threadIdx.x]);
}

程序中需要用到的生成100M随机数的函数：

void* big_random_block(int size)
{
	unsigned char* data = (unsigned char*)malloc(size);
	if (!data)
	{
		printf("Memery allocate failed!\n");
		return NULL;
	}
	for (int i=0; i<size; i++)
	{
		data[i] = rand();
	}

	return data;
}

程序编译是可能会遇到提示atomicAdd undefined的情况，这跟GPU的型号有关，可在项目属性--配置属性--CUDA RuntimeAPI--GPU--GPU Architecture选择相应的编译选项。

参考资源：

Jason Sanders, Edward Kandrot, CUDA By Example: An Introduction toGeneral-Purpose GPU Programming (2011).