一、前言

       本文介紹CUDA編程的共享內(nèi)存和同步。共享內(nèi)存中的變量（核函數(shù)中用__shared__聲明），在GPU上啟動(dòng)的每個(gè)線程塊，編譯器都創(chuàng)建該變量的副本，若啟動(dòng)N個(gè)線程塊，則有N個(gè)該變量副本，為每個(gè)線程塊私有；同步則是使線程塊中所有的線程能夠在執(zhí)行完某些語句后，才執(zhí)行后續(xù)語句。

二、線程塊、線程索引

以下為線程塊與線程的層次結(jié)構(gòu)圖

       每個(gè)線程均獨(dú)自執(zhí)行核函數(shù)，若在核函數(shù)中聲明共享變量，則每個(gè)線程塊均擁有該變量的一個(gè)副本，且該副本為該程序塊內(nèi)的所有線程所共享。

三、共享變量和同步例子

(1)以下程序?qū)崿F(xiàn)了點(diǎn)積運(yùn)算，計(jì)算公式為 f(n) = 1+2*2+ 3*3+ … (n-1)*(n-1)，使用共享變量計(jì)算各個(gè)程序塊內(nèi)所有線程的求和運(yùn)算結(jié)果。

#include <cuda_runtime.h>  #include <iostream>  //main1.cu#include "book.h"using namespace std;#define N 33*1024          //數(shù)組長度const int threadsPerBlock = 64;   //每個(gè)線程塊的線程數(shù)量const int blocksPerGrid = 64;     //第一維線程格內(nèi)線程數(shù)量__global__ void add(float *a, float *b, float *c){	__shared__ float cache[threadsPerBlock];         //__shared__聲明共享變量，每個(gè)線程塊均有自己的副本，被其所有	                                                 //線程共享,這里用于存放每個(gè)線程塊內(nèi)各個(gè)線程所計(jì)算得的點(diǎn)積和	int index =threadIdx.x + blockIdx.x *blockDim.x;    //將線程塊、線程索引轉(zhuǎn)換為數(shù)組的索引	int cacheIdx = threadIdx.x;	float temp = 0;	while (index < N){		temp += a[index] * b[index];		index += gridDim.x * blockDim.x;	}	cache[cacheIdx] = temp;      //存放每個(gè)線程塊內(nèi)各個(gè)線程所計(jì)算得的點(diǎn)積和	__syncthreads();               //cuda內(nèi)置函數(shù)，使所有線程均執(zhí)行完該命令前代碼，才執(zhí)行后面語句，也即保持同步	                               //目的為獲得各個(gè)cache副本，此時(shí)共有64個(gè)cache副本		//規(guī)約運(yùn)算，將每個(gè)cache副本求和，結(jié)果保存于cache[0]	int i = blockDim.x / 2;	while (i != 0){		if (cacheIdx < i){			cache[cacheIdx] += cache[i + cacheIdx];		}		__syncthreads();    //所有線程完成一次規(guī)約運(yùn)算，方可進(jìn)行下一次		i /= 2;	}	if (cacheIdx == 2)     //一個(gè)操作只需一個(gè)線程完成即可		c[blockIdx.x] = cache[0]; //所有副本的cache[0] 存放于數(shù)組c}int main(){	float a[N], b[N];	float *c = new float[blocksPerGrid];	float *dev_a, *dev_b, *dev_c;	//gpu上分配內(nèi)存	HANDLE_ERROR(cudaMalloc((void**)&dev_a, N*sizeof(float)));	HANDLE_ERROR(cudaMalloc((void**)&dev_b, N*sizeof(float)));	HANDLE_ERROR(cudaMalloc((void**)&dev_c, N*sizeof(float)));	//為數(shù)組a,b初始化	for (int i = 0; i < N; ++i){		a[i] = i;		b[i] = i;	}	//講數(shù)組a,b數(shù)據(jù)復(fù)制至gpu	(cudaMemcpy(dev_a, a, N*sizeof(float), cudaMemcpyHostToDevice));	(cudaMemcpy(dev_b, b, N*sizeof(float), cudaMemcpyHostToDevice));		add <<< blocksPerGrid, threadsPerBlock >> >(dev_a, dev_b, dev_c);	//將數(shù)組dev_c復(fù)制至cpu	HANDLE_ERROR(cudaMemcpy(c, dev_c, blocksPerGrid*sizeof(float), cudaMemcpyDeviceToHost));		//進(jìn)一步求和	double sums = 0.0;	for (int i = 0; i < blocksPerGrid; ++i){		sums += c[i];	}	//顯示結(jié)果	cout << "gpu dot compute result:" << sums << "/n";	sums = 0.0;	for (int i = 0; i < N; ++i){		sums += i*i;	}	cout << "cpu dot compute result:" << sums << "/n";	//釋放在gpu分配的內(nèi)存	cudaFree( dev_a);	cudaFree(dev_b);	cudaFree(dev_c);	delete c;	return 0;}
運(yùn)行結(jié)果
(2)以下程序使用二維程序塊共享變量計(jì)算圖像數(shù)據(jù)，生成圖像
//main2.cu#include <cuda_runtime.h>  #include <iostream>  #include "book.h"#include <opencv2/opencv.hpp>using namespace cv;using namespace std;#define PI 3.1415926#define DIM 1024       //灰度圖像的長與寬__global__ void kernel(uchar * _ptr ){	int x = threadIdx.x + blockIdx.x * blockDim.x;	int y = threadIdx.y + blockIdx.y * blockDim.y;	int idx = x + y *gridDim.x *blockDim.x;	__shared__ float shared [16][16] ;   //每個(gè)線程塊中每個(gè)線程的共享內(nèi)存緩沖區(qū)	const float period = 128.0f;	shared[threadIdx.x][threadIdx.y] = 255 * (sinf(x*2.0f*PI / period) + 1.0f)*(sinf(y*2.0f*PI / period) + 1.0f) / 4.0f;	__syncthreads();   //使所有shared副本均被計(jì)算完成	_ptr[idx] = shared[15 - threadIdx.x][15 - threadIdx.y];}int main(){	Mat src(DIM,DIM , CV_8UC1 , Scalar::all(0));	uchar *ptr_dev;	HANDLE_ERROR(cudaMalloc((void**)&ptr_dev, DIM * DIM*sizeof(uchar)));	dim3 blocks(DIM / 16, DIM / 16);	dim3 threads(16 ,16);	kernel << < blocks, threads >> >( ptr_dev );	HANDLE_ERROR(cudaMemcpy(src.data, ptr_dev, DIM * DIM*sizeof(uchar), cudaMemcpyDeviceToHost));	cudaFree(ptr_dev);	namedWindow("Demo",  0);	imshow("Demo" , src);	waitKey(0);	return 0;}運(yùn)行結(jié)果：