从图形渲染到通用计算的GPU基础知识(下)GPU与通用计算, GPGPU, NVIDIA GPU架构解析, NVIDIA

五. GPU与通用计算

HPC （High Performance Computing, HPC）高性能计算

高性能计算（High Performance Computing, HPC）是指通过集中使用多台计算机或多处理器系统来完成一项计算任务，以达到提高计算性能的方法。高性能计算的主要应用领域包括科学计算、工程计算、金融计算、医疗计算、气象计算、地球科学计算等。

在过去的几十年里，高性能计算主要依靠集中式计算机系统（如超级计算机）来提供大量的计算资源。然而，随着计算机硬件技术的不断发展，特别是图形处理器（GPU）技术的发展，高性能计算现在也可以通过并行计算来实现。并行计算是指同时进行多个任务，以提高计算效率的方法。

CUDA（Compute Unified Device Architecture）和OpenCL（Open Computing Language）是两种最常用的高性能并行计算技术，它们分别由NVIDIA公司和Khronos(柯罗诺斯) Group发展。这两种技术都提供了一种编程模型，以便于开发人员利用GPU进行高性能计算。

通用图形处理器（GPGPU）

通用图形处理器（GPGPU，General Purpose Graphics Processing Unit）最早由NVIDIA公司的Mark J. Harris(马克.J.哈里斯)于2002年提出。基于图形渲染管线的流水线特征，GPU本质上是一个可同时处理多个计算任务的硬件加速器。由于GPU中包含了大量的计算资源，Mark J. Harris自2002年就开始尝试在GPU上做通用并行计算方面的研究。不过，由于架构及编程平台的限制，研究人员还只能使用 OpenGL 或其他渲染语言编写通用计算程序，程序员需要将通用计算使用的数据结构翻译成图形结构（例如各种图元），因此开发调试困难较大。直到CUDA和OpenCL的出现，GPU才正式进入GPGPU时代。

图5-1 GPGPU的典型应用

从图形渲染到通用计算的转变

随着图形处理需求的日益复杂和硬件加速性能的不断完善,有越来越多的功能被添加到图形流水织中,形成了更为丰富的图形流水线操作步骤和流程。图形流水线中可编程的部分,则是GPU演化成GPGPU的基础。

固定功能流水线（Fixed Function Pipeline）

早期的图形硬件都是固定功能流水线，不可编程，GPU的每个阶段都是预定义的，用于执行特定的任务，例如变换和光照（T&L）、光栅化等。开发者不能编写自定义的着色器来控制这些阶段；所有的处理逻辑都是硬件决定的。这种架构在早期的GPU中非常常见。

可编程流水线（Programmable Pipeline）

在2001-2006年,可编程实时图形流水线的出现将顶点处理和片段处理移到了可编程处理器上。由于可编程性的引人，GPU不再是一个功能单一的设备，拥有了更好的可扩展性和适应性，GPGPU也正是在可编程图形流水线阶段开始发展起来的。在将GPU运用到科学计算上时，这些可编程的着色器和着色语言(为着色器编程的语言)就成了技术的核心。把算法用着色语言实现，再加载到着色器里，同时把原本的图形对象替换为科学计算的数据，这就实现了GPU对通用数据的处理。

图5-2 图形流水线，灰色部分是可编程的

图5-3 图形渲染流水线

分离着色器架构（Separate Shader Architecture）

随着图形技术的发展，GPU开始支持顶点着色器（Vertex Shaders）和片元着色器（Fragment Shaders）或称像素着色器（Pixel Shaders），这两种类型的着色器在早期是独立实现的。顶点着色器负责处理顶点数据，而片元着色器负责处理像素数据。这种架构允许更灵活的图形处理，因为开发者可以自定义这些着色器来控制图形渲染的详细过程，但顶点和片元（或像素）着色器依然是分开的，各自拥有专用的处理单元。

图 5-4 2006年 NVIDIA 7900GTX显卡的管线架构图

Vertex Engines: 顶点引擎（也称为顶点着色器）负责处理3D图形管线中的顶点数据。它们执行各种变换、照明和其他顶点级别操作，准备数据进行下一步的光栅化过程。
Z-Cull: Z-剔除是一种隐藏表面消除技术，用于提高渲染效率。在渲染过程中，它可以快速剔除那些不可见的像素（因为它们被其他几何体遮挡），从而减少不必要的像素着色器计算。
Triangle Setup/Raster: 三角形设置和光栅化是将顶点处理过的三维几何体转换为二维屏幕像素的过程。这涉及确定哪些像素属于给定的三角形，并准备这些像素进行后续的像素着色处理。
Shader Instruction Dispatch: 着色器指令分派是指GPU如何管理和分派着色器程序中的指令给不同的着色器核心。这是GPU在执行顶点和像素着色器程序时的一个关键组成部分。
24 Fragment Shaders: 片段着色器（也称为像素着色器）负责计算最终像素的颜色和其他属性。GeForce 7900 GTX具有24个片段处理单元，用于执行这些操作。
Fragment Crossbar: 片段交叉开关是一个数据交换机制，允许不同的片段着色器单元高效地交换和处理数据。这有助于提高GPU内部数据处理的灵活性和效率。
16 Raster Operation Pipelines: 光栅操作管线（ROP）用于处理渲染管线的最后阶段，包括混合、深度测试、模板测试等操作，最终将像素写入帧缓冲区。
Memory Partition: 内存分区涉及到GPU如何管理和分配其内存资源，包括显存的读写操作。这对于保持高性能和避免瓶颈非常重要。

统一着色器架构（Unified shader Architecture）

传统的GPU厂商通常采用固定比例的顶点着色器和像素着色器单元(比如经典的1:3黄金渲染架构),但这种做法常常会导致单元利用率低下的问题。比如,一段着色程序中包含10%的顶点着色器指令,剩下 90%都是像素着色器指令,那么顶点着色器一段时间内将处于空闲状态,反之同理。为解决这一问题,统一着色器架构整合了顶点着色器和像素着色器,这种无差别的着色器设计,使得GPU成了一个多核的通用处理器

在早期分离渲染结构下的GPU，顶点着色器（VS）和像素着色器（PS）的硬件结构是独立的（位置参考图4-1），它们各有各的寄存器、运算单元等部件。这样很多时候，会造成顶点着色器与像素着色器之间任务的不平衡。比如,一段着色程序中包含10%的顶点着色器指令,剩下 90%都是像素着色器指令,那么顶点着色器一段时间内将处于空闲状态,反之同理，对于顶点数量多的任务，像素着色器空闲状态多；对于像素多的任务，顶点着色器的空闲状态多，于是，为了解决VS和PS之间的不平衡，引入了统一着色器架构（Unified shader Architecture）。统一着色器架构整合了顶点着色器和像素着色器，此架构的GPU，VS和PS用的都是相同的Core。也就是，同一个Core既可以是VS又可以是PS。这种无差别的着色器设计，使得GPU成了一个多核的通用处理器。

****微软在设计Xbox360游戏机的时候，提出了统一渲染架构的概念。NVIDIA起初是比较抗拒，它认为传统的分离设计在硬件方面更加高效，但最后看到大势所趋、2006年，NVIDIA 公布了统一着色器架构(unified shader architecture)和其 GeForce8系列 GPU。AMD是2007年。统一渲染架构固然是微软为提升硬件效率所提出，但它进一步提升了GPU内部运算单元的可编程性，让GPU运行高密集度的通用计算任务就成为可能，这意味着GPU可以打破3D渲染的局限，迈向更为广阔的天地。

