OpenCL技术方案示例

OpenCL技术方案示例

NaplesPU

http://www.naplespu.com/

http://www.naplespu.com/doc/index.php?title=Main_Page

主页

NaplesPU是一个全面的开源多核加速器，涵盖了从计算核心到片上互连、一致性存储器层次结构和编译工具链的所有体系结构层。NaplesPU完全用System Verilog HDL编写，利用了现代计算体系结构中常见的三种并行形式，特别是在GPU设备等异构加速器中：矢量并行、硬件多线程和多核心组织。NPU开源项目配备了一个针对NaplesPU矢量ISA的完整的基于LLVM的编译器，可以让您尝试当今许多核心技术的所有风格。

NPU多核架构基于通过片上网络（NoC）连接的可配置瓦片的可参数化网格。每个瓦片都有一个缓存控制器和一个目录控制器，用于处理不同瓦片中不同内核之间的数据一致性。计算核心基于矢量流水线，具有轻量级控制单元，从而将大部分硬件资源用于数据并行内核的加速。内存操作和长延迟指令通过利用硬件多线程来屏蔽。每个硬件线程（大致相当于OpenCL术语中的波前或NVIDIA术语中的CUDA翘曲）都有自己的PC、寄存器文件和控制寄存器。NaplesPU系统中的线程数可由用户配置。

 

OpenCL 对 NaplesPU 的支持

OpenCL 将平台定义为主机连接到的一组计算设备。每个设备进一步划分为多个计算单元 （CU），每个计算单元都定义为处理元素 （PE） 的集合。目标平台在架构上是围绕单个内核设计的，其结构是一组最多八个硬件线程。每个硬件线程相互竞争，以访问 16 个硬件通道，对 32 位宽整数或浮点操作数执行标量和向量运算。

 

 计算设备抽象以物理方式映射在那不勒斯PU众核架构上。那不勒斯PU众核可以根据内核数量进行配置。每个 NPU 内核都映射在 OpenCL 计算单元上。在内部，NPU 内核由硬件线程组成，每个线程都代表 OpenCL 处理元素的抽象。

执行模型匹配

从执行模型的角度来看，OpenCL 依赖于一个 N 维索引空间，其中每个点代表一个内核实例执行。由于物理内核实例的执行是由硬件线程完成的，因此 OpenCL 工作项映射到 NPU 单个硬件线程上。因此，工作组被定义为一组硬件线程，工作组中的所有工作项都在单个计算单元上执行。

 

内存模型匹配

OpenCL 将内存划分为四个不同的空间：

  • 全局空间和常量空间：所有工作组中的所有工作项都可以访问这些空间中的元素。
  • 本地空间：仅对工作组中的工作项可见。
  • 私有空间：仅对单个工作项可见。

目标平台由 DDR 内存提供，即 OpenCL 命名法中的设备内存。因此，变量在物理上映射到此内存上。编译器本身通过查看地址空间限定符来验证是否满足 OpenCL 约束。

每个 NPU 内核还配备了一个 Scratchpad 存储器，这是每个内核独有的片上非相干存储器部分。此存储器符合 OpenCL 本地存储器功能

 

最后，NPU 内核中的每个硬件线程都有一个专用堆栈。此内存部分对于每个硬件线程（即 OpenCL 工作项）是私有的，其他线程无法寻址。因此，每个堆栈都充当 OpenCL 专用存储器。

编程模型匹配

OpenCL 支持两种编程模型：数据并行和任务并行。数据并行模型要求 OpenCL 索引空间的每个点都执行一个内核实例。由于每个点都代表一个工作项，并且这些点映射在硬件线程上，因此可以正确满足数据并行要求。请注意，实现的模型是一个宽松的版本，不需要严格的一对一数据映射。

任务并行编程模型要求内核实例在索引空间的任何点独立执行。在这种情况下，每个工作项都不受限制地执行其他工作项的相同内核实例。编译器前端定义了一组可用于此目的的内置函数。此外，每个 NPU 内核由 16 个硬件通道构建，可用于实现锁步执行。因此，实现了 OpenCL 支持，以允许使用向量类型。因此，以下数据类型支持向量执行：

  • charn、ucharn 分别映射在 vec16i8 和 vec16u8 上，其中 n=16。不支持 n 的其他值。
  • shortn、ushortn 分别映射在 vec16i16 和 vec16i32 上，其中 n=16。不支持 n 的其他值。
  • intn、uintn 分别映射在 vec16i32 和 vec16u32 上，其中 n=16。不支持 n 的其他值。
  • floatn，映射在 vec16f32 上，其中 n=16。不支持 n 的其他值。

OpenCL 运行时设计

OpenCL API 是一组用于协调和管理设备的功能，它们为运行应用程序和监控其执行提供支持。这些 API 还提供了一种检索设备相关信息的方法。

下图描述了 OpenCL 规范中定义的 OpenCL 运行时的 UML 类图。灰色填充框表示由于缺少硬件支持而不可用的功能。

 

自定义 OpenCL 运行时依赖于两个主要抽象：

  • 低级抽象（不完全依赖于硬件）提供设备-主机通信支持。
  • 根据 OpenCL API，高级抽象管理运行在设备上的内核的生命周期。

OpenCL 示例

以下代码显示了在 NPU 设备上运行的 OpenCL 中的向量矩阵乘法。

#include <opencl_stdlib.h>
#define WORK_DIM 4
 
__kernel void kernel_function(__global int16 *A, __global int16 *B, __global int16 *C, int rows, int cols)
{
   __private uint32_t threadId = get_local_id(0); 
 
   uint32_t nT = WORK_DIM; // number of threads
   uint32_t nL = 16;       // number of lanes
 
   uint32_t N = rows;
   uint32_t nC = N / nL;
   uint32_t ndivnT = N / nT;
   uint32_t tIdndivnT = threadId * ndivnT;
   uint32_t tIdndivnTnC = tIdndivnT * nC;
   for (uint32_t i = 0; i < ndivnT * nC; i++)
   {
       uint32_t col = (tIdndivnT + i) % nC;
       C[tIdndivnTnC + i] = 0;
       for (uint32_t j = 0; j < nC; j++)
       {
           for (uint32_t k = 0; k < nL; k++)
           {
               C[tIdndivnTnC + i] += A[tIdndivnTnC + i - col + j][k] * B[(nC * k) + (j * N) + col];
           }
       }
   }
}
posted @ 2024-05-24 04:22  吴建明wujianming  阅读(106)  评论(0)  编辑  收藏  举报