GPU 调度策略架构与CUDA运行机制（二）

1 GPU运行机制总述

市面上有很多GPU厂家，他们产品的架构各不相同，但是核心往往差不多，整明白了一个基本上就可以触类旁通了。

1.0 GPU计算流程（CPU协同GPU计算）

一个典型的计算流程是这样的:

• 数据从CPU的内存拷贝到GPU的内存
• CPU把计算指令传送给GPU
• GPU把计算任务分配到各个CUDA core并行处理
• 计算结果写到GPU内存里, 再拷贝到CPU内存里.

1.1 Host与Device

一个CUDA程序的可以分为两个部分（两者拥有各自的存储器）：

• 在CPU上运行的称为Host程序
• 在GPU上运行的称为Device程序

1.2 Kernel

GPU上运行的函数又被叫做Kernel函数。

Host程序在调用Device程序时，可以通过参数确定执行该Kernel的CUDA threads的数量。

每个Thread在执行Kernel函数时，会被分配一个thread ID，Kernel函数可以通过内置变量threadIdx访问。

1.3 图解说明从CPU到GPU的运行机制

CUDA在执行的时候是让Host程序里面的一个一个的Kernel函数按照Grid（线程网格）的概念在GPU上执行。

一个Kernel函数对应一个Grid。

每个Grid中的任务是一定的。当要执行这些任务的时候，每一个Grid又把任务分成一部分一部分的Block（线程块），Block中间有若干Thread（线程），再分成线程来完成。

1.4 Single Instruction Multiple Threads（SIMT）

• GPU中的SIMT体系结构相对于CPU的SIMD（单指令多数据，Single Instruction Multiple Data）。中文翻译：单指令多线程。SIMT对于可编程性的好处使得NVIDIA的GPU架构师为这种架构命名，而不是将其描述为 SIMD 。
• 为了有效地管理和执行多个单线程，流多处理器（SM）采用了SIMT架构。此架构在第一个unified computing GPU中由NVIDIA公司生产的GPU引入。
• GPU使用SIMT执行 32 个并行线程的 Warp ，实现单指令、多线程，这使得每个线程能够访问自己的寄存器，从不同的地址加载和存储，并遵循不同的控制流路径。CUDA编译器和GPU一起工作，以确保Warp的线程组尽可能频繁地被分配到SM中，一起执行相同的指令序列，从而最大限度地提高性能。
• 每个线程可以包含控制流指令（控制流指令为标量指令）
• 同组Warp中的这些线程可以执行不同的控制流路径
• 当一个Warp中的线程分支到不同的执行路径时，产生分支发散（Branch divergence）

优势

• 共享控制逻辑可以有更多的空间面基去分配给计算单元
• 大量的并行操作，不需要进行复杂的控制编程

SIMD VS SIMT

• CPU中通过SIMD来处理矢量数据；纯粹使用SIMD不能并行的执行有条件跳转的函数，很显然条件跳转会根据输入数据不同在不同的线程中有不同表现。
• GPU则使用SIMT，无需开发者费力把数据凑成合适的矢量长度，并且SIMT允许每个线程有不同的分支，利用SIMT 才能做到不同分支的并行操作。

GPU 工作原理

本章将从 GPU 硬件基础和英伟达 GPU 架构两个方面讲解 GPU 的工作原理。英伟达 GPU 有着很长的发展历史，整体架构从 Fermi 到 Blankwell 架构演变了非常多代，其中和 AI 特别相关的就有 Tensor Core 和 NVLink。

本节首先讲解 CPU 和 GPU 架构的区别，之后以$AX+Y$这个例子来探究 GPU 是如何做并行计算的，为了更好地了解 GPU 并行计算，对并发和并行这两个概念进行了区分。此外会讲解 GPU 的缓存机制，因为这将涉及到 GPU 的缓存（Cache）和线程（Thread）。

GPU 工作原理

基本工作原理

首先通过$AX+Y$这个加法运算的示例了解 GPU 的工作原理，$AX+Y$ 的示例代码如下：

void demo(double alpha, double *x, double *y)
{
    int n = 2000;
    for (int i = 0; i < n; ++i)
    {
        y[i] = alpha * x[i] + y[i];
    }
}

示例代码中包含 2 FLOPS 操作，分别是乘法（Multiply）和加法（Add），对于每一次计算操作都需要在内存中读取两个数据，$x[i]$ 和 $y[i]$，最后执行一个线性操作，存储到 $y[i]$ 中，其中把加法和乘法融合在一起的操作也可以称作 FMA（Fused Multiply and Add）。

在 O(n) 的时间复杂度下，根据 n 的大小迭代计算 n 次，在 CPU 中串行地按指令顺序去执行 $AX+Y$ 程序。以 Intel Exon 8280 这款芯片为例，其内存带宽是 131 GB/s，内存的延时是 89 ns，这意味着 8280 芯片的峰值算力是在 89 ns 的时间内传输 11659 个比特（byte）数据。$AX+Y$ 将在 89 ns 的时间内传输 16 比特（C/C++中 double 数据类型所占的内存空间是 8 bytes）数据，此时内存的利用率只有 0.14%（16/11659），存储总线有 99.86% 的时间处于空闲状态。

内存总线 99.86%时间处于空闲状态