图5-5 为NVIDIA 第一代统一着色器架构（Tesla架构）的架构图，这种统一的 GPU结构，它以多个可编程流处理器(Streaming Processor，SP)组成的并行阵列为基础，统一了顶点、几何、像素处理和并行计算，而不像早期的GPU那样对每种类型都有专用的分立处理器。这个架构基于NVIDIA GeForce8800 GPU构建。它将112个SP阵列组织成了14个流多处理器(StreamingMultiprocessor，SM)，14个SM 又组成了7个纹处理簇(Texture Processing Cluster，TPC)、共享纹理单元和纹理 L1 缓存。纹理单元会将过滤后的结果传给SM，因为对于连续的纹理请求来说,支持的过滤区域通常是重叠的,因此一个小的L1纹理缓存可以有效地减少存储器系统请求。在统一的SM及其 SP核上，既可以运行包括顶点、几何及片段处理的图形应用，也可以运行普通的计算程序。处理器阵列通过一个内部互连网络与光栅操作处理器、L2纹理缓存、动态随机存储器(DRAM)和系统存储器相连。

图5-9进一步展示了图5-2所示的逻辑图形流水线中各个阶段是如何映射到图5-5的统一架构上的。从图5-9中可以看出，专用的图形处理单元与统一的计算处理单元有机地结合在一起，顶点处理、几何处理和像素处理等可编程着色器的处理过程都是在统一的SM阵列和SP单元上完成的。图形数据在图形和计算处理单元之间不断循环，完成原有的逻辑图形流水线的处理。

图5-5 2007年 GF 8800 管架构图

图5-6 中文对照

图 5-7 纹理/处理器集群（TPC）

图 5-8 流式处理器（SM）

图 5-9 图形流水线向统一化着色器架构的映射

通用计算API

上面提到了，在将GPU运用到科学计算上时，是把算法用着色语言实现。但着色器编程语言是为复杂的图形处理任务设计的，而非通用科学计算，所以在使用时需要通过一系列非常规的方法来达到目的。这种方式要求编程人员不仅要熟悉自己需要实现的计算和并行算法，还要对图形学硬件和编程接口有深入的了解，开发难度很高。于是专门用于通用计算（GPGPU，即General-Purpose computing on Graphics Processing Units）。这类API不仅限于图形渲染，而是利用GPU进行高性能计算任务。

计算API	维护组织	适用平台
CUDA	NVIDIA	主要用于NVIDIA的GPU
OpenCL	Khronos Group	跨平台，支持多种厂商的GPU、CPU、FPGA等

CUDA 与 OpenCL

CUDA

2007年 NVIDIA 发布 CUDA（Compute Unified Device Architecture，统一计算设备架构），这是是一种新型的硬件和软件架构，它将 GPU 视为数据并行计算设备，在其上进行计算的分配和管理，而无需将其映射到图形API。它可用于 GeForce 8 系列、Tesla 解决方案和一些 Quadro 解决方案。操作系统的多任务机制负责管理多个并发运行的 CUDA 应用序和图形应用程序对 GPU的访问。

OpenCL

2008年，苹果公司向Khronos Group提交了一份关于跨平台计算框架的草案，该草案由苹果公司开发，并与AMD、IBM、Intel和NVIDIA公司合作逐步晚上。这个跨平台计算框架就是OpenCL。20088年12月8日，OpenCL1.0技术规范发布。2010年6月14日，OpenCL1.1发布，2011年11月19日，OpenCL1.2发布，2013年11月19日，OpenCL2.0分布。OpenCL是一个异构并行计算平台编写程序的工作标准，此异构计算可映射到CPU、GPU、DSP和FPGA等计算设备。OpenCL提供了底层硬件结构的抽象模型，旨在提供一个通用的开发的API。开发人员可以编写在GPU上运行的通用计算程序，而无需将其算法映射到OpenGL或DirectX的3D图形的API上。

二者的区别

CUDA和OpenCL的主要区别在于它们的设计目标和使用范围。CUDA是NVIDIA专门为其图形处理器（GPU）设计的并行计算编程模型，主要面向NVIDIA的GPU。而OpenCL是Khronos(柯罗诺斯) Group为多种类型的并行计算设备（如GPU、DSP、FPGA等）设计的通用并行计算编程模型，主要面向多种类型的并行计算设备。

各自的优缺点

CUDA的优点：

高性能：CUDA利用GPU的高性能并行计算能力，提供了高性能的高性能并行计算。
易用性：CUDA提供了简单易用的编程模型和API，使得开发人员可以快速上手并行计算编程。
丰富的库支持：CUDA提供了丰富的数学库和优化库，使得开发人员可以轻松地利用这些库来提高程序性能。

CUDA的缺点：

平台限制：CUDA主要面向NVIDIA的GPU，因此在其他类型的并行计算设备上可能无法运行。

OpenCL的优点：

跨平台兼容性：OpenCL为多种类型的并行计算设备设计，可以在不同类型的并行计算设备上运行。
通用性：OpenCL提供了通用的编程模型和API，使得开发人员可以在不同类型的并行计算设备上编程。
开源性：OpenCL是开源的，因此开发人员可以自由地使用和修改OpenCL技术。

OpenCL的缺点：

性能：由于OpenCL需要考虑多种类型的并行计算设备，因此在某些应用场景下可能无法实现与CUDA相同的性能。

六. GPGPU架构相关

下文的软硬件相关架构信息均以NVIDIA的GPGPU和CUDA 为主。

并行体系结构

GPGPU通常由成百上千个架构相对简易的基本运算单元组成。在这些基本运算单元中，一般不提供复杂的诸如分支预测、寄存器重命名、乱序执行等处理器设计技术来提高单个处理单元性能，而是采用极简的流水线进行设计。每个基本运算单元可同时执行一至多个线程，并由GPGPU中相应的调度器控制。GPGPU作为一个通用的众核处理器，凭借着丰富的高性能计算资源以及高带宽的数据传输能力在通用计算领域占据了重要的席位。虽然各个GPGPU厂商的芯片架构各不相同，但几乎都是采用众核处理器阵列架构，在一个GPU芯片中包含成百上千个处理核心，以获得更高的计算性能和更大的数据带宽。

GPU中执行的线程对应的程序通常成为内核（kernel），这与操作系统中的内核是完全不同的两个概念。除此之外，GPU中执行的线程与CPU或者操作系统中定义的线程也有所区别，GPU中的线程相对而言更为简单，所包含的内容也更为简洁。在GPU众核架构中，多个处理核心通常被组织成一个线程组调度执行单位，线程以组的方式被调度在执行单元中执行，如NVIDIA的流多处理器、AMD的SIMD执行单元。同一个线程组中的线程执行相同的程序指令，并以同步的方式执行，每个线程处理不同的数据，实现数据级并行处理。不同GPU架构对线程组的命名各不一样，如NVIDIA将线程组称为warp，AMD将线程组称为wavefront。线程组中包含的线程数量各不相同，从4个到128个不等。除此以为，线程组的组织执行模式也各不相同，常见的执行模式有SIMT（Single Instruction Multiple Threads，单指令多线程）执行模式和SIMD（Single Instruction Multiple Data，单指令流多数据流）执行模式两种。

SIMD

单指令流多数据流(Single Instruction Stream &.Multiple Data Stream,SIMD)SIMD是一种典型的并行体系结构，采用一条指令对多个数据进行操作。向量处理器就是SIMD的典型代表。很多微处理器也增加了 SIMD模式的指令集扩展。在实际应用中SIMD通常要求问题中包含大量对不同数据的相同运算(如向量和矩阵运算)。通常情况下,SIMD需要有高速I/O及大容量存储来实现高效并行。GPGPU也借鉴了 SIMD的方式,通过内置很多 SIMD处理单元和多线程抽象来实现强大的并行处理能力

SIMT

