cuda学习笔记(1)

1.基本流程

cuda代码（.cu）的目的是并行运算。只要在c/c++代码中调用以 __ global__为关键字修饰的函数( __ global __ void function( type *InputArrayA, type *InputArrayB, type *OutputArrayA) )，称为核函数，代码经nvcc编译，识别到核函数就会编译成gpu指令码; 调用该函数时，要在函数名称加上 <<<blocksPerGrid, threadsPerBlock>>> ( function<<<blocksPerGrid, threadsPerBlock>>>( type *InputArrayA, type *InputArrayB, type *OutputArrayA) )。不过，gpu 只能操作gpu上的变量，所以在调用 __ global __ 函数之前，先用 cudaMalloc 申请好在cuda变量内存（__global函数的参数：input array，output array)，并用 cudaMemcpy （cudaMemcpyHostToDevice) 赋值输入array。待函数执行完成后，执行结果保存在输出array中，用 cudaMemcpy （cudaMemcpyDeviceToHost) 把执行结果从gpu内存中copy到cpu中，并行计算完成，用 cudaFree 释放之前申请的cuda变量内存。 以上就是cpu代码中调用gpu的流程。

1.1 blocksPerGrid, threadsPerBlock 说明

调用 cuda 核函数需要指定调用多少个block,每个block包含多少个thread。其中，多个block组成一个 grid . 共调用了 blocksPerGrid*threadsPerBlock 个并行执行的线程，所以要在cuda核函数中明确的指定每个线程执行时对应的array index。注意：thread, block有.x, .y二维数据，但有时只用其中一维.x 。下面将给出一个简单的demo，执行c = a+b运算。

example1: simple_add.cu

#include "../common/book.h"

#define N   10

__global__ void add( int *a, int *b, int *c ) 
{
    int tid = blockIdx.x;    // this thread handles the data at its thread id
    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;

    // 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 ) );
    //N blocks, 1 thread per block for N length arrays parallel computation(add)
    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
    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaFree( dev_b ) );
    HANDLE_ERROR( cudaFree( dev_c ) );

    return 0;
}