CPU / GPU原理与 CUDA

作者: Teci | 来源:发表于2019-10-09 14:45 被阅读0次

1. CPU 并行

1.1 概述

之所以先从 CPU 聊起而不是直接切入正题讲 GPU, 是因为并行思想在CPU和GPU上是相通的, 而首先了解 CPU 的模式, 也更有助于之后熟悉 GPU 与 CPU 之间的交互。

因为 CPU 首先得取得数据, 才能进行运算, 所以很多时候, 限制我们程序运行速度的并非是 CPU 核的处理速度, 而是数据访问的速度, 下面罗列了一些硬件设备的访问速度 (1Gbps = 125 MB/s):

网卡 (NIC), 1Gbps
机械硬盘(HDD) , 1-2Gbps, 如果连接了 SATA3 的话, 最高可达到 6Gbps. 其实对于 HDD 而言, 最耗时的部分不是传送数据的时间, 而是寻找数据的时间, 因为 HDD 内部构造就像老式的唱片机, 一个旋转头和一个圆盘, 而定位某个数据需要将旋转头移动到数据储存在圆盘上的位置, 这个过程是较为耗时的, 所以如果数据是不规则储存的, 即分散储存, 那么找到所有数据的时间还需要更长, 因此虽然 SATA3 最高可支持 6Gbps 的速率, 但是 HDD 是很难达到这个最高速率的
USB 2.0 的速率是 0.48Gbps, 而 USB3.0 支持 5Gbps, USB3.1 支持 10Gbps
固态硬盘(SDD) 使用 SATA3 的话可支持 4-5Gbps
DDR4 内存和 CPU(Core i7) 之间的传输速率是 160-480Gbps
GPU 内部的内存传输速率在 800-8000Gbps

市面上卖的 CPU 大多都是 x核2x线程, 譬如 8核16线程, 10核20线程, 这是因为这些核都采用了 Hyper-Threading 技术, 该技术可以让一个 microprocessor(核) 表现出两个分离的 processor 的形式, 不过这两个 processor 需要共享相同的核资源, 譬如 cache memory 等。

但是如果你打开当前系统的资源管理器, 你可能会看到当前系统中存在大概 1000+ 个线程, 但是你的 CPU 是 10 核 20 线程, 很显然是不合理的。这是因为资源管理器中显示的线程不是正在活动的线程, 而其中绝大部分都在睡眠状态, 譬如某个网络监听窗口, 一直在沉睡, 直到某个网络包到来, 系统唤醒了该线程, 线程处理完这个网络包后继续睡眠。此处的10核20线程指的是能同时执行(execution) 的线程, 而某个线程被启动(launch), 但是在后台睡眠而不被执行也是可以的。

程序可以被分为以下三类:

core-intensive (核资源密集型)
memory-intensive (内存资源密集型)
I/O-intensive (I/O 资源密集型)

一个 4GHz 的单核 CPU 的能耗大概是一个 1GHz 的单核 CPU 的 16 倍, 而一个 4GHz 的单核 CPU 的能耗与一个双核的每个核 3GHz 的CPU能耗相当。

1.2 CPU 并行化程序设计

需要明确的点是:

当我们需要 launch 多个线程的时候, 确保大量的计算放在线程中, 而不要放在多线程执行完并同步后的代码中
launch 多个线程是有 overhead 的, 因为需要给每个线程分配资源, 句柄等
在 c 语言中, 一旦一个线程被成功 launch, 系统会给其分配一个 handler(句柄), 一个虚拟CPU (10核20线程相当于有20个虚拟CPU), 一个stack(栈)区, 并将 handler 返回给launch这个线程的主线程
总共有两种类型的核, 一种是 in-order(inO) 的, 另一种是 out-of-order(ooO) 的。
- 其中 inO 意味着其严格按照二进制码中的顺序来执行代码, 而 ooO 则是当前哪个 operand(操作元) 可以执行, 就执行哪个, 譬如对于一系列指令, inO 只能等一条指令执行完了后才能执行后一条指令, 而 ooO 可以在等待当前这条指令执行完之前, 先着手执行其后面的指令。
- 很显然, 在指令执行调度上 ooO 是比 inO 高效的, 然而, 因为 inO 较为简单, 所以制造 inO 核所需要的基础 chips(芯片) 更小, 所以一个 inO 核中可以放置更多的基础芯片, 因为基础芯片的数量增加了, 所以 inO 可以以更高的时钟频率来工作
- 同时, 因为 inO 更为简单, 所以耗能较低
如果一个线程需要使用大量的核的资源, 那么这个线程就被称为 thick 线程, 反之被称为 thin 线程, 如果一个线程被设计为尽可能少的需要核的资源, 那么由这种线程组成的多线程程序, 能最大地提升多线程的效率, 这也是为什么微软设计 windows 系统时, 确保其每个系统线程所占用的核的资源最少, 这样一来就不会影响用户的应用程序的执行

1.3 内核与内存

CPU架构以及 L1, L2, L3 缓存

1.jpg

上图是 i7-5930K 的 CPU 结构, 下图是每个核中详细的结构。

2.jpg

芯片组和DRAM

