【笔记】CUDA （一） - 介绍、架构、编程模型基础

• 文档：
  1. 官方 guide 👍 - https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
  2. 官方 api - https://docs.nvidia.com/cuda/cublas/index.html
  3. 官方 sample - https://docs.nvidia.com/cuda/cuda-samples/index.html#vector-addition
• 参考：
  1. CUDA编程入门极简教程 - https://zhuanlan.zhihu.com/p/34587739
     1. An Even Easier Introduction to CUDA - https://developer.nvidia.com/blog/even-easier-introduction-cuda/
     2. Unified Memory in CUDA 6 - https://developer.nvidia.com/blog/unified-memory-in-cuda-6/
     3. Maximizing Unified Memory Performance in CUDA - https://developer.nvidia.com/blog/maximizing-unified-memory-performance-cuda/
  2. 理解CUDA架构、编程与进阶使用 - https://blog.csdn.net/qq_37764141/article/details/122607968
  3. 扫地的小何尚 - CUDA编程
     1. CUDA编程手册中文版—CUDA简介 https://blog.csdn.net/kunhe0512/article/details/124120941
     2. CUDA 编程手册中文版—编程模型 https://blog.csdn.net/kunhe0512/article/details/124121001

GPU

GPU（Graphics Processing Unit）在相同的价格和功率范围内，比CPU提供更高的指令吞吐量和内存带宽。许多应用程序利用这些更高的能力，在GPU上比在CPU上运行得更快(参见 GPU应用程序)。其他计算设备，如FPGA，也非常节能，但提供的编程灵活性要比GPU少得多。

GPU 专门用于高度并行计算，因此设计时更多的晶体管用于数据处理，而不是数据缓存和流量控制。

将更多晶体管用于数据处理，例如浮点计算，有利于高度并行计算。
GPU可以通过计算隐藏内存访问延迟，而不是依靠大数据缓存和复杂的流控制来避免长时间的内存访问延迟，这两者在晶体管方面都是昂贵的。

CUDA 介绍

2006年，NVIDIA公司发布了CUDA。

CUDA（Compute Unified Device Architecture）是建立在NVIDIA的CPUs上的一个通用并行计算平台和编程模型。
（即显卡厂商 NVIDIA 推出的通用并行计算架构，该架构使 GPU 能够解决复杂的计算问题）

近年来，GPU最成功的一个应用就是深度学习领域，基于GPU的并行计算已经成为训练深度学习模型的标配。目前，最新的CUDA版本为CUDA 9。

CUDA 包含了 CUDA 指令集架构（ISA）以及 GPU 内部的并行计算引擎。 开发人员可以使用 C语言 来为 CUDA™架构编写程序，所编写出的程序可以在支持 CUDA™ 的处理器上以超高性能运行。CUDA3.0 已经开始支持 C++ 和 FORTRAN。

如下图所示，支持其他语言、应用程序编程接口或基于指令的方法，例如 FORTRAN、DirectCompute、OpenACC。

可扩展的编程模型

多核 CPU 和众核 GPU 的出现意味着主流处理器芯片现在是并行系统。挑战在于开发能够透明地扩展可并行的应用软件，来利用不断增加的处理器内核数量。就像 3D 图形应用程序透明地将其并行性扩展到具有广泛不同内核数量的多核 GPU 一样。
CUDA 并行编程模型旨在克服这一挑战，同时为熟悉 C 等标准编程语言的程序员保持较低的学习曲线。
.

其核心是三个关键抽象⭐️（它们只是作为最小的语言扩展集向程序员公开）

1. 线程组的层次结构
2. 共享内存
3. 屏障同步

