[CUDA]CUDA C并行编程

CUDA将GPU称为设备，将CPU称为主机，一般的计算方式是将数据从内存拷贝进GPU内存(显存)，通过GPU计算再拷回内存中。

下面的代码是一个经典的通过GPU进行的向量加法运算

代码来自http://bbezxcy.iteye.com/blog/2220819

#include<cuda_runtime.h>
#include<windows.h>
#include<iostream>
using namespace std;
const int nMax = 30000;
__global__ void addKernel(float *aaa,float *bbb, float *ccc)
{
	//int i = blockIdx.x;
	int i = threadIdx.x + blockIdx.x*blockDim.x;
	ccc[i] = 0;
	if (i < nMax)for (int j = 0; j < 500; j++)ccc[i] += aaa[i] * bbb[i];
}
void add(float *a, float *b,float *c,int i){
	for (int j = 0; j<500; j++) c[i] += a[i] * b[i];
}
int main(){
	float a[nMax], b[nMax], c[nMax];
	float *devA, *devB, *devC;
	clock_t startT, endT;
	for (int i = 0; i < nMax; i++){
		a[i] = i*1.010923;
		b[i] = 2.13*i;
	}
	startT = clock();
	cudaMalloc((void**)&devA, nMax*sizeof(float));
	cudaMalloc((void**)&devB, nMax*sizeof(float));
	cudaMalloc((void**)&devC, nMax*sizeof(float));
	endT = clock();
	cout << "分配设备空间耗时 " << endT - startT << "ms"<<endl;


	startT = clock();
	cudaMemcpy(devA, a,nMax*sizeof(float),cudaMemcpyHostToDevice);
	cudaMemcpy(devB, b, nMax*sizeof(float), cudaMemcpyHostToDevice);
	endT = clock();
	cout << "数据从主机写入设备耗时 " << endT - startT << "ms" << endl;

	startT = clock();

	cudaEvent_t start1;
	cudaEventCreate(&start1);
	cudaEvent_t stop1;
	cudaEventCreate(&stop1);
	cudaEventRecord(start1, NULL);

	addKernel<<<60,501>>>(devA, devB, devC);

	cudaEventRecord(stop1, NULL);
	cudaEventSynchronize(stop1);
	float msecTotal1 = 0.0f;
	cudaEventElapsedTime(&msecTotal1, start1, stop1);
	//cout << msecTotal1 << "ddd" << endl;
	endT = clock();
	cout << "GPU计算耗时 " << msecTotal1 << "ms" << endl;

	startT = clock();
	cudaMemcpy(c, devC, nMax*sizeof(float), cudaMemcpyDeviceToHost);
	endT = clock();
	cout << "数据从设备写入主机耗时 " << endT - startT << "ms" << endl;

	cout <<"GPU计算结果 "<< c[nMax - 1] << endl;
	for (int i = 0; i < nMax; i++){
		a[i] = i*1.010923;
		b[i] = 2.13*i;
		c[i] = 0;
	}
	startT = clock();
	for (int i = 0; i < nMax; i++){
		add(a, b, c, i);
	}
	endT = clock();
	cout << "CPU计算耗时 " << endT - startT << "ms" << endl;
	cout << "CPU计算结果 " << c[nMax - 1] << endl;

        //释放在设备上分配的空间
	cudaFree(devA);
	cudaFree(devB);
	cudaFree(devC);
	cin >> a[0];
	return 0;
}

 上面的代码中使用了一些通用模式

1，调用cudaMalloc();这个函数在设备（GPU）上分配内存。一般来说为了避免内存泄漏计算完成之后需要通过cudaFree来释放内存空间。

2，cudaMemcpy(devA, a,nMax*sizeof(float),cudaMemcpyHostToDevice);这个函数用来处理主机和设备之间的数据拷贝。最后的参数cudaMemcpyHostToDevice代码是从主机拷贝去设备，如果需要从设备拷贝数据到主机需要将这个参数改为cudaMemcpyDeviceToHost。

3，在定义函数的时候在前面加上__global__，则这个函数就是一个在主机上调用，在设备上运行的函数。在上面的代码里，调用__global__函数代码是

addKernel<<<60,1>>>(devA, devB, devC);

这里的<<<60,501>>>的意思是，调用函数的时候，开出60个线程格，每个线程格包含501个线程。在global函数中通过代码int i = threadIdx.x + blockIdx.x*blockDim.x;得到当前线程是第几个线程。

 

在调用global函数的时候，我们可以通过dim3类型变量修改调用函数的方式

例如dim3 grid(10,10);addKernel<<<grid,501>>>(devA, devB, devC);这样就可以把按照一维排列的线程块改为在二维空间内排布。函数内可以通过一下代码得到当前线程的标号

int x = blockIdx.x;

int y = blockIdx.y;

int threadId = x + y *gridDim.x;

 

CUDA为我们内建了一些变量用于访问线程格、线程块的尺寸和索引等信息，它们是：

      1. gridDim：代表线程格（grid）的尺寸，gridDim.x为x轴尺寸，gridDim.y、gridDim.z类似。拿上图来说，它的gridDim.x = 3，gridDim.y = 2，gridDim.z = 1。

      2. blockIdx：代表线程块（block）在线程格（grid）中的索引值，拿上图来说，Block（1,1）的索引值为：blockIdx.x = 1，blockIdx.y = 1。

      3. blockDim：代表线程块（block）的尺寸，blockDIm.x为x轴尺寸，其它依此类推。拿上图来说，注意到Block（1,1）包含了4 * 3个线程，因此blockDim.x = 4, blockDim.y = 3。

 

      4. threadIdx：线程索引，前面章节已经详细探讨过了，这里不再赘述。