承影Ventus GPGPU【一】简介

凳子花❀

已于 2025-01-03 13:43:06 修改

阅读量1.1k

点赞数 30

分类专栏： AI芯片数字IC设计承影Ventus GPGPU 文章标签：承影 GPGPU Ventus GPGPU

于 2025-01-03 11:46:16 首次发布

本文链接：https://blog.csdn.net/qq_42759162/article/details/144906741

版权

数字IC设计同时被 3 个专栏收录

39 篇文章

订阅专栏

AI芯片

29 篇文章

订阅专栏

承影Ventus GPGPU

5 篇文章

订阅专栏

乘影Ventus GPGPU

前言

本文是承影Ventus GPGPU系列第一篇，主要是介绍承影Ventus GPGPU的一些基础概念和设计理念。本文的所有参考文档都来源于官方网址：OpenGPGPU，除了官网外，github中也有一些文档：ventus-gpgpu。其他系列博客内容：

简介

“乘影”GPGPU概述

“乘影”是清华大学集成电路学院开发的一款基于RISC-V向量扩展（RVV）的开源通用GPU（GPGPU）。它旨在通过开放的指令集架构（ISA）打破传统GPU厂商对高性能计算芯片的垄断，促进国内GPU设计公司核心技术的发展。乘影的所有文档在：文档-OpenGPGPU。以下是“乘影”的主要特点：

开源性：基于RISC-V指令集，完全开源，支持社区参与和创新。
SIMT架构：采用单指令多线程（SIMT）架构，每个warp由32个线程组成，硬件将warp分时映射到RVV处理器的lane上执行。
编程模型：兼容OpenCL编程模型，支持NDRange、WorkGroup（CTA/Block）、WorkItem（Thread）等概念。
指令集扩展：在RISC-V基础指令集上进行了扩展，支持向量指令、自定义指令、异步数据拷贝指令等，以满足GPGPU的特殊需求。

硬件开发进展

微架构设计

SM（Streaming Multiprocessor）：每个SM是一个支持多warp调度的RISC-V向量处理器，能够执行多个warp的指令流。warp切换类似于超线程（hyper-threading），按周期级进行调度。
寄存器堆：
- 每个warp拥有64个32位标量寄存器（sGPR）和256个1024位向量寄存器（vGPR），其中vGPR的实际宽度为num_thread * 32位（即32个线程，每个线程32位）。
- 特殊寄存器包括x0（零寄存器）、x1（返回PC寄存器）、x2（栈指针寄存器）、x3（共享内存基址）、x4（私有内存基址）等。
- CSR（控制与状态寄存器）用于存储warp ID、workgroup ID、kernel metadata baseaddr等信息。
分支处理：使用SIMT-stack硬件管理分支，减少流水线停顿，嵌套分支和循环也能快速处理。分支-合并操作仅需四条指令即可完成。

功能单元

vALU（向量算术逻辑单元）、vFPU（向量浮点单元）、vSFU（向量特殊函数单元）、vMUL（向量乘法单元）：这些功能单元均为可折叠、全流水配置，可以通过num_lane参数配置硬件单元数目，结合num_thread自由控制每类向量指令的执行周期数。
典型延迟：
- ALU：1周期
- MUL：2周期
- FMUL：3周期
- FADD：3周期
- FMACC：5周期

缓存设计

连贯性指导的缓存一致性（RCC）：基于RISC-V弱存储模型（RVWMO），“乘影”实现了释放连贯性指导的缓存一致性（RCC），降低了L1-L2带宽开销和硬件复杂度，同时减少了编程框架的额外负担。
缓存策略：
- 私有内存（Private Memory）：写回（write-back）
- 全局内存（Global Memory）：直写（write-through）
- 缓存一致性操作通过清空MSHR（缺失状态记录器）、全局冲刷、全局无效化等微架构操作予以支持。

张量计算单元（Tensor Core）

支持484规模的矩阵/张量运算，主要用于卷积和矩阵乘加操作。底层硬件通过FPU阵列实现矩阵乘加，作为功能单元接入流水线。

数据异步拷贝机制（DMA Engine）

提供了多种异步数据拷贝指令，如cp_dma、cp_dma_bulk、cp_dma_tensor等，用于高效的数据搬移。例如，cp_dma_tensor可以从全局内存中拷贝一个张量类型的数据到共享内存，适用于深度学习中的张量操作。

官方描述

乘影 GPGPU 指令集以 RISC-V 向量扩展（后文简称为 RVV）为核心设计 GPGPU，相比 RISC-V 标量指令，具有更丰富的表达含义，可以实现访存特性表征、区分 workgroup 和 thread 操作等功能。核心思想是在编译器层面以 v 指令作为 thread 的行为描述，并将 thread->warp/workgroup 的公共数据合并为标量指令。硬件上一个 warp 就是一个 RVV 程序，通常向量元素长度为 num_thread，同时又将workgroup 中统一执行的公共地址计算、跳转等作为标量指令执行，即 Vector-Thread 架构。硬件将warp 分时映射到 RVV 处理器的 lane 上去执行。

一个warp中的多个线程组成一个vector；

但公共地址计算、跳转是标量指令；

相比其它SIMT架构，在硬件上的折中是⽆法实现完全的per thread per pc，仍然需要以workgroup（或分⽀状态下的warp_split）执⾏。RVV指令集在变⻓上有三个⽅⾯的体现：硬件vlen改变；SEW元素宽度改变；LMUL分组改变。本架构特点在于这三个参数在编译期都已固定，元素数⽬⼤部分情况也固定为num_thread，本架构本质上是SIMT。