单指令多线程（Single Instruction Multiple Threads）是SIMD的升级版，由NVIDIA公司生产的GPU引入。这个概念与SIMD非常相似，它也是一次性给多个线程发出相同的指令，让它们同时执行。但SIMT与SIMD的不同在于，SIMT可以更好地处理条件分支。在SIMD中，如果线程需要根据条件选择执行不同的操作，那么可能会浪费一些计算资源。但在SIMT中，即使线程需要根据条件选择执行不同的操作，GPU也能保证所有线程都能充分利用。

SIMD和SIMT的主要区别在于如何处理条件分支。在SIMD中，如果一条指令的执行依赖于某个条件，而这个条件在数据之间有所不同，那么这就会导致问题。这是因为所有的数据都需要执行相同的指令，所以如果有一些数据需要执行一个操作，而其他数据需要执行另一个操作，那么这就会导致一些核心空闲，浪费计算资源。

举个例子，假设我们有一个SIMD设备，有4个处理核心，我们需要对一个数组进行操作，如果数组元素大于10，我们就给它加1，小于等于10的就给它减1。假设数组是[8,12,7,15]，那么第一个和第三个元素需要执行减1的操作，第二个和第四个元素需要执行加1的操作。在SIMD中，我们不能同时执行加1和减1的操作，因此我们必须分两步来进行，先执行所有的减1操作，然后执行所有的加1操作。在执行减1操作的时候，第二个和第四个核心就闲置了，同样，在执行加1操作的时候，第一个和第三个核心就闲置了，这就浪费了计算资源。

而在SIMT中，我们可以让每个线程独立地处理自己的数据元素。在上述例子中，我们可以让第一个线程处理第一个元素，第二个线程处理第二个元素，以此类推。因此，我们可以在一个线程上执行减1的操作，同时在另一个线程上执行加1的操作。这样就避免了因为条件分支而导致的核心闲置，提高了计算效率。

NVIDIA与通用并行计算

SP（streaming Process），SM（streaming multiprocessor）是硬件（GPU hardware）概念。而thread（线层），block（线程块），grid（线程网格），warp（线程束）是软件上的（CUDA）概念。

CUDA

要理解NVIDIA 并行计算的硬件架构，就得先了解下CUDA。NVIDIA CUDA（Compute Unified Device Architecture，统一计算设备架构），是一种新型的硬件和软件架构，当通过CUDA编译时，我们可以将GPU视为一种强大的并行计算设备，它能够同时执行大量的计算线程。在这种情况下，GPU充当主CPU的辅助处理器。换句话说，那些在主机上运行、需要处理大量数据和进行密集计算的应用程序部分，可以转移到GPU上进行处理。具体来说，应用程序中需要多次执行且每次执行时数据不同的那部分，可以被拆分成多个线程，这些线程在GPU上并行执行。为了实现这一点，相关的函数会被编译成GPU能理解的指令集（称为kernel），然后这个kernel被传输到GPU上执行。同时，主机（CPU）和设备（GPU）各自拥有独立的DRAM内存，分别称为主机内存和设备内存。数据可以通过使用高性能的直接内存访问（DMA）引擎，从一个内存（例如主机内存）复制到另一个内存（例如设备内存）。这种数据传输方式利用了GPU的高速DMA接口，可以有效提升数据处理的效率。

图6-1 CUDA支持相关

图6-2 CUDA软件堆栈

CUDA的软件堆栈(图6-1)由三层构成：CUDA Library、CUDA Runtime API、CUDA Driver API。 CUDA runtime API和CUDA driver API提供了实现设备管理（Device management)，上下文管理（Context management），存储器管理（Memory Control），代码块管理（Code Module management），执行控制（Excution Control），纹理索引管理（Texture Reference management）与OpenGL和Direct3D的互操作性（Interoperity with OpenGL and Direct3D）的应用程序接口。CUDA Library 提供了简单高效的常用函数。

相较于串行代码，在CUDA编程模型中引入了主机端和设备端的概念，其中CPU作为主机端（Host），GPU作为设备端（Device）。CPU负责任务的调度，数据的传输、逻辑处理以及运算量少的计算， GPU硬件主要通过“CUDA核（SP）”进行并行计算。将在 CPU 上执行的代码称为主机代码，在 GPU 上运行的代码称为设备代码。设备端代码又称为核函数（kernel）。Host端的串行程序和Device端的并行程序可以各自独立运行，如图6-3所示。GPU程序可以异步执行，当CUDA程序中在GPU开始执行后，程序的流程控制权立刻交还给Host端串行程序，即CPU可以在GPU进行大规模并行运算时进行串行运算，提高异构设备的运行效率。

图6-3-1 CUDA程序执行流程

图6-3-2 CUDA程序执行流程

图6-4 CUDA线程组织模型

图6-4中可以看到CUDA的简单模型。其中图里的Host指的是CPU端，Device指的是GPU端，Kernel指的是具体执行的程序算法或者说是一个函数。数个线程(Thread)组成一个线程块(Block)，数个线程块组成一个线程网格(Grid)。

内核函数(Kernel)

内核函数是CUDA每个线程执行的函数

kernel：GPU 上运行的代码称为设备代码。设备端代码又称为核函数（kernel）,下面是一段简单的内核函数，两个大小为 N 的向量 A 和 B 相加，并将结果存储到向量 C 中：

// Kernel definition

global void VecAdd(float* A, float* B, float* C)

{

    int i = threadIdx.x;

    C[i] = A[i] + B[i];

}
int main()

{

    ...

    // Kernel invocation with N threads 执行 VecAdd() 的 N 个线程中的每一个线程都会执行一个加法。

    VecAdd<<<1, N>>>(A, B, C);

    ...

}

在CUDA中，线程索引以三维形式表示，对应于线程块（Block）内部线程的布局也是三维的。每个线程都有一个三维索引（threadIdx.x,threadIdx.y,threadIdx.z），用于在其所属的线程块内标识其位置。同时，线程块本身在网格（Grid）中也有一个三维索引（blockIdx.x,blockIdx.y,blockIdx.z），用于在整个网格中标识线程块的位置。这样，结合网格和块的维度信息（blockDim和gridDim），可以为网格中的每个线程计算出一个全局唯一的索引。

threadIdx.x是线程在所处线程块中的X方向的索引。由于本例中是定义的1维排布，因此X方向索引即为线程的ID。

由于GPU中的每个线程都会执行相同的VecAdd函数，因此不同的线程需要使用自己独有的ID来区分彼此，来获取不同的数据。这就是SIMT的概念，即“相同指令，不同线程”。

线程层次

为方便起见，threadIdx 是一个 3 分量(3-component)向量，因此可以使用一个一维、二维或三维的线程索引(thread index)来识别线程，形成一个具有一个维度、两个维度或三个维度的、由线程组成的块，我们称之为线程块(thread block)。

线程的索引和它的线程 ID 以一种直接的方式相互关联：对于一维块，它们是相同的；对于大小为(Dx, Dy)的二维块，索引为(x, y)的线程的线程ID为(x + yDx)；对于大小为 (Dx, Dy, Dz) 的三维块，索引为 (x, y, z) 的线程的线程 ID 为 (x + yDx + zDxDy)。

例如，下面的代码将两个大小为NxN的矩阵A和B相加，并将结果存储到矩阵C中:

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
                       float C[N][N])
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}
int main()
{
    ...
    // Kernel invocation with one block of N * N * 1 threads
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

每个块的线程数量是有限制的，因为一个块中的所有线程都应当驻留在同一个处理器核心上，并且共享了该核心有限的内存资源。在当前的GPU中，一个线程块最多包含1024个线程。但是，一个内核可以由多个形状相同的线程块执行，因此线程的总数等于每个块的线程数乘以块数。块被组织成一维、二维或三维的线程块网格(grid)，网格中的每个块都可以由一个具有一维、二维或三维的唯一索引进行识别，该索引可以通过内置的blockIdx变量在内核中访问。线程块的维度可以通过内置的blockDim变量在内核中访问。

扩展前面的MatAdd()示例对多个块进行处理，代码如下所示。

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}
int main()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

