[菜鸟每天来段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).