CUDA C的并行编程

1.首先我们可以先看一下传统的CPU向量编程

#define N   10
void add( int *a, int *b, int *c ) {
    int tid = 0;    // this is CPU zero, so we start at zero
    while (tid < N) {
        c[tid] = a[tid] + b[tid];
        tid += 1;   // we have one CPU, so we increment by one
    }
}
int main( void ) {
    int a[N], b[N], c[N];
    // fill the arrays 'a' and 'b' on the CPU
    for (int i=0; i<N; i++) {
        a[i] = -i;
        b[i] = i * i;
    }
    add( a, b, c );    // display the results
    for (int i=0; i<N; i++) {
        printf( "%d + %d = %d\n", a[i], b[i], c[i] );
    }
    return 0;
}

上面的add函数针对的是单核CPU编程，如果遇到了多核CPU，那么可以修改add函数，进而每个核可以分配不同的任务，比如说是双核CPU编程，add函数可以修改如下所示

void add( int *a, int *b, int *c ) 
{
    int tid = 0;
    while (tid < N) {
        c[tid] = a[tid] + b[tid];
        tid += 2;
    }
}
void add( int *a, int *b, int *c ) 
{
    int tid = 1;
    while (tid < N) {
        c[tid] = a[tid] + b[tid];
        tid += 2;
    }
}

但是，这样的编程带来了很多缺点，比如说要多编很多的代码。这里引入GPU编程，那么add函数就可以编写成设备函数

首先，是main函数：

#define N   10
int main( void ) {
    int a[N], b[N], c[N];
    int *dev_a, *dev_b, *dev_c;
    // allocate the memory on the GPU
    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) ) );
    // fill the arrays 'a' and 'b' on the CPU
    for (int i=0; i<N; i++) {
        a[i] = -i;
        b[i] = i * i;
    }
    // copy the arrays 'a' and 'b' to the GPU
    HANDLE_ERROR( cudaMemcpy( dev_a, a, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );
    HANDLE_ERROR( cudaMemcpy( dev_b, b, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );
    add<<<N,1>>>( dev_a, dev_b, dev_c );
    // copy the array 'c' back from the GPU to the CPU
    HANDLE_ERROR( cudaMemcpy( c, dev_c, N * sizeof(int),
                              cudaMemcpyDeviceToHost ) );
    // display the results
    for (int i=0; i<N; i++) {
        printf( "%d + %d = %d\n", a[i], b[i], c[i] );
    }
    // free the memory allocated on the GPU
    cudaFree( dev_a );
    cudaFree( dev_b );
    cudaFree( dev_c );
    return 0;
}

其次，是add功能函数：

__global__ void add( int *a, int *b, int *c ) {
    int tid = blockIdx.x;    // handle the data at this index
    if (tid < N)
        c[tid] = a[tid] + b[tid];
}

1>在设备上执行的add函数也是用C语言实现的，同时在函数的前面添加__global__

2> <<< >>>内的参数和内核自己的参数

<<<N，1 >>>内的参数用于在运行时描述如何加载内核（也就是CUDA的C函数），第一个参数是相应设备执行内核的block数。通过内置变量blockIdx可以知道目前哪个block正在运行。

为什么选用blockIdx.x？因为CUDA C设置了一组二维的block。但是注意，luanch的block数量不要超过65535，这是个上限。

3.和C函数相应的CUDA  C 函数

malloc() -> cudamalloc()

free -> cudafree()

memcpy() -> cudaMemcpy()   最后一个参数指明拷贝方向，cudaMemcpyDeviceToHost；cudaMemcpyHostToDevice；cudaMemcpyDeviceToDevice

像是其他的内核函数如，cudaGetDeviceCount(&count)用于检查设备数量，cudaDeviceProp是设备信息结构，cudaGetDeviceProperties(cudaDeviceProp &,int i)用于得到第i个设备信息。