九、GPU体系结构

xiaozongpeng · 发表于 2022-5-15 09:09

0.参考资料
1.Streaming Processor
CUDA core，CUDA内存模型最基本处理单元，具体指令是在SP上处理的。GPU进行并行计算，就是很多个SP同时做处理。
每个SP会分配相应的寄存器和局部内存，寄存器和局部内存只能被自己访问。不同SP之间彼此独立。
2.Stream Multiprocessor
多个SP加上其他一些资源组成一个SM，也叫GPU大核。其他资源如：warp scheduler、register资源、shared memory、L1 cache、SFU(Special Function Unit)、LD/ST单元等。
L1缓存为SM内的数据，SM内的运算单元能够共享，但跨SM之间的L1不能相互访问。register资源均分给SM内所有SP。shared memory被SM上正在执行的线程块的线程所共享。
SM的数量由芯片决定的。例如，在Tegra系列中，一个GPU中通常只有2个SM，每一个SM中包含4个WARP，每一个WARP中有32个SP(thread)。因此，一个SM中有128个SP。
每个SM包含的SP数量依据GPU架构而不同。1.x硬件，一个SM包含8个SP，2.0是32个，2.1是48个，3.0和3.5是192个。SM目前也称为MP，在KEPLER架构(SM3.0和3.5)下也称为SMX。
3.WARP
最小的硬件执行单位。SM先将所有的SP在物理上分成几个WARP，WARP中SP执行相同的指令，同时工作。即取指令单元取一条指令同时发射给WARP中的所有的SP。
如下图所示，Tegra中一个WARP由32个SP(thread)组成，32个SP(thread)是一起工作的，执行相同的指令。申请线程数尽量为32的倍数。如果线程数量不是32的倍数，假如是1，则WARP会生成一个掩码，当一个指令控制器对一个WARP单位的线程发送指令时，32个线程中只有一个线程在真正执行，其他31个进程会进入静默状态。

4.单指令多线程(Single-Instruction Multiple-Thread，SIMT)
(1)线程束是最基本的执行单元，一个线程束包含32个并行SP(Thread)，这些SP以不同数据资源执行相同的指令，线程束采用的是SIMT架构。
(2)由于线程束的大小为32，所以Block包含thread的个数一般设置为32的倍数。

5.device
多个多核处理器(SM)和其他一些资源构成GPU。其他资源如：global memory、local memory、L2缓存、constant memory、texture memory等。
6.存储结构
(1)全局内存
全局内存(global memory)是数据存储的常用内存，它能被设备内所有线程访问、全局共享，为片下(off chip)内存。跟CPU架构一样，运算单元不能直接使用全局内存的数据，需要经过缓存，过程如下图所示：

全局内存申请使用cudaMalloc。
(2)局部内存(local memory)
局部内存(local memory)是线程独享的内存资源，线程之间不可以相互访问，硬件位置是片下(off chip)状态，所以访问速度跟全局内存一样。局部内存主要用来解决当寄存器不足时的场景，即在线程申请的变量超过可用的寄存器大小时，nvcc会自动将一部数据放置到局部内存里面。
注意，局部内存设置的过程是在编译阶段就会确定。
(3)寄存器(register)
寄存器(register)是线程能独立访问的资源，它所在的位置与局部内存不一样，是在片上(on chip)的存储，用来存储一些线程的暂存数据。寄存器的速度是访问中最快的，但是它的容量较小。以目前最新的Ampere架构的GA102为例，每个SM上的寄存器总量256KB，使用时被均分为了4块，且该寄存器块的64KB空间需要被warp中线程平均分配，所以在线程多的情况下，每个线程拿到的寄存器空间相当小。
(4)L1/L2缓存
L1/L2缓存(Cache)数据缓存，这个存储跟CPU架构的类似。L2为所有SM都能访问到，速度比全局内存块，所以为了提高速度有些小的数据可以缓存到L2上面；L1为SM内的数据，SM内的运算单元能够共享，但跨SM之间的L1不能相互访问。
对于开发者来说，需要注意L2缓存能够提速运算，比如CUDA11 A100上面L2缓存能够设置至多40MB的持续化数据(persistent data)，L2上面的持续化数据能够拉升算子kernel的带宽和性能。
(5)共享内存(shared memory)
共享内存(shared memory)存储硬件位于芯片上(on chip)，访问速度较快，共享内存主要是缓存一些需要反复读写的数据。是一种在block内能访问的内存。
注：