不同处理器计算 $AX+Y$ 时的内存利用率，不管是 AMD Rome 7742、Intel Xeon 8280 还是 NVIDIA A100，对于 $AX+Y$ 这段程序的内存利用率都非常低，基本 ≤0.14%。
AMD Rome 7742Intel Xeon 8280NVIDIA A100Memory B/W(GB/sec)2041311555DRAM Latency(ns)12289404Peak bytes per latency24,88811,659628,220Memory Efficiency0.064%0.14%0.0025%
由于上面的 $AX+Y$ 程序没有充分利用并发和线性度，因此通过并发进行循环展开的代码如下：

void fun_axy(int n, double alpha, double *x, double *y)
{
    for (int i = 0; i < n; i += 8)
    {
        y[i + 0] = alpha * x[i + 0] + y[i + 0];
        y[i + 1] = alpha * x[i + 1] + y[i + 1];
        y[i + 2] = alpha * x[i + 2] + y[i + 2];
        y[i + 3] = alpha * x[i + 3] + y[i + 3];
        y[i + 4] = alpha * x[i + 4] + y[i + 4];
        y[i + 5] = alpha * x[i + 5] + y[i + 5];
        y[i + 6] = alpha * x[i + 6] + y[i + 6];
        y[i + 7] = alpha * x[i + 7] + y[i + 7];
    }
}

每次执行从 0 到 7 的数据，实现一次性迭代 8 次，每次传输 16 bytes 数据，因此同样在 Intel Exon 8280 芯片上，每 89 ns 的时间内将执行 729（11659/16）次请求，将程序这样改进就是通过并发使整个总线处于一个忙碌的状态，但是在真正的应用场景中：

• 编译器很少会对整个循环进行超过 100 次以上的展开；
• 一个线程每一次执行的指令数量是有限的，不可能执行非常多并发的数量；
• 一个线程其实很难直接去处理 700 多个计算的负荷。

由此可以看出，虽然并发的操作能够一次性执行更多的指令流水线操作，但是同样架构也会受到限制和约束。

将 $Z=AX+Y$ 通过并行进行展开，示例代码如下：

void fun_axy(int n, double alpha, double *x, double *y)
{
    Parallel for (int i = 0; i < n; i++)
    {
        y[i] = alpha * x[i] + y[i];
    }
}

通过并行的方式进行循环展开，并行就是通过并行处理器或者多个线程去执行 $AX+Y$ 这个操作，同样使得总线处于忙碌的状态，每一次可以执行 729 个迭代。相比较并发的方式：

• 每个线程独立负责相关的运算，也就是每个线程去计算一次 $AX+Y$；
• 执行 729 次计算一共需要 729 个线程，也就是一共可以进行 729 次并行计算；
• 此时程序会受到线程数量和内存请求的约束。

GPU 线程原理

GPU 整体架构和单个 SM（Streaming Multiprocessor）的架构，SM 可以看作是一个基本的运算单元，GPU 在一个时钟周期内可以执行多个 Warp，在一个 SM 里面有 64 个 Warp，其中每四个 Warp 可以单独进行并发的执行，GPU 的设计者主要是增加线程和增加 Warp 来解决或者掩盖延迟的问题，而不是去减少延迟的时间。

GPU 整体架构与 SM 架构

为了有更多的线程处理计算任务，GPU SMs 线程会选择超配，每个 SM 一共有 2048 个线程，整个 A100 有 20 多万个线程可以提供给程序，在实际场景中程序用不完所有线程，因此有一些线程处于计算的过程中，有一些线程负责搬运数据，还有一些线程在同步地等待下一次被计算。很多时候会看到 GPU 的算力利用率并不是非常的高，但是完全不觉得它慢是因为线程是超配的，远远超出大部分应用程序的使用范围，线程可以在不同的 Warp 上面进行调度。
Pre SMA100Total Threads2048221,184Total Warps646,912Active Warps4432Waiting Warps606,480Active Threads12813,824Waiting Threads1,920207,360

小结

本节首先从架构层面分析了 CPU 和 GPU 的主要区别，因为 CPU 的设计目标是尽可能在低的延迟下执行任务，GPU 的设计目标是最大化吞吐量，因此决定了 CPU 适合处理顺序执行的任务，GPU 适合处理大规模并行计算的任务。

以 $AX+Y$ 为例讲解了并发和并行的区别以及和串行的区别，在串行计算时内存利用率很低，当把程序用并发的方式展开，并发操作的多流水线会受到 CPU 架构的限制，当程序用并行的方式循环展开时，程序的执行则只受到线程数量和内存请求的约束，因此 CPU 跟 GPU 的本质区别是并行的问题而不是并发的问题，GPU 通过大量的线程提供并行的能力。

为了提供并行的能力，GPU 通过多级缓存、多级流水、多级 Cache 提供并行的机制，同时可以尽可能减少内存的时延。为了将数据充分利用起来，引入了 GPU 线程的原理，GPU 里面提供了大量超配的线程去完成对不同层级数据的搬运和计算。

1 GPU的软件抽象

软件资源的抽象即为GPU的线程模型，可以分为Grid、Block、Thread和Warp。

Grid、Block、Thread是一种软件组织结构，是线程组织的三个层次，并不是硬件的，因此理论上我们可以以任意的维度（一维、二维、三维）去排列Grid，Block，Thread；在硬件上就是一个个的SM或者SP，并没有维度这一说，只是软件上抽象成了具有维度的概念。