硬件实现

NVIDIA GPU 架构通过多线程流式多处理器 (SM: Streaming Multiprocessors) 可扩展阵列构建。当主机 CPU 上的 CUDA 程序调用内核网格（grid）时，网格的块(block)被枚举并分发到具有可用执行能力的多处理器(SM)。一个线程块(block)的线程(thread)在一个 SM 上同时执行，同时多个线程块也可以在一个 SM 上同时执行(一个block只能调度到一个Streaming Multiprocessor上运行。一个Streaming Multiprocessor可以同时运行多个block)。当线程块终止(block)时，新块(block)在空出的 SM 上启动。SM 旨在同时执行数百个线程。为了管理如此大量的线程，SM 采用了一种称为 SIMT（Single-Instruction, Multiple-Thread: 单指令多线程）的独特架构（NVIDIA第一代Tesla架构首次使用了SIMT）。

SM 多线程流式处理器

SP（streaming processor）： 最基本的处理单元，也称为CUDA core。负责执行浮点运算和整型运算等指令。在NVIDIA的术语中，SP可以看作是GPU中负责执行具体计算任务的核心。具体的指令和任务都是在SP上处理的，GPU进行并行计算，也就是很多个SP同时做处理。

图6-5-1早期SM的结构

图6-5-2 早期的cuda core内部分单精度浮点型运算单元和整型运算单元，后来被拆分

图6-5-3 Volta架构下SM被重新设计，CUDA core被拆分

SM（streaming multiprocessor）： 多个SP加上其他的一些资源组成一个SM，也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。图6-6是一个SM的基本组成，其中每个绿色小块代表一个SP。

图6-6-1 软硬件层级对比

图 6-6-2 Pascal 架构下的SM

Warp

SIMD各执行单元会共享register file等资源（还有cache、memory），SIMT则对于每个执行单元（thread）具有独立的memory、register file等资源。 ···

Warp是SM的基本执行单元。一个Warp包含32个并行thread，这32个thread执行于SIMT模式。也就是说所有Thread执行同一条指令，并且每个Thread会使用各自的数据执行该指令。

SM的 SIMT 单元以 32 个并行线程为一组来创建、管理、调度和执行线程，这样的线程组称为 warp （线程集束）。一个warp中的线程必然在同一个block中，如果block所含线程数目不是warp大小的整数倍，那么多出的那些thread所在的warp中，会剩余一些inactive的thread，也就是说，即使凑不够warp整数倍的thread，硬件也会为warp凑足，只不过那些thread是inactive状态，需要注意的是，即使这部分thread是inactive的，也会消耗SM资源。由于warp的大小一般为32，所以block所含的thread的大小一般要设置为32的倍数。

当一个 SM 执行一个或多个线程块时，它将线程块拆分为 Warp，每个 Warp 都由 Warp 调度程序来调度执行。每个块被分割成 Warp 的方式是一样的；每个 warp 包含连续的线程，这些线程 ID 是不断增加的，第一个 Warp包含线程0。线程层次结构描述了线程 ID 如何对应到块中的线程索引。

如图6-7-1所示，一个 Warp 一次执行一条公共指令，因此当一个 Warp 中的 32 个线程都执行相同的执行路径时，Warp 效率才完全发挥出来。如果 Warp 中的线程由于数据依赖而发生条件分支发散，则 warp 会执行每个需要的分支路径，同时禁用不在该路径执行的线程。分支发散只发生在一个 Warp 内；不同 Warp 之间独立执行，无论是执行带有分叉的还是不带有分叉的代码路径。

在SIMT架构中，尽管每个线程有能力独立执行指令，但它们仍然被组织在warps中进行协同工作，以提高效率。分支（decision-making）是SIMT架构处理的关键特点，它允许在同一个warp中的线程遵循不同的执行路径，而不是像SIMD那样所有元素都必须执行完全相同的操作。但是，分支会导致效率下降，因为不是所有线程都在每个时刻都在执行有用的工作。因此，在SIMT架构中编程时，开发者会尽量减少这种分支，以最大化效率。

图 6-7-1

Warp的执行和调度：

Warp: 在NVIDIA的CUDA架构中，一个Warp是由32个线程组成的组，这些线程可以同时执行相同的指令序列。
Warp Scheduler: 图6-7-1 显示了两个Warp调度器，每个调度器负责一组Warp的执行。调度器将指令分配给Warp中的线程，并管理执行的时间安排。

指令的执行：

从图6-7-1中可以看到，多个Warps被同时调度并执行指令。每个Warp按顺序处理指令，但所有线程在执行指令时必须同步，即一个Warp内的所有线程同时执行同一指令。
如果Warp中的某些线程需要执行不同的执行路径（如if/else结构所示），会导致执行分岔，称作“divergence”。这会导致整个Warp的执行效率下降，因为所有线程必须等待最慢的路径完成后才能继续同步执行。

warp不应该是SIMT体系结构的么，理论来说它下面的每一个进程不应该都可以单独执行某个分支么，为什么还出现上图的分批执行和等待？

确实，SIMT允许每个线程（在Warp中的一个线程）独立地选择执行路径，但这并不意味着没有成本。下面详细解释为什么会出现分批执行和等待的情况：

SIMT架构的工作原理

在NVIDIA的SIMT架构中，一个Warp是32个线程的集合，这些线程同时执行相同的指令序列。理论上，每个线程可以独立地选择其执行路径（例如，if-else语句中的不同分支）。然而，Warp中所有线程同步执行同一指令的特性导致了一些限制：

分支发散（Branch Divergence）

当Warp中的线程遇到一个分支指令（如if-else结构）时，不同的线程可能会选择不同的执行路径。这被称为“分支发散”。在SIMT架构中，尽管每个线程理论上可以单独执行，但由于硬件的限制和效率考虑，GPU仍然会尝试将这些线程同步执行。

执行和等待的过程

当发生分支发散时，Warp调度器必须处理所有可能的执行路径。GPU会先执行一个分支的所有线程，然后保存这部分线程的状态，再回来执行另一个分支的所有线程。这意味着即使某些线程已经完成了它们的执行路径，它们也必须等待其他线程完成其路径，直到所有线程都能再次同步。这种执行方式确保了硬件的高效利用，但也可能导致等待和执行效率下降。

性能影响

分支发散会导致线程执行效率下降，因为不同的执行路径需要被顺序处理，而非真正并行执行。这种效率下降尤其在Warp中有大量分支发散时显著。因此，在编写GPU代码时，通常建议尽量减少分支发散，以提高执行效率和减少线程等待时间。

总的来说，尽管SIMT允许每个线程理论上可以独立执行，但实际上为了保证执行的效率和硬件资源的合理利用，Warp中的线程执行可能会因为分支发散导致必须顺序处理不同的执行路径。这就是为什么会出现分批执行和等待现象。

很多文章会说SP对应一个线程，然而并不是这样的，官方论坛的人回复，SP只是负责执行浮点运算和整型运算等指令，它可能都不会完整的执行完一个线程。SP 最直接指的是浮点 ALU。它处理浮点加法和乘法，但一般情况下不处理其他指令。例如，如果有一个整数加法指令，SP 将不会被调度浮点 ALU来处理该指令，而是整数 ALU。

