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