Compute Unified Device Architecture

如何编译？使用 nvcc 将 cu 编译为 o 文件；（Q：链接器使用什么？）

学习资料：

1. 祖传代码：https://github.com/GZhonghui/cuda_samples
2. CUDA 编程入门（HPC Wiki） https://hpcwiki.io/gpu/cuda/
3. CUDA Tutorial https://cuda-tutorial.readthedocs.io/
4. xiaoxiong（关于CUDA的博客，很有帮助） https://littlebearsama.github.io/tags/CUDA/
5. 张也冬 https://zyddora.github.io/
6. GPU高性能编程CUDA实战
7. https://github.com/CodedK/CUDA-by-Example-source-code-for-the-book-s-examples-
8. https://github.com/ischintsan/cuda_by_example

1. GT：Graphics Technology
2. GTX：Giga Texel Shader eXtreme
3. RTX：引入了RT Core、Tensor Core

关于硬件的理解（细节处可能有错误）：原本GPU是为了绘制图形所设计的硬件，主要工作是计算vert和frag等着色器，主要也是进行向量矩阵等运算，后来NV把这部分计算硬件设计成了通用的，也就是CUDA中的SM/CUDA Core（这里补充一下，SM=流多处理器，SP=流处理器=CUDA Core，1个SM包含多个SP），原本的功能不变，同时也可以运行自定义的通用计算代码了。后来，因为深度学习的需要，NV加了Tensor Core，因为实时光线追踪的需要，NV加了RT Core，这两种都是单独的新硬件，（目前）都是进行其对应类型的计算，（应该）还不能做通用类型的计算，所以不在CUDA的范围内。

微架构和硬件有关，微架构决定Compute Capability（计算能力），计算能力写为8.0（不是具体这个数字，为了展示格式举个例子而已），有时候也叫80，sm_80。根据微架构（or计算能力），安装CUDA的时候要选择合适的版本，同样，根据CUDA的版本，安装驱动的时候，也要选择合适的驱动版本（一个驱动有最高支持的cuda版本，一个cuda有最低需要的驱动版本）

NV官方提供的API or samples，有些也需要对应的计算能力范围，比如sm50 - sm90

Grid（对应一个global函数，是逻辑概念） > SM（Stream Multiprocessor，SM是硬件概念。有时候也被叫做“GPU大核”）>= Block（1个SM上，运行1个或者多个Block，取决于Block的资源需求，所以不用担心，一个Block的Thread分配少了，导致SM的资源被浪费；SM资源可分，但是1个Block不能被分开在不同的SM上运行） > Thread Warp（线程束，参考内存管理的部分） > Thread

关于Nvidia GPU硬件架构和对应的CUDA编程模型：
CUDA是在Nvidia的GPU上编写并行程序的模型，在CUDA模型中，Nvidia的GPU由多个SM流式多处理器组成，SM包括了多个CUDA核心，CUDA核心是用于计算的最小单位，每一个核心上运行一个线程。
在启动并行程序的时候，我们会产生一个Grid，Grid包括多个SM的计算任务，但是我们不需要考虑运行的硬件到底有几个SM，GPU驱动会自动分配SM任务，如下图

程序是自动并行运算的，在编程层面，我们创建的概念是一个Grid，一个Grid包含多个Blocks（线程块，可指定数量），每一个Block包含多个Threads（计算的最小单位，单个线程，可指定数量），所有的Threads一起运行，如下图

那么我们就可以得到硬件和编程概念上的联系了：

1. GPU设备：对应Grid，Grid可能比较小，分配到部分GPU的资源，可能比较大，分配到全部GPU的资源并排队（串行）执行，分配的结果与计算任务的大小和GPU的计算资源相关，由驱动自动决策，不需要我们考虑。
2. SM流式多处理器：对应Block，一个SM承担一个Block的计算任务。
3. CUDA计算核心：对应Thread，一个计算核心运行一个Thread，多个计算核心同时运行实现并行运算。

注意：无论是Block还是Thread，并行执行，执行的先后顺序都是不确定的，有依赖关系的计算任务不能并行！

host指主机端（CPU），device指设备端（GPU），3种函数类型如下：

1. __host__: 这类函数与正常的函数没有区别。其只能被 host 上执行的函数（__host__）调用，并在 host 上执行。使用修饰符__host__修饰，在CPU端调用，也在CPU端执行，也就是最普通的函数，不加修饰符默认就被看作是Host函数
2. __global__: 这类函数可以被任何函数调用，并在 device 上执行。使用修饰符__global__修饰，在CPU端调用，在GPU端运行，是CPU给GPU下达命令、沟通交互的关键（CUDA Kernel，也就是核函数）
3. __device__: 这类函数只能被 device 上执行的函数（__device__ 或 __global__）调用，并在 device 上执行。使用修饰符__device__修饰，只能在GPU端调用（也就是在Global函数或者Device函数中调用），也是在GPU端运行

关于调用的方式：

1. 在启动 CUDA Kernel 时，<<<x,y>>> 中的第一个数字是每一个 Grid 中的 Block 数量，第二个数字是每一个 Block 中的 Thread 数量（每一个 Block 中含有的 Thread 数量是相等的）。那么，Grid的数量是？Grid应该只有1个，1个global函数对应1个Grid
2. Host函数和Device函数就是正常的调用方式，但是Global函数涉及到了CPU和GPU的交互，在调用的时候就是启动了一个并行计算的任务，所以需要指定GPU计算需要的尺寸，使用<<<,>>>（3层的尖括号）来指定Block的数量、每个Block中Thread的数量
3. 调用Global函数是异步的，也就是调用Global函数是不会阻塞主进程的，GPU的计算在后台并行进行，所以也需要一些同步手段

指定Grid和Block的维度：
Gird是包含N个Block的集合，Block是包含N个Thread的集合，这些集合可以有维度，比如Block包含N个Thread，这N个Thread可以是(N)，也可以是(2,N>>1)，也可以是(2,2,N>>2)，只要Thread的总数在范围之内就可以（每个Block内最多包含的Thread数量是一定的，这个数字是由硬件决定的）。另外，每个维度的范围也有一个上限，比如，通过 cudaDeviceProp 结构体中的 maxThreadsDim[] 数组，可以查询每个 Block 在各个维度中支持的最大线程数。
维度最多有3个，以Block举例来说，在只有1维的时候，每个Thread通过ID.x区分，2维的时候通过ID.x和ID.y区分。正确的选择维度可以提升性能。在调用global函数的时候，<<<x,y>>>内的x和y不一定必须是标量，也可以是CUDA提供的向量类型，比如dim3就是CUDA提供的3维向量（x,y,z），通过dim3的形式就可以指定Grid和Block的维度。

1. 环境配置
2. 内存管理
3. CUDA SDK
4. NVCC
5. 核函数
6. 设备属性
7. 修饰符
8. 同步
9. 调试

1. 内存管理.odt
2. 函数.odt
3. 数据类型.odt
4. 环境配置.odt
5. 硬件系列.odt
6. 线程束.odt
7. 编程模型.odt
8. 编译.odt
9. 设备属性.odt