这些抽象提供了细粒度的数据并行和线程并行，嵌套在粗粒度的数据并行和任务并行中。它们指导程序员将问题划分为可以由线程块并行独立解决的粗略子问题，并将每个子问题划分为可以由块内所有线程并行协作解决的更精细的部分。
这种分解通过允许线程在解决每个子问题时进行协作来保留语言表达能力，同时实现自动可扩展性。实际上，每个线程块都可以在 GPU 内的任何可用multiprocessor上以乱序、并发或顺序调度，以便编译的 CUDA 程序可以在任意数量的多处理器上执行，如下图所示，并且只有运行时系统需要知道物理multiprocessor个数。
这种可扩展的编程模型允许 GPU 架构通过简单地扩展multiprocessor和内存分区的数量来跨越广泛的市场范围：高性能发烧友 GeForce GPU ，专业的 Quadro 和 Tesla 计算产品 (有关所有支持 CUDA 的 GPU 的列表，请参阅 支持 CUDA 的 GPU）。

⚠️注意：GPU 是围绕一系列流式多处理器 (SM: Streaming Multiprocessors) 构建的（有关详细信息，请参 阅硬件实现）。 多线程程序被划分为彼此独立执行的线程块，因此具有更多multiprocessor的 GPU 将比具有更少多处理器的 GPU 在更短的时间内完成程序执行。

CUDA 架构

GPU并不是一个独立运行的计算平台，而需要与CPU协同工作，可以看成是CPU的协处理器，因此当我们在说GPU并行计算时，其实是指的基于CPU+GPU的异构计算架构。在异构计算架构中，GPU与CPU通过PCIe总线连接在一起来协同工作：

• CPU所在位置称为为主机端（host）
• GPU所在位置称为设备端（device）

如下图所示

可以看到：

• GPU包括更多的运算核心（ALU：算数逻辑单元），其特别适合数据并行的计算密集型任务，如大型矩阵运算；
• 而CPU的运算核心较少，但是其可以实现复杂的逻辑运算，因此其适合控制密集型任务

另外：

• CPU上的线程是重量级的，上下文切换开销大；
• 但是GPU由于存在很多核心，其线程是轻量级的。

因此，基于CPU+GPU的异构计算平台可以优势互补，CPU负责处理逻辑复杂的串行程序，而GPU重点处理数据密集型的并行计算程序，从而发挥最大功效。

CUDA是NVIDIA公司所开发的GPU编程模型，它提供了GPU编程的简易接口，基于CUDA编程可以构建基于GPU计算的应用程序。CUDA提供了对其它编程语言的支持，如C/C++，Python，Fortran等语言，这里我们选择CUDA C/C++接口对CUDA编程进行讲解。开发平台为Windows 10 + VS 2013，Windows系统下的CUDA安装教程可以参考这里。

CUDA 编程模型基础

在给出CUDA的编程实例之前，这里先对CUDA编程模型中的一些概念及基础知识做个简单介绍。

# host、device：

CUDA编程模型是一个异构模型，需要CPU和GPU协同工作。在CUDA中，host和device是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。CUDA程序中既包含host程序，又包含device程序，它们分别在CPU和GPU上运行。同时，host与device之间可以进行通信，这样它们之间可以进行数据拷贝。

典型的CUDA程序的执行流程如下：

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

# kernel：

上面流程中最重要的一个过程是调用CUDA的核（kernel）函数来执行并行计算。

• kernel 是CUDA中一个重要的概念。
  （CUDA C++ 通过允许程序员定义称为 kernel 的 C++ 函数来扩展 C++，当调用内核时，由 N 个不同的 CUDA 线程并行执行 N 次，而不是像常规 C++ 函数那样只执行一次。）
• kernel是在device上线程中并行执行的函数，核函数用 __global__ 符号声明，在调用时需要用 <<<grid, block>>> 来指定 kernel 要执行的线程数量。
  （请参阅 C++ 语言扩展）
  每个执行内核的线程都有一个唯一的线程 ID，可以通过内置变量在内核中访问。
• 在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID，这个ID值可以通过核函数的内置变量threadIdx来获得。

由于GPU实际上是异构模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词开区别host和device上的函数，主要的三个函数类型限定词如下：

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

要深刻理解 kernel，必须要对 kernel 的线程层次结构有一个清晰的认识。

# Guid、Block：（1-dim、2-dim、3-dim 结构）

首先，GPU上很多并行化的轻量级线程。kernel在device上执行时实际上是启动很多线程，一个 kernel 所启动的所有线程称为一个 网格（grid） ，同一个网格上的线程共享相同的全局内存空间；
grid是线程结构的第一层次，而网格又可以分为很多 线程块（block），一个线程块里面包含很多 线程，这是第二个层次。

线程两层组织结构如上图所示，这是一个 gird 和 block 均为2-dim（dimension）的线程组织。

grid 和 block 都是定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量；在定义时，缺省值初始化为1。
因此 grid 和 block 可以灵活地定义为1-dim，2-dim以及3-dim结构，对于图中结构（主要水平方向为x轴），定义的 grid 和 block 如下所示，kernel 在调用时也必须通过 执行配置<<<grid, block>>> 来指定 kernel 所使用的线程数及结构。

dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<< grid, block >>>(prams...);

所以，一个线程需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置

如图中的 Thread (1,1)：

threadIdx.x = 1
threadIdx.y = 1
blockIdx.x = 1
blockIdx.y = 1

# 全局ID：（blockDim、gridDim 变量）

一个线程块上的线程是放在同一个流式多处理器（SM)上的，但是单个SM的资源有限，这导致线程块中的线程数是有限制的，现代GPUs的线程块可支持的线程数可达1024个。

有时候，我们要知道一个线程在 block 中的全局ID，此时就必须还要知道block的组织结构（1-dim？2-dim？3-dim？），这是通过线程的内置变量 blockDim 来获得。它获取线程块各个维度的大小。对于一个2-dim的 block $(D_x,D_y)$ ，线程 $(x,y)$ 的ID值为 $(x+y\ast D_x)$ ，如果是3-dim的block ( D