thread，block，gird在不同维度的大小根据算力不同是有限制的：所以在不同CUDA版本或在编译时没有指定架构的情况下，可能CUDA版本也会对thread，block，grid在不同维度的大小产生影响。

1.1 Grid（线程网格）

一个Kernel函数对应一个Grid。

一个Grid中会分成若干个Block。同一Grid下的不同Block可能会被分发到不同的SM上执行。

Grid跑在GPU上的时候，可能是独占一个GPU，也可能是多个kernel函数并发占用一个GPU（后面这种实现需要fermi及更新的GPU架构支持）。

1.2 Block

数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过shared memory通信

1.3 Thread

一个CUDA的并行程序会被以许多个Thread来执行

每个Thread中的局域变量被映射到SM的寄存器上，而Thread的执行则由CUDA核心也就是SP来完成。

1.4 Warp

Warp是GPU执行程序时的调度单位，同一个Warp里的线程执行相同的指令，即SIMT。

一个SM的CUDA core会分成几个Warp（即CUDA core在SM中分组)，由Warp scheduler负责调度。尽管Warp中的线程从同一程序地址，但可能具有不同的行为，比如分支结构。因为GPU规定同一Warp中所有线程在同一周期执行相同的指令，Warp发散分支过多会导致有效分支减少性能下降。

一个SM同时并发的Warp是有限的，因为资源限制，SM要为每个线程块分配共享内存，也要为每个线程束中的线程分配独立的寄存器，所以SM的配置会影响其所支持的线程块和Warp并发数量。

一个Warp中的线程必然在同一个block中，如果block所含线程数目不是Warp大小的整数倍，那么多出的那些thread所在的Warp中，会剩余一些inactive的thread，也就是说，即使凑不够Warp整数倍的thread，硬件也会为Warp凑足，只不过那些thread是inactive状态，需要注意的是，即使这部分thread是inactive的，也会消耗SM资源。由于warp的大小一般为32，所以block所含的thread的大小一般要设置为32的倍数。

例：如果一个块中有128个线程，那么线程0-31将在一个Warp中，32-63将在下一个Warp中

Warp非常重要，原因如下：

• Warp中的线程是被绑定在一起的。如果Warp中的一个线程沿着if-else块的if侧走，而其他线沿着else侧走，那么实际上所有32条线程都会沿着两侧走。在执行功能上是没有问题的，那些不应该被执行分支的线程会被禁用，因此始终获得正确的结果，但是如果双方都很长，那么性能损失就很重要。
• Warp内的线程（实际上是半纠缠的（self-warp））一起从内存中获取数据，是一起访问共享内存中的同一段数据同一段的。也就是说如果可以确保Warp中的所有线程都从同一段内获取数据，就只需要实现一次内存转换。
• 如果它们都从随机地址获取数据，那么就需要排队去实现32次内存转换。

2 软件抽象和硬件结构的一一对应关系

硬件结构可以参考之前的一篇文章

2.1 Block对应于SM

• SM上可以同时存在多个Block被执行，这些Block不一定来自同一个kernel函数。
• SM设备有Device Limit，Warp和Block的数量不能超过对应的上限。
• 除了受到设备定义的限制之外，还受到硬件资源的限制： - SP的寄存器数量- 线程块消耗的共享内存量

每个线程会占用一定数量的寄存器和Shared Memory，因此SM上同时存活的Block数目不应当超过这些硬件资源的限制。由于SM上可以同时有来自不同kernel的Block存在，因此有时候即便SM上剩余资源不足以再容纳一个kernel A的Block，但却仍可能容纳下一个kernel B的Block。

• 一个线程块的thread只能在一个SM上调度

2.2 Block与Thread之间的联系Warp 对应于 SM与SP之间的联系

• 软件抽象里，认为任务分配到Block之后，所有的线程是并行执行的，这只是个逻辑上无懈可击的抽象，事实上我们不可能对一个任意大小的Block都给出一个同等大小的CUDA核心阵列去推动它的并行计算，来真正并行的执行它们。因而有了Warp这个概念。物理上，Block被划分成一块块的warp分别映射到CUDA核心阵列上执行，每一个warp就都可以理解为是一个线程的集装箱，为的是线程数量固定统一可以给他分配统一的硬件资源，每个集装箱只装一种货物，也就是下面同步执行的意思。
• 目前，CUDA中的Warp都是从threadIdx = 0开始，以threadIdx连续的32个线程为一组划分得到，即便最后剩下的线程不足32个，也将其作为一个Warp。CUDA kernel的配置中，我们经常把Block的size设置为32的整数倍，正是为了让它能够精确划分为整数个Warp（更深刻的原因和存储器访问性能有关，但这种情况下仍然和Warp的size脱不了干系）。
• Warp是SM调度和执行的基础概念。Block被划分成32个线程组成的Warp。这样，大量的Warp生存在SM上，等待被调度到CUDA核心阵列去执行。
• Warp中的活动线程由Warp Scheduler驱动。每一块SM中有单独的一个或者多个Warp Scheduler（举例：GM204中32个CUDA核心共享一个Warp Scheduler），以及多个CUDA核心。
• 当一个Warp执行中出现等待（存储器读写延迟等）后，Warp Scheduler就迅速切换到下一个可执行的Warp，对其发送指令直到这个Warp又一次出现等待，周而复始。这就是常说“用多线程掩盖延迟”。SM会从驻留在SM上的所有Warp中进行指令调度。（这里的驻留表示已经可以被执行的Warp，会从这里挑选，这时候挑选出来的Warp能来自于驻留在SM上的任何线程块）。
• 通常一个SM中的SP会分成几个Warp（也就是SP在SM中是进行分组的，物理上进行的分组）。
• 同步执行：Warp中的32个SP是一起工作的，执行相同的指令，如果没有这么多thread需要工作，那么这个Warp中的一些SP是不工作的，处于闲置状态。