共享内存与L1的差异。共享内存与L1的位置、速度极其类似(在很多显卡上，共享内存和L1 缓存使用的是同一块硬件)，区别在于共享内存的控制与生命周期管理与L1不同，共享内存的使用受用户控制，L1受系统控制，shared memory更利于block之间数据交互。
共享内存大小只有几十K，过度使用共享内存会降低程序的并行性。

共享内存有两种空间申请方式，静态申请和动态申请。静态申请指共享内存大小明确；动态申请指共享内存大小不明确，使用时确定。
申请：__shared__关键字修饰，大小有限制，不能过分使用。申请方式分为静态申请和动态申请。
使用：每个线程从全局内存索引位置读取元素，将它存储到共享内存之中。在使用共享内存时，如果块内线程同步函数(__syncthreads())使用不当，会导致数据交叉。
(6)常量内存(constant memory)
常量内存(constant memory)为片下存储的只读内存，訪問速度與local memory、global memory相同。所有线程都能访问的只读存储器，host可讀可寫。
为什么要设立单独的常量内存，直接用global memory或shared memory不行吗？主要解决一个WARP内多线程访问相同数据的速度太慢的问题。运算的thread要同时访问一个constant_A常量，存储介质上面constant_A数据只保存了一份，会出现先后访问的问题，使得并行计算的thread出现了运算时差。针对于此，提出常量内存。
如下图所示，常量内存对应的cache位置产生多个副本，让thread访问时不存在冲突，从而提高并行度。

使用__constant__限定符修饰，編譯時將变量存儲在常量内存中；定义常量内存时，如果定义在所有函数之外，則作用于整个文件。
使用cudaMemcpyToSymbol()將数据拷贝到常量内存中。

__constant__ int devVar;
cudaMemcpyToSymbol(devVar,hostVar,sizeof(int),0,cudaMemcpyHostToDevice);
cudaMemcpyFromSymbol(hostVar,devVar,sizeof(int),0,cudaMemcpyDeviceToHost);(7)图像/纹理内存(texture memory)
图像/纹理内存(texture memory)为片下存储的只读内存，訪問速度與local memory、global memory、constant memory相同。所有线程都能访问的只读存储器，host可讀可寫。
图像/纹理内存通常用来处理1D/2D/3D结构数据(相邻数据之间存在一定关系，或者相邻数据之间需要进行相同的运算)。texture在运算之前能进行一些处理(或者说它本身就是运算)，比如聚合、映射等；texture memory进行图像类数据加载时，warp内的thread访问的数据地址相邻，从而减少带宽的浪费。
如图所示，是一个P100显卡的SM架构，包含了四个Tex。Tex是专门用来处理texture的单元，进行数据拿取(fetch)的时候，能够在一个clock时钟内完成对数据的一些预处理。

(8)对比

7.GPU宏观结构
由于纳米工艺的引入，GPU可以将数以亿记的晶体管和电子器件集成在一个小小的芯片内。从宏观物理结构上看，现代大多数桌面级GPU的大小跟数枚硬币同等大小，部分甚至比一枚硬币还小。

当GPU结合散热风扇、PCI插槽、HDMI接口等部件之后，就组成了显卡。

显卡不能独立工作，需要装载在主板上，结合CPU、内存、显存、显示器等硬件设备，组成完整的PC机。

8.GPU微观架构发展历程
GPU微观结构因不同厂商、不同架构会有所差异，但核心部件、概念、以及运行机制大同小异。

9.NVidia Tesla架构
Tesla微观架构总览图如下。下面将阐述它的特性和概念：
(1)拥有7组TPC(Texture/Processor Cluster，纹理处理簇)
(2)每个TPC有两组SM(Stream Multiprocessor，流多处理器)
(3)每个SM包含：

6个SP(Streaming Processor，流处理器)
2个SFU(Special Function Unit，特殊函数单元)
L1缓存、MT Issue(多线程指令获取)、C-Cache(常量缓存)、共享内存

(4)除了TPC核心单元，还有与显存、CPU、系统内存交互的各种部件。

10.NVidia Fermi架构
Fermi架构如下图，它的特性如下：
(1)拥有16个SM
(2)每个SM

2个Warp(线程束)
两组共32个Core
16组加载存储单元(LD/ST)
4个特殊函数单元(SFU)

