实现GPU编程模型：Rigel架构与RCUDA框架解析

### 实现GPU编程模型：Rigel架构与RCUDA框架解析 #### 1. GPU编程模型基础 GPU编程模型中，二维网格由线程块组成，而每个一维、二维或三维的线程块又进一步分解为线程。线程块内的线程协同执行，它们并发启动且执行过程相互交织，可通过显式屏障控制这种交织。不过，线程块并不保证能并发运行，这限制了它们之间的交互。 内核函数采用单程序多数据（SPMD）形式编写，由网格中的每个线程执行。内核函数中的变量默认是每个线程私有的，但也可以注释为共享变量。对于共享变量，每个线程块都有一个私有实例，其所有组成线程都可以访问。共享变量可用的内存容量有限，但延迟较低。全局变量对所有线程可见，存储在大容量DRAM中，访问延迟较长。 基于GPU架构的编程模型特点和性能影响，程序员需遵循特定的软件设计原则： - 应共同调度使用相同数据的任务或操作，以利用局部性。小数据集可以快速且重复地访问，而对大数据集的流式或随机访问会限制性能。 - 需要足够数量的线程块来充分利用片上并行性，同时线程块大小会影响执行开销和负载均衡。 #### 2. 相关框架与架构 ##### 2.1 MCUDA MCUDA是一个公开可用的CUDA源到源翻译框架，之前用于将CUDA内核转换为CPU的并行C代码。在MCUDA中，线程块内的CUDA线程在循环中组合并序列化，创建遍历各个CUDA线程索引的代码。其翻译过程通过使线程块成为最小的并行任务来增加工作粒度，执行时线程块映射到单独的操作系统线程并行执行。 ##### 2.2 Rigel架构 Rigel是一款拥有1024个核心的多指令多数据（MIMD）计算加速器，针对任务和数据并行的视觉计算工作负载，可扩展到数千个并发任务。其设计目标是提供高计算密度，同时支持易于使用的传统编程模型。 Rigel的基本处理元素是经过面积优化的双发射顺序执行核心，具有类似RISC的指令集架构（ISA）、单精度浮点运算单元（FPU）和独立的取指单元。八个核心和一个共享缓存组成一个Rigel集群，集群通过双向树状互连逻辑分组为一个瓦片。芯片上分布着八个瓦片，每个瓦片包含16个集群，通过多级互连连接到32个全局缓存库。最后一级全局缓存为8个高带宽GDDR内存控制器提供缓冲。 Rigel的架构参数如下表所示： | 参数 | 值 | 单位 | | ---- | ---- | ---- | | 核心数 | 1024 | - | | 内存带宽 | 192 | GB/s | | DRAM通道数 | 8 | - | | L1I大小 | 2 | kB | | L1D大小 | 1 | kB | | L2集群缓存（总计） | 8 | MB | | L3全局缓存（总计） | 4 | MB | 应用程序使用基于任务的API为Rigel开发，一个任务映射到一个Rigel核心。任务长度可变，且不同步执行。任务的生成和分配是动态的，由软件处理，硬件仅实现全局和集群级别的原子操作。这种软件方法允许灵活的执行模型，可用于将CUDA映射到该架构。 #### 3. RCUDA框架 RCUDA是一个允许在Rigel上执行CUDA代码的框架，它包含两个主要组件： ##### 3.1 源代码转换 CUDA内核源代码需要进行转换，以适应Rigel的MIMD执行模型。在集群内，线程可以动态映射到核心，并在核内的同步点之间的循环中串行执行。每当出现同步点时，集群上的线程队列会重置，以便集群可以再次遍历每个线程。 共享变量存储为每个集群的数据结构，每个核心可以通过集群缓存读写共享数据。由于允许CUDA线程在同步点后在集群内的核心之间迁移，局部变量存储在集群级数据结构中。但在同步点之间产生和消耗的局部CUDA线程变量不需要复制，因为当CUDA线程移动到另一个核心时不会使用这些变量。 原本针对x86通用CPU的主机代码必须手动编辑，以避免使用Rigel不支持的编程接口。此外，Rigel上的主机代码由同一内存空间中的单个核心执行，而不是在单独的主机处理器上执行。主机代码所需的更改包括合并单独的主机和设备内存分配，并去除在Rigel的单地址空间中不必要的复制操作。 ##### 3.2 运行时库 RCUDA框架的第二个主要组件是软件运行时库，它提供工作分配和CUDA内置函数（如`syncthreads()`和原子操作）的实现。 RCUDA在全局芯片级别和本地集群级别进行分层工作分配。CUDA使用线程块网格来定义工作，线程块内的线程同时执行，允许线程块内的线程同步。在RCUDA框架中，一个线程块在单个Rigel集群上执行。只有当核心等待工作时，集群中的一个核心才会尝试获取一个线程块。按需获取允许线程块动态分配到集群。一个核心一次只获取一个块，这以需要更多获取操作为代价提高了负载均衡。 在本地集群级别，RCUDA控制代码通过在集群内的核心之间划分线程来处理工作分配。每个集群包含八个具有独立指令流和共享缓存的核心。CUDA线程可以静态映射，即每个核心执行固定部分的线程；也可以动态映射，即核心按需分配线程，以提高负载均衡，但代价是更频繁且可能有竞争的出队操作。 ##### 3.3 内核执行 当调用内核时，一个核心初始化运行时并将内核函数参数写入全局内存。然后，每个集群中的一个核心尝试通过原子递减剩余线程块的全局计数器来获取一个线程块标识符（ID）。如果块ID为非负，该核心初始化集群任务队列，每个核心开始执行线程块中的任务。完成一个线程块后，重复该过程。当集群中的所有核心完成执行，且没有更多线程块可用时，核心进入屏障，等待芯片上的所有其他核心完成。所有核心进入屏障后，控制返回给主机代码。 #### 4. RCUDA优化 ##### 4.1 内核代码转换 一些CUDA构造在Rigel上的映射效果不佳，主要是共享内存和线程同步。 **共享内存移除**：在NVIDIA GPU上，使用共享内存对于良好的性能至关重要，因为GPU的缓存和预取能力有限。通过使用共享内存，程序员可以将数据放在共享暂存内存中，利用时间局部性，避免内存带宽瓶颈。许多CUDA内核仅使用共享内存来保存全局数据的高带宽只读副本，在计算前将全局内存内容填充到共享内存中，计算时使用共享内存，最后将结果写回全局内存。 例如，下面是一个使用共享内存的转置内核代码： ```c __global__ void transpose(float odata[][], float idata[][], int width, int height) { __shared__ float block[BLOCK_DIM][BLOCK_DIM+1]; // read the matrix tile into shared memory unsigned int xIndex = ```