CUDA 学习笔记
主要参考《GPU高性能编程CUDA实战》
一般将CPU及其系统的内存称为主机，将GPU及其内存称为设备。
在GPU设备上执行的函数通常称为核函数(Kernal);
一个简单的程序:
1
2
3
4
5
6
7
8
9
10
| #include <stdio.h>
__global__ void kernel(void){
}
int main(void){
kernel<<<1,1>>>();
printf("Hello world\n");
return 0;
}
|
global修饰符告诉编译器，函数应该被编译为在设备上而不是在主机上运行。于是，kernel()函数将被交给编译设备代码的编译器，而main()函数将被交给主机编译器。
尖括号表示要将一些参数传递给运行时系统，用以告诉运行时如何启动设备代码。
再来看一个更复杂的程序:
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
| #include <stdio.h>
#include <cuda_runtime.h>
__global__ void add( int a, int b, int *c ) {
*c = a + b;
}
int main( void ) {
int c;
int *dev_c;
cudaMalloc( (void**)&dev_c, sizeof(int) );
add<<<1,1>>>( 2, 7, dev_c );
cudaMemcpy( &c, dev_c, sizeof(int),cudaMemcpyDeviceToHost ) ;
printf( "2 + 7 = %d\n", c );
cudaFree( dev_c );
return 0;
}
|
cudaMalloc()用于分配内存，但与malloc()不同的是，它是在GPU上分配内存。由于分配的地址时在GPU上，所以程序员一定不能在主机代码中对这个指针进行解引用。同样的，也不能用free()函数来释放这些内存，只能用cudaFree()。
为了让主机代码能得到GPU计算的结果，可以用cudaMemcpy()函数，它的第四个参数数据移动的方向，是从主机到设备、设备到主机还是设备到设备。
上面说的这些内容好像对加快程序的执行速度没什么帮助，我们用下面这个向量求和的例子来了解CUDA如何让程序变得更快:
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
| #include<stdio.h>
#include<book.h>
#define N 10
__global__ void add(int* a,int* b,int* c){
int tid=blockIdx.x;
if(tid<N)c[tid]=a[tid]+b[tid];
}
int main(void){
int a[N],b[N],c[N];
int* dev_a,*dev_b,*dev_c;
HANDLE_ERROR(cudaMalloc((void**)&dev_a,N*sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**)&dev_b,N*sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**)&dev_c,N*sizeof(int)));
for(int i=0;i<N;i++)a[i]=-i,b[i]=i*i;
HANDLE_ERROR(cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_c,c,N*sizeof(int),cudaMemcpyHostToDevice));
add<<<N,1>>>(dev_a,dev_b,dev_c);
HANDLE_ERROR(cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost));
for(int i=0;i<N;i++)printf("%d %d\n",i,c[i]);
cudaFree(dev_a),cudaFree(dev_b),cudaFree(dev_c);
return 0;
}
|
注意到，我们用了add<<<N,1>>>。N代表了设备在执行核函数时使用的并行线程块(Block)的数量。当启动核函数时，我们告诉运行时，我们想要一个一维线程格(Grid)，其中包含N个线程块。
如何在代码中知道当前正在运行的是哪个线程块？可以通过blockIdx.x得到，它是一个内置变量。这样，我们就把向量求和并行化，达到了加速的效果。
在上面的例子中，我们启动了多个线程块，但每个线程块只有1个线程。要想在每个线程块里启动多个线程，可以这样写:
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
| #include<stdio.h>
#include<book.h>
#define N 100
__global__ void add(int* a,int* b,int* c){
int tid=threadIdx.x+blockIdx.x*blockDim.x;
if(tid<N)c[tid]=a[tid]+b[tid];
}
int main(void){
int a[N],b[N],c[N];
int* dev_a,*dev_b,*dev_c;
HANDLE_ERROR(cudaMalloc((void**)&dev_a,N*sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**)&dev_b,N*sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**)&dev_c,N*sizeof(int)));
for(int i=0;i<N;i++)a[i]=-i,b[i]=i*i;
HANDLE_ERROR(cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_c,c,N*sizeof(int),cudaMemcpyHostToDevice));
add<<<(N+127)/128,128>>>(dev_a,dev_b,dev_c);
HANDLE_ERROR(cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost));
for(int i=0;i<N;i++)printf("%d %d\n",i,c[i]);
cudaFree(dev_a),cudaFree(dev_b),cudaFree(dev_c);
return 0;
}
|
注意到启动核函数时用的是<<<(N+127)/128,128>>>,(N+127)/128是对N/128向上取整，得到线程块的数量。第二个128代表一个线程块有128个线程。这样子可能会导致线程数量大于N，但没有关系，我们会在核函数里判断线程索引是否越界。
在核函数中，线程索引通过$tid=threadIdx.x+blockIdx.x*blockDim.x$得到。
为了避免使用过多的线程块，我们直接指定线程块的数量，然后在核函数中用while循环对下表递增，如下所示:
1
2
3
4
| __global__ void add(int* a,int* b,int* c){
int tid=threadIdx.x+blockIdx.x*blockDim.x;
while(tid<N)c[tid]=a[tid]+b[tid],tid+=blockDim.x*gridDim.x;
}
|