在SIMT这种架构下，线程被组织成更小的集合，称为“workgroup”（在OpenCL中）或“warp”（在CUDA中）。每个workgroup或warp中的线程执行相同的指令，但可以有不同的执行路径，例如在遇到分支指令时。

“仍然需要以workgroup（或分⽀状态下的warp_split）执⾏”意味着即使在这种架构下，线程并不是完全独立的，它们仍然需要以这些小集合的形式来执行指令。在分支发生时，可能需要将一个warp分割成更小的部分（warp_split），以便每个部分可以独立执行不同的分支路径。这是为了在保持SIMT架构的效率的同时，处理程序中的条件分支。

硬件vlen改变：vlen（Vector Length）指的是向量的长度，即一次可以处理的数据元素的数量。在RVV指令集中，vlen可以在硬件层面上改变，这意味着可以根据需要处理不同长度的向量，提供更大的灵活性。

SEW元素宽度改变：SEW（Scalar Element Width）是指单个元素的宽度，即每个元素占用的位数。在RVV指令集中，SEW可以改变，允许处理不同数据宽度的元素，例如，可以处理8位、16位、32位或64位的数据。

LMUL分组改变：LMUL（Lane Multiple）是指在一次操作中可以并行处理的向量通道数。在RVV指令集中，LMUL可以改变，这意味着可以调整并行处理的数据量，以适应不同的计算需求。

CTA = workgroup = block，Compute Thread Array

软件工具链

整体架构

编译器：基于LLVM实现，支持OpenCL C 2.0的完整编译流程，包括编译、链接、优化等。
运行时环境（Runtime）：实现了OpenCL API接口，负责任务分配、内存管理、同步和事件管理等功能。运行时环境生成动态库libpocl.so，并与硬件驱动程序交互。
设备驱动程序（Driver）：提供统一的设备接口，屏蔽底层硬件差异，支持不同类型的GPGPU设备。驱动程序生成动态库libspike_driver.so，并实现了内存分配、命令队列管理、设备控制等功能。

关键功能

任务分配：Host端通过CTA Scheduler将任务以WorkGroup为单位发送给GPGPU，线程块调度器（CTA Scheduler）负责管理SM资源，以warp为单位将任务分配给SM执行。
内存管理：支持OpenCL内存模型，包括全局内存、共享内存、私有内存等。驱动程序负责分配和释放设备内存，并将数据从主机端搬移到设备端或从设备端搬移到主机端。
同步和事件管理：支持OpenCL中的同步和事件管理函数，确保任务之间的正确顺序和依赖关系。例如，clEnqueueNDRangeKernel用于启动内核执行，clWaitForEvents用于等待事件完成。

测试套件

使用OpenCL兼容性测试套件（OpenCL-CTS）进行测试，总体通过率约为80%。具体测试项目包括基本功能、编译器、API、计算信息、内存分配、事件管理等，部分项目的通过率较高，如basic（93%）、api（91%）、computeinfo（100%）等。

指令集架构

RISC-V向量扩展（RVV）

“乘影”GPGPU指令集基于RISC-V向量扩展（RVV），并在其基础上进行了裁剪和修改，以适应SIMT编程模型。支持的基本指令集包括RV32I、RV32M、RV32A、zicsr、zfinx等。
向量指令集zve32f子集用于处理单精度浮点数（FP32）的向量运算，支持加载、存储、加法、减法、乘法、除法、平方根等常见浮点运算。为了简化硬件实现，省略了一些复杂的特性，如变长元素数目或宽度、向量间的数据交换（gather, scatter, shuffle等）。

自定义指令

分支控制指令：使用B型指令格式实现线程束分支控制，支持条件分支（如vbeq、vbge等）和重汇聚（如setrpc）。分支-合并操作通过SIMT-stack硬件管理，减少流水线停顿。
异步数据拷贝指令：如cp_dma、cp_dma_bulk、cp_dma_tensor等，用于高效的数据搬移。
计算指令：添加了12位立即数的加法指令vadd12，降低寻址、立即数装载/加减、跳转时的立即数扩展频率。还支持矩阵乘加（MMA）指令和特殊函数（如指数函数）指令，支持不同规模和精度的矩阵运算。
寄存器扩展指令：如regext和regexti，用于扩展寄存器编码和立即数，允许更灵活的指令表达。

缓存一致性

RCC（Release Consistency-directed Cache Coherence）：基于RISC-V弱存储模型（RVWMO），“乘影”实现了连贯性指导的缓存一致性（RCC），降低了L1-L2带宽开销和硬件复杂度。通过显式指令（如FENCE、AMO原子指令）控制严格的数据同步行为，确保线程间的正确通信。

术语表

在这里插入图片描述

参数表

在这里插入图片描述

总结

“乘影”GPGPU是一款基于RISC-V向量扩展的开源通用GPU，具有丰富的指令集扩展和高效的硬件设计。它通过SIMT架构实现了高性能并行计算，并且兼容OpenCL编程模型，支持多种存储空间的访问和管理。软件工具链提供了完整的编译、运行时和驱动支持，确保了良好的开发体验和性能表现。未来，“乘影”将继续发展，进一步优化硬件架构和指令集，推动开源GPGPU生态的繁荣。