2.3 Thread对应于SP

• Thread在SP也就是CUDA Cores上执行
• Thread会被分配Register/Local Memory，数据存在这里
• SM上的CUDA核心是有限的，它们代表了能够在物理上真正并行的线程数（也就是优化到最佳情况下所能最大达到同一时刻在运行的并行数量）
• 每一个线程都有自己的寄存器内存和local memory，一个warp中的线程是同时执行的，也就是当进行并行计算时，线程数尽量为32的倍数，如果线程数不上32的倍数的话；假如是1，则warp会生成一个掩码，当一个指令控制器对一个warp单位的线程发送指令时，32个线程中只有一个线程在真正执行，其他31个 进程会进入静默状态。

3 软件抽象和硬件结构对应关系的例子

把GPU跟一个学校对应起来，学校里有教学楼、操场、食堂，还有老师和学生们；很快有领导（CPU）来检查卫生（需要执行的任务Host程序），因此这个学校的学生们要完成打扫除的工作（Device程序）。

• 软件抽象资源包括Thread、Warp、Block和Grid
• 硬件资源包括SP和SM

3.1 软件抽象

Grid对应的是年级

是抽象的划分组织方式

根据年级划分任务，Grid可以分为多个不同的班级

Block对应的是班级

是抽象的划分组织方式

每个班级有若干的同学（线程），可能一个两个不同的年级会出现在同一层楼（SM），或者一层楼只有一个班级，或者没有班级，但是每一层楼的班级最大数量是固定的

Warp对应的是兴趣小组

每个小组有32个学生；（同一时间他们一定是一个班级下的小组）

并且数量固定，即使凑不满这么多学生需要加进来不干活的学生，凑够一个小组

只要求他们有着一样的兴趣爱好（能执行相同的任务）

Thread对应的是学生

一个Thread对应一个SP

每个学生都有个课桌 ，放自己的物品，不能让别人用，表示每个Thread在软件上都有自己的空间（寄存器等）

3.2 硬件资源

SM对应的是教学楼的一个楼层

是实际存在的资源

一个楼层上可以有多个班级，年级和楼层并没有确定的对应关系，一个楼层中可以有很多来自不同的年级的Block

SM中的SP会被分成兴趣小组，承接不同的任务

SP对应的是学生

一个SP对应一个Thread

是实际存在的资源

每个学生都有个课桌 ，放自己的物品，不能让别人用，表示每个SP在硬件上都有自己的空间（local memory + registers)；

在楼层中，有公共的空间（走廊、厕所等），这一层楼的所有同学都可以停留，表示一个SM中有shared memory，这个SM上的Block都可以访问；（shared memory是不是所有的block都可以访问）

学校里的公共区域，比如操场、食堂等，所有同学都可以去运动、吃饭，表示GPU中有一些公共的存储空间供所有的Grid访问。

3.3 执行任务

虽然GPU是并行运行，但也并不是我们理想中所有的Thread一起工作，在打扫卫生时，并不是所有学生一起干活，学生经过老师（这里我们理解为Wrap Scheduler）安排后，分为一组一组的小组，每一个小组都只会做一件一样的事情，如果有人先做完了或者不需要做，那么他也会在旁边等他的组员，处于等待状态idle。

4 用多线程掩盖延迟

Global Memory访存延迟可以达到数百个时钟周期，即便是最快的Shared Memory和寄存器在有写后读依赖时也需要数十个时钟周期。这似乎和CUDA强大的处理能力完全相悖。

为什么GPU具有这么高的计算能力？如果连寄存器都这么慢，怎么会有高性能呢？难道这不会成为最大的瓶颈吗？

因为这个高延迟的开销被掩盖了，掩盖在大量线程之下。更清楚的说，控制单元（Warp Scheduler）在多组线程之间快速切换，当一组线程Warp（一个线程组，在CUDA里叫做Warp）因为访存或其他原因出现等待时，就将其挂起，转而执行另一组线程，GPU的硬件体系允许同时有大量线程存活于GPU的SM（流多处理器）之中，这种快速切换保证资源的最大利用率——控制单元始终有指令可以发放，执行单元始终有任务可以执行，仍然可以保持最高的指令吞吐，每个单元基本都能保持充分的忙碌。

这就是GPU硬件设计中非常有特色的基本思想：用多线程掩盖延迟。这一设计区别于CPU的特点是，大量高延迟寄存器取代了少量低延迟寄存器，寄存器的数量保证了可以有大量线程同时存活，且可以在各组线程间快速切换。尽管每个线程是慢的，但庞大的线程数成就了GPU的数据吞吐能力。

下面图片可以说明：GPU用多个Warp掩盖延迟 / 与CPU计算模式的对比