(3)每个Warp

16个Core
Warp编排器(Warp Scheduler)
分发单元(Dispatch Unit)

(4)每个Core

1个FPU(浮点数单元)
1个ALU(逻辑运算单元)

11.NVidia Kepler架构
Kepler除了在硬件有了提升，有了更多处理单元之外，还将SM升级到了SMX。SMX是改进的架构，支持动态创建渲染线程(下图)，以降低延迟。

12.NVidia Maxwell架构
采用了Maxwell的GM204，拥有4个GPC，每个GPC有4个SM，对比Tesla架构来说，在处理单元上有了很大的提升。

13.NVidia Turing架构(2018)
下图是采纳了Turing架构的TU102 GPU，它的特点如下：

6 GPC(图形处理簇)
36 TPC(纹理处理簇)
72 SM(流多处理器)
每个GPC有6个TPC，每个TPC有2个SM
4,608 CUDA核
72 RT核
576 Tensor核
288 纹理单元
12x32位GDDR6内存控制器(共384位)

单个SM的结构图如下：

每个SM包含：

64 CUDA核
8 Tensor核
256 KB寄存器文件

TU102 GPU芯片实物图：

14.thread
thread相当于硬件中的SP，是最小的逻辑单位，WARP是最小的硬件执行单位。
15.block
若干个thread(典型值是128~512个)组成一个block，blcok执行在SM上，一个SM可能有一个或多个 blocks。block中的thread通过shared memory进行通信。block中的thread使用SIMT模式执行相同命令不同数据。一般，一个block里有很多个WARP。
16.GRID
多个block则会再构成grid，每个grid会有自己的global memory、constant memory 和 texture memory。
一个kernel对应一个grid，该grid又包含若干个block，block内包含若干个thread。如图5所示，thread以warp为单位被SM的scheduler发射到SP或者其他单元(如SFU、LD/ST  unit执行相关操作)，需要等待的warp会被切出(依然是resident状态)，以空出执行单元给其他warps。

17.索引计算
一个Grid可以包含多个Blocks，Blocks的组织方式可以是一维的、二维或三维；Bolock包含多个threads，这些thread的组织方式也是一维、二维或三维。

(1)grid划分成1维，block划分成1维
int threadId=blockIdx.x*blockDim.x+threadIdx.x;(2)grid划分成1维，block划分为2维
int threadId=blockIdx.x*blockDim.x*blockDim.y+threadIdx.y*blockDim.x+threadIdx.x;(3)grid划分成1维，block划分为3维
int threadId=blockIdx.x*blockDim.x*blockDim.y*blockDim.z
         +threadIdx.z*blockDim.y*blockDim.x+threadIdx.y*blockDim.x+threadIdx.x;(4)grid划分成2维，block划分为1维
int blockId=blockIdx.y*gridDim.x+blockIdx.x;
int threadId=blockId*blockDim.x+threadIdx.x;(5)grid划分成2维，block划分为2维
int blockId=blockIdx.x+blockIdx.y*gridDim.x;
int threadId=blockId*(blockDim.x*blockDim.y)+(threadIdx.y*blockDim.x)+threadIdx.x;(6)grid划分成2维，block划分为3维
int blockId=blockIdx.x+blockIdx.y*gridDim.x;
int threadId=blockId*(blockDim.x*blockDim.y*blockDim.z)
         +(threadIdx.z*(blockDim.x*blockDim.y))
         +(threadIdx.y*blockDim.x)+threadIdx.x;(7)grid划分成3维，block划分为1维
int blockId=blockIdx.x+blockIdx.y*gridDim.x+gridDim.x*gridDim.y*blockIdx.z;
int threadId=blockId*blockDim.x+threadIdx.x;(8)grid划分成3维，block划分为2维
int blockId=blockIdx.x+blockIdx.y*gridDim.x+gridDim.x*gridDim.y*blockIdx.z;
int threadId=blockId*(blockDim.x*blockDim.y)+(threadIdx.y*blockDim.x)+threadIdx.x;(9)grid划分成3维，block划分为3维
int blockId=blockIdx.x+blockIdx.y*gridDim.x+gridDim.x*gridDim.y*blockIdx.z;
int threadId=blockId*(blockDim.x*blockDim.y*blockDim.z)
         +(threadIdx.z*(blockDim.x*blockDim.y))
         +(threadIdx.y*blockDim.x)+threadIdx.x;18.GPU编程流程
