- host CPU和内存 (host memory)
- Device GPU和显存 (device memory)
SIMT模型与SIMD模型的区别
SIMD(Single Instruction Multi Data，单指令多数据)模型要求同一个向量中的所有元素要在统一的同步组中一起执行，而SIMT则允许属于同一个线程束的多个线程独立执行，这几个线程可以有不同的行为。因此SIMT(Single Instruction Multi Thread)允许线程级并发，也就是在统一线程束下的线程可以同时做不同的事情。
SIMT有SIMD所不具备的三个特征：
- 每个线程都有自己的指令地址计数器
- 每个线程都有自己的寄存器状态
- 每个线程可以有一个独立的执行路径
硬件基础
系统每一bit的输入与输出，包括磁盘、网络控制器、键盘、鼠标、USB设备以及GPU的输入输出，都要通过芯片组。直到最近，芯片组被一分为二：一个是连接大多数外围设备和系统的“南桥”，另一个是包含图形总线(加速图形端口，后被PCIe接口取代)和内存控制器(通过前端总线与内存相连)的“北桥”。
外部数据总线是中央处理器CPU(Central Processing Unit)的一部分，是CPU与外部数据传输的通道。外部数据总线一次可传输二进制数据的位数越大，CPU与外部交换数据的能力越强。
GPU内存控制器总是和GPU集成，它是在一个与CPU内存控制器完全不同的约束集下设计。
CPU的缺陷
CPU 与 GPU 架构的一个主要区别就是 CPU 与 GPU 映射寄存器的方式。 CPU 通过使用寄存器重命名和栈来执行多线程。为了运行个新任务, CPU 需要进行上下文切换,将当前所有寄存器的状态保存到栈(系统内存)上,然后从栈中恢复当前需要执行的新线程上次的执行状态。这些操作通常需要花费上百个 CPU 时钟周期。如果在 CPU 上开启过多的线程,时间儿乎都将花费在上下文切换过程中寄存器内容的换进/换出操作上。因此,如果在 CPU 开启过多的线程,有效工作的吞吐量将会快速降低。
然而, GPU 却恰恰相反。 GPU 利用多线程隐藏了内存获取与指令执行带来的延迟。因此,在 GPU 上开启过少的线程反而会因为等待内存事务使 GPU 处于闲置状态。此外, GPU 也不使用寄存器重命名的机制,而是致力于为每一个线程都分配真实的寄存器。因此,当需要上下文切换时,所需要的操作就是将指向当前寄存器组的选择器(或指针)更新,以指向下一个执行的线程束的寄存器组,因此儿乎是零开销。
一个线程束即同时调度的一组线程。在当前的硬件中,一个线程束包含32个线程。因此,在一个 SM 中,每次换进/换出、调度都是32个线程同时执行。
每个 SM 能调度若干个线程块。在 SM 层,线程块即若干个独立线程束的逻辑组。编译时会计算出每个内核线程需要的寄存器数目。所有的线程块都具有相同的大小,并拥有已知数目的线程,每个线程块需要的寄存器数目也就是已知和固定的。因此, GPU 就能为在硬件上调度的线程块分配固定数目的寄存器组。
线程层次
每32线程分为一个warp
GigaThread/Grid：GPU级别的调度单位，共享L2级别的Cache
SM(Streaming Multiprocessor)/block：SM调度block(线程块)，共享L1级别的Cache
Warp/Threads：一个Warp包含32个hread，Warp是CUDA core级别(一个thread对应一个CUDA core，但一个CUDA core可以通过分时复用的方式实现多个thread)的调度单位，允许同一个warp中的thread读取其他thread的值，共享L0级别的Cache
每一个warp共享一个contex执行上下文 取指译码器
Thread : sequential execution unit
- 所有线程执行相同的核函数
- 并行执行
Block : a group of threads
- 执行在同一个SM(Streaming Multiprocessor)
- 同一个Block中的线程共享一块Shared Memory（L1级别的Cache）
- CUDA不支持超过1024个线程的线程块
Grid : a collection of thread blocks
- 一个Grid当中的Block可以在多个SM中执行，一个SM可以有很多Block
每个SM所能容纳的block数量受到以下两方面限制：
- 每个SM可提供的寄存器空间固定。
如果每个线程所需的寄存器过多，则每个SM能够调度的线程数量就受到限制。 - 每个SM能够调度的线程束的数量有限制。
CUDA的执行流程
- 加载核函数
- 将Grid分配到一个Device(GPU)
- 根据<<<…>>>内的第一个参数，Giga thread eigine将block分配到SM中。一个Block内的线程一定会在同一个SM内，一个SM可以有很多Block。
- 根据<<<…>>>内的第二个参数，Warp调度器将调用线程
- Warp调度器为了提高运行效率，会将每32个线程分为一组，称作一个warp。block(x,y,z)按先X方向，再Y方向，再Z方向优先组合成warp
- 每个warp会被分配到32个Core上运行。
CPU与GPU交互、多GPU间交互原理
Page-locked or pinned memory transfers attain the highest bandwidth between the host and the device.
Pinned memory is allocated using the cudaHostAlloc() functions in the Runtime API.
cudaHostRegister() 可以将host memory变为锁页内存.
计算和内存拷贝异步执行
cudaMemcpy()是主机内存与设备内存间阻塞式数据传输的API，只有当数据传输完成后才会将控制权返回主机线程。
cudaMemcpyAsync()是主机锁页内存与设备内存间非阻塞式数据传输的API，会立即将控制器返回主机线程。比cudaMemcpy()多一个参数Stream ID.
- 锁页主机内存：GPU可以直接访问的CPU内存；
- 命令缓冲区：由CUDA驱动程序写入命令，GPU从此缓冲区读取命令并控制其执行；
- CPU/GPU同步：指的是CPU如何跟踪GPU的进度。
CPU/GPU并发
CUDA内核的所有启动都是异步的：CPU通过将命令写入命令缓冲区来请求启动内核，然后直接返回，而不检查GPU进度。内存复制也可以采用异步方式，这使CPU/GPU并发以及可能使内存复制与内核处理并发执行。
命令缓冲区和同步位置都位于锁页主机内存上，被CUDA驱动程序和GPU使用，以跟踪相应进程。主要的GPU操作后都跟着一个将新进度值写入共享同步位置的命令。如下图所示，进度值将一直为3,直到GPU完成当前命令的执行并将4写人同步位置中为止。
CUDA既隐式又显式地暴露了这些硬件功能。上下文范围的同步通过简单调用cuCtxSynchronize()、cudaThreadSynchronize()等函数来检查GPU请求的最近同步值，并且一直等待，直到同步位置获得该值。例如，在下图中，如果由CPU写人的命令8之后紧接着cuCtxSynchronize()或cudaThreadSynchronize()函数，则驱动程序将一直等待，直到共享同步值大于或等于8为止。
CUDA 事件则更明确地暴露了这些硬件能力。cuEventRecord()函数的作用是将一个命令加入队列使得一个新的同步值写人共享同步位置中,cuEventQuery()和 cuEventSynchronize()则分别用于检查和等待这个事件的同步值。早期版本的 CUDA 只是简单地轮询共享的同步位置,反复地读内存,直到等待准则满足为止。但是这种方法代价很大,且只有当应用程序不必等待太久时才有用(即同步位置不一定要被读取很多遍,就可以因等待标准已经得到满足而退出)。对大多数应用程序来说,基于中断的方案( CUDA 公开称为“阻塞同步”)更好,因为它们使 CPU 等待线程挂起,直到 GPU 发出中断信号为止。驱动程序将 GPU 中断映射到一个特定的线程同步原语平台,如Win32事件或 Linux 的信号。
通过指定 CU_CTX _BLOCKING_SYNC 到 cuCtxCreate()或指定 cudaDeviceBlockingSync 到 cudaSetDeviceFlags()，应用程序可以强制使用上下文范围的同步进人阻塞状态。然而,使用阻塞的 CUDA 事件(指定 CU_EVENT_BLOCKING_SYNC 到 cuEventCreate()或指定 cudaEventBlockingSync 到 cudaEventCreate() )更可取，因为它们粒度更细且可以与任何类型的 CUDA 上下文进行无缝互操作。
敏锐的读者可能会关注 CPU 和 GPU 在不使用原子操作或其他同步原语的情况下读取和写人这种共享的内存位置。但由于 CPU 只读取共享位置,竞争条件并不是关注点。这样,最坏的情况是 CPU 读取了一个“过时的”值,从而导致其等待时间要比实际的长。
CPU接口与内部GPU同步
GPU 可能包含多个引擎,以使内核执行和内存复制并发进行。在这种情况下, GPU 的驱动程序将写人一些命令,这些命令被分发到同时运行的不同引擎中。每个引擎都有自己的命令缓冲区和共享同步值。图2-28显示的是两个复制引擎和一个计算引擎并行工作时的情形。其中,主机接口负责读取命令并将其调度到相应引擎。在图2-28中,一个主机到设备的内存复制操作和两个相关操作(一个内核启动和一个设备到主机的内存复制)已被提交给硬件。依据 CUDA 编程抽象模型,这些操作都是在同一个流上进行的。这个流就像是一个 CPU 线程,在内存复制之后接受内核启动的提交。因此, CUDA 驱动程序为了实现内部 GPU 同步必须将命令插入主机接口的命令流中。
如上图所示，主机接口在协调流同步需求上起着核心作用。例如,在完成所需的内存复制之前,内核是启动不了的。此时,DMA单元可以停止给既定引擎传递命令,直到同步位置获得一个特定值为止。此操作与 CPU / GPU 同步类似,但 GPU 是对它内部进行同步的。
此硬件机制之上的软件抽象模型是一个 CUDA 流。它在该操作中就像 CPU 线程一样都是串行排队,为了并发执行需要多个流。由于引擎间共享命令缓冲区,应用程同的流里以软件的方式流水线化它们的请求。
GPU间同步
由于图2-26-图2-28中的同步位置都是在主机内存上,所以它们可以被系统中的任何一个 GPU 访问。其结果是,在CUDA4.0中,英伟达能够在cudaStreamWaitEvent()和 cuStreamWaiEvent() 函数的形式中添加 GPU 之间的同步。这些 API调用导致驱动程序为主机接口将等待命令插人当前 GPu 的命令缓冲区中,使得 GPU 一直等待,直到事件的给定同步值被写入为止。从CUDA4.0开始,事件不一定会被等待中的同一个 GPU用信号唤醒。流原先只能在单个 GPU 的硬件单元之间同步执行,现在已提升到可以在 GPU 之间同步执行了。