SP与线程的关系
- 一个SP在任意时刻可以执行来自某个线程的指令，但它并不是“一个线程”。GPU上的线程调度单位是warp，每个warp包含固定数量的线程（例如，在很多NVIDIA架构中，一个warp包含32个线程）。
- Warp执行模型：当一个warp被选中执行时，它的线程会被分配到可用的SP上。如果一个SM（流式多处理器）有足够的SP，一个warp中的线程可以同时执行。否则，这些线程的执行可能会分成几个步骤完成，这取决于可用的硬件资源和具体的指令。
- 因此，一个SP不是直接与一个单独的线程绑定的。它根据调度策略执行warp中线程的指令，这些指令可以来自不同的线程。
为什么不能将SP等同于线程
- 资源共享与调度：在GPU内部，资源（如SP）是被多个线程共享的。线程的执行并不是绑定到特定的SP上，而是根据GPU调度器的策略动态分配的。
- 执行效率：通过将线程以warp为单位组织和执行，GPU能够高效地管理和执行并行计算。这种设计允许在硬件层面上实现指令级并行和数据级并行。

CUDA为什么要分线程块和线程网格？

主要是硬件的性能限制（通信和共享内存）和并发计算的考虑。

图6-7-2 CUDA内存模型

图6-7-3 SM内存模型

存储器	片上/片外	缓存	存取	范围	生命周期
寄存器	片上	n/a	R/W	一个线程	线程
本地	片外	1.0以上有	R/W	一个线程	线程
共享	片上	n/a	R/W	块内所有线程	块
全局	片外	1.0以上有	R/W	所有线程+主机	主机配置
常量	片外	Yes	R	所有线程+主机	主机配置
纹理	片外	Yes	R	所有线程+主机	主机配置

首先，GPU在硬件层次上并不知道什么“三维”线程。在硬件看来所有东西都是一维的。设计成三维的索引方式是因为许多并行计算的程序的本身具有维度这个概念。例如图像处理程序通常是在一个二维平面上做并行计算，许多科学计算则是在一个三维的空间里并行，如下图。

这样的话，写程序的时候你就不用写这样的代码了。直接用threadIdx.x, threadIdx.y就可以了。

x = threadIdx % width;y = threadIdx / width;

图6-8

共享内存和同步

共享内存：在一个块内的所有线程可以访问相同的共享内存，这使得它们可以高效地共享数据和协调操作。如果没有块这一层次，线程间的这种高效通信将变得困难。
同步：块内的线程可以通过同步原语（如__syncthreads()）进行同步，这对于保证数据一致性和执行正确性非常重要。

2 . 资源分配和管理

CUDA将GPU资源（如寄存器、共享内存）分配给线程块，而不是单个线程。块的存在使得资源管理更为高效，因为它允许集中管理和优化资源的使用。

并行执行的层次化

硬件映射：线程块的概念与GPU硬件的物理结构紧密相关。在GPU中，多处理器可以执行一个或多个块，每个块内的线程可以映射到多处理器上的核心。
负载平衡：块为负载平衡提供了一个自然的单位。GPU可以在块之间分配工作负载，使得多处理器可以同时处理多个块。

编程灵活性

块提供了一个中间层级，允许程序员根据应用需求和GPU的硬件特性来优化线程的组织和调度。这种灵活性是高效利用GPU的关键。

性能优化

在块内部，线程可以通过共享内存进行高速数据交换，这比通过全局内存进行通信要快得多。缺少块层次会限制这种优化手段的使用。如果去掉块这一层次，CUDA编程模型将无法有效地利用GPU的并行计算能力，同时会失去很多关于资源管理、线程通信和同步的优化手段。
从程序员的角度看，貌似grid-thread两级层次就够了，不同grid表征不同任务，不同thread表征不同执行流，这么思考很合理。block这个层次，更多的是软件对硬件性能的妥协，因为硬件如果想要性能好，势必要有局部性。有了DDR，为什么还有cache呢？因为想利用局部性，提升性能。CACHE可以做到对程序员不可见，BLOCK很难做到程序员不可见。（除非牺牲大量性能。）（如果想追求极致性能，程序员也需要看见CACHE。）

浮点运算单元的浮点型

FP32（单精度浮点数）

精度：32位
结构：1位符号位，8位指数位，23位尾数位（小数部分）
用途：FP32是最常用的浮点数格式之一，广泛应用于科学计算、3D图形处理以及深度学习等领域。它提供了足够的精度，满足大多数应用的需求，同时保持了较好的性能。

FP64（双精度浮点数）

精度：64位
结构：1位符号位，11位指数位，52位尾数位
用途：FP64提供了更高的计算精度，适用于需要极高数值精度和范围的应用，如科学研究、天文学计算、精确工程模拟等。由于其较高的计算开销，一般只在精度非常关键的场合使用。

FP16（半精度浮点数）

精度：16位
结构：1位符号位，5位指数位，10位尾数位
用途：FP16提供了较低的数值精度，但在某些场合下足以满足需求，尤其是在深度学习训练和推理中。使用FP16可以减少内存使用、加速计算以及降低能耗，对于在资源受限的设备上运行模型特别有用。

FP8（8位浮点数）

精度：8位
结构：一种常见的分配方式是1位符号位，4位指数位，3位尾数位，但具体的位分配可能根据实现有所不同。
用途：FP8是一种较新的格式，主要用于深度学习领域，尤其是模型推理阶段，以进一步减少模型的大小和计算需求。它使得在移动设备和边缘计算设备上运行复杂的神经网络模型成为可能。

七. NVIDIA Tesla系列GPU

NNVIDIA GPU 三大系列

自1999年起，NVIDIA发布了世界上第一块可以被称作GPU的GeForce 256，逐渐形成了以家庭娱乐的GeForce、专业绘图工作站的Quadro以及服务器高性能计算的Tesla为核心的三大GPU产品系列。其中GeForce与Quadro系列类似，它们主要面向图形处理，而Tesla系列专注于计算，即使它们的核心采用同代架构，但在细节上会因计算场景的不同而会有一些差异，下图是Ampere架构下的分别用于计算以及图形的SM架构图。

命名规则

GeForce系列

在NVIDIA显卡中，从2004年开始就有【GT】的代号，代表着中低端显卡，到了2005年引入了【GTX】的代号，直接代表着高端或者顶级显卡。10系列显卡开始, 没有GT和GTS的代号, 仅保留GTX，在新的20系列显卡中，由于新增了光线追踪的特性，因此由原来的GTX前缀改为RTX，这里的【RT】就代表着光线追踪（ray tracing的缩写）

Tesla系列

NVIDIA一般以历史上一些著名科学家的名字命名自己的GPU微架构，上面8种微架构分别是：特斯拉（Tesla），费米（ Fermi ），开普勒（ Kepler ），麦克斯韦（ Maxwell ），帕斯卡（ Pascal ），伏打（ Volta ），图灵（ Turing ），安培（ Ampere ）, 赫柏（ Hopper ）, 布莱克韦尔（Blackwell）。

Quadro系列

NVIDIA已经开始逐步将Quadro系列产品线整合到了其RTX系列之下，并逐渐停止了Quadro品牌的使用。这一变化主要发生在2020年左右，标志着NVIDIA对其产品线进行了重组和品牌策略的调整。Quadro系列长期以来一直是面向专业工作站市场的高性能图形卡，特别是在CAD（计算机辅助设计）、3D渲染、视频编辑和计算密集型应用程序中广受欢迎。然而，随着技术的发展，NVIDIA推出了RTX系列，该系列不仅包括面向游戏市场的高性能图形卡，也包括了面向专业市场的产品，这些专业产品继承了Quadro系列的定位，提供了高精度的计算能力、更大的显存以及针对专业应用优化的特性。

NNVIDIA GPU 市场占有率

