简介
CPU 和 GPU - 异构计算的演进与发展世界上大多数事物的发展规律是相似的，在最开始往往都会出现相对通用的方案解决绝大多数的问题，随后会出现为某一场景专门设计的解决方案，这些解决方案不能解决通用的问题，但是在某些具体的领域会有极其出色的表现。
GPU 的架构；内存管理；任务管理；数据类型。
内存将数据传输到CPU大概每秒大概传输200GB（也就是每秒25G FP64），cpu 计算能力大概是2000 GFLOPs FP64，这两者的比值就是设备的计算强度。也就是cpu 每秒不对一个数据处理80次，cpu 就会空闲，但也没有什么算法需要对一个数据处理80次。当增加FLOPs速度比增加内存带宽的速度快的时候，计算强度就会上升。为什么？物理上光速 3亿米/秒，电脑时钟30亿Hz，在一个时钟周期，光只传播了10cm，电流在硅中的传播速度只有光的五分之一（6万公里/秒），物理实际上很复杂，根据经验，一个时钟周期内，电流的移动只有20mm。 CPU 的期望是一个线程基本完成所有工作，将这些线程从一个切换到另一个是非常昂贵的（上下文切换），cpu设计者把所有资源都投入到延迟上了。GPU 设计师将所有资源的都投入到增加线程中，而不是减少延迟。此外，gpu 使用寄存器缓存来解决高延迟问题，以及通过靠近数据来减少延迟（内存传输数据慢，cpu 没有忙起来，内存也没有忙起来）。gpu 选择增加线程，每个sm2048个线程，一次跑一批（warp）线程，当一些线程因为等待延迟关闭时，其它线程大概已经load好数据了准备运行了。这就是gpu工作的秘密，它可以在不同的warp之间切换，并且在一个时钟周期内完成，所以根本没有上下文开销。gpu 是一个吞吐量系统，总是超量分配线程，超量分配意味着总是在内存快的（指的是数据ready？）时候工作。（a grid represents all work to be done）所有要做的工作都被分解成线程块（the grid comprises many blocks with an equal number of threads），每个块都有并行线程，保证线程同时运行，这样它们就可以共享数据（threads within a block run independently but may synchrionze to exchange data）（机器学习有一些计算，比如all-to-all）。但所有的块都是在超量分配模式下独立调度的，这样才能两全其美。但它也允许一定数量的线程相互交互，这就是gpu编程的本质。PS：是不是可以认为，内存延迟高，这时为了cpu和内存提速（都别闲着），做大带宽，比如虽然10s传一次数据，但一次传10M，gpu 准备大量线程同时干活儿。后续有数据就干活儿，没干活儿就让位给另一批数据ready的线程。
GPU
各种游戏里面的人物的脸，并不是那个相机或者摄像头拍出来的，而是通过多边形建模（Polygon Modeling）创建出来的。而实际这些人物在画面里面的移动、动作，乃至根据光线发生的变化，都是通过计算机根据图形学的各种计算，实时渲染出来的。
图像进行实时渲染的过程，可以被分解成下面这样 5 个步骤：
- 顶点处理（Vertex Processing）。构成多边形建模的每一个多边形呢，都有多个顶点（Vertex）。这些顶点都有一个在三维空间里的坐标。但是我们的屏幕是二维的，所以在确定当前视角的时候，我们需要把这些顶点在三维空间里面的位置，转化到屏幕这个二维空间里面。这个转换的操作，就被叫作顶点处理。这样的转化都是通过线性代数的计算来进行的。可以想见，我们的建模越精细，需要转换的顶点数量就越多，计算量就越大。而且，这里面每一个顶点位置的转换，互相之间没有依赖，是可以并行独立计算的。
- 图元处理。把顶点处理完成之后的各个顶点连起来，变成多边形。其实转化后的顶点，仍然是在一个三维空间里，只是第三维的 Z 轴，是正对屏幕的“深度”。所以我们针对这些多边形，需要做一个操作，叫剔除和裁剪（Cull and Clip），也就是把不在屏幕里面，或者一部分不在屏幕里面的内容给去掉，减少接下来流程的工作量。
- 栅格化。我们的屏幕分辨率是有限的。它一般是通过一个个“像素（Pixel）”来显示出内容的。对于做完图元处理的多边形，把它们转换成屏幕里面的一个个像素点。每一个图元都可以并行独立地栅格化。
- 片段处理。在栅格化变成了像素点之后，我们的图还是“黑白”的。我们还需要计算每一个像素的颜色、透明度等信息，给像素点上色。
- 像素操作。把不同的多边形的像素点“混合（Blending）”到一起。可能前面的多边形可能是半透明的，那么前后的颜色就要混合在一起变成一个新的颜色；或者前面的多边形遮挡住了后面的多边形，那么我们只要显示前面多边形的颜色就好了。最终，输出到显示设备。
经过这完整的 5 个步骤之后，完成了从三维空间里的数据的渲染，变成屏幕上你可以看到的 3D 动画了。称之为图形流水线（Graphic Pipeline）。这个过程包含大量的矩阵计算，刚好利用了GPU的并行性。
现代 CPU 里的晶体管变得越来越多，越来越复杂，其实已经不是用来实现“计算”这个核心功能，而是拿来实现处理乱序执行、进行分支预测，以及高速缓存部分。而在 GPU 里，这些电路就显得有点多余了，GPU 的整个处理过程是一个流式处理（Stream Processing）的过程。因为没有那么多分支条件，或者复杂的依赖关系，我们可以把 GPU 里这些对应的电路都可以去掉，做一次小小的瘦身，只留下取指令、指令译码、ALU 以及执行这些计算需要的寄存器和缓存就好了。
于是，我们就可以在一个 GPU 里面，塞很多个这样并行的 GPU 电路来实现计算，就好像 CPU 里面的多核 CPU 一样。和 CPU 不同的是，我们不需要单独去实现什么多线程的计算。因为 GPU 的运算是天然并行的。无论是对多边形里的顶点进行处理，还是屏幕里面的每一个像素进行处理，每个点的计算都是独立的。
一方面，GPU 是一个可以进行“通用计算”的框架，我们可以通过编程，在 GPU 上实现不同的算法。另一方面，现在的深度学习计算，都是超大的向量和矩阵，海量的训练样本的计算。整个计算过程中，没有复杂的逻辑和分支，非常适合 GPU 这样并行、计算能力强的架构。
为什么深度学习需要使用GPU
为什么深度学习需要使用GPU？相比cpu，gpu
- gpu核心很多，比如CPU来讲它多少core呢？我看过前几天发布会有至强6的E系列最高288core的，那GPU呢，上一代H100是1万8。
- gpu内存带宽更高，速度快就贵，所以显存容量一般不大。因为 CPU 首先得取得数据， 才能进行运算， 所以很多时候，限制我们程序运行速度的并非是 CPU 核的处理速度， 而是数据访问的速度。
- 控制流，cpu 控制流很强，alu 只占cpu的一小部分。gpu 则要少用控制语句。CPU需要很强的通用性来处理各种不同的数据类型，同时又要逻辑判断又会引入大量的分支跳转和中断的处理。这些都使得CPU 里的晶体管变得越来越多，越来越复杂，其实已经不是用来实现“计算”这个核心功能，而是拿来实现处理乱序执行、进行分支预测，以及高速缓存。GPU 专门用于高度并行计算，面对的则是类型高度统一的、相互无依赖的大规模数据和不需要被打断的纯净的计算环境。因此设计时更多的晶体管用于数据处理，而不是数据缓存和流量控制。GPU 只有 取指令、指令译码、ALU 以及执行这些计算需要的寄存器和缓存。CPU 上不同线程一般是执行不同任务，GPU同一个block的线程执行的则是相同的kernel函数。PS： 将更多晶体管用于数据处理，例如浮点计算，有利于高度并行计算。我们一般习惯将cpu的控制单元和计算单元视为一个整体，而gpu 一般会独立看待控制单元和计算单元，所以觉得它们差别很大。
- 编程，cpu 是各种编程语言，编译器成熟。
如图所示，CPU在芯片领域中主要用于降低指令时延的功能，例如大型缓存、较少的算术逻辑单元（ALU）和更多的控制单元。与此相比，GPU则利用大量的ALU来最大化计算能力和吞吐量，只使用极小的芯片面积用于缓存和控制单元，这些元件主要用于减少CPU时延。
CPU / GPU原理与 CUDAGPU 一开始是没有“可编程”能力的，程序员们只能够通过配置来设计需要用到的图形渲染效果（图形加速卡）。在游戏领域, 3D 人物的建模都是用一个个小三角形拼接上的, 而不是以像素的形式, 对多个小三角形的操作, 能使人物做出多种多样的动作, 而 GPU 在此处就是用来计算三角形平移, 旋转之后的位置。为了提高游戏的分辨率, 程序会将每个小三角形细分为更小的三角形，每个小三角形包含两个属性, 它的位置和它的纹理。在游戏领域应用的 GPU 与科学计算领域的 GPU 使用的不同是, 当通过 CUDA 调用 GPU 来进行科学计算的时候, 计算结果需要返回给 CPU, 但是如果用 GPU 用作玩游戏的话, GPU 的计算结果直接输出到显示器上, 也就不需要再返回到 CPU。
深度学习的模型训练，指的是利用数据通过计算梯度下降的方式迭代地去优化神经网络的参数，最终输出网络模型的过程。在这个过程中，通常在迭代计算的环节，会借助 GPU 进行计算的加速。
GPU 架构
- GPU的core不能做任何类似out-of-order exectutions那样复杂的事情，总的来说，GPU的core只能做一些最简单的浮点运算,例如 multiply-add(MAD)或者 fused multiply-add(FMA)指令，后来经过发展又增加了一些复杂运算，例如tensor张量(tensor core)或者光线追踪(ray tracing core)相关的操作。
- GPU的编程方式是SIMD(Single Instruction Multiple Data)意味着所有Core的计算操作完全是在相同的时间内进行的，但是输入的数据有所不同。如果这个工作给到CPU来做，需要N的时间才可以做完，但是给到GPU只需要一个时钟周期就可以完成。
- 多个core之间通讯：在图像缩放的例子中，core与core之间不需要任何协作，因为他们的任务是完全独立的。然而，GPU解决的问题不一定这么简单，假设一个长度为8的数组，在第一步中完全可以并行执行两个元素和两个元素的求和，从而同时获得四个元素，两两相加的结果，以此类推，通过并行的方式加速数组求和的运算速度。如果是长度为8的数组两两并行求和计算，那么只需要三次就可以计算出结果。如果是顺序计算需要8次。如果GPU想要完成上述的推理计算过程，显然，多个core之间要可以共享一段内存空间以此来完成数据之间的交互，需要多个core可以在共享的内存空间中完成读/写的操作。我们希望每个Cores都有交互数据的能力，但是不幸的是，一个GPU里面可以包含数以千计的core，如果使得这些core都可以访问共享的内存段是非常困难和昂贵的。出于成本的考虑，折中的解决方案是将各类GPU的core分类为多个组，形成多个流处理器(Streaming Multiprocessors )或者简称为SMs。
- SM块的底部有一个96KB的L1 Cache/SRAM。每个SM都有自己的L1缓存，SM间不能互相访问彼此的L1。L1 CACHE拥有两个功能，一个是用于SM上Core之间相互共享内存（寄存器 也可以），另一个则是普通的cache功能。存在全局的内存GMEM，但是访问较慢，Cores当需要访问GMEM的时候会首先访问L1,L2如果都miss了，那么才会花费大代价到GMEM中寻找数据。
和CPU对比
CPU和GPU的主要区别在于它们的设计目标。CPU的设计初衷是执行顺序指令，一直以来，为提高顺序执行性能，CPU设计中引入了许多功能。其重点在于减少指令执行时延，使CPU能够尽可能快地执行一系列指令。这些功能包括指令流水线、乱序执行、预测执行和多级缓存等（此处仅列举部分）。而GPU则专为大规模并行和高吞吐量而设计，但这种设计导致了中等至高程度的指令时延。这一设计方向受其在视频游戏、图形处理、数值计算以及现如今的深度学习中的广泛应用所影响，所有这些应用都需要以极高的速度执行大量线性代数和数值计算，因此人们倾注了大量精力以提升这些设备的吞吐量。我们来思考一个具体的例子：由于指令时延较低，CPU在执行两个数字相加的操作时比GPU更快。在按顺序执行多个这样的计算时，CPU能够比GPU更快地完成。然而，当需要进行数百万甚至数十亿次这样的计算时，由于GPU具有强大的大规模并行能力，它将比CPU更快地完成这些计算任务。
计算架构
GPU架构总体如下图所示：
两级线程层次结构（带上grid也有说三层的，比较新的Hooper 架构 引入了Thread Block Clusters 层次），可以分为两个粒度来看 GPU：
- 以SM 为基本单元来看GPU 整体架构，GPU由多个SM组成，而在SM之外，仅仅有global memory和L2 cache两个组件。PS：gpu sm 更类似于cpu 里的core，不同sm执行不同的指令单元
- SM的硬件架构：核心组件包括内存、计算单元和指令调度。每个SM包含多个核心（在 Fermi 架构之前，处理核心被称为 Stream Processor，每个 SP 可以执行一个线程的计算任务，在 Fermi 架构之后，英伟达将处理核心更名为 CUDA 核心），它们共享一个指令单元，但能够并行执行不同的线程。每个SM中的共享内存允许线程之间进行有效的数据交换和同步。 在Fermi 架构中，每个 SM 包含 2 个线程束（Warp），一个 Warp 中包含 16 个 Cuda Core，共 32 个 CUDA Cores。随着 Volta 架构的推出，V100 GPU 每个SM配备了 8 个 Tensor Core。
流式多处理器（Streaming Multiprocessor、SM）是 GPU 的基本单元，每个 GPU 都由一组 SM 构成，SM 中最重要的结构就是计算核心 Core
- 线程调度器（Warp Scheduler）：线程束（Warp）是最基本的单元，每个线程束中包含 32 个并行的线程，GPU 控制部件面积比较小，为了节约控制器，一个 Warp 内部的所有 CUDA Core 的 PC（程序计数器）一直是同步的，但是访存地址是可以不同的，每个核心还可以有自己独立的寄存器组，它们使用不同的数据执行相同的命令，这种执行方式叫做 SIMT（Single Instruction Multi Trhead）。调度器会负责这些线程的调度；
-
一个 Warp 中永远都在执行相同的指令，如果分支了怎么处理呢？其实 Warp 中的 CUDA Core 并不是真的永远都执行相同的指令，它还可以不执行。这样会导致 Warp Divergence，极端情况下，每一个Core的指令流都不一样，那么甚至还可能导致一个 Warp 中仅有一个 Core 在工作，效率降低为 1/32.
- GPU 需要数据高度对齐，一个 Warp 的内存访问是成组的，一次只能读取连续的且对齐的 128byte（正好是WarpSize 32 * 4 byte），CPU 是一个核心一个 L1，GPU 是两个 Warp 一个 L1 Cache，整个Warp 有一个核心数据没准备好都执行不了。
- GPU 的线程切换不同于 CPU，在 CPU 上切换线程需要保存现场，将所有寄存器都存到主存中，GPU 的线程切换只是切换了寄存器组（一个 SM 中有高达 64k 个寄存器），延迟超级低，几乎没有成本。一个 CUDA Core 可以随时在八个线程之间反复横跳，哪个线程数据准备好了就执行哪个。 这是 GPU 优于 CPU 的地方，也是为了掩盖延迟没办法的事情。
-
- CUDA Core：向量运行单元 ，在Fermi 架构中，每一个 Cuda Core 由 1 个浮点数单元 FPU 和 1 个逻辑运算单元 ALU 组成。
- Tensor Core：张量运算单元（FP8、FP16、BF16、TF32、INT8、INT4），2017 年提出的 Volta 架构，引入了张量核 Tensor Core 模块，一种专为 AI 训练和推理设计的可编程矩阵乘法和累加单元。TensorCore及其相关的数据路径是定制的，以显著提高浮点计算吞吐量。每个TensorCore提供一个4x4x4矩阵处理数组，它执行操作D=A*B+C，其中A、B、C和D是4×4矩阵。每个TensorCore每个时钟周期可以执行64个浮点FMA混合精度操作，而在一个SM中有8个TensorCore，所以一个SM中每个时钟可以执行1024(8x64x2)个浮点操作。 Tensor Core
- 特殊函数的计算单元（Special Functions Unit、SPU），（超越函数和数学函数,反平方根、正余弦啥的）
- Dispatch Unit：指令分发单元
与个人电脑上的 GPU 不同，数据中心中的 GPU 往往都会用来执行高性能计算和 AI 模型的训练任务。正是因为社区有了类似的需求，Nvidia 才会在 GPU 中加入张量（标量是0阶张量，向量是一阶张量， 矩阵是二阶张量）核心（Tensor Core）18专门处理相关的任务。张量核心与普通的 CUDA 核心其实有很大的区别，CUDA 核心在每个时钟周期都可以准确的执行一次整数或者浮点数的运算，时钟的速度和核心的数量都会影响整体性能。张量核心通过牺牲一定的精度可以在每个时钟计算执行一次 4 x 4 的矩阵运算。PS：就像ALU 只需要加法器就行了（乘法指令转换为多个加法指令），但为了提高性能，直接做了一个乘法器和加法器并存。
内存架构
与线程层次对应的是显存层次，不同层次的线程可以访问不同层次的显存。
- Multi level Cache：多级缓存（L0/L1 Instruction Cache、L1 Data Cache & Shared Memory）。GPU的特点是有很多的ALU和很少的cache. 缓存的目的不是保存后面需要访问的数据的，这点和CPU不同，而是为thread提高服务的。如果有很多线程需要访问同一个相同的数据，缓存会合并这些访问，然后再去访问dram（因为需要访问的数据保存在dram中而不是cache里面），获取数据后cache会转发这个数据给对应的线程，这个时候是数据转发的角色。但是由于需要访问dram，自然会带来延时的问题。GPU的控制单元（左边黄色区域块）可以把多个的访问合并成少的访问。
- 每个SM还有一个L1缓存，缓存从L2缓存中频繁访问的数据
- 所有SM都共享一个L2缓存，缓存全局内存中被频繁访问的数据，以降低时延。需要注意的是，L1和L2缓存对于SM来说是公开的，也就是说，SM并不知道它是从L1还是L2中获取数据。SM从全局内存中获取数据，这类似于CPU中L1/L2/L3缓存的工作方式。
- 存储和缓存数据的寄存器文件（Register File）。每个SM有大量的寄存器，被SM内的核心（Core）之间共享。
- 常量内存 (Constants Caches)：：用于SM上执行的代码中使用的常量数据, Constant 声明的变量就会在这里存。仅可由 CPU 写入，但可被所有 GPU 线程读取。适合存储小规模的、不变的数据（如配置信息、系数等）。
__constant__ float constData[256]; // 常量内存 - 访问存储单元（Load/Store Queues）：在核心和内存之间快速传输数据；
- 共享内存（Shared Memory）。每个SM有一块共享内存，SRAM内存，供运行在SM上的线程块共享使用。
__shared__ float sharedA[TILE_SIZE][TILE_SIZE]; // 共享内存 -
Global memory（也就是常说的显存）:我们用nvidia-smi命令得到的就是显存的大小，也叫全局内存，or 片外全局内存，存取的时延比较高
// 通常而言，全局内存主要适用于存储程序的大部分输入输出数据，尤其是需要 GPU 和 CPU 共享的大容量数据。 // 示例：在矩阵乘法中，两个矩阵的元素可以存储在全局内存中，以便所有线程都可以访问。 __global__ void matrixMultiplication(float *A, float *B, float *C, int N) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; float sum = 0.0; // 本地内存（Local Memory） for (int i = 0; i < N; ++i) { sum += A[row * N + i] * B[i * N + col]; } C[row * N + col] = sum; }
不同层次的显存访问延迟不同，Ampere 架构的 GPU一些重要的运算延迟周期时间
- 访问全局内存（高达80GB）：约380个周期
- 二级缓存（L2 cache）：约200个周期
- 一级缓存或访问共享内存（每个流式多处理器最多128KB）：约34个周期
- 乘法和加法在指令集层面的结合（fused multiplication and addition，FFMA）：4个周期
- Tensor Core（张量计算核心）矩阵乘法运算：1个周期
由于不同的存储器访问延迟差距较大，如果我们在编程的时候可以利用片内存储器降低访问延迟，就可以提升 Kernel 的性能。庆幸的是，在 GPU 编程中，CUDA 为 Shared Memory 提供编程接口，这使得开发者在设计 Kernel 实现时，可以利用 Shared Memory 访问延迟低的特点加速 Kernel 的性能。所以在 GPU 编程中，Kernel 的设计是以 Thread Block 这个粒度展开的。但这样会导致两个问题：
- 单个 Thread Block 处理的数据规模有限，原因是 Shared Memory 的容量有限。
- SM 利用率较低。单个 Thread Block 可配置的最大线程数为 1024，每个 Thread Block 会分配到一个 SM 上运行。假如每个 Thread Block 处理较大规模的数据、计算，Kernel 一次仅发射很少的 Thread Block，可能导致某些 SM 处于空闲状态，计算资源没有被充分挖掘，这样同样会限制 Kernel 的整体性能。例如在 LLM 长文本推理 进行 Decoding Attention时， 𝐾、𝑉 长度较长，此时由于显存上限问题， batch size 会小，这导致单个 Thread Block 访问的数据量、计算量较大，同时发射的 Thread Block 的数量较少，导致某些 SM 处于空闲状态，限制 Kernel 性能。 按 Thread Block 这个粒度划分子任务已经难以处理一些场景，限制了 Kernel 运行效率。解决这个问题的最直接的方式是：提供更大粒度的线程组Thread Block Clusters。 Hopper 架构特性：Distributed Shared Memory
CPU 与GPU
CPU 与GPU 协作
GPU 无法自己独立工作，其工作任务还是由 CPU 进行触发的。整体的工作流程可以看做是 CPU 将需要执行的计算任务异步的交给 GPU，GPU 拿到任务后，会将 Kernel 调度到相应的 SM 上，而 SM 内部的线程则会按照任务的描述进行执行。
大多数采用的还是分离式结构，AMD 的 APU 采用耦合式结构，目前主要使用在游戏主机中，如 PS4。
- 锁页：GPU 可以直接访问 CPU的内存。出于某些显而易见的原因，cpu 和gpu 最擅长访问自己的内存，但gpu 可以通过DMA 来访问cpu 中的锁页内存。锁页是操作系统常用的操作，可以使硬件外设直接访问内存，从而避免过多的复制操作。”被锁定“的页面被os标记为不可被os 换出的，所以设备驱动程序在给这些外设编程时，可以使用页面的物理地址直接访问内存。PS：部分内存的使用权暂时移交给设备。
- 命令缓冲区：CPU 通过 CUDA 驱动写入指令，GPU 从缓冲区 读取命令并控制其执行，
- CPU 与GPU 同步：cpu 如何跟踪GPU 的进度
对于一般的外设来说，驱动程序提供几个api接口，约定好输入和输出的内存地址，向输入地址写数据，调接口，等中断，从输出地址拿数据。输出数据地址 command_operation(输入数据地址)。gpu 是可以编程的，变成了输出数据地址 command_operation(指令序列,输入数据地址)
系统的三个要素: CPU，内存，设备。CPU 虚拟化由 VT-x/SVM 解决，内存虚拟化由 EPT/NPT 解决，设备虚拟化呢？它的情况要复杂的多，不管是 VirtIO，还是 VT-d，都不能彻底解决设备虚拟化的问题。除了这种完整的系统虚拟化，还有一种也往往被称作「虚拟化」的方式: 从 OS 级别，把一系列的 library 和 process 捆绑在一个环境中，但所有的环境共享同一个 OS Kernel。
不考虑嵌入式平台的话，那么，GPU 首先是一个 PCIe 设备。GPU 的虚拟化，还是要首先从 PCIe 设备虚拟化角度来考虑。一个 PCIe 设备，有什么资源？有什么能力？
- 2 种资源: 配置空间；MMIO(Memory-Mapped I/O)
- 2 种能力: 中断能力；DMA 能力
一个典型的 GPU 设备的工作流程是:
- 应用层调用 GPU 支持的某个 API，如 OpenGL 或 CUDA
- OpenGL 或 CUDA 库，通过 UMD (User Mode Driver)，提交 workload 到 KMD (Kernel Mode Driver)
- Kernel Mode Driver 写 CSR MMIO，把它提交给 GPU 硬件
- GPU 硬件开始工作… 完成后，DMA 到内存，发出中断给 CPU
- CPU 找到中断处理程序 —— Kernel Mode Driver 此前向 OS Kernel 注册过的 —— 调用它
- 中断处理程序找到是哪个 workload 被执行完毕了，…最终驱动唤醒相关的应用
本质上GPU 还是一个外设，有驱动程序（分为用户态和内核态）和API，用户程序 ==> API ==> CPU ==> 驱动程序 ==> GPU ==> 中断 ==> CPU.
Grid—> Block—>threads
CUDA里另外一个不次于kernel的概念就是三级线程管理：Grid—> Block—>threads。
- grid：kernel 在 device上跑，实际上启动一大堆线程，一个 kernel 所启动的所有线程称为一个Grid，一个Grid的所有线程是共享一大段内存，也就是相同的全局内存(显存)空间。
- Grid再分下去就是block层级，block里面才是装的thread，也就是线程。虽然一个Grid里面的所有线程，都是共享全局显存地址空间，但是，block之间都是隔离的，自己玩自己的，并行执行（注意，不是并发），每个 block自己的共享内存（Shared Memory），里面的Thread 共享，别的block的thread不能来访问。
- block 内部的 threads，怎么玩都可以了，可以同步，也可以通过 shared memory通信。
Grids是跑在Device（GPU）层级，block相当于包工头对应SM层级，而真正的干活的CUDA core呢，对应的就是拿到现成thread去执行了。
GPU架构与计算入门指南CUDA是NVIDIA提供的编程接口，用于编写运行在其GPU上的程序。
- 在CUDA中，你会以类似于C/C++函数的形式来表达想要在GPU上运行的计算，这个函数被称为kernel。
- kernel在并行中操作向量形式的数据，这些数字以函数参数的形式提供给它。 一个简单的例子是执行向量加法的kernel，即接受两个向量作为输入，逐元素相加，并将结果写入第三个向量。
要在GPU上执行kernel，我们需要启用多个线程，这些线程总体上被称为一个网格（grid），但网格还具有更多的结构。一个网格由一个或多个线程块（有时简称为块）组成，而每个线程块又由一个或多个线程组成。
CUDA为什么要分线程块和线程网格？
- grib,block,thread这套编程模型对应着硬件结构的抽象，以适应不同的硬件结构。GPU相对硬件结构简单，以堆砌达到预期性能。所以堆砌核心之间和核心之内的内存机制、指令发射等等都是不一样的。
- 这套编程模型强迫你优化线程通信。越快的存储越贵，存储空间也就越小，越不可能全局共享。共享内存是一种专门供单个CUDA线程块内的线程进行通信的内存，与全局内存相比，优势在于可以显著提高纯线程间的通信速度。但共享内存中的数据不能被主机直接访问，必须通过内核函数将其复制到全局内存。
线程块和线程的数量取决于数据的大小和我们所需的并行度。例如，在向量相加的示例中，如果我们要对256维的向量进行相加运算，那么可以配置一个包含256个线程的单个线程块，这样每个线程就可以处理向量的一个元素。如果数据更大，GPU上也许没有足够的线程可用，这时我们可能需要每个线程能够处理多个数据点。
编写一个kernel需要两步。第一步是运行在CPU上的主机代码，这部分代码用于加载数据，为GPU分配内存，并使用配置的线程网格启动kernel；第二步是编写在GPU上执行的设备（GPU）代码。在GPU上执行Kernel的步骤
- 将数据从主机复制到设备。 在调度执行kernel之前，必须将其所需的全部数据从主机（即CPU）内存复制到GPU的全局内存（即设备内存）。在最新的GPU硬件中，我们还可以使用统一虚拟内存直接从主机内存中读取数据。
- SM上线程块的调度。当GPU的内存中拥有全部所需的数据后，它会将线程块分配给SM。同一个块内的所有线程将同时由同一个SM进行处理。为此，GPU必须在开始执行线程之前在SM上为这些线程预留资源。在实际操作中，可以将多个线程块分配给同一个SM以实现并行执行。由于SM的数量有限，而大型kernel可能包含大量线程块，因此并非所有线程块都可以立即分配执行。GPU会维护一个待分配和执行的线程块列表，当有任何一个线程块执行完成时，GPU会从该列表中选择一个线程块执行。
- 单指令多线程 (SIMT) 和线程束（Warp）。众所周知，一个块（block）中的所有线程都会被分配到同一个SM上。但在此之后，线程还会进一步划分为大小为32的组（称为warp），并一起分配到一个称为处理块（processing block）的核心集合上进行执行。SM通过获取并向所有线程发出相同的指令，以同时执行warp中的所有线程。然后这些线程将在数据的不同部分，同时执行该指令。在向量相加的示例中，一个warp中的所有线程可能都在执行相加指令，但它们会在向量的不同索引上进行操作。由于多个线程同时执行相同的指令，这种warp的执行模型也称为单指令多线程 （SIMT）。这类似于CPU中的单指令多数据（SIMD）指令。 PS： 有点类似hadoop 任务调度的意思，cpu 是driver 进程，gpu是 worker 进程。 SIMT编程模型由Thread和Block组成，block偏软件概念（一个block多少个thread 可配），warp偏硬件概念（32个thread）。
CUDA编程体系和GPU硬件的关系
warp（gpu的一个单位）是典型的单指令多线程（SIMT，SIMD单指令多数据的升级）的实现，也就是32个线程同时执行的指令是一模一样的，只是线程数据不一样，这样的好处就是一个warp只需要一个套逻辑对指令进行解码和执行就可以了，芯片可以做的更小更快，之所以可以这么做是由于GPU需要处理的任务是天然并行的。
CUDA有host和device的概念，在 CUDA程序构架中，Host 代码部分在CPU上执行，就是一般的C。当遇到程序要进行并行处理的，CUDA就会将程序编译成GPU能执行的程序，并传送到GPU，这个被编译的程序在CUDA里称做核（kernel），Device 代码部分在 GPU上执行。kernel是CUDA中的核心概念之一。CUDA执行时最重要的一个流程是调用CUDA的核函数来执行并行计算。Kernel 程序通常以一种拷贝和计算（copy and compute）模式执行，即，首先从全局内存中获取数据，并将数据存储到共享内存中，然后对共享内存数据执行计算，并将结果（如果有）写回全局内存。
| CUDA编程抽象 | 线程层次结构 | 存储层次结构 | 缓存 |
|---|---|---|---|
| 线程/thread | ALU/cuda core | Local Memory。每个线程都有本地内存，存储临时变量。 | |
| 线程块/thread block | 流多处理器 (SM) | 共享内存/Shared Memory。同一个Block内的线程可以用共享内存共享数据。 | 每个流式多处理器（SM）都有自己的L1 Cache |
| 网格/grid | GPU device | 全局内存/Global Memory。可以被所有块上的所有线程访问 | 所有 SM 共享L2 Cache |
不同层次的显存访问延迟不同，以 PCIE 80GB 的 H800为例，其 Global Memory 的访问延迟约为 478 个时钟周期，Shared Memory 的访问延迟约为 30 个时钟周期，Register 约为 1 个时钟周期。由于不同的存储器访问延迟差距较大，如果我们在编程的时候可以利用片内存储器降低访问延迟，就可以提升 Kernel 的性能。庆幸的是，在 GPU 编程中，CUDA 为 Shared Memory 提供编程接口，这使得开发者在设计 Kernel 实现时，可以利用 Shared Memory 访问延迟低的特点加速 Kernel 的性能。所以在 GPU 编程中，Kernel 的设计是以 Thread Block 这个粒度展开的。
从大到小来谈：CUDA编程就是在GPU硬件上启动了线程集合，为了更好的调度线程，GPU采用了分层的架构，在最高层的Grid负责将Block分配到哪些SM硬件上，在SM内部将由Warp调度那些线程来执行当前的任务：SM在实际计算时，会把block中的thread进一步拆分为warp，一个warp是32个thread，同一个warp里的thread，会以不同的数据，执行同样的指令，SM 一次只会执行一个warp。为了使SM忙碌起来，当一个warp遇到IO时，指令调度器会让SM执行另一个warp，这样就可以使SM保持忙碌，从而提高效率。因此，在编程时，最好是保证SM有足够多的warp进行切换。但是，warp驻留在SM是有代价的，即占据了内存，包括寄存器和共享内存等。为此，CUDA中有一个重要的概念叫做Occupancy（占用率），占用率是每个多处理器的活动warp与可能的活动warp的最大数量的比值，占用率跟寄存器数量、共享内存使用等因素有关。例如，在计算能力为7.0的设备上，每个多处理器有65,536个32位寄存器，最多可以有2048个线程同时驻留（64个warps）。PS：这不就是进程（wrap）等cpu排队嘛，只有64个wrap在同一时间能跑
从小到大来谈：CUDA 编程主打一个多线程 thread，多个 thread 成为一个 thread block，同一个 block 内的 thread 共享Shared Memory/L1 cache/SRAM，而 thread block 就是由这么一个 Streaming Multiprocessor (SM) 来运行的。
- 一个 SM 里面有多个 subcore，每个 subcore 有一个 32 thread 的 warp scheduler 和 dispatcher, 在一个 warp 中的所有线程都会同时执行相同的指令，但是输入的数据不同，这种机制也被称为 SIMD（单指令多数据）或 SIMT（单指令多线程）模型。
- GPU 的调度单元以 warp 为单位进行调度，而不是单个线程。这意味着整个 warp 会被分配到一个流多处理器（SM）上并一起执行。在 CUDA 中，占用率是一个重要的性能指标，表示每个 SM 上激活的 warps 与 SM 可以支持的最大 warp 数量的比例。更高的占用率通常意味着更好的硬件利用率。
- 如果 warp 中的所有线程都采取相同的分支路径（例如，都满足某个条件语句），则它们会继续同步执行。但是，如果线程在分支上有不同的路径（即分歧），则 warp 会执行每个路径，但不是所有线程都会在每个路径上活跃。这可能导致效率下降，因为即使某些线程在特定路径上没有工作，整个 warp 也必须等待该路径完成。为了确保高效执行，开发人员可能需要确保他们的代码减少 warp 分歧。
- Global memory 就是我们常说的 显存 (GPU memory)，其实是比较慢的。Global memory 和 shared memory 之间是 L2 cache，L2 cache 比 global memory 快。每次 shared memory 要到 global memory 找东西的时候, 会去看看 l2 cache 里面有没有, 有的话就不用去 global memory 了. 有的概率越大, 我们说 memory hit rate 越高, CUDA 编程的一个目的也是要尽可能提高 hit rate. 尤其是能够尽可能多的利用比较快的 SRAM (shared memory).但是因为 SRAM 比较小, 所以基本原则就是: 每次往 SRAM 移动数据的, 都可能多的用这个数据. 避免来来回回的移动数据. 这种 idea 直接促成了最近大火的 FlashAttention. FlashAttention 发现很多操作计算量不大, 但是 latency 很高, 那肯定是不符合上述的 “每次往 SRAM 移动数据的”. 怎么解决呢?Attention 基本上是由 matrix multiplication 和 softmax 构成的. 我们已经知道了 matrix multiplication 是可以分块做的, 所以就剩下 softmax 能不能分块做? softmax 其实也是可以很简单的被分块做的. 所以就有了 FlashAttention.
GPU的线程相对于CPU来讲属于十分轻量级的线程，创建和切换的开销都很小，而并行执行的数量以千计。但是另外一方面，GPU的线程并不能像CPU的线程那样自由。GPU的线程在执行的时候是分块（block）执行的，所以块（block）内的线程其实是共享pc寄存器。因此，虽然在编程的时候（编程模型当中），GPU的线程与CPU的线程类似，单独执行一段代码（称为kernel），但是实际上在GPU硬件上执行的时候，其实是将使用同一个kernel的多个线程归并在一个块（block）当中，用SIMD的方式去执行的。这种执行方式就隐含了，在任何一个时刻，一个块当中的所有线程，会进行一模一样的动作：如果是读内存，那么大家一起读；如果是写内存，那么大家一起写。从而，对于这种情况，相较于各个线程有自己的一片内存区域（按照线程组织数据），按照块组织数据效率会更高。所以我们可以看到，一般给CPU用的数据，都是线性排列的。而给GPU用的数据，基本上都是按照块（对应着GPU编程模型当中的线程块）来组织的。例如，从全局内存中加载的粒度是32*4字节，恰好是32个浮点数，每个线程束中的每个线程恰好一个浮点数。同样的原因，在片上存储空间以及多层高速缓存（cache）的组织方面，GPU也是突出了一个分块交换的概念，对线程组（block）的尺寸是非常敏感的。而这些在CPU上就不是那么明显。GPU深度学习性能的三驾马车：Tensor Core、内存带宽与内存层次结构 PS：线程多 ==> 共享pc/线程按block组织 ==> simd ==> 内存按块组织。
与模型的关系（不完整）：在计算模型推理时，我们通常会将模型视为单个块（block），但实际上模型由许多矩阵组成。当我们运行推理时，每个矩阵都被加载到内存中。具体来说，每个矩阵的块被加载到设备内存中，即共享内存单元（在A100上只有192KB）。然后，该块用于计算批次中每个元素的结果。需要注意的是，这与GPU RAM（即HBM）不同，A100具有40GB或80GB的HBM，但只有192KB的设备内存。因为我们不断地在设备内存中搬运数据，所以这在执行数学运算时会导致一个内存带宽瓶颈。我们可以通过计算模型大小/内存带宽比来近似传输权重所需的时间，并通过模型FLOPS/GPU FLOPS来近似计算所需的时间。
CUDA——GPU编程
NVIDIA 率先在 GPU 中引入了通用计算能力，使得开发者能利用 CUDA 编程语言来驱动。这时候 GPU 的核心都是 CUDA Core。由于一个 GPU 里面有大量的 CUDA Core，使得并行度高的程序获得了极大的并行加速。但是，CUDA Core 在一个时钟周期只能完成一个操作，矩阵乘法操作依然需要耗费大量的时间。NVIDIA 为了进一步加速“加乘运算”，在 2017 年推出了 Volta 架构的 GPU，从这个架构开始 Tensor Core 被引入。它可以在一个时钟周期完成两个 4×4x4 半精度浮点矩阵的乘法（64 GEMM per clock）。
基本概念
CUDA编程指北：从入门到实践 未读完，可继续。 CUDA 程序一般使用 .cu 后缀，编译 CUDA 程序则使用 nvcc 编译器。一般而言，一个 CUDA 程序的结构如下：
int main() {
主机代码; // 负责 CPU 和 GPU数据传输、GPU内存管理、以及启动 GPU 内核（内核启动参数指定了 GPU 上线程的数量和分布方式）等
核函数调用; // 每个内核函数在 GPU 的众多 CUDA 核心上并行执行，在 GPU 的多个线程上同时执行
主机代码;
核函数调用;
......
return 0;
}
__global__ void 核函数1(parameters) {
// 在设备代码中，计算任务被分解为多个线程，这些线程组成线程块（Block），多个线程块组成一个线程网格（Grid）。CUDA 提供了 threadIdx、blockIdx 等内置变量来获取线程的索引，从而让每个线程在数据中找到属于自己的计算任务。
......
}
__global__ void 核函数2(parameters) {
......
}
前缀__global__用来定义一个核函数，在 CUDA 中，核函数只能返回 void 类型（无返回值），这意味着当我们需要写计算结果时，应该在参数列表中传入一个用来存放计算结果的指针，然后将计算结果写回到这个指针指向的存储空间中。CUDA 核函数传入的参数必须是指向设备内存，因此，我们必须预先在主机代码中分配设备内存并初始化。分配设备内存可以使用 cudaMalloc 函数，初始化设备内存则可以将一段已经初始化好的主机内存拷贝到刚分配好的设备内存中，这可以使用 cudaMemcpy 函数实现，这两个函数的函数原型如下：
cudaError_t cudaMalloc(void** d_ptr, unsigned int size);
cudaError_t cudaMemcpy(void* d_ptr, void* h_ptr, unsigned int size, enum cudaMemcpyKind)
PS：所以在推理框架中，显存管理是推理框架负责。核函数 都是被封装后，注册到算子里，被类似op.forward 触发执行。 核函数是无状态的。 核函数的调用语法（内核启动语法）如下所示：
// CUDA 使用特殊的语法 <<<Grid, Block>>> 启动内核函数。
kernel_function<<<grid_size, block_size>>>(parameters)
// 也可以认为是
kernel<<<numBlocks, threadsPerBlock>>>(parameters)
// numBlocks 表示线程块的数量，threadsPerBlock 表示每个线程块中包含的线程数。
// 通过指定线程块数和线程数，内核启动控制了 GPU 的并行粒度。较大的数据集通常需要更多的线程和线程块来充分利用 GPU 的并行能力。
CUDA 的核函数设计一般遵循如下范式：data1，data2 … 表示需要处理的数据指针，index1 和 index2 … 用来定位需要计算的数据的位置，some_operation 对这些数据进行指定的计算操作，然后写回到参数列表中传入的用于记录结果的 result 指针中。总结下来就是两部曲：确定线程和数据的对应； 对需要处理的数据执行操作。PS：CUDA 最难的是并行思想。并行思想其中难上加难的东西是数据分组。并行计算中，最重要的一点是为数据分组成多个小数据块，每个线程（进程）再去实现SPMD或者SIMD/T。而这个数据分组的方式，存储方法等等直接的影响到你这个并行程序最终的性能。大部分的并行程序，解决了数据分组问题，其本身的问题就解决了，算法本身的优化反倒是不是那么的重要了。
__global__ void kernel_function(data1, data2, ..., result) {
index1, index2, ... = get_index(thread_info)
result = some_operations(data1[index1], data2[index2], ...)
}
内核启动后，GPU 可以异步执行任务，CPU 继续进行其他操作，直至需要等待 GPU 完成。开发者可以利用这种异步特性，使程序在 CPU 和 GPU 间并行执行，达到更高的并行效率。此外，CUDA 提供了同步函数（如 cudaDeviceSynchronize），确保 CPU 在需要时等待 GPU 完成所有操作，避免数据不一致的问题。
编译 CUDA 程序：编译 CUDA 程序需要使用 Nvidia 官方提供的编译器 nvcc。nvcc 会先将所有源代码先分离成主机代码和设备代码，主机代码完整支持 C++ 语法，设备代码只部分支持 C++ 语法。nvcc 先将设备代码编译为 PTX（parallel thread execution）伪汇编代码，再将 PTX 代码编译为二进制的 cubin 目标代码。CUDA 中核函数也因此不能直接作为类的成员函数，如果希望使用面向对象，我们一般通过包装函数调用核函数，然后将这个包装函数作为成员函数。
cuda的寄存器存储和共享内存
// 静态声明shared memory
__global__ void my_kernel() {
__shared__ int i;
}
// 动态声明shared memory
my_kernel(grid_dim, block_dim, 8) // 调用的时候指定动态内存的大小。这里8就是动态分配的。
__global__ void my_kernel() {
extern __shared__ int arr[];
}
// kernel函数中寄存器分配的场景一般有如下几种：临时变量，循环中分配的变量，函数调用和函数返回值。
__device__ int do_something(int x) {
return x*100; //返回值会被存放在register上。
}
__global__ void my_kernel() {
int a; //临时变量.a会被存放在register中
// for中的变量i会被存放在register中
for(int i=0;i<100;++i){
...
}
// do_something()的返回值会被存放在register中
int b = do_something(a);
// 大数组不会被存放在register上
int arr[1000];
}
大数组、结构体、动态分配的数组不会被存放在register上，而是存放在local memory中。local memory是global memory中的一块区域，由cuda自动分配，专门用来存放线程私有的数据。很显然，它的访问速度会比register 和shared memory慢很多。
cuda的同步有两种：
- 系统级同步：同步host和device的工作，用cudaDeviceSynchronize()接口。这个接口会阻塞所有host的工作，直到cuda端的工作完成。
- block线程同步。同步同一个block内的线程，用__synthreads()接口。同一个block内的线程用register和shared memory 进行通信。 cuda不同block之间的线程无法同步。如果需要，只能使用系统级同步方式，，使用cudaDeviceSynchronize()进行等待，在不同block的线程达到checkpoint后结束当前的kernel，开启新的kernel。
执行过程
CUDA C++ 编程指北-第一章：入门以及编程模型 未读。
device 函数和global函数因为需要在GPU上运行，因此不能调用常见的一些 C/C++ 函数（因为这些函数没有对应的 GPU 实现）。
| 限定符 | 执行 | 调用 | 备注 |
|---|---|---|---|
| global | 设备端执行 | 可以从主机调用也可以从某些特定设备调用 | 异步操作，host 将并行计算任务发射到GPU的任务调用单之后，不会等待kernel执行完就执行下一步 |
| device | 设备端执行 | 设备端调用 | |
| host | 主机端执行 | 主机调用 |
典型的CUDA程序的执行流程如下：
- 分配host内存，并进行数据初始化；
- 分配device内存，并从host将数据拷贝到device上；
- 调用CUDA的核函数在device上完成指定的运算；
- 将device上的运算结果拷贝到host上；
- 释放device和host上分配的内存。
矩阵加法示例
// __global__ 表示在device上执行，从host中调用
// 两个向量加法kernel，grid和block均为一维
__global__ void add(float* x, float * y, float* z, int n){
// 获取全局索引
int index = threadIdx.x + blockIdx.x * blockDim.x;
// 步长
int stride = blockDim.x * gridDim.x;
for (int i = index; i < n; i += stride){ # cuda编程经常会跨步长，每个thread 只计算数据的局部和/积/xx。
z[i] = x[i] + y[i];
}
}
int main(){
int N = 1 << 20;
int nBytes = N * sizeof(float);
// 申请host内存
float *x, *y, *z;
x = (float*)malloc(nBytes);
y = (float*)malloc(nBytes);
z = (float*)malloc(nBytes);
// 初始化数据
for (int i = 0; i < N; ++i){
x[i] = 10.0;
y[i] = 20.0;
}
// 申请device内存
float *d_x, *d_y, *d_z;
cudaMalloc((void**)&d_x, nBytes);
cudaMalloc((void**)&d_y, nBytes);
cudaMalloc((void**)&d_z, nBytes);
// 将host数据拷贝到device
cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);
cudaMemcpy((void*)d_y, (void*)y, nBytes, cudaMemcpyHostToDevice);
// 定义kernel的执行配置
dim3 blockSize(256);
dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
// 执行kernel
add << < gridSize, blockSize >> >(d_x, d_y, d_z, N); # 第一个数字指明改程序分配多少个block，第二个数字程序指明每个block中的thread个数
// 将device得到的结果拷贝到host
cudaMemcpy((void*)z, (void*)d_z, nBytes, cudaMemcpyDeviceToHost);
// 检查执行结果
float maxError = 0.0;
for (int i = 0; i < N; i++)
maxError = fmax(maxError, fabs(z[i] - 30.0));
std::cout << "最大误差: " << maxError << std::endl;
// 释放device内存
cudaFree(d_x);
cudaFree(d_y);
cudaFree(d_z);
// 释放host内存
free(x);
free(y);
free(z);
return 0;
}
如何在 CPU 之上调用 GPU 操作？可以通过调用 __global__ 方法来在GPU之上执行并行操作。我的第一份CUDA代码 - xcyuyuyu的文章 - 知乎 kernel在调用时也必须通过执行配置<<<grid, block>>>来指定kernel所使用的线程数及结构。一个thread需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识，其中blockIdx指明block在grid中的位置，而threaIdx指明线程所在block中的位置。
矩阵计算优化（没太懂）
GPU存储可分为物理内存（硬件真实存在的）和逻辑内存（由cuda做抽象的）。 为什么要这么分呢？因为各个GPU的物理内存架构是不一样的，如果你写代码时还要考虑每个GPU的独特性，那可太痛苦了。所以cuda在这里帮了大忙：它对内存架构做了一层抽象，你只要按照它抽象后的框架写代码就可以。
每个thread占用一个SP（cuda core），即1个warp会占用1个SM上的32个SP。 有了这些前置知识，现在我们可以来看cuda矩阵优化的过程了。 假设矩阵
- A = (M,K) = (512,512)
- B = (K,N) = (512,512)
- C = AB = (M,K) * (K,N) = (512,512)
每个thread负责读取A矩阵的一行和B矩阵的一列，去计算C矩阵的一个元素。则一共需要M*N个thread。 矩阵A和矩阵B都存储在global memory，每个thread直接从global memory上进行读数，完成计算：
- 为了计算出C中的某个元素，每个thread每次都需要从global memory上读取A矩阵的一行（K个元素），B矩阵的一列（K个元素），则每个thread从global memory上的读取次数为2K
- C中共有M*N个thread，则为了计算出C，对global memory的总读取次数为： 2MNK
Naive GEMM的代码见下（完整代码见 sgemm_naive.cu ）：
// 将二维数组的行列索引转成一维数组的行列索引，这样可以更高效访问数据
// row, col：二维数组实际的行列索引，ld表示该数组实际的列数
// 例：二维数组实际的行列索引为(1, 3)，即第二行第四个元素，二维数据的总列数 = 5
// 返回的一位数组形式的索引为: 1*5 + 3 = 8
#define OFFSET(row, col, ld) ((row) * (ld) + (col))
// 定义naive gemm的kernel函数
__global__ void naiveSgemm(float * __restrict__ a, float * __restrict__ b, float * __restrict__ c, const int M, const int N, const int K) {
// 当前thread在C矩阵中的row
int m = blockIdx.y * blockDim.y + threadIdx.y;
// 当前thread在C矩阵中的col
int n = blockIdx.x * blockDim.x + threadIdx.x;
if (m < M && n < N) {
float psum = 0.0;
// 告知编译器自动展开循环体，这样可以减少循环控制的开销（循环次数小的时候可以这么做）
#pragma unroll
// 取出A[row]和B[col]，然后逐个元素相乘累加，得到最终结果
for (int k = 0; k < K; k++) {
// a[OFFSET(m, k, K)]: 获取A[m][k]
// b[OFFSET(k, n, N)]: 获取B[k][n]
psum += a[OFFSET(m, k, K)] * b[OFFSET(k, n, N)];
}
c[OFFSET(m, n, N)] = psum;
}
}
const int BM = 32, BN = 32;
const int M = 512, N = 512, K = 512;
dim3 blockDim(BN, BM);
dim3 gridDim((N + BN - 1) / BN,(M + BM - 1) / BM);
GPU/CUDA/驱动和机器学习训练框架的关系
显卡是硬件，硬件需要驱动，否则不能调用其计算资源。CUDA又是什么？
- 在2007年之前，GPU由CPU操作，CPU把一些图形图像的计算任务交给GPU执行。程序员不需要与GPU打交道。随着GPU计算能力的发展，越来越多的计算场景由GPU完成效果会更好。但现有的程序无法直接自由控制GPU的处理器。当然程序员也可以直接写代码与显卡驱动对接，从而直接控制GPU的处理器，但这样代码恐怕写起来要让人疯掉。nvidia当然会有动力提供一套软件接口来简化操作GPU的处理器。nvidia把这一套软件定义为CUDA。
- 多核 CPU 和众核 GPU 的出现意味着主流处理器芯片现在是并行系统。挑战在于开发能够透明地扩展可并行的应用软件，来利用不断增加的处理器内核数量。CUDA 并行编程模型旨在克服这一挑战，同时为熟悉 C 等标准编程语言的程序员保持较低的学习曲线。CUDA 编程手册系列第一章：CUDA 简介
gpu 和 cuda 和 gpu driver 之间的关系：比如 TX3090需要Compute Capability在8.6以上的cuda，而满足这个要求的cuda又只有11.0以上的版本。而cuda11版本又需要版本号>450的显卡驱动。
显卡
==> Compute Capability 查看显卡支持的Compute Capability, https://developer.nvidia.com/cuda-gpus
==> cuda
==> GPU driver 查看cuda对驱动的要求 (https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html
==> tf/pytorch version
Compute Capability的数值和GPU的计算速度无关，但是和GPU可执行的任务种类有关。The compute capability of a device is represented by a version number, also sometimes called its “SM version”. This version number identifies the features supported by the GPU hardware and is used by applications at runtime to determine which hardware features and/or instructions are available on the present GPU.
从模型到算子，以卷积计算为例
- 从卷积到矩阵乘 将输入数据（Feature Map）和卷积核数据进行重排，卷积操作本质上可以等效理解为矩阵乘操作。卷积操作的过程大概可以描述为按照约定的窗口大小和步长，在 Feature Map 上进行不断地滑动取数，窗口内的 Feature Map 和卷积核进行逐元素相乘，再把相乘的结果累加求和得到输出 Feature Map 的每个元素结果。
- 矩阵乘分块 Tilling。卷积转换后的矩阵乘的维度非常大，而芯片里的内存空间往往是有限的（成本高），表现为越靠近计算单元，带宽越快，内存越小。为了平衡计算和内存加载的时间，让算力利用率最大化，AI 芯片往往会进行由远到近，多级内存层级的设计方式，达到数据复用和空间换时间的效果。根据这样的设计，矩阵乘实际的数据加载和计算过程将进行分块 Tilling 处理。Tiling（平铺）是一种优化技术，它涉及将大的矩阵分解成更小的块或“瓦片”（tiles），这些小块的大小通常与CPU或GPU的缓存大小相匹配，以便可以完全加载到缓存中。
- 矩阵乘的库。矩阵乘作为 AI 模型中的重要性能算子，CPU 和 GPU 的平台上都有专门对其进行优化实现的库函数。比如 CPU 的 OpenBLAS, Intel MKL 等，GPU 的 cuBLAS, cuDNN 等。实现的方法主要有 Loop 循环优化 (Loop Tiling)和多级缓存 (Memory Hierarchy)。
- 矩阵乘的优化。在具体的 AI 芯片或其它专用芯片里面，对矩阵乘的优化实现主要就是减少指令开销，可以表现为两个方面：
- 让每个指令执行更多的 MACs 计算。比如 CPU 上的 SIMD/Vector 指令，GPU 上的 SIMT/Tensor 指令，NPU 上 SIMD/Tensor,Vector 指令的设计。
- 在不增加内存带宽的前提下，单时钟周期内执行更多的 MACs。比如英伟达的 Tensor Core 中支持低比特计算的设计，对每个 cycle 执行 512bit 数据的带宽前提下，可以执行 64 个 8bit 的 MACs，大于执行 16 个 32bit 的 MACs。