cuda之内存共享案例

1.代码

#include<iostream>
#include<cstdlib>
using namespace std;
#define N 64

__global__ void inverseArray(int *arr,int n) {
	__shared__ int s[N];  /* 共享数组变量 */
	int ind = threadIdx.x;
	int tr = n-ind-1;
	s[ind] = arr[ind];
	__syncthreads();  /* 同步所有线程 */
	arr[ind] = s[tr];
}

void initializeArray(int *arr) {
	for(int i=0;i<N;i++)
		arr[i] = i;
}

int main() {
	int *a,*dev_a;
	int *r;
	r = (int *)malloc(sizeof(int)*N);
	a = (int *)malloc(sizeof(int)*N);
	initializeArray(a);
	for(int i=0;i<N;i++)
		cout<<a[i]<<" ";
	cout<<endl;
	cudaMalloc((void **)&dev_a,sizeof(int)*N);
	cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice);
	inverseArray<<<1,N>>>(dev_a,N);
	cudaMemcpy(r,dev_a,N*sizeof(int),cudaMemcpyDeviceToHost);
	for(int i=0;i<N;i++)
		cout<<r[i]<<" ";
	cout<<endl;
	return 0;
}

2.描述

上面代码中，__shared__为一个block内所有线程的共享内存，其生命周期即该block的生命周期，其只能被该block内的线程所访问。

__syncthreads()像是一个壁垒，它会迫使一个block内的所有线程等待执行，等到所有线程都被阻塞之后再开始执行。