英伟达目前在数据中心AI市场拥有98% 的市场份额，而AMD仅有1.2%的市场份额，英特尔则只有不到1%。
英伟达在全球离散GPU市场份额中占据了80.3%的份额，是市场的绝对领导者。 AMD排名第二，占据了19.7%的份额，是英伟达的主要竞争对手。英特尔则在整个GPU市场中占据了1.4%的份额，主要是由于其集成显卡在PC市场中的份额。
英伟达数据中心业务占公司收入达到83%

NVIDIA GPU关键技术

Core 计算核心

CUDA Cores（计算单元）

CUDA Cores，或称为计算单元（早期架构中叫SP），是NVIDIA GPU的基础处理单元，用于处理图形和一般计算任务。它们是执行着色、物理模拟、像素处理等操作的核心组件。CUDA Cores适用于广泛的并行计算任务，包括图形渲染、视频处理、科学计算和深度学习等。CUDA编程模型允许开发者利用这些核心来加速计算密集型应用。

RT Cores（光线追踪单元）

RT Cores是专门设计用于加速光线追踪计算的处理单元。光线追踪是一种高级渲染技术，能够模拟光线如何在真实世界中与物体互动，产生逼真的视觉效果。RT Cores能够加速光线与场景中物体相交的计算，大幅提高光线追踪性能。RT Cores主要用于图形渲染，特别是在游戏、视觉效果和建筑可视化中，为用户提供更加逼真的光照和阴影效果。

图7-2-1 未开启光追

图7-2-2 开启光追

Tensor Cores（张量单元）

Tensor Cores是专门为加速深度学习算法而设计的处理单元。本质上， “Tensor core" 是加速矩阵乘法的处理单元。它们能够高效地执行深度学习所需的矩阵乘法和累加运算，特别是在处理混合精度计算时显著提高性能。这使得Tensor Cores在训练和推理深度神经网络时非常有效。Tensor Cores主要用于加速AI和深度学习应用，包括图像和语音识别、自然语言处理、推荐系统等。它们是现代AI加速计算的关键组件。

DLSS

DLSS 全称Deep Learning Super Sampling（深度学习超级采样）。它最初的目的是为了弥补在使用光线追踪功能时的显卡性能损失，借助这项独特的AI技术，DLSS可以在维持显示质量的同时显著提高游戏的性能，正因为此，后续迭代发布的DLSS 2.0也成为了NVidia显卡的一项独立功能。Nvidia 使用神经网络超级计算机，通过机器深度学习来训练其人工智能，具体方式是不断地让人工智能自动比对超高分辨率静态图像（SSAA超级采样抗锯齿版本）和低分辨率图像之间的差异，使人工智能有能力推理出把图像从较低的分辨率提升到较高的分辨率的方法。也就是说，这项技术以较低的分辨率渲染图形以保持较高的性能，然后应用各种效果输出高分辨率图像的整体效果。

这项技术对于需要高帧率（大于60 FPS）和高分辨率（4K） 运行的游戏特别有用。例如，使用 DLSS功能，以1080p分辨率渲染游戏帧，这种负载更容易获得更高的帧速率，然后放大图形以4K分辨率输出，从而带来超高清图像显示质量。依赖于RTX 显卡的特殊硬件Tensor核心。所以目前只有RTX 20系列和RTX 30系列才有这项功能。不过 RTX 30系列配置的是 Nvidia 第二代 Tensor 核心，DLSS性能更好。

分布式通信

NVLink

在进入大模型时代后，大模型发展已是人工智能的核心，但训练大模型实际上是一项比较复杂的工作，因为它需要大量的 GPU 资源和较长的训练时间。此外，由于单个 GPU 工作线程的内存有限，并且许多大型模型的大小已经超出了单个 GPU 的范围。所以就需要实现跨多个 GPU 的模型训练，这种训练方式就涉及到了分布式通信。在传统的方案中，GPU互联通常采用PCIe。

了解决这个问题，nvidia开发了一个全新的互联构架nvlink。单条nvlink是一种双工双路信道，第一版的实现被称为nvlink 1.0，与P100 GPU一同发布。一块P100上，集成了4条nvlink。每条link具备双路共40GB/s的带宽，整个芯片具备整整160GB/s的带宽。

当然，nvlink不仅仅只是限定在GPU之间互联上。IBM将nvlink 1.0添加到他们基于Power8+微架构的Power处理器上，这一举措使得P100可以直接通过nvlink于CPU相连，而无需通过pcie。通过与最近的power8+ cpu相连，4GPU的节点可以配置成一种全连接的mesh结构。

图7-3-1 传统互联方式

图7-3-2 NVLNK GPU与GPU互联方式

图7-3-3

NVSwitch

NVSwitch芯片是一种物理芯片，类似于交换机ASIC，可通过高速的NVLink接口连接多个GPU，提高服务器内的通信和带宽，允许多个GPU以全连接的方式互连，从而实现任意GPU对之间的高速通信。NVSWITCH设计用于构建较大的GPU集群，其中每个GPU都需要与集群中的其他所有GPU高效通信。

图7-5-1

###### 区别

设计目标：NVLINK主要针对单个GPU与GPU或GPU与CPU之间的高速通信，而NVSWITCH旨在实现多GPU系统中的全互联，支持更复杂的多GPU配置。
通信拓扑：NVLINK提供点对点的通信路径，NVSWITCH则实现了全连接网络，支持任意GPU之间的直接通信。
应用场景：NVLINK适合需要高速点对点通信的场景，而NVSWITCH适用于需要大规模GPU协作的高性能计算和深度学习训练场景

MIG(GPU虚拟化)

NVIDIA Multi-Instance GPU (MIG) 技术是 NVIDIA 推出的一种 GPU 虚拟化技术，允许一块物理 GPU 被分割成多个独立的 GPU 实例，每个实例可以被分配给不同的虚拟机、容器或用户。这种技术有助于更有效地利用 GPU 资源，提高 GPU 的共享性和多租户支持。

MIG 技术通常需要硬件和软件支持，包括支持 MIG 的 NVIDIA GPU 和相应的驱动程序。这使得 MIG 技术成为数据中心和云计算环境中更好地管理 GPU 资源的有力工具。它有助于提高 GPU 利用率，降低成本，并更好地满足不同应用程序和用户的需求。

MIG 技术关键特点

资源划分：MIG 允许将一块物理 GPU 分割成多个 GPU 实例，每个实例具有自己的 GPU 核心、GPU 内存、NVLink 带宽等资源。这样可以更好地控制和划分 GPU 资源。
多租户支持：MIG 技术可以用于虚拟化 GPU，以便不同用户或应用程序可以共享同一块物理 GPU 而不会相互干扰。
动态资源调整：管理员可以根据工作负载的需求动态地重新配置 MIG 实例的资源，从而实现更好的资源利用和性能。
容错性：MIG 技术支持 GPU 实例的隔离，这意味着一个 GPU 实例中的问题不会影响到其他实例，从而提高了系统的容错性。
部署灵活性：MIG 技术可以用于云计算、虚拟化环境、容器化应用程序等多种情境，为不同的部署需求提供了灵活性。

图7-6-1 MIG

图7-6-2 支持MIG的显卡

#### NVIDIA GPGPU : Tesla系列

架构演进

图 7-7-1 Tesla系列显卡架构演进

图7-7-2 Tesla系列显卡算力版本

##### 里程碑架构

Fermi架构

Fermi 是第一个为高性能计算(High Performance Computing,HPC)应用提供所需功能的架构，支持符合 IEEE 754-2008 标准的双精度浮点，融合乘加运算(FusedMultiply-Add,FMA)，提供从寄存器到 DRAM 的 ECC(Error Correcting Code)保护，具有多级别的缓存,并支持包括 C、C++、FORTRAN、Java、MATLAB 和 Python 等编程语言通常，Fermi架构被认为是第一个完整的 GPGPU计算架构,它实现了图形性能和通用计算并重,为通用计算市场带来前所未有的变革。