GPU因为多个Warp可以快速切换来掩盖延迟，而CPU用快速的寄存器来减小延迟。两者的重要区别是寄存器数目，CPU的寄存器快但少，因此Context Switch代价高；GPU寄存器多而慢，但寄存器数量保证了线程Context Switch非常快。同时也是因为GPU对高延迟的容忍度比较高，他只追求在长时间内比较稳定的较大吞吐量，而不在意响应时间。

4.1 多少线程才能够掩盖掉常见的延迟呢？

对于GPU，最常见的延迟大概要数寄存器写后读依赖，即一个局域变量被赋值后接着不久又被读取，这时候会产生大约24个时钟周期的延迟。为了掩盖掉这个延迟，我们需要至少24个Warp轮流执行，一个Warp遇到延迟后的空闲时间里执行其余23个Warp，从而保持硬件的忙碌。在Compute Capability 2.0，SM中有32个CUDA核心，平均每周期发射一条指令的情况下，我们需要24 ∗ 32 = 768 24*32 = 76824∗32=768个线程来掩盖延迟。
保持硬件忙碌，用CUDA的术语来说，就是保持充分的Occupancy，这是CUDA程序优化的一个重要指标。

5 关于现代GPU如此进行软件抽象和硬件设计的一些思考

整个设计逻辑关系我觉得可以归结为如下的情况

• 目标是实现任务
• 发现任务具有如下的特性：允许一定的延迟；需要大吞吐量；有大量同样的操作或者计算
• 所以设计了现有的硬件体系架构，软件抽象模型

那么为什么这样的计算或者说任务可以被如上所说的硬件软件更好的完成呢？

其实是因为我们是在已知任务特性的情况下（我们实际使用中所需要完成的任务大概率属于这些，或者说这些任务在CPU上比较容易有掣肘），才把结构设计成这样的。

• 第一方面： - 现实世界中应用在大规模数据上的计算，通常都涵盖在这一计算模式之中，因而考虑更复杂的模式本质上是不必要的。

比如计算大气的流动，每一点的风速仅仅取决于该点邻域上的密度和压强分布；

比如计算图像的卷积，每一个输出像素都仅是对应源点邻域和一个卷积核的内积。

• 从这些例子中我们可以看到，除了各个数据单元上进行的计算是一样的，计算中数据之间的相互影响也具有某种“局域性”，一个数据单元上的计算最多需要它某个邻域上的数据。这一点意味着线程之间是弱耦合的，邻近线程之间会有一些共享数据（或者是计算结果），远距离的线程间则独立无关。

这个性质反映在CUDA里，就是Block划分的两重天地：Block内部具有Shared Memory，线程间可以共享数据、通讯和同步，Block外部则完全独立，Block间没有通讯机制，相互执行顺序不影响计算结果。这一划分使得我们既可以利用线程间通讯做一些复杂的应用和算法加速，又可以在Block的粒度上自由调度计算任务，在不同计算能力的硬件平台上自适应的调整任务安排。

• 第二方面：

多个线程同步执行一致的运算，使得我们可以用单路指令流对多个执行单元进行控制，大幅度减少了控制器的个数和系统的复杂度

• 第三方面：

把注意力放在“几乎一致”这里。最简单的并行计算方案是多路数据上同时进行完全一致的计算，即SIMD（单指令多数据流）。这种方案是非常受限的。事实上我们可以看出，“完全一致”是不必要的。只要这些计算在大多数时候完全一致，就可以对它们做类似于SIMD的加速，不同点是在计算分叉时候，各个线程不一致的特殊情况下，只需要分支内并行，分支间串行执行即可，毕竟这些只是很少出现的情况。 这样，把“完全一致”这个限制稍微放松，就可以得到更广阔的应用范围和不输于SIMD的计算性能，即SIMT（单指令流多线程）的一个重要环节，这是GPU强大处理能力的原因。

LAST 参考文献

GPU 初理解 - 简书

GPU架构之处理模块 - 知乎

GPU中的基本概念 - 云+社区 - 腾讯云

CUDA, 软件抽象的幻影背后 之二 | 奇点视觉

CUDA, 软件抽象的幻影背后 | 奇点视觉

GPU编程1–GPU中的基本概念 - 知乎

(3条消息) gpu的单位表示_GPU中的基本概念_weixin_39717121的博客-CSDN博客

CUDA的thread,block,grid和warp - 知乎

GPU编程3–GPU内存深入了解 - 知乎

GPU架构之Hierarchy Memory多级存储 - 知乎

cuda编程(一)：GPU概念与架构 - 知乎

GPU计算 – GPU体系结构及CUDA编程模型

Nvidia GPU架构 - Cuda Core，SM，SP等等傻傻分不清？_咚咚锵的博客-CSDN博客_cuda sm

Fermi威力完美呈现，GeForce GTX 580全球同步评测 - 超能网
————————————————

                        版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
                     

原文链接：https://blog.csdn.net/qq_41554005/article/details/119760698

示例：一维数组的求和计算

代码中注释的一、二处究竟该怎么来写？

---

线程参数设置 情况1：一维grid，一维block （线程分配）

grid(1,1,1): block数量=111

block(length,1,1): thread数量=length11

总thread数量 = (111)(length1*1)

-------------------------------------------------------------------------------------------

线程参数设置 情况二2：一维grid，二维block （线程分配）