芯片组和CPU的构建模块主要是MOS晶体管，但DRAM的构建模块是存储电荷的极小电容器。CPU 通过将 MOS 晶体管缩小使 CPU 中能添加更多的晶体管, 而 DRAM通过将更多的电容器转入连续的区域来获得更大的储存。
由于电荷存储在极小的电容器中，因此在一定时间后（例如50 ms）会消耗（即泄漏）电荷。由于这种泄漏，必须读取数据并将其放回DRAM（即刷新）。考虑到刷新的缺点，允许一次一个字节地访问数据是没有意义的。因此，每次访问数据最好是大块大块地存取（例如，每次4KB）。
DRAM一次访问一行，每行是DRAM内存的最小可访问量。现代DRAM中的行大约为2-8 KB。要访问某行，需要一定的时间（延迟）。但是，一旦访问该行（由DRAM内部进入行），该行实际上可以自由访问。CPU访问DRAM的延迟大约为200-400个周期，而访问同一行中的后续元素只需几个周期。即获取一行数据有很长的延迟，但一旦读取，对该行的访问速度非常快。除了可以访问行之外，DRAM还具有各种其他延迟，例如行到行延迟等。这些参数中的每一个都由存储器接口标准指定，这些标准由协议定义。
内存带宽的增加应允许更多的硬件线程更快地将数据带入内核, 从而当更多数量的已启动软件线程同时需要来自DRAM主内存的数据时, 不会造成阻塞, 避免内存带宽饱和。

L1, L2, L3 缓存和执行单元

L1, L2, L3 缓存是 SRAM(static random access memory), 而不是 DRAM, 其访问速度远远大于 DRAM。
L1 缓存
- 总共有 64KB, 其中 32KB 用来储存数据, 即上图中的 L1D, 剩下 32KB 用来储存指令, 即 L1I, L1I 中储存的是最常用的一些指令
- 访问 L1 缓存的速度很快 (4 周期的 load-to-use 延迟)
- 每个核都有独有的 L1 缓存, 因为每个核共有两个线程, 所以这两个线程需要共享 L1 缓存, 而这两个线程之间可以彼此之间可以通过 L1 缓存收/发数据
L2 缓存
- 总共有 256KB, 其不像 L1 缓存将数据和指令分开, 而是将指令和数据储存在一起
- 访问 L2 的缓存速度较块 ( 11-12 周期的 load-to-use 延迟 )
- 当系统决定某个数据或某个指令不再常用时, 会将其从 L1 缓存中移除, 因为 L2 缓存比 L1 缓存大, 所以从 L1缓存中移除的数据或指令可能仍然存在与 L2 缓存中
- 每个核都有其独有的 L2 缓存
L3 缓存
- 总共有 15MB
- 访问 L3 缓存比 DRAM 块, 但是比 L2 缓存稍慢 (约等于 22 周期的 load-to-use 延迟)
- 每个核都有其独有的 L3 缓存
L1, L2, 和 L3 中数据的出入完全由CPU控制，不受程序员的控制。但是，通过将数据操作保持在小循环中，程序员可以很大程度上影响高速缓冲存储器的效率。
数据会首先进入 L3 缓存, 然后从 L3 缓存进入 L2 缓存, 最后再从 L2 缓存进入 L1 缓存, 所以, 为了在最大程度上利用缓存提高效率, 我们应该尽可能做到:
1. 每个线程重复访问32 KB数据区域
2. 尝试在可能的情况下将更广泛的访问限制为256 KB，
3. 在考虑所有已启动的线程时, 尝试在L3 (例如，15 MB) 储存累积的数据
4. 如果必须超过L3的大小，请确保在超过此区域之前大量使用 L3
执行单元
- 执行单元分为两类：ALU（算术逻辑单元）负责整数运算，逻辑运算如OR，AND，XOR等.FPU（浮点单元）负责浮点（FP）操作，如FP ADD和FP MUL（乘法）。除（整数或FP）除法比加法和乘法更复杂，因此有一个单独的除法单位。但是整数除法显著地比整数乘法要慢。所有这些执行单元都由一个核中的两个线程共享。
- 在每一代中，更复杂的计算单元可用作共享执行单元。但是，多个单元被合并用于可能由两个线程 (例如ALU) 执行的常见操作, 不过每一代的确切细节可能会改变。但在过去3到3年的CPU设计中，ALU-FPU功能分离从未改变过。
- 计算两个线程生成的地址，以将两个线程中的数据写回内存。对于地址计算，加载和存储地址生成单元（LAGU和SAGU）由两个线程共享，以及正确排序目标存储器地址 (MOB) 的单元。
- 指令只被预取并解码一次并传递到所有者线程。因此，两个线程共享预取器和解码器。
- 我们在写代码时需要考虑某个计算所消耗的资源, 譬如计算 √(x²+y²) 就是计算密集型的指令, 其需调用的运算单元包括
  - 两次 FP-MUL (浮点相乘)
  - 一次 FP-ADD (浮点相加)
  - 一次开方操作 (用的应该是卡马克算法)

3.jpg

假设我们现在要左右翻转一张图片, 即上图的狗, 我们初步的想法是将该 2MB 的图片整个储存在全局内存中, 然后一行行迭代, 每迭代一行则将该行左右相对的像素点交换, 那么按照这种逻辑设计出来的程序如下:

...
for(row=ts; row<=te; row++) {
  col=0;
  while(col<ip.Hpixels*3/2){ 
    // example: Swap pixel[42][0] , pixel[42][3199] 
    pix.B = TheImage[row][col];   // TheImage 使用的是全局内存(DRAM), 储存了图片数据
    pix.G = TheImage[row][col+1]; 
    pix.R = TheImage[row][col+2]; 
    TheImage[row][col] = TheImage[row][ip.Hpixels*3-(col+3)]; 
    ...

上述程序如果想要翻转一个像素, 其需要访问6次内存(TheImage), 前三次读取该像素点的三通道的值, 后三次一一将读取的值放到对应的翻转后的位置上, 因为图片的一行总共有3200个像素点, 所以翻转一行总共需要 3200×6=19200 次内存访问

很显然, 上述程序设计十分低效, 因为:

我们每次都读取全局内存, 而且每次读取的内容都不一样, 这样系统根本没办法使用 L1, L2, L3 缓存将常用的数据缓存下来用以加速
当我们访问 DRAM 时, 尽可能地一次存取较大并且连续的块, 譬如 1KB, 4KB 之类大小的块, 而不是一个字节一个字节地存取, 那样相当耗时, 而上述程序读取 TheImage 时, 都是以像素点通道为单位读取, 一个像素点通道一个字节, 读取粒度太小

那么改进后的代码为:

unsigned char Buffer[16384]; // This is the buffer to use to get the entire row 
...
for(row=ts; row<=te; row++) {
// bulk copy from DRAM to cache
  memcpy((void *) Buffer, (void *) TheImage[row], (size_t) ip.Hbytes); col=0;
  while(col<ip.Hpixels*3/2){
     pix.B = Buffer[col];
     pix.G = Buffer[col+1];
     pix.R = Buffer[col+2];
     Buffer[col] = Buffer[ip.Hpixels*3-(col+3)];
     Buffer[col+1] = Buffer[ip.Hpixels*3-(col+2)];
     Buffer[col+2] = Buffer[ip.Hpixels*3-(col+1)];
     Buffer[ip.Hpixels*3-(col+3)] = pix.B;
     Buffer[ip.Hpixels*3-(col+2)] = pix.G;
     Buffer[ip.Hpixels*3-(col+1)] = pix.R;
     col+=3;
}
// bulk copy back from cache to DRAM
memcpy((void *) TheImage[row], (void *) Buffer, (size_t) ip.Hbytes); 
...

上述代码的优点是:

通过一个 Buffer, 每次将一行的像素数据缓存到 L1, L2, L3 缓存中, 这样一来有两个好处, 一是不再需要访问内存, 并提高了效率, 而且 buffer 相当是连续的储存, 不像之前访问 TheImage 那样离散地访问像素点
因为所处理的图片是 22MB 的, 假设我们的 L3 缓存是 15MB, 那么在程序运行的过程中, L3 缓存会不断的将不常用的数据删除, 并读取常用的像素点, 如此周而复始

上述说法有一点不准确, 即不是我们主动地将其放入 L1, L2, L3 缓存中, 在我们创建 Buffer 的时候, 其还是在 DRAM 中的, 只是操作系统发现 Buffer 比较常用, 因此主动地将其放入 L1, L2, L3 缓存中

程序设计所用内存

4.jpg

当我们执行一个程序时, 它需要几个内存空间

一个栈内存区
- 其用来储存传入函数的的参数变量和从函数返回的参数变量及指针, 在函数中定义的一些基本类型的变量和对象的引用变量都在函数的栈内存中分配。当在一段代码块定义一个变量时，编译器就在栈中为这个变量分配内存空间，当超过变量的作用域后，编译器会自动释放掉为该变量所分配的内存空间，该内存空间可以立即被另作他用。
- 栈的优势是，存取速度比堆要快，仅次于寄存器，栈数据可以共享。但缺点是，存在栈中的数据大小与生存期必须是确定的，缺乏灵活性。栈中主要存放一些基本类型的变量（,int, short, long, byte, float, double, boolean, char）和对象句柄。栈有一个很重要的特殊性，就是存在栈中的数据可以共享。
- 在上下文切换的时候(context switch), 原有线程的栈区会被保存下来
一个堆内存区
- 在标准C语言上，使用malloc等内存分配函数获取内存即是从堆中分配内存, 从堆中分配的内存需要手动释放，如果不释放，而系统内存管理器又不自动回收这些堆内存的话, 那就一直被占用。如果一直申请堆内存，而不释放，内存会越来越少，很明显的结果是系统变慢或者申请不到新的堆内存。而过度的申请堆内存 (譬如在函数中申请一个1G的数组), 会导致堆被压爆
- 我们掌握堆内存的权柄就是返回的指针，一旦丢掉了指针，我们便无法主动释放它。这便是内存泄露。而如果在函数中申请一个数组，在函数体外调用使用这块堆内存，结果是未定义的。我们知道在c/c++ 中定义的数组大小必需要事先定义好，他们通常是分配在静态内存空间或者是在栈内存空间内的，但是在实际工作中，我们有时候却需要动态的为数组分配大小，这时就要用到堆内存分配的概念。
- 在堆内存分配时首先应该知道操作系统有一个记录空闲内存地址的链表，当系统收到程序的申请时，会遍历该链表，寻找第一个空间大于所申请空间的堆结点，然后将该结点从空闲结点链表中删除，并将该结点的空间分配给程序，另外，对于大多数系统，会在这块内存空间中的首地址处记录本次分配的大小，这样，代码中的delete语句才能正确的释放本内存空间。另外，由于找到的堆结点的大小不一定正好等于申请的大小，系统会自动的将多余的那部分重新放入空闲链表中。
- 堆内存是向高地址扩展的数据结构，是不连续的内存区域。这是由于系统是用链表来存储的空闲内存地址的，自然是不连续的，而链表的遍历方向是由低地址向高地址。堆内存的大小受限于计算机系统中有效的虚拟内存。
- 由此可见，堆内存获得的空间比较灵活，也比较大。堆内存是由new分配的内存，一般速度比较慢，而且容易产生内存碎片,不过用起来最方便
一个储存代码的区域
- 储存程序代码和在程序中定义的 constant 变量, 这个区域不会被修改, 包括这个区域内储存的常量变量

我们需要知道的是, 一个程序设计为 12 个线程执行并不意味着速度加快了 12 倍, 在不同线程之间切换以及上下文切换, 或多线程多线程抢占内存带宽造成的阻塞等等, 会产生一个 Parallelization Overhead, 我们可以通过下述公式计算 Parallelization Overhead

$Parallelization Overhead=1- \frac{单个线程执行时间}{N个线程总共执行时间\times N}$

物理内存与虚拟内存

系统将所请求的内存大小看做一系列的 pages 的组合, 即页, 每个页是 4KB, 如果一个用户请求 1MB 的内存大小, 那么其请求 256 页
而系统会将当前常用的页放置在内存中, 不常用的放在磁盘中, 而仅仅会提供给程序一个虚拟地址, 这个虚拟地址可能在内存中, 也可能在磁盘中
因为虚拟内存的缘故, 所以假设我们的物理内存是 8GB, 但我们却可以分配超过 8GB 的虚拟内存
malloc() 会分配一个虚拟内存地址, 该地址指向的位置既可能在磁盘中, 也可能在内存中

锁页内存与可分页内存

锁页内存是分配的内存地址就是物理内存内的地址, 而不像虚拟内存一样可能在磁盘中, 不过若分配过多的锁页内存, 会导致其他程序可用的内存减小。
- 假设物理内存是 8GB, 而虚拟内存是 64GB, 为某个程序分配了 2GB 的锁页内存后, 此时可用物理内存还剩 6GB, 不过虚拟内存还是 64GB, 虽然虚拟内存还是 64GB, 但是其灵活性下降, 因为可用的真实物理内存从 8GB 下降到 6GB
malloc()将分配标准的，可分页的主机内存。而cudaHostAlloc()将分配页锁定的主机内存。页锁定的主机内存也称为固定内存或不可分页内存，它的重要属性就是：操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此，操作系统能够安全的使用应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位。
当使用可分页内存进行从CPU到GPU的复制时，复制操作将执行两遍，第一遍从可分页内存复制到一块"临时的"锁页内存，然后再从这个锁页内存复制到GPU上。因此，当在GPU和主机间复制数据时，这种差异会使也锁定主机内存的性能比标准可分页内存的性能要高大约2倍。
cudaHostAlloc 和 malloc 操作都会在系统内存中分配一块区域, 但其区别是, cudaHostAlloc 除了像 malloc 分配内存外, 还额外有一个锁页操作, 而这锁页内存会消耗额外的时间
固定内存是一把双刃剑，当使用固定内存时，将失去虚拟内存的功能。特别是，应用程序中使用每个锁页内存时都需要分配物理内存，因为这些内存不能交换到磁盘上。这意味着，与使用标准的malloc调用相比，系统将更快的耗尽内存。因此，建议仅对 cudaMemcpy() 调用中的源内存或者目标内存，才使用锁页内存，并且在不需要的时候立即释放。
锁页内存允许GPU上的DMA控制器请求主机传输，而不需要CPU主机处理器的参与
CPU仍然可以访问上述锁页内存，但是此内存是不能移动或换页到磁盘上的
在GPU上分配的内存默认都是锁页内存，这只是因为GPU不支持将内存交换到磁盘上
在主机上分配锁页内存有以下两种方式：
- 使用特殊的cudaHostAlloc函数，对用的释放内存使用cudaFreeHost函数进行内存释放
- 使用常规的malloc函数，然后将其注册为(cudaHostRegister)锁页内存,注册为锁页内存只是设置一些内部标志位以确保内存不被换出，并告诉CUDA驱动程序，该内存为锁页内存，可以直接使用而不需要使用临时缓冲区

使用锁页内存需要注意以下几点：

锁页操作会消耗额外的时间, 随着所分配的内存区域的增大, 其耗时也增长, 所以不能分配太多，太多的话会降低系统整体性能
锁页内存和显存之间的拷贝速度是6G/s，普通的内存和显存之间的拷贝速度是3G/s（显存之间的拷贝速度是30G/s,CPU之间的速度是10G/s）
使用cudaHostAlloc函数分配内存，其内的内容需要从普通内存拷贝到锁页内存中，因此这种拷贝会带来额外的CPU内存拷贝时间开销，CPU需要把数据从可分页内存拷贝到锁页，但是采用cudaHostRegister把普通内存改为锁页内存，则不会带来额外的cpu内存拷贝时间开销，因为cudaHostAlloc的做法是先分配锁页内存，这时里面是没有数据的，那么需要将一般的内存拷贝过来，而对于cudaHostRegister内存，他是之间就使用malloc分配好的，cudaHostRegister只是设置一些内部标志位以确保其不被换出，相当于只是更改了一些标志位，就不存在前面说的数据拷贝
在某些设备上，设备存储器和主机锁页存储器之间的数据拷贝和内核函数可以并发执行
在某些设备上，可以将主机的锁页内存映射到设备地址空间，减少主机和设备之间的数据拷贝，要访问数据的时候不是像上面那那样将数据拷贝过来，而是直接通过主机总线到主机上访问，使用cudaHostAlloc分配时传入cudaHostAllocMapped，或者使用cudaHostRegister时传入cudaHostRegisterMapped标签
默认情况下，锁页内存是可以缓存的。在使用cudaHostAlloc分配时传入cudaHostAllocWriteCombined标签，将其标定为写结合，这意味着该内存没有一级二级缓存，这样有利用主机写该内存，而如果主机读取的话，速度将会极其慢，所以这种情况下的内存应当只用于那些主机只写的存储器

1.4 线程管理与同步

5.jpg

如果当前 CPU 是 10核20线程的, 那么在某一时刻, 不能超过 20个线程出于 Running 状态
当我们调用 pthread_create() 创建一个进程时, 操作系统则会去查看, 是否有足够的资源来新建一个线程? 如果有的话, 那么一个 handle 会被分配给该线程, 并为其创建栈区和必要的内存空间, 此时, 这个新建立的线程就进入了 Runnable 队列, 等待被调度
一旦某个在 Runnable 队列中的线程准备执行, 那么一个虚拟的 CPU 则会分配给它, 而其此时进入了 Running 状态
当某个 Running 状态的线程因为等待某个资源(譬如调用了 scanf 函数, 等待用户输入), 而卡住时, 其会被重新放入到 Runnable 队列中, 同时, 其在 Running 状态得到的寄存器信息等等, 即其运行状态, 会被保存到一个区域, 当其等待的资源准备好后, 重新运行该线程, 并将其寄存器信息恢复, Running ==> Runnable 这个过程称为 context switch(上下文切换)
当某线程等待的资源暂时无法得到或无法确定大概需要多久能得到时, 那么该线程则会被放到 Stopped 队列中, 当该资源准备好后, 则该线程又会回到 Runnable 队列中
如果一个线程执性完成, 那么其就会进入 Terminated 状态, 被放入 Terminated 中的线程无法再回到 Runnable 队列中来
假设一个 CPU 最多只能支持 20核, 如果某个程序launch了25个线程, 有可能其效率还是比launch 20个线程高, 这是因为 context switch(上下文切换) 的缘故

2. GPU 并行

2.1 GPU / CUDA 概述

在游戏领域, 3D 人物的建模都是用一个个小三角形拼接上的, 而不是以像素的形式, 对多个小三角形的操作, 能使人物做出多种多样的动作, 而 GPU 在此处就是用来计算三角形平移, 旋转之后的位置, 如下图
- 6.jpg
  * E2C991444715105B658510DADAB9FF7D.jpg
- 而为了提高游戏的分辨率, 程序会将每个小三角形细分为更小的三角形
- 每个小三角形包含两个属性, 它的位置和它的纹理
在游戏领域应用的 GPU 与科学计算领域的 GPU 使用的不同是, 当通过 CUDA 调用 GPU 来进行科学计算的时候, 计算结果需要返回给 CPU, 但是如果用 GPU 用作玩游戏的话, GPU 的计算结果直接输出到显示器上, 也就不需要再返回到 CPU
因为 CPU 和 GPU 的硬件架构上的不同, 譬如 CPU 的内存是用 DDR4 而 GPU 的内存是用 GDDR5, 它们的 ISAs (Instruction set architecture) 应该完全不同, 但是 nvcc 编译器解决了这一问题, 其编译的 .cu 文件(类 c 语言), 将其中涉及 CPU 的代码编译为 CPU 指令, 而将涉及 GPU 的代码编译为了对应的 GPU 指令, 大大简化了 CUDA 编程的难度。

2.1 GPU 通用架构

GPU 软件架构

每 32 个线程组成一个线程束 (warp), 一个线程束可以被看做是程序执行的一个最小单元, 而 GPU 也是基于线程束来执行程序, 假如某个代码只需要 12 个线程, 那么 GPU 也会启动 32 个线程来执行这个代码, 其中的 12 个线程就被完全浪费了
但一般来说, 一个仅包含 32 个线程的 warp 实在是太小了, 因此我们就将多个线程束合在一起, 组成 block, 常用的 block 大小是 32, 64, 128, 256, 512, or 1024 threads/block
- 每个 warp 相当于是 code execution 的最小单元
- 每个 block 相当于是 code launch 的最小单元
而多个 block 在一起组成一个 grid, grid 可以使 1D, 2D或3D 的 blocks 的排列
- 虽然GPU可以用3D grid, 譬如 x, y , z 每一维度都可以 launch 假设200个block, 总共200^3=8000000 个block, 但不代表这 800000 可以在用 1D grid 的时候全部用一个维度 launch, 譬如对于 GT630 GPU 而言其 x 维度最大只支持 65535 个block
block 之间不能有资源上的相互依赖
假设一个 block 有 256 个线程, 不代表每个线程都会被即刻执行, 其首先会被分为 ceil(256/32)=8 个 warp, 然后将这 8 个 warp 按从 0 到 7 编号, 然后按 warp0, warp1, …, warp7 的书序执行
- 正常情况下, 在写 CUDA 程序时, 我们是不用考虑 warp 的 ID 的, 仅当我们写底层的 CUDA assembly language (PTX) 编码时才需要考虑 warp ID
CPU 和 GPU 之间的任何交互都会通过 NRE (Nvidia Runtime Engine), NRE 存在于显卡驱动中
- 当 NRE 检测到代码中有错误的地址访问时, 会终止程序, 但是需要注意的是, 程序是运行在 CPU 中的, 所以仅仅操作系统才可以终止程序, 所以 NRE 会通知操作系统, 然后操作系统将该程序终止
每一代的 Parallel Thread Execution (PTX) 和 Instruction Set Architecture (ISA) 都会改进
- PTX 是 Intermediate Representation (IR) 并且对于不同操作系统而言都是一样的, 我们在用 nvcc 编译 .cu 文件时, 可以加上 --ptx 来生成 PTX代码, 也可加上 --cubin 来生成 CUDA 的二进制文件, CUBIN 针对不同系统是不同的

GPU 硬件架构

7.jpg

Giga Thread Scheduler (GTS) 模块用来管理安排 block, 即每当一个 SM 完成当成工作后, GTS 会给其安排一个新的 block
一个 GPU 有越多的 SM, 其速度就越快
虽然每个 SM(Streaming Multiprocessor) 一次只能执行一个 block, 但其可以一次接收多个 block, 并将其放入一个队列中缓存起来
假设 launch 了 166656 个 block, 但是一个 Pascal 显卡的 SM 最多接收 32 个 block, 而该显卡有 60 个 SM, 32*60=1920, 远远小于 launch 的 166656 个block, 那么 GTS 就让剩下的 block 就排在队列之中等待之前的 SM 执行完毕后空出来, 同时需要注意的是, 分配给每个 SM 的 32 个 block 中, 一次只能执行一个 block, 其他的 31 个 block 都需要等待, 同时 GTS 也会将 blockDim, blockId 以及 gridDim 等参数随着 block 传入到 SM 中, 用来帮助计算每个 thread 的 ID
每一个 block 都会得到一个当前执行的 CUDA 命令代码的二进制版本, 即 CUBIN, 而每个 SM 在收到一个 block 时会将其包含的 CUBIN 缓存在 SM 中的指令缓存中
假设现在有 6 个 SM, 那么加载 blocks 的顺序就是:
- Block0→SM0, Block1→SM1, Block2→SM2, Block3→SM3, Block4→SM4, Block5→SM5
- Block6→SM0, Block7→SM1, Block8→SM2, Block9→SM3, Block10→SM4, Block11→SM5
- 那么每个 SM 中缓存 Block 的队列即为:
- SM0 =⇒ [ Block0, Block6, Block12, Block18, Block24, Block30, Block36, Block42 ]
- SM1 =⇒ [ Block1, Block7, Block13, Block19, Block25, Block31, Block37, Block43 ]
- …...
- SM5 =⇒ [ Block5, Block11, Block17, Block23, Block29, Block35, Block41, Block47 ]
当 SM 中当前执行的 block 因为访问某些资源而阻塞时, SM 可以选择执行队列中的其他的 block, 以避免 SM 空闲下来, 因此, 对于 SM1而言, 有可能出现的情况是, Block7 率先完成了工作, 此时 GTS 就会将下一个 Block 传入到 SM1 中原来 Block7 所在的位置, 因为下一个需要执行的是 Block48 (Block47 已经传给 SM5 了), 那么此时 SM1 的队列就是
- SM1 =⇒ [ Block1, Block48, Block13, Block19, Block25, Block31, Block37, Block43 ]
当 GTS 给一个 SM 分配一个 block 后, 其就与这个 block 没关系了, 就该轮到该 SM 给这个 block 分配 thread ID, 缓存资源, 寄存器等

2.2 GPU 不同系列架构

Fermi 架构

8.jpg

9.jpg

SFU 是 Special Function Unit, 其用来执行一些特殊的函数, 譬如 sin(), cos(), log() 等
LD/ST 是 Load/Store, 对于内存数据的读写请求都被放置在该模块的队列中, 当读写完成后, 该请求从队列中移除
每个 Core 都包含一个 Floating Point (FP) 和一个 Integer (INT) 执行单元用来执行浮点或整型的指令
L1 缓存用来储存那些常用的数据, 其与 Shared Memory 共用 64 KB 的大小, 64 KB 被分为 (16 KB+48 KB) 或 (48 KB+16 KB)

Kepler 架构

10.jpg

11.jpg

SM 在 Kepler 架构中被称为 SMX
DPU 是 Double Precision Unit, 即可以更高效地计算 double 类型数据

Maxwell 架构

12.jpg

13.jpg

SM 在 Maxwell 中被称为 SMM
一个 SMM 包含四个 sub-uint, 其中每个 sub-unit 包含 32 个 core, 这四个 sub-unit 共用一个 Instruction Cache, 而每个 sub-unit 独占一个 Instruction Buffer

Pascal 架构

14.jpg

15.jpg

High Bandwidth Memory (HBM2), 其通过使用4096-bit 内存总线带宽, 最高可支持 720 GBps 传输速率
GTS 在此处称为 GIGA THREAD ENGINE (GTE)

模块介绍

FPU (Floating Point Unit)
- FPU 可以用来计算 double 类型数据, 其通过不断地循环来计算
- 对于那些没有 DPU 模块的 GPU, 在使用 FPU 来计算 double 类型数据时 (譬如 24 位或 32 位的double), 其耗时是计算单精度浮点类型数据的 24 或 32 倍
DPU (Double Precision Unit)
- DPU 往往会比 Core 在物理上大一些, 因为随着尾数的增加, 乘法器的大小呈平方扩大, 即一个单浮点类型需要 23 位尾数, 而一个双精度浮点需要 52 位尾数, 故与 FPU 相比一个 GPU 大概是其四倍的大小 (注意 FPU 是包含在 Core 中的)
SFU (Special Function Unit)
- 其用来计算一些特殊的运算, 如 sin(), cos(), exp(), log(), sqrt() 等
RF (Register FIle)
- 假设现在程序中有如下变量
  - double R, G, B
  - unsigned int ThrPerBlk, MYbid, MYtid, MYgtid,
  - unsigned int BlkPerRow, RowBytes, MYrow, MYcol,
  - unsigned int MYsrcIndex, MYpixIndex
- 假设每个寄存器是 32-bit 的, 那么每个 double 类型都需要 2 个寄存器, 而每个 unsigned int 需要 1 个, 因此为了储存这些变量总共需要耗费 16 个寄存器, 而编译器可能还需要消耗一些寄存器储存临时运算的值, 那么假设该程序总共需要耗费 24 个寄存器
- 假设我们每个 block 有 128 个线程, 那么每个 block 就需要 24*128=3072≈3K 个寄存器, 而假设 Pascal 架构的 GPU 每个 SM 的队列中最多可容纳 32 个 block, 那么总共需要 32*3K=96K 个寄存器, 而若每个 SM 只有 32K 个寄存器(总共 128KB), 因此对于 GTE 而言, 因为寄存器(RF)数量的限制, 其最多只能向每个 SM 中加载 10 个 block
- 所以在设计 CUDA 程序时, 尽可能地限制 RF 的使用
LDST (Load/Store Queues)
- 该模块用来在 core 和 memory 之间传输数据, 当一个 core 请求从 memory 中读取或写入数据时, 该请求则会被储存在 LD/ST 中, 并等待其完成, 在等待其完成的过程中, 另外一个 warp 会被启动并执行
L1 和 Texture Cache
- L1 缓存是硬件控制的缓存, 即代码无法控制 L1 缓存中储存的数据, L1 缓存用以储存常用的数据
- Texture 缓存用来储存之前提到的游戏建模中小三角形的纹理
Shared Memory
- 共享内存是代码控制的储存, 通过 CUDA 编程, 可以显式地规定哪部分数据储存在共享内存中
- 当使用 Shared Memory 的时候要注意内存大小, 不然有可能比没用 Shared Memory 速度还慢, 因为当 shared memory 用完的时候, 没有资源供新的 block 执行, 故新的 block 会被阻塞, 而对于那些没有用 shared memory 的, 虽然没有 shared memory 稍显耗时, 但其 block 不会被阻塞, 所以最终可能速度更快
Constant Cache
- 其用来保存不可变的值, 即常量
- 其只被写入一次(即初始值), 但可被多次读取
Instruction Cache
- 其用来保存当前 SM 所执行 Block 中的指令, 每个 block 中都包括其所需执行的指令
Instruction Buffer
- 其用来保存当前 SM 的局部指令, 即从 Instruction Cache 中复制来的指令, Instruction Buffer 相对于 Instruction Cache 的关系就像 L1 相对于 L2 的关系
Warp Schedulers
- 当每个 block 被传递给 SM 时, warp scheduler 用来将每个 block 转化成线程束的形式, 即假设我们启动了 256 threads/block 那么就会转化成 8 warps/block, 即转化为如下形式
  - schedule warp0: gridDim.x=166656, blockDim.x=256, blockIdx.x=0
  - schedule warp1: gridDim.x=166656, blockDim.x=256, blockIdx.x=0
  - …...
  - schedule warp7: gridDim.x=166656, blockDim.x=256, blockIdx.x=0
- 需要注意的是此处只是 schedule 了, 而不是 dispatch, 所以此处没有分配 thread ID,仅当所有资源都准备好时, 才会 dispatch
Dispatch Units
- 该模块在每个 warp 已经被 schedule 后, 并且所有的资源都准备完毕时启动
- 其会给每个线程分配 threadIdx.x, threadIdx.y, 和 threadIdx.z, 即如下形式
  - gridDim.x=166656, blockDim.x=256, blockIdx.x=0, threadIdx.x=0…32
  - 32 个 core, 每个 core 负责执行一个线程
GLOBAL MEMORY
- 全局内存与 DDR4 内存一样, 一次读取一块内存的效率比一个比特一个比特读取的效率更高

2.3 CPU / GPU 传输

一些定义
- latency(延迟), 第一个包从发出到收到的时间差
- throughput(吞吐量), 在一段时间内的平均传输速率, 即传输的总量除以总时间
- bandwidth(带宽), 最大的 throughput(吞吐量)
- UpStream 带宽是 CPU->GPU 的带宽, 而 DownStream 带宽是 GPU->CPU 的带宽, 而 PCIe 的好处是支持 UpStream和DownStream 的同时传输
当数据通过 PCIe 传输时, 无论是从 CPU->GPU 还是 GPU->CPU, CPU 始终通过虚拟内存页来参与到这个过程中
仅当 CPU, GPU 以及主板都支持 PCIe 的某个特定版本时, 譬如 PCIe 3.0, 才能充分发挥 PCIe 的传输速率, 如果其中某一个不支持, 那么则无法使用 PCIe 3.0
16.jpg
- 图中的内存总线是 L3 与 DDR4 之间的连接部分, 68 GB/s, 即 CPU 缓存与内存的传输速率
- PCIe 3.0 是 PCI EXPRESS BUSS, 传输速率为 16 GB/s
- GPU 内部的储存, 被称作 GPU global memory, 即图中的 GDDR5, 其与 GPU L2 缓存交互的速率是 336 GB/s , 需要注意的是 GPU 最多只有 L2, 没有 L3 缓存, 即 GPU 中的 L2 缓存与 CPU 中的 L3 缓存都是最后一层缓存 Last Level Cache (LLC), LLC 直接与内存交互
17.jpg
- 上图是不同内存的带宽

2.4 限制 SM 速度的因素

Shared Memory
- 如果程序被设计为需要大量依赖共享内存, 那么当共享内存耗尽时, SM 就会被阻塞, 而不能在不同 block 之间切换执行
Register File
- 使用的寄存器的数量与代码中定义的变量数量有关, 其用来保存 kernel 中变量的值
- 假设每个 kernel 中可使用的寄存器是 255 个, 如果使用的寄存器超出了这个部分, 则其就会用内存当做寄存器, 而内存当做寄存器很显然会慢, 而每个 block 中最多有 32K 或 64K 个寄存器, 而对每个 block 允许的寄存器的数量的限制, 直接影响了每个 block 中所能容纳的 thread 数量
  - if your kernels require 240 registers, launching 512 threads/block will require 512 × 240 = 122,880 = 120 K registers inside the register ﬁle. However, even if your Compute Capability allows 64 K registers, this is still more than the 120 K you need. This means that you cannot launch more than 256 threads/block before hitting the register ﬁle limitation, which would require 256 × 240 = 61,440 = 60 K registers
  - 假设当前的 kernel 需要 240 个寄存器, 那么启动 512 threads/block 就需要 512 × 240 = 122,880 = 120 K 个在RF中的寄存器, 然而, 假设根据我们当前的 compute capability, 我们只能允许 64K 个寄存器, 很显然是不够的, 因此我们不被允许启动 512 threads/block, 而最多只能启动 256 threads/block, 因为 256 × 240 = 61,440 = 60 K < 64 K
Threads / Block
- 如果该值太小, 即每个 block 过小, 那么我们就需要更多的 block 完成特定任务, 但是每个 SM 一次启动的 block 数量是有上限的
- 如果该值太大, 那么有可能会造成浪费, 譬如我们总共需要 5220 个线程, 而我们的 block 大小是 512 threads / block, 那么我们就需要 11 个 block, 11 * 512 - 5220=412, 即我们浪费了 412 个线程
Warp
- 每个 SM 中可容纳的 warp 数量, 假设该值是 64, 那么一个 warp 中有 32 个线程, 那么也就意味着每个 block 最多可容纳 32*64=2048 个线程
我们将上述情况汇总, 以一个例子来看一下:
- 现在有以下条件
  1. 假设 block 大小是 320 threads/block (10 warps/block)
  2. 假设 GTS 给当前 SM 分配了 6 个 block
  3. 每个 SM 中可容纳 64 个 warp
  4. 假设每个 block 都需要 10 KB 的 shared memory 大小
  5. 当前 SM 可允许的最大 shared memory 大小是 48 KB
- 根据前三个条件, 因为 SM 分配了 6 个 block, 故总共启动了 6*10=60 个 warp, 也就是说每个 SM 浪费了 64-60=4 个warp, 此时的占用率是 60/64≈94%
- 根据后两个条件, 因为 shared memory 的限制, 我们最多只能启动 4 个 block, 因为 4*10=40<48KB, 此时占用率为 4*10(warps/block) / 64 ≈ 63%
- 假如我们将 block 的大小改为 512 threads/block (16 warps/block), 此时每个 block 依然只消耗 10 KB 的 shared memory, 那么以上条件不变的话, 我们启动 4 个 block, 就可以达到 4*16/64=100% 的占用率

2.5 设计 CUDA 代码时的一些注意事项

通过如下代码我们在 main 函数中定义了一个 buffer 变量, 并给其分配了 GPU 上的内存
- unsigned char *buffer;
- cudaMalloc((void**)&buffer, SIZE);
- 你可能会感到困惑, 从 buffer 的定义方式上来看, 其很显然是一个 CPU 上的指针变量, 指向 GPU 上的某块内存区域, 这说法看起来很让人困惑, 因为在常理中, CPU 上的指针变量指向的应该是 CPU 的内存地址
- 我们首先来看一个指针变量是什么, 一个指针变量无非就是一个 64位的integer , 这个 integer 代表了内存上某块地址, 正因为指针变量是一个 integer, 所以其可以进行加减操作, 通过加减操作可以指向当前地址前后的不同地址
- 那么对于这个指针变量本身, 其只知道自己储存的是一个 64位的integer, 而无从了解其是内存中的还是GPU内存中的地址, 而判断这个的是 nvcc 编译器, 其通过 buffer 这个 64位integer 的值, 判断出其是属于内存还是GPU内存, 然后根据其所属地, 分配出一块空间
GPU 对于 a*(b+c) 有特殊的运算操作符, 但是对于 a+b+c 没有, 这就可能导致 a+b+c 反而比 a*(b+c) 还慢
OR, AND, 移位操作 (>>) 等在 GPU 中计算速度很快, 可以在代码中尽可能使用, 譬如如下代码
- //NOW: A=[B1,R0,G0,B0] B=[G2,B2,R1,G1] C=[R3,G3,B3,R2], 其中 A, B, C 都是 32 bit 变量, 这三个变量每个都包含 4 个内部值, 每个值都是 8 bit
- // D=[B2,R3,G3,B3]
- D = (C >> 8) | ((B << 8) & 0xFF000000); //其中 0xFF000000 转化为二进制为 11111111,00000000,00000000,00000000
- // E=[G1,B1,R2,G2]
- E = (B << 24) | (B >> 24) | ((A >> 8) & 0x00FF0000) | ((C << 8) & 0x0000FF00); // 0x0000FF00 转化为二进制为 00000000,00000000,11111111,00000000
- // F=[R0,G0,B0,R1]
- F=((A << 8) & 0xFFFF0000) | ((A >> 16) & 0x0000FF00) | ((B >> 8) & 0x000000FF);
线程束分化 (thread divergence), 当一个 warp 中, 因为数据不同, 不同线程运行同一份代码时, 根据 if 而给出不同的 TRUE/FALSE 结果, 并因为结果不同而执行不同的操作时, 就会导致线程束分化, 线程束分化会影响并行效率, 因为仅当一个 warp 中所有线程都在做相同的事时, 效率才最高

本文主要内容来自书籍: GPU Parallel Program Development Using CUDA (Tolga Soyata)
[http://www.hds.bme.hu/~fhegedus/00%20-%20Numerics/B2018%20GPU%20Parallel%20Program%20Development%20Using%20CUDA.pdf]