CUDA C的并行编程

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

#define N   10void 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   10int 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个设备信息。