grid(1,1,1): block数量=111

block(8,2,1): thread数量=821

总thread数量 = 16

我们一定要有并行思想，这里有16个线程，kernel启动后，每个线程都有自己的索引号，比如某个线程位于grid中哪个维度的block（即blockIdx.x,blockIdx.y,blockIdx.z），又位于该block的哪个维度的线程（即threadIdx.x,threadIdx.y,threadIdx.z），利用这些线程索引号映射到对应的数组下标，我们要做的工作就是将保证这些下标不重复（如果重复的话，那就惨了），最初那种一维的计算方式就不行了。因此，通过使用threadIdx，blockDim来进行映射（偏移）。blockDim.x=8,blockDim.y=2

--------------------------------------------------------------------------------------

线程参数设置 情况3：一维grid，一维block （block分配）

---------------------------------------------------

线程参数设置情况4: block和thread都分配

---

线程参数设置 情况5：二维grid，二维thread

示例：倒推其线程参数设置

它的线程参数设置是怎样的？线程索引怎么计算？

参数设置为：

总Thread数量： 84182*1 = 512

一维数组的线程索引计算方法：

二维数组的线程索引计算方法：

根据CUDA算力不同thread，block，gird在不同维度的大小是有限制的:

Cuda Wrap的限制：

发布于 2024-06-14 10:33・IP 属地浙江

前言

说到GPU估计大家都不陌生，但是提起gpu底层的一些架构以及硬件层一些调度策略的话估计大部分人就很难说的上熟悉了。当然这个不是大家的错，主要是因为Nv gpu的整个生态都是闭源的，所以大家了解起来就会有一些障碍。最近这半年笔者有幸参与了一些gpu的项目，在这个过程当中也花了一些时间去理了一下gpu相关的东西，故借这篇文章给大家简单介绍一下。下面的行文将基于以下三个层面进行阐述：

• CUDA编程模型
• GPU 底层硬件架构
• 硬件层的调度策略

CUDA编程模型

为了让习惯了以cpu为计算主体的广大开发者也能够快速的开发出基于gpu来进行计算的应用程序，英为达在2007年发布了一种新的编程模型框架cuda。简单来讲，cuda里面提供了基于gpu来进行并行计算的编程范式以及大量的api。这里需要强调的是基于cuda的应用程序，它的程序主体仍然运行在cpu上，开发者可以通过其提供的api将相关代码offload到gpu上去执行比如一些矩阵运算等。从大的层面来讲cuda编程主要可以分为下面三大步：

1. 通过cuda api(比如cudaMemcpy)将input data 从host memory copy到device memory
2. 通过cuda api将gpu code load到gpu上去执行
3. device将执行之后的结果dma到host memory注：host-> cpu server device->gpu为了让大家更好地去理解相关的流程，这里给大家先介绍一下cuda编程模型当中的一些核心概念。

kernel

对于操作系统同学来说此kernel非彼kernel，这里的kernel更准确的来说是叫核函数，在概念上跟大家熟悉的c++/c函数差不多，只不过它是在gpu上执行的。

图1 the kernel function execute gpu

每一个cuda 核函数的开始处都有一个__global__的关键字来进行申明。具体例子可以参考一下图1所示，可能大家会对图中的threadID感到有些摸不到头脑。下面我们就来介绍一下thead相关的概念。

thread blocks and Grid

为了能够更好地讲清楚thread blocks和grid的概念，我们先上一个基于gpu进行矩阵计算的代码sample。

// Kernel - Adding two matrices MatA and MatB
___global__ void MatAdd(float MatA[N][N], float MatB[N][N], float MatC[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        MatC[i][j] = MatA[i][j] + MatB[i][j];
}
 
int main()
{
    ...
    // Matrix addition kernel launch from host code
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks((N + threadsPerBlock.x -1) / threadsPerBlock.x, (N+threadsPerBlock.y -1) / threadsPerBlock.y);
    //核心函数launch
    MatAdd<<<numBlocks, threadsPerBlock>>>(MatA, MatB, MatC);
    ...
}

从上面的示例当中

MatAdd

是一个核函数，可以看到其具体实现里面有

blockId

和

threadId

的代码段，也就是说核函数会被拆分为多个thread去gpu上执行。接下来我们看一下

block

和

thread

的声明：

dim3 threadsPerBlock(16, 16);
dim3 numBlocks((N + threadsPerBlock.x -1) / threadsPerBlock.x, (N+threadsPerBlock.y -1) / threadsPerBlock.y);

cuda里面用关键字

dim3

来定义block和thread的数量，以上面来为例先是定义了一个

16*16

的2维threads也即总共有256个thread，接着定义了一个2维的blocks。因此在在计算的时候，需要先定位到具体的block，再从这个bock当中定位到具体的thread，具体的实现逻辑见

MatAdd

函数。再来看一下grid的概念，其实也很简单它是多个block组成的一个集合。thread、block 和grid的关系具体见下图：

图2 grid、block and thread

stream

中文翻译为"流"，它主要是通过提升kernel函数的并发性来提升整个计算的运行效率。下面我们来看一下在cuda编程模型当中具体是如何使用stream的。

cudaStream_t stream[nStreams];  
for (int i = 0; i < nStreams; i ++)  
{  
    checkCuda(cudaStreamCreate(&stream[i]));  
}  
for (int i = 0; i < nStreams; i ++)  
{  
    checkCuda(cudaStreamDestroy(stream[i]));  
}

上面所展示的是stream的创建和销毁，接下来我们来看一下如何使用stream

for (int i = 0; i < nStreams; i ++)   
{  
    int offset = i * streamSize;  
    checkCuda(cudaMemcpyAsync(&d_a[offset], &a[offset], streamBytes, cudaMemcpyHostToDevice, stream[i]));  
    kernel_function<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);  
    checkCuda(cudaMemcpyAsync(&a[offset], &d_a[offset], streamBytes, cudaMemcpyDeviceToHost, stream[i]));  
}

