并行编程实战——CUDA入门编程的函数

一、CUDA的函数类别

已经分析过CUDA的开发方式有两种情况，一种是使用Runtime,另外一种是设备驱动的形式。这里只对前者进行分析，后者有机会再进行说明。
对于开发框架来说，所谓定义的函数，一定是分成两大类，对外的框架接口和内部实现的重要功能函数。由于此处不是分析CUDA框架底层的实现，所以这里重点介绍一下框架的接口函数的功能和应用。
一般来说，CUDA函数分成几大类：
1、内存管理函数
用于设备内存的分配和释放；设备与主机间数据的复制、设置；统一内存的分配和管理；异步内存的复制；相关缓存管理的函数；
2、设备管理函数
用来获取可用设备数量、当前设备的设定和获取以及设备属性的管理等相关函数
3、同步/异步并发函数
主要提供块内线程同步、主机与设备同步以及流同步等；同时提供异步并发的相关接口函数
4、纹理内存函数
纹理的引用、获取及绑定纹理内存等函数
5、图形互操作函数
它主要用来与OpenGL 和 Direct3D两个图形库进行互操作
6、流和事件函数
主要提供流的创建、销毁以及相关的事件管理函数
7、错误处理函数
主要用于检查运行时生成的错误相关的接口函数
8、其它函数
如数学函数以及原子操作等函数

CUDA框架的接口函数都是以cuda为前缀的，这也是一个非常好的编程风格。在后面文章中将根据这几类不断的对相关的常见函数进行引入并分析说明。

二、通过例程引入

为了能更好的理解常用的接口函数，先看一下在Windows平台上的默认的kernel.cu这个文件中的代码：

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

int main()
{
    const int arraySize = 5;
    const int a[arraySize] = { 1, 2, 3, 4, 5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    // Add vectors in parallel.
    cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }

    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
        c[0], c[1], c[2], c[3], c[4]);

    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
    cudaError_t cudaStatus;

    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // Allocate GPU buffers for three vectors (two input, one output)    .
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    // Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    // Launch a kernel on the GPU with one thread for each element.
    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

    // Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

    // Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);

    return cudaStatus;
}

下面将针对上面代码中的接口函数进行分析说明。

三、函数说明

在上面的代码中，既然是默认的程序，一般代表着，内部出现的函数是高概率出现的，下面就针对函数出现的先后进行初步的说明：
1、cudaSetDevice()
它是一个设备管理函数，用来设备在哪个设备（GPU）上运行，默认一般是cudaSetDevice(0)
2、cudaMalloc()
主要用于分配显存空间，也就是设备的内存分配
3、cudaMemcpy()
用于在主机和设备间数据的复制
4、核函数
核函数使用__global__标志，其具体的内容在前面分析过，不再赘述
5、cudaGetLastError()
用于获取最新的运行时调用错误，针对所有的CUDA错误，均可通过函数cudaGetErrorString来获取错误的具体信息
6、cudaDeviceSynchronize()和cudaThreadSynchronize()
这个函数比较重要，它有点类似一些异步框架中的同步等待函数，比如线程中的join()函数。这个函数要确保以前调用的与当前设备相关的主机线程中CUDA API都已执行完成。
在前面分析过，主机和设备间是异步执行的，为了确保主机程序退出时，CUDA相关调用已经完成，它必须阻塞主机线程直到相关CUDA的执行完成（包括数据的反向复制、错误的检查和处理。这个操作的结查就是性能一定会受损。
cudaThreadSynchronize()只处理当前的主机线程，而前者则包括所有与当前设备连接的主机线程。一个粒度小，一个粒度大。
7、cudaFree()
释放设备的内存和前面的cudaMalloc()匹配使用
8、cudaDeviceReset()
此函数用来重置设备，即释放相关的设备内存并重置相关状态

其实通过上面的初步分析，可以明显感觉到一些主流框架使用的味道，不管是同步处理还是内存管理以及相关的错误检查，只要有一定的开发经验的都会感到很熟悉。所以大家不必有畏难情绪。

四、总结

上面只是针对Runtime在Windows平台上的初步的接口函数进行了说明，让大家有一个特定的靶子进行分析学习。不同版本的CUDA可能在细节上有所差异，函数的可能有增有减，这个需要大家根据自己的实际版本查看相应的函数接口。特别是在一些最新版本或较老的版本中间，差异有可能相当大。大家不必计较这些不同，反而可以从不同函数的版本的迭代过程中，寻找到CUDA版本升级的趋势和方向。