[菜鸟每天来段CUDA_C]基于共享内存的位图与syncthreads的使用

本文使用CUDA实现基于共享内存的位图显示。位图中每个位置的像素值由每个线程计算，计算结果保存到缓冲区(共享内存)中。

结果为一个由多个绿色球形构成的网格(如下图)。

图中可以看出：没用同步(syncthreads)的运行结果是错误的，原因在于一个线程块的线程没有全部计算结束就对共享内存赋值。

Syncthreads的主要功能是对线程块中的线程进行同步，确保线程块中的每个线程都执行完__syncthreads()函数之前的语句才会执行下一条语句。

部分代码如下：

/********************************************************************
*  sharedMem.cu
*********************************************************************/

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cutil_inline.h>
#include "CPUBitmap.h"

#define DIM 512
#define PI 3.1415926535897932f

/************************************************************************/
/* Init CUDA                                                            */
/************************************************************************/
bool InitCUDA(void)
{
  ......
}

/************************************************************************/
__global__ void kernel(unsigned char* ptr)
{
	int x = threadIdx.x + blockIdx.x * blockDim.x;
	int y = threadIdx.y + blockIdx.y * blockDim.y;
	int offset = x + y * blockDim.x * gridDim.x;

	__shared__ float sharedMem[16][16];
	const float period = 128.0f;
	sharedMem[threadIdx.x][threadIdx.y] = 
		255 * (sinf(x*2.0f*PI/period) + 1.0f) * 
		      (sinf(y*2.0f*PI/period) + 1.0f) / 4.0f;

	__syncthreads();
	
	ptr[offset*4 + 0] = 0;
	ptr[offset*4 + 1] = sharedMem[15-threadIdx.x][15-threadIdx.y];
	ptr[offset*4 + 2] = 0;
	ptr[offset*4 + 3] = 255;
}
/************************************************************************/

int main(int argc, char* argv[])
{

	if(!InitCUDA()) {
		return 0;
	}

	CPUBitmap bitmap(DIM, DIM);
	unsigned char* devBitmap;

	cutilSafeCall(cudaMalloc((void**)&devBitmap, bitmap.image_size()));

	dim3 grids(DIM/16, DIM/16);
	dim3 threads(16, 16);

	kernel<<<grids, threads>>>(devBitmap);

	cutilSafeCall(cudaMemcpy(bitmap.get_ptr(), devBitmap, bitmap.image_size(), cudaMemcpyDeviceToHost));

	bitmap.display_and_exit();

	cudaFree(devBitmap);

	return 0;
}

完整代码链接： http://download.csdn.net/detail/jonny_super/6606341

(VS2008 + OpenGL + CUDA)

参考资源：

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