stream

具体用法如上面sample所示，如果你不显示的申请stream的话系统也会有一个default的

stream0

。大家可以从下面的这张图比较直观地看到两者在执行效率上的区别：

图3 cuda stream 串行和并行执行

GPU底层硬件架构

上面所讲的都是cuda层面的概念，下面我们来讲一下GPU的底层硬件架构。GPU最核心的东西就是SM(stream multiprocessor)，上面所讲的thread最终的执行体就是SM。因为nv每一代gpu的SM都有一些小的差别，所以这里我们从nv最近5代的gpu来简单讲一下。

Pascal架构之P100

Nvidia 在2016年发布了他的第一代数据中心专用gpu P100，也是在p100这代gpu当中首次引入了float16的支持。下面我们来看p100的sm的底层微架构，具体如下图所示：

图4 P100 SM

每张p100 gpu卡总共有56个SM，每个SM包含了两个SP(stream processor)，下面我们来详细介绍一下每个SP的相关组成。

• core也称之为cuda core，主要用来进行FP和INT的计算
• DP Unit主要是在HPC场景用来进行double precison 计算，而机器学习场景基本上不会用到
• SFU也是一个计算单元，它主要负责 sine, cosine, log and exponential等函数的计算
• LD/ST即load uint和store unit即内存控制器的常用组件
• Register File即寄存器组
• Tex即图形渲染时需要用到的内存

上面还有一些组件比如warp scheduler、dispatch unit，这些都会在后面的调度章节进行详细介绍。

Volta架构之V100

2018年NV推出了v100 gpu卡，每张卡拥有80个SM，每个SM包含了4个SP。v100 SM底层微架构如下图所示：

图5 v100 SM

相关的组件这里就不详细介绍了，从图中可以看到v100 sp跟p100 sp的区别主要有1)去掉了DP Unit从而为增加更多的FP、INT unit腾出位置。2)SFU 数量减少，多增加了TENSOR CORE。 因为是首次引入tensor core，这里我们来详细介绍一下tensor core的作用。它主要用来做矩阵的MAC运算即两个矩阵的乘积与另外一个矩阵的和。

图6 tensor core 4x4 Matrix Multiply and Accumulate

从图6可以看到tensor core MAC运算是支持混合精度运算的，这里需要强调的是MAC操作是在一个cycle里面完成的。具体来说gpu主要是通过FMA(Fused multiply-add)指令在一个运算周期内完成一次先乘再加的浮点运算。

图7 Multiplication and addition happen in one clock cycle also known as FMA

Turing架构之TU102/TU104/TU106

Turing系列的gpu卡也是在2018年发布的，与Volta系列不同的是Turing主打的是游戏加速场景。当然它也可以用在AI场景，比如T系列的tensor core除了支持FP16之外还支持INT8和INT4。

图8 turing SM

从图8可以看到T系列SM跟V系列SM不同之处在于引入了RT CORE，从turing spec里面可以知道它主要是用来加速3D场景ray tracing。

Ampere 架构之A100

Nvidia 在2020发布了Ampere系列gpu卡，从功能上来说它是V系列的继承者。A100 GPU拥有108 颗SM。

图9 a100 sm

A100 GPU上引入了第三代tensor core，新的TC支持了从INT4、INT8、FP16、TF32到FP64所有类型的数据运算。同时在性能上也要比V100 TC增加了很多，具体如下图所示：

图10 A100 tc throughput and Efficiency

Hopper 架构之H100

Nvidia在2022年3月发布了Hopper系列gpu也即H100，每张H100 gpu 卡拥有144 颗SM，其中每个SM包含4颗SP。

图11 H100 SM

如上图所示H100 的 SM当中引入了第四代的TENSOR CORE，与A100 的第三代tensor core相比H100的第四代tensor core在性能上又有了新的增强，具体如下：

图12

1. GPU Resource Management

GPU channel是GPU与CPU之间的桥接接口，通过CPU向GPU发送GPU指令的唯一通道，GPU channel包含了两类用于存储GPU指令的buffer：

GPU command buffer (也称之为FIFO push buffer)
Ring buffer (也称之为indirect buffer)，从上图中看出，这个buffer是环形结构的，即其容量是固定的，这也是为什么叫Ring buffer的原因吧
当GPU指令被写入到GPU command buffer时，系统还会向Ring buffer中写入与此指令所对应的packet，packet包含了此指令在GPU command buffer中的偏移位置与长度数据。

在执行指令的时候，GPU不是直接从GPU command buffer中读取数据，而是先经过Ring buffer读取出当前待处理指令的相关信息，再据此读取GPU command（这也是为什么Ring buffer被称之为indirect buffer的原因）。
————————————————

                        版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
                     

