并行编程实战——CUDA编程基础入门-CSDN博客

本文链接：https://blog.csdn.net/fpcc/article/details/149149118

一、概述

通过前面的学习，大家基本明白了什么是CUDA和如何安装其开发环境。本文将对CUDA编程的基本内容与大家共同学习一下，明白CUDA编程的流程和相关的一些整体的框架，并对一些具体的细节和关键字进行说明。

二、CUDA的基础框架说明

CUDA编程中，是通过CPU和GPU共同来完成的，CPU居于主导的控制地位而GPU用于协助进行大规模的数据计算和处理。CPU与GPU通过PCIE总线连接到一起（当然，未来会不会有更先进的其它总线或通信方式也不好说），GPU之间通信通过PCIe或者是NVLink来实现。CPU被称为主机端（host）而GPU端被称为设备端（device），这在后面的关键字中大家可以清晰的看出来。
在这里插入图片描述

上图中左侧为host端，右侧为device端，它们之通过PCIE总线相连。
CUDA存在着两种开发方式，一种是使用Runtime API 另外一种是使用 Driver API，两种 API的应用场景各有不同。不过由于这次安装在Windows上的为runtime API，所以会以 runetime API 为主进行开发，另外一种如果用到会在文中指出说明。
CUDA支持多种不同类型的显卡进行共同应用，是一种可扩展的编程模型.
在这里插入图片描述

CUDA核心是三个关键的抽象概念即线程组的层次结构、共享内存和屏障同步。
在CUDA编程模型中，重点要明白以下几部分：
1、Kernels
在CUDA模型中对C++进行了扩展，它允许开发者定义一些函数，这些函数允许在调用时被N个不同的CUDA线程并行执行。在CUDA中这个函数被称为Kernels，由关键字__global__定义实现。
2、线程的层次结构
在CUDA编程中线程的控制处理，经历了从Thread Block到Thread Block Clusters的过程，也就是说线程是划分成不同的块或块簇来处理的，与线程块中的线程在流式多处理器上保证共同调度的方式类似，块簇中的线程块也保证在GPU中的GPU处理集群（GPC）上共同调度。
无论是块或块簇，都会使用网格（Grid）来管理，可以划分为一维、二维和三维等情况，一个块中最多有1024个线程，块簇可由开发者自主定义块的数量，但在CUDA支持的一些环境中，它会有一定的限制数量。
在这里插入图片描述

3、内存的层次结构
CUDA内存层次结构相对于其它比较容易理解，特别是有内存库底层开发或相关经验的开发者来说，更容易理解。其实就是CUDA线程执行的过程中可以处理多种类型的内存，如私有内存、共享内存或全局内存，另外还提供了只读内存如常量内存空间和纹理内存（Texture memory）
在这里插入图片描述

三、主要的关键字

在CUDA编程中主要的关键字有：
1、global
此关键字用来定义核函数，由<<< >>>配置在执行核心的执行线程数，每个线程都有唯一一个线程ID并可在内核中进行访问。它的特点包括：
主机端或计算能力在大于等于3.2的设备端调用，在设备端执行；
其返回值必须为void类型且不能是类成员函数；
必须指定其执行配置即上面提到的<<<>>>;
函数的调用是异步的即其在设备执行完成前返回；
2、device
它被定义为只能由设备端调用并且在设备端执行的函数，需要注意的是，它不能与关键字__global__一起使用

3、host
这个关键字用来定义在主机上执行的函数，其特点为：
函数只能在主机端调用和执行；
它不可以和关键字__global__一起使用，但可以与__device__一起使用。

四、CUDA开发流程

在CUDA编程中，设备端的大部分程序代码可以独立于主机端的程序代码执行，当一个Kernel函数被调用后，控制权会返回到主机端的CPU控制。一般来说，CPU负责串行的工作而由GPU进行并行的操作。
CUDA的开发流程虽然因为不同的接口方式有所不同，但一般分以下几步：
1、分配CPU和GPU内存
2、复制数据从CPU端到GPU端
3、调用CUDA kernel来执行运算
4、运算结束将结果数据从GPU拷贝至CPU
5、释放CPU和GPU内存

五、CUDA的开发例程

看一个官网上的例程就明白了：

// Device code
__global__ void VecAdd(float* A, float* B, float* C, int N)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N)
        C[i] = A[i] + B[i];
}

// Host code
int main()
{
    int N = ...;
    size_t size = N * sizeof(float);

    // Allocate input vectors h_A and h_B in host memory
    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);
    float* h_C = (float*)malloc(size);

    // Initialize input vectors
    ...

    // Allocate vectors in device memory
    float* d_A;
    cudaMalloc(&d_A, size);
    float* d_B;
    cudaMalloc(&d_B, size);
    float* d_C;
    cudaMalloc(&d_C, size);

    // Copy vectors from host memory to device memory
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // Invoke kernel
    int threadsPerBlock = 256;
    int blocksPerGrid =
            (N + threadsPerBlock - 1) / threadsPerBlock;
    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // Copy result from device memory to host memory
    // h_C contains the result in host memory
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // Free device memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // Free host memory
    ...
}