GPU编程自学3 —— CUDA程序初探

深度学习的兴起，使得多线程以及GPU编程逐渐成为算法工程师无法规避的问题。这里主要记录自己的GPU自学历程。

《GPU编程自学1 —— 引言》
《GPU编程自学2 —— CUDA环境配置》
《GPU编程自学3 —— CUDA程序初探》
《GPU编程自学4 —— CUDA核函数运行参数》
《GPU编程自学5 —— 线程协作》
《GPU编程自学6 —— 函数与变量类型限定符》
《GPU编程自学7 —— 常量内存与事件》
《GPU编程自学8 —— 纹理内存》
《GPU编程自学9 —— 原子操作》
《GPU编程自学10 —— 流并行》

三、 CUDA程序初探

3.1 主机与设备

通常将CPU及其内存称之为主机，GPU及其内存称之为设备。

如下图所示，新建一个NVIDIA CUDA工程，并命名为 “1-helloworld”

之后发现项目里多了一个 “kernel.cu”的文件，该文件内容是一个经典的 矢量相加 的GPU程序。

可以暂时全部注释该代码，并尝试编译运行下面的我们经常见到的编程入门示例：

#include <iostream>

int main()
{
    std::cout<<"Hello, World!"<<std::endl;
    system("pause");
    return 0;
}

这看起来和普通的C++程序并没什么区别。这个示例只是为了说明CUDA C编程和我们熟悉的标准C在很大程度上是没有区别的。同时，这段程序直接运行在主机上。

接下来，我们看看如何使用GPU来执行代码。如下：

#include <iostream>

__global__  void mkernel(void){}

int main()
{
    mkernel <<<1,1>>>();
    std::cout<<"Hello, World!"<<std::endl;
    system("pause");
    return 0;
}

与之前的代码相比，这里主要增加了

一个空的函数mkernel()，并带有修饰符 global
对空函数的调用，并带有修饰符 <<<1,1>>>

_global_ 为CUDA C为标准C增加的修饰符，表示该函数将会交给编译设备代码的编译器(NVCC)并最终在设备上运行。而 main函数则依旧交给系统编译器(VS2013)。

其实，CUDA就是通过直接提供API接口或者在语言层面集成一些新的东西来实现在主机代码中调用设备代码。

3.2 第一个GPU程序：矢量相加

下面主要通过代码解读的形式来进行我们的第一个GPU程序。

程序遵循以下流程：
主机端准备数据 -> 数据复制到GPU内存中 -> GPU执行核函数 -> 数据由GPU取回到主机

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

// 接口函数： 主机代码调用GPU设备实现矢量加法 c = a + b
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; // 获取线程ID
    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 };

    // 并行矢量相加
    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]);

    // CUDA设备重置，以便其它性能检测和跟踪工具的运行，如Nsight and Visual Profiler to show complete traces.traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}

// 接口函数实现： 主机代码调用GPU设备实现矢量加法 c = a + b
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;

    // 选择程序运行在哪块GPU上，(多GPU机器可以选择)
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // 依次为 c = a + b三个矢量在GPU上开辟内存 .
    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;
    }

    // 将矢量a和b依次copy进入GPU内存中
    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;
    }

    // 运行核函数，运行设置为1个block，每个block中size个线程
    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

    // 检查是否出现了错误
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

    // 停止CPU端线程的执行，直到GPU完成之前CUDA的任务，包括kernel函数、数据拷贝等
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

    // 将计算结果从GPU复制到主机内存
    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;
}