图7-2-3 Fermi架构

图7-2-4 Fermi SM

图7-2-5 Fermi架构下的逻辑渲染关系

Fermi架构下的图形渲染管线

程序在图形API中创建一个drawcall。之后经过一些验证后到达驱动，然后用GPU可读的编码格式将指令插入到Pushbuffer（推送缓存）。这个阶段会产生许多跟CPU侧相关的瓶颈，这就是为什么程序员用好API以及所用技术能充分利用上现代GPU的性能十分重要。
经过一段时间或者直接的“刷新”调用后驱动上已经缓存了足够多的工作到Pushbuffer并将它发送给GPU进行处理（以及一些操作系统的参与）。GPU的主接口（Host Interface）获取到经过前端（Front End）处理的指令。
之后我们开始在基元分配器（Primitive Distributor）中开始分配的工作。在这里会对索引缓冲中的序列进行处理，并且批量生成之后我们会将其发送给多个GPC上的三角形。

Kepler 架构

Kepler 是为了高性能科学计算而设计的，相比Fermi架构效率更高，性能更好,其突出优点是双精度浮点运算能力高并且更加强调功耗比。但是双精度能力在深度学习训练上作用不大，所以 NVIDIA 又推出了 Maxwell 架构来专门支持神经网络的训练。Maxwell架构支持统一虚拟内存技术，允许CPU直接访问显存和 GPU访问主存。随后,NVIDIA又推出了 Pascal架构，进一步增强了 GPGPU 在神经网络方面的适用性。

Fermi，Kepler，Maxwell架构演进：主要还是在SM内堆SP

图7-2-6 架构演进

Volta 架构

Volta 对 GPU的核心，即流多处理器(SM)的架构进行了重新设计 。Volta 架构比前代 Pascal设计能效高 50%。在同样的功率范围下,单精度浮点(FP32)和双精度浮点(FP64)性能有大幅提升。Volta架构还新增了专门为深度学习而设计的张量核心(tensor core)单元。Volta架构的另一个重要改动是独立的线程调度。之前 NVIDIA 一直使用SIMT 架构，即一个线程束(warp)中的 32 个线程共享一个程序计数器(Program Counter,PC)和栈(stack)。Volta架构中每个线程都有自己的程序计数器和堆栈，使得线程之间的细粒度控制成为可能.

Volta架构的CUDA Core开始区分不同类型的CUDA核心（如专门的FP32 Core、INT32 Core等），在前几代架构中，一个CUDA Core在每个时钟周期里只能为一个线程执行一条浮点或整数指令。但是从Volta架构开始，将一个CUDA Core 拆分为两部分：FP32 Core 和 INT32 Core，好处是在同一个时钟周期里，可以同时执行浮点和整数指令，提高计算速度。当CUDA程序执行FP32运算时，GPU会通过其FP32 Core来执行这些操作，因为这些核心专门针对处理32位浮点数进行了优化。相对应地，如果程序包含双精度浮点（FP64）运算或整数（INT32）运算，那么对应的FP64 Core或INT32 Core将会被利用。

图7-8-1 Pascal架构的SM

图7-8-2 Volta架构的SM

Turing 架构

Turing 架构最重要的新特性是加入了专门用于加速光线追踪的 RT(Ray-Tracing)核心，实现了计算机图形学的一大突破,使得实时光线追踪成为可能。另外，深度学习超采样(Deep Learning Super Sampling,DLSS)使用专为游戏而设的深度神经网络,使用超高质量的64倍超级采样图像或真实画面进行训练，进而通过张量核心来推断高质量的抗锯齿结果。在 Turing 架构上,张量核心不仅可以加速 DLSS 等特性，也可以加速某些基于 AI的降噪器，以清理和校正实时光线追踪渲染的画面。

Ampere 架构

Ampere 架构是 NVIDIA 在 2020 年新推出的 GPGPU 架构,其旗舰产品 Ampere A100中张量核心单元的性能比 Volta架构中张量核心单元的性能提高了2.5倍,比传统的CUDA 核心单元执行单精度浮点乘加的性能提高了20倍。

计算能力	架构	发布年代	Cores/SM	总 SM 数	CUDACores	L1 Cache(KB)	L2 Cache(KB)
1.0	Tesla	2006
2.0	Fermi	2009	32	16 SM	512	48	768
3.0	Kepler	2012	192	15 SMX	2880	48	1536
4.0	–
5.0	Maxwell	2014	128	24 SMM	3072	96	2048
6.0	Pascal	2016	64	60 SM	3840	64	4096
7.0	Volta	2018	648 个 Tensor Core	80 SM	5120	与共享内存共用 128（最多 96）	6144
7.5	Turing	2018	648 个 Tensor Core	72 SM	4608	与共享内存共用 128（最多 96）	6144
8.0	Ampere	2020	644 个 Tensor Core	108 SM	6912	与共享内存共用 192（最多 164）	40960
9.0	Hopper	2022	1284 个 Tensor Core	144 SM	18432	与共享内存共用 256	61440

Ampere 架构

以GA100为主，介绍Ampere架构

常见术语

流处理器（stream multiprocessors，SM） ：从英文名字就可以理解出来，流处理器的集合，是GPU架构中的基本计算单元，可以理解为工厂架构中的车间，，由SP，DP，SFU等运算单元组成。
SMX： Kepler架构中的SM。
SMM： Maxwell架构中的SM。
流处理器（stream processor，SP） ：也称为core，是GPU运算的最基本单元，类似于计算机组成课程中的CPU内部的ALU（不严谨），是执行计算的。
CUDA core： SP的另一个名称，或者说是SP的升级版，称为CUDA core，始于Fermi架构，每个 CUDA 内核都有一个完全流水线化的整数算术逻辑单元 (ALU) 和浮点单元 (FPU)。
- FP32 Core ： 是可以执行单精度浮点运算的CUDA核心。这是一种常用的精度，平衡了计算的准确度和性能。在图形处理、游戏开发和某些科学计算中，FP32提供了足够的精度。
- INT32 Core ： 是专门执行32位整数运算的CUDA核心。整数运算在计算机图形、视频处理以及某些类型的数据处理任务中非常重要。与FP32运算相比，INT32运算通常用于不需要或不适合使用浮点数表示的数据和算法。在一些较新的NVIDIA GPU架构中，如Turing和Ampere，INT32运算能力得到了显著的增强，以支持更复杂的图形和计算任务。
- FP64 Core： 是可以执行双精度浮点运算的CUDA核心。FP64使用64位来表示浮点数，提供了更高的精度但以牺牲一定的性能为代价。双精度运算对于科学计算、工程模拟和任何需要极高数值准确度的应用至关重要。不是所有的NVIDIA GPU都提供大量FP64核心，因为这些核心的面积和能源消耗较大。
双精度浮点运算单元（DP） ：专用于双精度浮点运算的处理单元，一种特殊工种，只能用于双精度浮点运算。
特殊功能单元（special function unit，SFU） ：用来执行超越函数指令，如sin,cos 倒数，平方根等函数。
线程处理器簇（thread processing cluster，TPC） ：由SM控制器，多个SM和L1cache（L1缓存）组成。
图形处理器簇（graph processing cluster，GPC） ：类似与TPC，但不是TPC的替代品，在Pascal架构中，同时出现了GPC和TPC，且GPC包含TPC，有的架构中没有TPC，有的架构中没有GPC，有的架构中TPC，GPC都没有，有的架构中TPC，GPC都有。。
流处理器阵列（scalable streaming processor array，SPA） ：所有处理器核心和高速缓存的综合，包含所有的SM，TPC，GPC.与存储系统共同构成GPU架构。
存储控制器（memory controller,MMC） 顾名思义，控制存储访问单元。
存取单元（load/store unites，LD/ST）
HBM (High Bandwidth Memory DRAM) ：高带宽存储器