原文链接：https://blog.csdn.net/denglin12315/article/details/122728991

GPU indirect buffer

IB (Indirect Buffer)间接缓冲特定引擎的命令缓冲区。与直接向队列中写入命令不同，您可以将命令写入一块内存，然后将指向该内存的指针放入队列中。然后，硬件将跟随指针并执行内存中的命令，然后返回到环中的其余命令。

如图

作者：人间正道是沧桑a
链接：https://www.jianshu.com/p/e835413c2e1f
来源：简书
著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。

GPU 底层调度

再聊调度之前，我们还是先来重点介绍几个相关的概念：channel、tsg、runlist、pbdma。

• channel

  这是nv driver层的才有的概念，每一个gpu应用程序会创建一个或者多个channel。而channel也是gpu硬件(在gpu context 层面来说)操作的最小单位。

• tsg

  全称为timeslice group，通常情况下一个tsg含有一个或者多个channel，这些channel 共享这个tsg的timeslice。

• runlist

  多个tsg或者channel的集合，gpu硬件就是从runlist上选取channel来进行任务执行。

• pbdma

  全称为pushbuffer dma。push buffer可以简单的理解为一段主机内存，这段内存主要有cpu写然后gpu来读。gpu通过从pushbuffer 里面拿到的数据生成相应的

command(也叫methods)

和

data(address)

。而上面讲到的channel里面包含有指向pushbuffer的指针。

图13

结合图13再给大家理一下上面几个概念之前的一些关联。首先，runlist里面的每个entry就是一个channel，每个channel里面有

Inst Blk Ptr

也即instance块指针，这些指针分别指向保存gpu上下文的内存和push buffer也即上图当中的PB seg。

接着我们先来简单的描述一下gpu应用是如何通过channel来提交任务的，具体流程如下：

    Submitting new work to a channel involves the following steps:

     1. Write methods to a pushbuffer segment
     2. Construct a new GP entry pointing to that pushbuffer segment
     3. Update GP_PUT in USERD( User-Driver Accessible RAM) to indicate the 
         new GP entry is ready
     4. Request the doorbell handle from RM, given the channel ID
     5. Write the channel's handle to the NOTIFY_CHANNEL_PENDING register

相信大家结合上面的一些讲述应该比较容易看懂上面的提交流程这里就不再赘述了，接下来我们回到调度正题上来。上面说到了应用提交work的相关流程，那这个work提交之后呢？这就涉及到如何将这些任务进行调度和执行了，下面我们先上一个整体调度架构图

图14 gpu scheduler

gpu的整个调度结构如图14所示，从左到右依次为Application scheduler、stream scheduler、thread block scheduler和warp scheduler。下面我们来一一对他们进行介绍。

K8S

User scheduler

Application scheduler

通常情况下两个不同的gpu应用是不能同时占用gpu的计算单元的，他们只能通过时分复用的方法来使用gpu。具体来讲就是gpu按照FIFO的策略依次从runlist上拿取channel进行执行，每一个channel只能运行一定的时间，等时间片用完之后就会进行切换来运行其他的channel。但是这种时分复用的调度算法有一个缺陷就是如果App每次提交的任务都比较小就无法占满gpu SM从而导致了gpu 整体使用率比较低。为了解决这个问题，nvidia 又提出了一另外一种调试算法叫

Multi-Process Service

，我们也叫空分。在MPS的场景下它允许两个不同的应用能够在同一时刻去占用不同的gpu sm，从而来提高gpu的使用率。

图15 MPS

stream scheduler

当gpu从runlist里面取出channel之后会生成相应的command和数据，而每个stream里面包含了一系列的commands。由于不同的应用的stream是可以设置不同的优先级的，所以stream scheduler主要负责不同应用的stream的调度和抢占。

Thread Block scheduler

它主要负责将thread block assign给gpu的sm，完成thread block跟gpu sm之间的一一映射。通常能不能将一个 kernel的thread block assign给某个sm主要看SM上的计算能力。举个例子，假如说一个sm支持 2048 threads和32 blocks，那么如果某个kernel有64个threads和64个blocks则scheduler也只能选这个kernel一半的blocks去运行。

warp scheduler

通常情况下一个warp包含了32个thread，warp scheduler的主要作用就是从wrap中获取准备好的待执行的instruction，并把这些instruction分配给sm上的Disaptch Unit。接着Dispatch Unit会把这些指令发送到SM的SIMD core上执行。

总结

gpu上的其他细节还有很多，笔者这篇文章就当作抛砖引玉了。如果大家想更加深入的研究的话可以去看看nv的一些open gpu doc，另外就是官方放出来的一些开源代码。这些都是非常重要的研究材料，仔细研读之后应该会有一些启发

本节视频

深入GPU原理：线程和缓存关系【AI芯片】GPU原理01_哔哩哔哩_bilibili​www.bilibili.com/video/BV1bm4y1m7Ki/?spm_id_from=333.999.0.0&vd_source=bf331b9ca4fb3b040bf9d1e87899c075​编辑

内容参考

AISystem/02Hardware/03GPUBase at main · chenzomi12/AISystem (github.com)​github.com/chenzomi12/AISystem/tree/main/02Hardware/03GPUBase

特别说明：本文是对开源项目AISystem的内容贡献

@ZOMI酱