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类型变量修改调用函数的方式
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:线程索引,前面章节已经详细探讨过了,这里不再赘述。