(1)主机端计算任务

申请显存、拷贝
核函数调用
转移、释放

(2)设备端计算任务

读写线程寄存器
读写block中共享内存
读写Grid中全局内存

19.核函数调用
核函数在设备端执行，在主机端调用，调用时必须申明执行参数。调用形式如下：
Kernel<<<Dg,Db,Ns,S>>>(param list);其中，<<<>>>内参数形式为<<<Dg,Db,Ns,S>>>，这些参数告诉编译器如何启动核函数。
(1)参数Dg表示grid的维度(一个grid中有多少个block)，类型为dim3。Dim3 Dg(Dg.x,Dg.y,1)表示grid中每行有Dg.x个block，每列有Dg.y个block，第三维恒为1(目前一个核函数只有一个grid)。整个grid中共有Dg.x*Dg.y个block，Dg.x和Dg.y最大值为65535。
(2)参数Db表示block的维度(一个block中有多少个thread)，类型为dim3。Dim3 Db(Db.x, Db.y, Db.z)表示block中每行有Db.x个thread，每列有Db.y个thread，高度为Db.z。整个block中共有Db.x*Db.y*Db.z个thread，Db.x和Db.y最大值为512，Db.z最大值为62。计算能力为1.0、1.1的硬件最多有768个thread，计算能力为1.2、1.3的硬件最多有1024个thread。
(3)参数Ns是可选参数，设置每个block静态分配shared Memory之外，动态分配时shared memory的大小，单位为byte。不需要动态分配时，该值为0或省略不写。
(4)参数S是cudaStream_t类型的可选参数，初始值为零，表示该核函数处在哪个流之中。
20.核函数调用注意事项
(1)在GPU上执行的函数。
(2)一般通过标识符__global__修饰。
(3)调用时必须声明内核函数的执行参数。
(4)调用通过<<<参数1,参数2>>>，参数1表示线程块参数，参数2表示线程参数。网格是由线程块和线程组成的。
(5)以网格(Grid)的形式组织，每个线程格由若干个线程块(block)组成，而每个线程块又由若干个线程(thread)组成。
(6)在编程时，必须先为kernel函数中用到的数组或变量分配好足够的空间，再调用kernel函数，否则在GPU计算时会发生错误。
21.CUDA编程的标识符号
核函数标识符号共有三种，分别是__device__、__global__和__host__。
(1)不同的标识符号对应不同的工作地点和被调用地点。
(2)使用__global__标识，必须返回void。
(3)__device__和__host__可以一起用。
(4)__device__和__global__函数的文件后缀为.cu。

22.cuda提供函数
(1)cudaMalloc()
申请显存。

cudaError_t cudaMalloc(void **devPtr,size_t size)
在设备全局内存中分配对象。
两个参数：地址、申请内存大小。

(2)cudaFree()
释放显存。

cudaError_t cudaFree(void* devPtr)
从设备全局内存中释放对象。
参数：指向释放对象的指针。

(3)cudaMemcpy()
内存与显存数据互相复制。

cudaError_t cudaMemcpy(void dst,const void src,size_t count,cudaMemcpyKind kind)
内存数据复制传递。
第四个参数：cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyDeviceToDevice、cudaMemcpyDefault(默认，一般不会使用)。
调用cudaMemcpy()传输内存是同步的。

以上三个函数是CUDA自带的，调用需要先：
#include &#34;cuda.h&#34;23.CUDA编程流程
23.1 编译
Linux上CPU编程，使用g++或gcc进行编译，再通过link生成可执行程序。
GPU上编程，编译器为NVCC(NVIDIA Cuda compiler driver)。设备端程序头文件放在.h或.cuh文件里，执行程序(__global__定义的函数)放在.cu文件里，然后使用NVCC来进行编译。主机端程序放在.h和.cpp(.c)里，然后使用g++或gcc进行编译。
23.2 多文件编译方法
(1)逐个文件编译(GPU和CPU的程序都编译成.o文件，然后把它们汇总到一起，并link成一个可执行文件.exe)，但这只适用于文件数较少的情况，当文件数较多时，这种办法就显得比较复杂。
(2)使用cmake方式编译，写一个CMakeLists.txt，下文有介绍。

24.cudaMalloc函数理解

		自动登录	找回密码
密码			立即注册

九、GPU体系结构

本帖子中包含更多资源