使用 nvcc 编译器

CUDA 编程入门

逻辑模型

[!NOTE]
Host 为 CPU 端，Device 为 GPU 端

1. 分配 host 内存，并进行数据初始化；
2. 分配 device 内存，并从 host 将数据拷贝到 device 上；
3. 调用 CUDA 的核函数在 device 上完成指定的运算；
4. 将 device 上的运算结果拷贝到 host 上；
5. 释放 device 和 host 上分配的内存。

[!NOTE] 运行机制

• __global__：在 device 上执行，从 host 中调用（一些特定的 GPU 也可以从 device 上调用），返回类型必须是 void，不支持可变参数参数，不能成为类成员函数。注意用 __global__ 定义的 kernel 是异步的，这意味着 host 不会等待 kernel 执行完就执行下一步。
• __device__：在 device 上执行，单仅可以从 device 中调用，不可以和 __global__ 同时用。
  • 在 __global__ 和 __device__ 函数中调用
• __host__：在 host 上执行，仅可以从 host 上调用，一般省略不写，不可以和 __global__ 同时用，但可和 __device__，此时函数会在 device 和 host 都编译。

• 核函数：在 GPU 上并行执行的函数，__global__ 进行声明
  • 调用方式：kernel_fun<<< grid, block >>>(prams...);
• 一个 kernel 所启动的所有线程称为一个网格（grid）
  • 网格又可以分为很多线程块（block），一个线程块里面包含很多线程
• 这样每个线程都有对应的坐标 threadIdx，就可以进行并行运算
  • 比如一个二维数组运算，可以复制到 GPU 中，然后在线程中根据坐标来操纵二维数组的运算，这样每个线程只需要管自身的运算即可
  • 线程有内置变量，参考 shader
    • 可以获取 block 的 id 和 线程的 id

硬件模型

• SM 流式多处理器，是执行线程的物理层
  • 含有多个 CUDA 核心
  • 线程会分配给 SM 进行计算
  • 执行单元是线程束（warps），包含 32 个线程
    • Block 分配到 SM 会划分为多个线程束，然后并行执行
    • 同一个线程束执行同一个指令，但是每个线程有自己的寄存器状态，所以行为可能不同
  • SM 实际上的并发会收到资源分配影响，所以逻辑层的并行并不会完全进行
    • 技巧是 block 设置为 32 的倍数