共享內(nèi)存“__share__”

線程同步機(jī)制“__syncthreads()”

#include "cuda_runtime.h"    #include <highgui.hpp>    using namespace cv;#define DIM 600   //圖像長(zhǎng)寬#define PI 3.1415926535897932f  __global__ void kernel(unsigned char *ptr){	// map from blockIdx to pixel position    	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 * 3 + 0] = 0;	ptr[offset * 3 + 1] = sharedMem[15 - threadIdx.x][15 - threadIdx.y];	ptr[offset * 3 + 2] = 0;}// globals needed by the update routine    struct DataBlock{	unsigned char   *dev_bitmap;};int main(void){	DataBlock   data;	cudaError_t error;	Mat image = Mat(DIM, DIM, CV_8UC3, Scalar::all(0));	data.dev_bitmap = image.data;	unsigned char    *dev_bitmap;	error = cudaMalloc((void**)&dev_bitmap, 3 * image.cols*image.rows);	data.dev_bitmap = dev_bitmap;	dim3    grid(DIM / 10, DIM / 10);	dim3   block(10, 10);	//DIM*DIM個(gè)線程塊  	kernel << <grid, block >> > (dev_bitmap);	error = cudaMemcpy(image.data, dev_bitmap,		3 * image.cols*image.rows,		cudaMemcpyDeviceToHost);	error = cudaFree(dev_bitmap);	imshow("__share__ and __syncthreads()", image);	waitKey();}如果線程間不加入__syncthreads()同步機(jī)制，同一線程塊內(nèi)不同線程訪問(wèn)sharedMem，獲取的結(jié)果可能是不一樣的，生成的圖像如下，有散亂的雜點(diǎn)：加入__syncthreads()同步機(jī)制，保證了同一線程塊中不同的線程都執(zhí)行完成__syncthreads()這個(gè)集合點(diǎn)之前的部分之后，才繼續(xù)往下執(zhí)行，所以不同的線程訪問(wèn)sharedMem獲取的結(jié)果是一致的，圖像無(wú)雜散點(diǎn)，是一個(gè)規(guī)律的排布：