A100 GPU 硬件结构

NVIDIA GA100 GPU 由多个 GPU 处理群集(GPC)，纹理处理群集(TPC)，流式多处理器 (SM) 和 HBM2 显存控制器组成。

GA100 GPU 的完整实现包括以下单元：

- 每个完整 GPU 有 8 个 GPC，每个 GPC 有 8 个 TPC，每个 TPC 有 2 个SM，每个 GPC 有 16 个SM，总共 128 个 SM；
- 每个 SM 有 64 个 FP32 CUDA 核，总共 8192 个 FP32 CUDA 核；
- 每个 SM 有 4 第三代 Tensor Core，总共 512 个第三代 Tensor Core；
- 总共 6 个 HBM2 堆栈，12 个 512 位显存控制器；

基于 GA100 GPU 实现的 A100 Tensor Core GPU 包括以下单元：

- 每个 GPU 有 7 个 GPC，每个 GPC 有 7 个或 8 个 TPC ，每个 TPC 有 2 个 SM，每个 GPC 最多 16 个 SM，总共 108个SM；
- 每个 SM 有 64 个 FP32 CUDA 核，总共 6912 个 FP32 CUDA 核；
- 每个 SM 有 4 个第三代 Tensor Core，总共 432个第三代 Tensor Core；
- 总共 5个HBM2 堆栈，10 个 512 位显存控制器；

图7-9-1 具有128 个 SM 的 GA100 GPU

GA100 SM 被划分为四个处理块（或分区），每个处理块都有一个 64 KB 的寄存器文件、一个 L0 指令缓存、一个warp线程束调度器、一个Dispatch unit分配单元以及其他单元。这四个分区共享L1 数据缓存/共享内存子系统。

图7-9-1 GA100的SM

参考资料

Tesla 架构资料

【G80架构论文链接】：主要介绍Tesla1.0 （一代）架构的特性；
【GT200介绍链接】：介绍Tesla2.0 架构的特点，以及与第一代的差异；
【8800系列第三方分析】：分析Tesla1.0 架构显卡的一些特性；
【GT200的架构原理分析】：详细讲解了Tesla2.0 的架构特点以及工作原理；

Fermi 架构资料

【Fermi架构白皮书】：介绍Fermi架构的官方文档；
【Fermi特点总结】：对Fermi架构的创新点进行总结；
【Fermi架构分析文章】：从CPU讲到GPU，再讲到Fermi架构；
【CUDA Floating Point & IEEE 754】：讲述Fermi推出的CUDA core支持的FMA计算的精度和性能；

Kepler 架构资料

【Kepler架构PPT】：GTC2012发布会上面，对KeplerGK110特点进行展示与介绍；
【Kepler架构一代白皮书】：对第一代kepler GK110 架构参数介绍；
【Kepler架构第二代白皮书】：对GK110和GK210统一介绍，包含参数对比；
【Kepler图像卡GeForce GTX 680】：介绍kepler在图形卡680上面的应用；
【GK106产品的参数对比】：介绍kepler GK106芯片在 GTX 6xx 系列的产品；

Maxwell 架构资料

【GTX-750-Ti 白皮书】：主要介绍显卡750和GM107（Maxwell第一代）芯片；
【GTX 980 白皮书】：介绍显卡980，同时介绍GM204(第二代) 芯片【Maxwell第二代架构GM204】：
【Maxwell架构特点介绍】Nvidia官网的一篇介绍Maxwell的blog；
【对Maxwell架构的5点认识】：总结文章，里面包含了和上一代Kepler的对比分析；
【MaxwellGDC15】：产品介绍PPT；
【Maxwell的CUDA应用】：简单介绍maxwell在CUDA上面的新特性；

Pascal 架构资料

【Pascal特性总结】：官宣blog，简要地介绍了架构特点；
【Pascal架构 P100 白皮书】：介绍Tesla p100（HPC产品）的特点、参数、新特性；
【GeForce_GTX1080白皮书】：介绍1080（图形产品）显卡的特性；
【HPC产品的对比】：对比P100、K80、K40、M40的特点；
【GeForce GTX TITAN X】：图形卡Titan x的具体参数特性；

Volta架构资料

【Volta架构 V100白皮书】: 以Tesla V100产品为主介绍Volta架构；
【Volta的特性评估】: 对比v100与前几代架构的性能差异；
【CUDA 9对Volta的使用】：介绍CUDA9对Volta架构的新硬件特性的支持；
【TITAN V图形卡】：介绍Titan V （ceo 版本）的参数。
【论文-Volta的benchmark分析】：分析Volta指令、内存、NVlink方面的性能；

Turing 架构资料

【Turing架构白皮书】：介绍架构特性，同时介绍了Turing在几个工程应用方向的性能；
【GTX20系列参数】：基于Turing架构的显卡GTX20系列的参数对比；
【Turing- Real Time Ray Tracing】：图灵卡的实时光追踪技术的介绍；

Ampere 架构资料

【Ampere架构 GA100白皮书】：以GA100为主，介绍Ampere架构；
【Ampere架构 GA102白皮书】：以GA102为主，介绍Ampere架构，着重讲图形方向的性能；
【A100特性手册】：“NVIDIA A100 TENSOR CORE GPU”；
【A100 80G pcie产品】：A100 80GB 的产品介绍；
【A100 40GB规格pcie产品手册】：产品介绍；

Hopper 架构资料

【H100白皮书】：介绍H100的特性；
【H100介绍】：官方介绍文档；
【H100特性对比】：与前几代架构的特性对比；

从图形渲染到通用计算的GPU基础知识(下)

五. GPU与通用计算

HPC （High Performance Computing, HPC） 高性能计算

通用图形处理器（GPGPU）

从图形渲染到通用计算的转变

固定功能流水线（Fixed Function Pipeline）

可编程流水线（Programmable Pipeline）

分离着色器架构（Separate Shader Architecture）

统一着色器架构（Unified shader Architecture）

通用计算API

CUDA 与 OpenCL

CUDA

OpenCL

二者的区别

各自的优缺点

六. GPGPU架构相关

并行体系结构

SIMD

SIMT

NVIDIA与通用并行计算

CUDA

内核函数(Kernel)

线程层次

硬件实现

SM 多线程流式处理器

Warp

CUDA为什么要分线程块和线程网格？

浮点运算单元的浮点型

FP32（单精度浮点数）

FP64（双精度浮点数）

FP16（半精度浮点数）

FP8（8位浮点数）

七. NVIDIA Tesla系列GPU

NNVIDIA GPU 三大系列

命名规则

GeForce系列

Tesla系列

Quadro系列

NNVIDIA GPU 市场占有率

NVIDIA GPU关键技术

Core 计算核心

CUDA Cores（计算单元）

RT Cores（光线追踪单元）

Tensor Cores（张量单元）

DLSS

分布式通信

NVLink

NVSwitch

MIG(GPU虚拟化)

MIG 技术关键特点

架构演进

Fermi架构

Kepler 架构

Volta 架构

Turing 架构

Ampere 架构

Ampere 架构

常见术语

A100 GPU 硬件结构

参考资料

图形编程相关资料

通用计算相关资料

NVIDIA Tesla 系列资料

分离渲染架构

Tesla 架构资料

Fermi 架构资料

Kepler 架构资料

Maxwell 架构资料

Pascal 架构资料

Volta架构资料

Turing 架构资料

Ampere 架构资料

Hopper 架构资料

HPC （High Performance Computing, HPC）高性能计算