ECE408 Lecture 21 GPU as part of the PC Architecture

这次回顾ECE408 Lecture 21，这次介绍GPU作为PC架构的一部分。

课程主页：

• https://wiki.illinois.edu/wiki/display/ECE408

搬运视频：

• https://www.youtube.com/playlist?list=PL6RdenZrxrw-UKfRL5smPfFFpeqwN3Dsz

目标

• 了解将GPU用作协处理器时数据传输对性能的影响；
  • 传统CPU的speeds和feeds；
  • 使用GPU的speeds和feeds；
• 开发用于现代GPU性能调整的知识库；

回顾：规范的CUDA程序结构

• 全局变量声明
• 内核函数；
  • __global__ void kernelOne(…)
• Main () // 主机代码
  • 在设备上分配内存空间；
    • cudaMalloc(&d_GlblVarPtr, bytes)
  • 将数据从主机传输到设备；
    • cudaMemcpy(d_GlblVarPtr, h_Gl ...)
  • 执行配置设置；
  • 内核调用；
    • – kernelOne<<<execution configuration>>>( args… );
  • 将结果从设备传输到主机；
    • cudaMemcpy(h_GlblVarPtr,…)
  • 可选：与主机计算解决方案进行比较；

带宽：现代计算机系统的引力

关键组件之间的带宽最终决定了系统性能：

• 特别适用于处理大量数据的GPU；
• 在某些情况下，缓冲、重新排序、缓存等技巧可能会暂时违背规则；
• 最终，性能会回落到“speeds和feeds”的要求；

经典（历史）PC架构

• 北桥连接3个必须高速通信的组件
  • CPU、DRAM、视频；
  • 视频需要优先的DRAM访问权限；
  • 以前的NVIDIA卡连接到AGP，传输速度高达2GB/s；
• 南桥用作较慢I/O设备的集中器；

（原创）PCI总线规范

• 南桥的连接
  • 最初为33 MHz，32 位宽，132 MB/秒的峰值传输速率；
  • 后来为66 MHz，64 位，528 MB/秒峰值；
  • 设备的上行带宽仍然很慢（~256MB/s 峰值）；
  • 带arbitration的共享总线；
    • arbitration获胜者成为总线主控，可以通过南桥和北桥连接到CPU或DRAM；

PCI作为内存映射I/O

• PCI设备寄存器映射到CPU的物理地址空间；
  • 通过加载/存储访问
    （内核模式）；
• 地址在启动时分配给PCI设备；
  • 所有设备都监听它们的地址；

PCI高速(PCIe)

交换式点对点连接：

• 每张卡都有专门的“链接”到中央交换机，没有arbitration；
• 分组交换机：消息形成虚拟通道；
• 用于QoS的优先数据包（例如用于实时视频流）；

PCIe Generations

• 在一代内，链路中的通道数量可以缩放；
  • 使用不同的物理通道（更多位/更宽的传输）；
  • ×1, ×2, ×4, ×8, ×16, ×32, …；
• 每个新一代都旨在将速度提高一倍；
  • 当前一代是PCIe 5.0，但它仅在非常有限的一组系统上受支持，例如 IBM Power10；
    • 32GT/秒；
  • 现代AMD、英特尔和IBM系统支持PCIe 4.0；
  • 然而，PCIe Gen 3仍然被广泛使用；

PCIe Gen 3链路和通道

• 每个链路由一个或多个通道组成；
  • 每个通道为1位宽（4条线，每2线pair可以在一个方向上传输8Gb/s）；
    • 2线pair对用于差分信号；
    • 上下游同步且对称；
  • 每个链路可以组合1、2、4、8、12、16条通道；
• 每个字节数据被128b/130b 编码成130 位，1 和0的个数相等；单路每条通道的净数据速率为7.8768 GB/s；
  • 因此，净数据速率为985 MB/s (x1) 1.97 GB/s (x2)、3.94 GB/s (x4)、7.9 GB/s (x8)、15.8 GB/s (x16)；

基础：8/10位编码

• 目标是保持DC平衡，同时为时钟恢复提供足够的状态转换；
• 20位流中1和0的数量差应$\le 2$；
• 任何流中的连续1或0不应超过5个；
• 00000000、00000111、11000001，这样的形式不太好；
• 01010101, 11001100，这样的形式比较好；
• 在10位的1024个模式中找到256个好的模式来编码一个8位数据；
• 20%的开销；

当前：128/130位编码

• 相同的目标：保持DC平衡，同时为时钟恢复提供足够的状态转换；
• 1.5% 的开销而不是20%；
• 每66位至少变化一位；

模式包含许多0和1

• 一个有趣的问题：
  • 如果我们需要$2^{128}$个字；
  • 从所有$2^{130}$个130位模式中选择；
  • 我们必须考虑包括多少个 0/1？
    • 答案：每个类型63-67；
• 因此128b/130bcode words非常平衡，并且有很多0-1 转换（用于时钟恢复）；

最近的PCIe PC架构

• PCIe构成了PC内的互连主干；
• 北桥和南桥是PCIe开关；
• 如何支持PCI？
  • 需要一个PCI-PCIe 桥，这是
    • 有时作为南桥的一部分包含在内；
    • 或者
      可以作为单独的PCIe I/O 卡添加；
  • 当前系统将PCIe控制器直接集成在带有CPU的芯片上；

现代英特尔PCIe PC架构

现代AMD PCle PC架构

GeForce GTX 1080(Pascal)

使用DMA的PCIe数据传输

DMA（Direct Memory Access) 用于充分利用I/O总线的带宽：

• DMA使用物理地址作为源和目标；
• 传输操作系统请求的字节数；
• 需要固定内存；

固定(pinned)内存

• DMA使用物理地址；
• 操作系统可能会意外地将DMA正在读取或写入的数据调出，并将另一个虚拟页面中的页面调出到同一位置；
• 固定内存无法调出；
• 如果主机内存中cudaMemcpy的源或目标没有被固定，它需要首先被复制到一个固定的内存——额外的开销；
• cudaMemcpy使用固定的主机内存源或目标要快得多；

分配/释放固定内存（又名页面锁定内存）

• cudaHostAlloc()
  • 三个参数；
  • 指向已分配内存的指针地址；
  • 已分配内存的大小（以字节为单位）；
  • 可选项：
    • 现在使用cudaHostAllocDefault；
• cudaFreeHost()
  • 一个参数；
  • 指向要释放的内存的指针；

使用固定内存

• 以与malloc()返回的相同的方式使用分配的内存及其指针；
• 唯一的区别是分配的内存不能被操作系统分页；
• cudaMemcpy函数在固定内存的情况下应该快大约2倍；
• 固定内存是一种有限的资源，其过度使用会产生严重后果；

重要趋势

• 了解昨天、今天和明天
  • PC世界变得越来越扁平；
  • CPU和GPU融合在一起；
  • 计算外包变得越来越容易；