承影Ventus GPGPU【四】硬件结构

乘影Ventus GPGPU

前言

本文是承影Ventus GPGPU系列第四篇，主要是介绍承影Ventus GPGPU的硬件结构。本文的所有参考文档都来源于官方网址：OpenGPGPU，除了官网外，github中也有一些文档：ventus-gpgpu。其他系列博客内容：

1. 承影Ventus GPGPU【一】简介
2. 承影Ventus GPGPU【二】指令集
3. 承影Ventus GPGPU【三】软件工具链
4. 承影Ventus GPGPU【四】硬件结构
5. 承影Ventus GPGPU【五】LLVM编译器配置

硬件方面:

⽬前NDRange拆分为workgroup在驱动上进⾏， workgroup拆分为warp在硬件上进⾏。

Vector-Thread架构: 将每个warp视为一段RVV程序，并采用周期级调度，类似于hyper-threading。

寄存器堆: 4-Bank SRAM结构，支持256个num_thread * 32bit vGPR和64个32bit sGPR，并可根据warp实际使用情况进行动态分配。

乘影GPGPU微架构介绍

支持多线程束（warp）调度的向量处理器

“乘影”GPGPU设计中，每个流多处理器单元（SM）能够处理多个线程束（warp），每个warp包含32个线程。这种多线程束调度机制允许硬件在同一时间执行多个线程，从而显著提高了并行度和计算效率。在《基于Chisel HDL的硬件开发进展》文件中详细描述了这一点。

计算任务的启动、分配、执行

计算任务从主机（Host）到设备（Device）的传递过程包括任务的启动、分配和执行。具体来说，主机将计算任务以工作组（WorkGroup）的形式发送给GPGPU，并附带元数据信息。线程块调度器（CTA Scheduler）负责管理SM资源，并以warp为单位将任务分发给具体的SM。线程束调度器则管理warp的基本信息和执行情况。这个流程在《乘影硬件开发进展》文件中有详细说明，并且在《乘影GPGPU架构文档手册v2.02》第16页有关互联网络的部分也有所提及。

自定义指令的支持

“乘影”GPGPU不仅支持标准的RISC-V指令集，还添加了特定于GPGPU的功能指令，如分支控制指令、线程同步指令等。这些自定义指令增强了“乘影”的功能，使其更适合通用图形处理任务。例如，《乘影：开源通用GPU指令集架构介绍》文件中的4.2章节详细解释了自定义指令的种类和用途，包括分支控制指令、寄存器扩展指令、寄存器对拼接指令、同步和任务控制指令、自定义计算指令、自定义立即数访存指令和自定义私有内存访存指令。

任务分配：Host - CTA Scheduler - SM

任务分配过程中，主机首先将任务以WorkGroup为单位发给GPGPU，这其中包括了必要的元数据信息。接着，线程块调度器进行SM资源的管理，以warp为单位发给SM。最后，线程束调度器管理warp基本信息和执行情况。这个流程在《乘影硬件开发进展》文件中有所阐述，并在《乘影GPGPU架构文档手册v2.02》第17页有关SIMT-stack补充的内容中得到了进一步的解释。

任务执行

每个SM可视为一个支持多warp调度的RISC-V向量处理器

每个SM可以被看作是一个支持多warp调度的RISC-V向量处理器。这意味着每个SM内部都有多个执行单元，可以在同一时间处理多个warp。这种设计使得“乘影”能够在硬件层面上实现高效的并行计算。在《乘影硬件开发进展》文件中对此有详细的介绍。

每个warp可视为一段RVV程序

每个warp相当于一段RVV（RISC-V Vector Extension）程序，这些程序会经历取指、译码、发射、执行和写回寄存器的过程。RVV指令集提供了丰富的向量操作能力，使得“乘影”能够高效地处理大规模并行计算任务。这部分内容在《乘影GPGPU指令集架构》文件的4.1章节中有详细介绍。

warp切换是类似hyper-threading的周期级调度

warp之间的切换类似于超线程技术（hyper-threading），是周期级的调度。这意味着在一个周期内，如果一个warp遇到等待事件，系统可以快速切换到另一个warp继续执行，以此掩盖延迟。这种设计确保了即使某些线程处于等待状态，其他线程仍然可以充分利用硬件资源，提高整体性能。在《乘影GPGPU架构文档手册v2.02》第20页有关发射仲裁的内容中也有相关描述。

乘影GPGPU微架构的具体组件

寄存器堆

寄存器堆为4-bank SRAM，具有1r1w（每bank一个读端口和一个写端口），并且是交织（interleave）和unified结构。根据warp的实际使用数目进行分配，每个warp至多有256个 num_thread* 32bit vGPR 和 64个 32bit sGPR。此外，还有CSR（Control and Status Registers）。寄存器位宽扩展指令（如regext和regexti）可以在译码时暂存寄存器扩展信息，并在后续有效指令译码时取出合并处理。这部分内容在《乘影硬件开发进展》文件中有详细描述。

分支处理

“乘影”采用SIMT-stack硬件管理分支，不占用通用寄存器，减少流水线停顿，嵌套分支和循环也能快速处理。仅用四条指令完成一次完整分支-合并操作。硬件调度发射执行时，ibuffer暂存warp信息，Scoreboard记录依赖，每周期切换，Round-Robin发射。这部分内容在《乘影硬件开发进展》文件中有详细描述。

双发射

为了进一步发掘并行度，“乘影”增加了额外的线程束调度器，不同线程束的标量指令和向量指令之间并行调度，理想情况下每周期发射一条标量指令和一条向量指令。在保持低硬件复杂度的情况下提高了整体性能。这部分内容在《乘影硬件开发进展》文件中有详细描述。

执行单元

“乘影”实现了张量计算TensorCore和exp函数，支持transformer运算。TensorCore可用于卷积和矩阵运算，支持8 * 4 * 8 fp32的运算。底层硬件通过fpu阵列实现矩阵乘加，作为功能单元接入流水线。这部分内容在《乘影硬件开发进展》文件中有详细描述。

缓存设计

乱序机制（指令级并行/ILP）和多级缓存对性能至关重要，但会对线程间同步（线程级并行/TLP）产生阻碍。连贯性问题涉及流水线中的乱序因素，一致性问题涉及缓存内容的不同步。“乘影”通过连贯性指导的缓存一致性（RCC）解决了这些问题，相比CPU硬件一致性和GPU软件一致性，降低了L1-L2带宽开销和硬件复杂度，减少了编程框架的额外负担。这部分内容在《乘影硬件开发进展》文件中有详细描述。

综上所述，“乘影”GPGPU通过其独特的微架构设计实现了高效的任务管理和并行计算能力，支持丰富的自定义指令，确保了对复杂计算任务的有效处理。此外，它还结合了RISC-V的向量扩展特性，使得warp级别的调度和执行更加灵活和高效。

寄存器堆

寄存器堆设计

“乘影”GPGPU的寄存器堆采用4-bank SRAM结构，每个bank支持1个读端口和1个写端口，并且是交织（interleave）的。这种设计使得寄存器堆可以在一个周期内同时处理多个读写请求，从而提高了访问效率。寄存器堆是统一（unified）的，这意味着它可以动态地根据warp的实际使用情况分配寄存器资源，确保资源的高效利用。具体来说，每个warp最多可以拥有256个32位宽的向量寄存器（vGPR）和64个32位宽的标量寄存器（sGPR），并且还包含一些控制状态寄存器（CSR）。这些寄存器的设计在《乘影硬件开发进展.pdf》第19页中有详细描述。

动态分配与统一结构

寄存器堆的统一结构允许根据warp的实际使用情况进行动态分配。这意味着，如果某个warp使用的寄存器较少，那么剩余的寄存器可以被其他warp使用，从而最大化了硬件资源的利用率。例如，假设一个warp只需要使用128个vGPR和32个sGPR，那么剩下的128个vGPR和32个sGPR可以被其他warp使用。这种动态分配机制不仅提高了资源利用率，还使得“乘影”GPGPU能够更好地适应不同的工作负载。

什么是Interleave（交织）？

Interleave（交织） 是一种内存或寄存器访问技术，通过将数据分散存储在多个独立的存储单元中，使得这些存储单元可以并行访问。具体来说，交织意味着将一个大的地址空间划分为多个较小的子空间，并将这些子空间映射到不同的存储模块（bank）上。这样，当多个访问请求同时到达时，每个请求可以被分配到不同的bank，从而避免了竞争和等待时间，提高了整体的访问效率。

在“乘影”GPGPU中，寄存器堆采用4-bank SRAM结构，并且是交织的。这意味着每个warp的寄存器被分散存储在这4个bank中，每个bank支持1个读端口和1个写端口。通过交织设计，即使在一个周期内有多个寄存器访问请求，这些请求也可以被分配到不同的bank，从而实现并行访问，减少了访问延迟。

示例说明：

假设一个warp需要在同一周期内读取4个不同的寄存器（v0, v1, v2, v3），并且这些寄存器分别位于4个不同的bank中。由于寄存器堆是交织的，这4个读取请求可以同时进行，而不会发生冲突。如果没有交织设计，所有请求都必须排队等待，导致访问延迟增加。

# 交织设计下的寄存器访问
vadd.v v0, v1, v2  # v0, v1, v2 分别位于不同的bank中
vmul.v v3, v4, v5  # v3, v4, v5 分别位于不同的bank中

在这个例子中，v0、v1、v2 和 v3、v4、v5 的读取操作可以在同一周期内并行完成，因为它们位于不同的bank中。这种设计显著提高了寄存器访问的吞吐量。

什么是Unified（统一）结构？

Unified（统一）结构 是指寄存器堆可以根据warp的实际使用情况进行动态分配，而不是为每个warp固定分配一定数量的寄存器。在传统的GPU设计中，每个warp通常会预先分配固定的寄存器资源，无论这些寄存器是否被充分利用。这种方式虽然简单，但会导致资源浪费，尤其是在某些warp只需要少量寄存器的情况下，其他warp可能会因为资源不足而无法高效运行。

“乘影”GPGPU的寄存器堆采用了统一结构，这意味着所有的寄存器资源（包括vGPR和sGPR）都是共享的，硬件会根据warp的实际需求动态分配寄存器。例如，如果一个warp只需要使用128个vGPR和32个sGPR，那么剩下的128个vGPR和32个sGPR可以被其他warp使用。这种设计不仅提高了资源利用率，还使得“乘影”GPGPU能够更好地适应不同的工作负载。

示例说明：

假设有两个warp，warp A 和 warp B。warp A 需要使用128个vGPR和32个sGPR，而warp B 需要使用256个vGPR和64个sGPR。在统一结构下，warp A 只会占用128个vGPR和32个sGPR，剩下的128个vGPR和32个sGPR可以被warp B 使用。因此，warp B 可以获得所需的全部256个vGPR和64个sGPR，而不会因为资源不足而受到影响。

# 统一结构下的寄存器分配
warp A: 128 vGPR + 32 sGPR
warp B: 256 vGPR + 64 sGPR

在这个例子中，warp A 和 warp B 根据实际需求动态分配了寄存器资源，确保了资源的高效利用。如果没有统一结构，warp B 可能会因为资源不足而无法获得足够的寄存器，导致性能下降。

4-Bank SRAM结构

什么是Bank？

Bank 是指内存中的一个独立存储区域，可以并行访问。在4-Bank SRAM结构中，SRAM被划分为4个这样的独立存储区域。

为什么需要Bank？

并行性：通过将内存划分为多个Bank，可以在同一时间访问不同的Bank，从而实现并行数据访问，增加带宽。

减少冲突：如果所有数据都存储在一个单一的Bank中，连续的访问可能会导致Bank冲突（即多个请求同时访问同一个Bank），从而降低效率。多Bank结构可以减少这种冲突。

4-Bank SRAM结构的特点：

256个num_thread * 32bit vGPR:

vGPR（Vector General-Purpose Registers）是向量寄存器，每个线程都有自己的vGPR。

在4-Bank SRAM结构中，每个线程可以有多个vGPR，总共支持256个num_thread（线程数）乘以32位宽的vGPR。

这些vGPR被分配到4个Bank中，每个Bank可以存储一定数量的vGPR。

64个32bit sGPR:

sGPR（Scalar General-Purpose Registers）是标量寄存器，通常被所有线程共享。

在这个结构中，有64个32位宽的sGPR，这些寄存器也是分布在4个Bank中。

动态分配:

寄存器堆可以根据warp（一组同时执行的线程）的实际使用情况进行动态分配。

这意味着不是所有的线程都需要同时使用所有的寄存器，因此可以根据需要将寄存器分配给不同的线程，从而提高寄存器的使用效率。

硬件结构：

Bank组织：每个Bank是一个独立的SRAM存储单元，可以并行地读取或写入数据。

交叉存储：为了提高访问效率，vGPR和sGPR的数据在Bank之间是交叉存储的。这意味着一个线程的寄存器数据可能分布在不同的Bank中。

数据路径：每个Bank都有自己的数据路径，允许同时访问不同的Bank，从而增加了数据吞吐量。

控制逻辑：控制逻辑负责管理寄存器的分配，确保线程请求的数据可以正确地从相应的Bank中读取或写入。

优势：

提高带宽：由于可以并行访问多个Bank，所以整体的数据带宽得到了提升。

减少延迟：通过并行访问和减少Bank冲突，可以减少数据访问的延迟。

灵活分配：动态分配机制允许更高效地使用寄存器资源，特别是在线程需求不同的情况下。

综上所述，4-Bank SRAM结构是一种优化的内存设计，它通过将寄存器划分为多个Bank，并允许并行访问，提高了GPGPU设备在执行并行任务时的性能和效率。

Operand Request:

这个部分负责接收操作数请求，即从外部输入需要处理的数据地址和数据类型等信息。

Read Arbiter 0-3:

这些仲裁器用于管理对四个不同bank的读访问。它们确保在多个并发读请求时能够公平地分配资源，避免冲突。

Vector Bank 3 和 Scalar Bank 3:

这两个bank分别用于存放向量数据和标量数据。Bank 3应该是指最后一个bank，bank是个长方体，所以应该是4个bank叠到一起的画法。

Collector Allocation:

接收IBuffer中的信息：Collector Allocation模块接收来自IBuffer的信息，这些信息可能包括线程将要执行的操作指令以及这些操作需要访问的寄存器索引。

数据访问：模块根据IBuffer中的寄存器索引信息，从4-Bank SRAM结构的寄存器堆中读取相应的数据。

合并和处理：一旦数据被读取，Collector Allocation模块会负责将这些数据合并并准备传递到后续的处理阶段。

为什么不直接从bank接收数据

指令与数据的分离：在GPU架构中，通常将指令的执行与数据的访问分离。IBuffer处理指令流，而寄存器堆处理数据流。

灵活性：通过使用IBuffer，可以更灵活地处理指令和操作数的映射，允许更复杂的调度和优化。

减少冲突：Collector Allocation模块可以根据IBuffer中的信息来优化数据访问模式，减少bank冲突，提高寄存器堆的访问效率。

IBuffer:

IBuffer（Instruction Buffer）通常用于暂存指令流，以便于顺序执行或优化指令调度。是一个中间缓冲区，它用于存储即将被执行的线程的指令以及相关的操作数信息。这些操作数信息可能包括寄存器索引，而不是直接的数据。

Ctrl Sig:

Ctrl Sig代表控制信号，这些信号用于控制整个系统的运行状态和行为，如启动/停止操作、同步等。

ALU Src 0, ALU Src 1, ALU Src 2, ALU Src 3:

这些是算术逻辑单元（ALU）的操作源，表示不同的输入端口，可以接受来自不同bank的数据进行运算。

data[num_thread-1:0][31:0]:

这是数据路径的一部分，其中num_thread表示线程的数量，[31:0]表示32位宽度的数据总线。

Write Back:

Write Back模块负责将处理完成的结果写回到相应的内存位置或寄存器中。

EX:

EX可能指的是Execution Unit（执行单元），它是实际执行数据处理的地方，包括算术运算、逻辑运算和其他操作。

总结来说，这个设计通过多bank的结构来提高并行性和效率，同时利用仲裁器和collector allocation机制来有效地管理和分发数据和控制信号。

步骤

步骤 1: 接收IBuffer中的信息

功能：Collector Allocation模块首先接收来自IBuffer的信息。这些信息通常包括指令、操作码、寄存器索引以及线程标识等。

执行方式：IBuffer将指令和相关数据推送到Collector Allocation模块。这个过程通常是顺序的，每个时钟周期可能推送一条或多条指令。

步骤 2: 解析指令

功能：模块解析接收到的指令，确定需要读取哪些寄存器以及执行哪些操作。

执行方式：模块中的解码逻辑会分析指令的操作码和操作数，识别出需要访问的寄存器地址。

步骤 3: 寄存器地址映射

功能：将解析出的寄存器索引映射到具体的4-Bank SRAM结构中的物理地址。

执行方式：通过查找表或地址计算逻辑，将逻辑寄存器索引转换为物理bank地址和偏移量。

步骤 4: 读取寄存器数据

功能：根据映射得到的物理地址，从寄存器堆的相应bank中读取数据。

执行方式：

无冲突访问：如果请求的寄存器位于不同的bank，可以并行读取，以最大化带宽。

冲突处理：如果请求的寄存器位于相同的bank，需要串行访问或使用bank冲突解决策略。

步骤 5: 数据合并

功能：将来自不同bank的数据合并起来，为后续的处理阶段做准备。

执行方式：数据合并逻辑可能会使用数据选择器（如多路复用器）来根据指令的需要组合或选择数据。

步骤 6: 数据传递

功能：将合并后的数据传递到后续的处理阶段，如执行单元（Execution Units）或下一级的流水线阶段。

执行方式：通过内部数据总线或缓冲区，将准备好的数据发送到下一个处理模块。

步骤 7: 流水线管理

功能：管理数据流和指令流，确保流水线的有效运行，处理延迟和依赖问题。

执行方式：使用各种控制逻辑来监控流水线的状态，调整数据流动，处理 hazard（如数据冒险、控制冒险和结构冒险）。

这些步骤是高度并行的，并且可能在不同的时钟周期内同时进行。Collector Allocation模块的设计必须确保高效的数据访问和传递，以支持GPU的高吞吐量和并行处理能力。通过这种方式，GPU能够处理大量的线程和数据，同时保持高性能和效率。

寄存器扩展指令

简介

为了支持更复杂的指令集，“乘影”GPGPU引入了寄存器位宽扩展指令（regext和regexti）。这些指令允许在编译时扩展寄存器编码的宽度，从而支持更多的寄存器或更大的立即数。具体来说：

• regext指令：扩展四个5位寄存器编码至8位。这意味着原本只能编码32个寄存器的5位字段现在可以编码256个寄存器。这为编译器提供了更多的寄存器选择，尤其是在需要大量寄存器的情况下。

• regexti指令：扩展两个5位寄存器编码至8位，并扩展5位立即数至11位。这不仅增加了寄存器的选择范围，还允许使用更大的立即数，从而减少了对内存的依赖。

这些扩展指令的工作原理如下：当遇到寄存器扩展指令时，译码器会将扩展信息暂存在译码级中。随后，在接续的有效指令进行译码时，译码器会取出扩展信息并与当前指令合并处理。这种设计避免了额外的寄存器占用，同时也简化了指令格式。

具体实现：

1. 寄存器扩展信息的暂存：当遇到寄存器扩展指令（如regext或regexti）时，译码器会将扩展信息（例如，额外的寄存器编码或立即数扩展）暂存在译码级中。这些信息不会立即应用到寄存器文件中，而是保存在一个临时缓冲区中，等待后续指令的处理。

2. 扩展信息的合并：当接续的有效指令（如加法、乘法等）进行译码时，译码器会从临时缓冲区中取出扩展信息，并将其与当前指令的寄存器编码或立即数合并。例如，如果当前指令使用了5位寄存器编码，而之前的regext指令扩展了这个编码至8位，那么译码器会将8位编码应用到当前指令中。

3. 执行阶段：经过合并后的指令会被发送到相应的执行单元（如ALU或FPU）进行执行。此时，扩展后的寄存器编码或立即数已经准备好，可以直接用于计算。

示例说明：

假设有一个kernel，其中涉及到大量的向量运算，需要使用超过32个vGPR来存储中间结果。如果没有寄存器扩展指令，编译器可能会将超出的部分寄存器溢出到栈中，导致频繁的访存操作。然而，使用regext指令后，编译器可以将这些寄存器映射到更高的编号（例如，从v32到v255），而不需要频繁地将数据保存到内存中。这显著减少了访存开销，提升了执行效率。

# 未使用寄存器扩展指令的情况
add.v v0, v1, v2  # 使用v0-v31
add.v v0, v1, v2  # 使用v0-v31
...
spill v32  # 将v32溢出到栈中
load v32   # 从栈中加载v32
...

# 使用寄存器扩展指令的情况
regext  # 扩展寄存器编码
add.v v32, v33, v34  # 使用v32-v255
add.v v35, v36, v37  # 使用v32-v255
...

在这个例子中，使用regext指令后，编译器可以直接使用v32-v255寄存器，而不需要将数据溢出到栈中，从而减少了访存开销，提升了执行效率。

硬件结构

这张图展示了寄存器位宽扩展指令（如 regext 和 regexti）在处理器中的处理流程。具体来说，这张图解释了寄存器扩展指令如何在译码阶段被处理，并如何与后续的有效指令合并。

详细解释

1. 寄存器扩展信息暂存：

   • Decode（译码）：当处理器遇到寄存器扩展指令（如 regext 或 regexti）时，这些指令包含的寄存器扩展信息会被暂存在译码级中。这些信息包括扩展的寄存器编号（ext_rs1, ext_rs2, ext_rs3, ext_rd）和扩展的立即数（ext_imm）。
   • 暂存区域：这些扩展信息被暂存在一个临时缓冲区中，等待后续指令的处理。

2. 有效指令的合并处理：

   • Instruction Buffer（指令缓冲区）：当接续的有效指令（如加法、乘法等）进行译码时，译码器会从暂存区域取出扩展信息，并将其与当前指令合并处理。
   • 合并处理：合并后的指令会被发送到指令缓冲区，准备执行。

具体步骤

1. 寄存器扩展指令的处理：

   • 当处理器遇到寄存器扩展指令时，译码器会识别出这是一个扩展指令，并将扩展信息（如扩展的寄存器编号和立即数）暂存在译码级中。
   • 这些信息不会立即应用到寄存器文件中，而是保存在一个临时缓冲区中，等待后续指令的处理。

2. 有效指令的处理：

   • 当接续的有效指令进行译码时，译码器会从临时缓冲区中取出扩展信息，并将其与当前指令合并处理。
   • 合并后的指令会被发送到指令缓冲区，准备执行。

总结

这张图展示了寄存器扩展指令在处理器中的处理流程，包括扩展信息的暂存和与后续有效指令的合并处理。通过这种方式，处理器可以在不增加硬件复杂度的情况下，支持更多的寄存器和更大的立即数，从而提升了代码的紧凑性和执行效率。

控制状态寄存器（CSR）

除了vGPR和sGPR，“乘影”GPGPU还包含了一些控制状态寄存器（CSR）。这些寄存器用于存储与warp、workgroup、kernel等相关的控制信息。例如，CSR可以存储warp ID、workgroup ID、kernel metadata baseaddr等信息，这些信息在分支控制、同步操作以及内存访问中起着重要作用。

具体功能：

• warp ID：标识当前warp的唯一ID，用于区分不同的warp。
• workgroup ID：标识当前warp所属的workgroup，用于在多线程调度中识别不同的workgroup。
• kernel metadata baseaddr：存储kernel元数据的基地址，用于访问kernel的相关配置信息。
• 其他控制信息：还包括一些用于分支管理、同步操作的控制寄存器，确保warp之间的正确协作。

总结

“乘影”GPGPU的寄存器堆设计采用了4-bank SRAM结构，支持1r1w操作，并且是交织和统一的。交织设计通过将寄存器分散存储在多个bank中，实现了并行访问，减少了访问延迟；统一结构则允许根据warp的实际使用情况进行动态分配，确保了资源的高效利用。此外，通过引入寄存器扩展指令（regext和regexti），“乘影”GPGPU能够在不增加硬件复杂度的情况下，支持更多的寄存器和更大的立即数，从而提升了代码的紧凑性和执行效率。最后，控制状态寄存器（CSR）为warp、workgroup和kernel的控制提供了必要的支持，确保了系统的正确性和高效性。

分支管理和双发射

SIMT-stack分支管理

为什么采用SIMT-stack可以不占用通用寄存器？

在传统的GPU设计中，分支处理通常依赖于通用寄存器来保存和恢复线程的状态。然而，这种方法会导致寄存器资源的浪费，并且在频繁的分支操作中会增加流水线停顿的时间。为了克服这些问题，“乘影”GPGPU采用了SIMT-stack（Single Instruction, Multiple Threads Stack）硬件管理分支的方法。

SIMT-stack的工作原理：

SIMT-stack是一种专门用于管理分支和同步的硬件机制，它通过维护一个栈结构来跟踪每个warp中的分支状态。具体来说，当遇到分支指令时，SIMT-stack会记录当前的程序计数器（PC）和掩码信息（mask），并根据分支条件决定执行哪个路径。这种设计使得分支管理不再依赖于通用寄存器，从而释放了宝贵的寄存器资源，使得更多的寄存器可以用于实际的计算任务。

为什么可以减少流水线停顿？

传统的分支处理方法会在遇到分支指令时暂停流水线，直到所有线程完成分支判断并进入同一路径后才继续执行。这会导致流水线停顿，降低执行效率。而SIMT-stack通过硬件管理分支，可以在分支发生时立即记录当前状态并继续执行未分歧的路径，从而减少了流水线停顿的时间。

具体过程：

1. 分支指令执行时：当遇到分支指令（如vbeq）时，SIMT-stack会计算出当前分⽀的掩码情形，并将else对应的掩码压入栈中。
2. 执行if段：随后，带有掩码的if段指令会被立即执行，而不必等待所有线程完成分支判断。
3. 合并路径：当if段末尾遇到join指令时，SIMT-stack会从栈中弹出else段及其对应掩码，继续执行else段。
4. 嵌套分支支持：SIMT-stack支持嵌套分支，最多可以嵌套到单warp中的线程数（32个）。如果每次总选择最多的方向去压栈而非默认压if，那么栈深度只需要5即可。

这种设计使得“乘影”GPGPU能够在分支发生时快速切换到未分歧的路径，减少了流水线停顿的时间，提高了整体执行效率。

SIMT-stack是什么？如何进行硬件管理分支？

SIMT-stack的定义：

SIMT-stack是“乘影”GPGPU中用于管理分支和同步的硬件机制。它通过维护一个栈结构来跟踪每个warp中的分支状态，确保在分支发生时能够快速恢复正确的执行路径。SIMT-stack的主要功能包括：

• 分支嵌套控制流管理：维护分支嵌套控制流，确保程序运行的正确性。
• 快速跳过不必要的程序段：在实际没有分支分歧发生时，跳过不必要的程序段，减少不必要的指令执行。
• 隐式mask设置：由SIMT-stack设置的隐式mask会在该warp执行过程中一直生效，直到有其他分支管理支持对其进行修改。这个mask与RVV软件形式的mask可以叠加生效。

硬件管理分支的具体过程：

1. 分支指令启动：当遇到分支指令（如vbeq）时，SIMT-stack会启动一个split，计算出当前分支的掩码情形。
2. 压栈操作：将else对应的掩码压入SIMT-stack，然后带有掩码的if段指令会被立即执行。
3. 执行if段：if段指令被执行，期间SIMT-stack会保持当前的分支状态。
4. 合并路径：当if段末尾遇到join指令时，SIMT-stack会从栈中弹出else段及其对应掩码，继续执行else段。
5. 嵌套分支支持：SIMT-stack支持嵌套分支，最多可以嵌套到单warp中的线程数（32个）。如果每次总选择最多的方向去压栈而非默认压if，那么栈深度只需要5即可。

示例说明：

假设有一个包含多个分支的代码段，如下所示：

if (condition) {
    // if段
} else {
    // else段
}

在“乘影”GPGPU中，这段代码会被编译成以下汇编指令：

vbeq vs2, vs1, label_else  # 如果vs2[i] == vs1[i]，则跳转到label_else
# if段指令
...
join                       # 合并if和else段
label_else:
# else段指令
...

当执行到vbeq指令时，SIMT-stack会计算出当前分⽀的掩码情形，并将else对应的掩码压入栈中。随后，带有掩码的if段指令会被立即执行。当if段末尾遇到join指令时，SIMT-stack会从栈中弹出else段及其对应掩码，继续执行else段。这种设计使得分支处理更加高效，减少了流水线停顿的时间。

SIMT堆栈的功能

• 维护分支嵌套控制流：
  • 在SIMT模型中，多个线程（通常组成一个warp）同时执行相同的指令，但每个线程可能有不同的执行路径（例如，由于分支条件）。SIMT堆栈用于跟踪每个线程在分支控制流中的位置，确保即使在不同路径上也能正确执行。
• 保障程序运行的正确性：
  • 当遇到分支指令时，SIMT堆栈会保存必要的信息，以便在分支结束后能够恢复正确的执行路径。
• 跳过不必要的程序段：
  • 如果一个warp中的所有线程都遵循相同的分支路径（没有分歧），SIMT堆栈可以用来跳过不执行的路径，从而提高效率。

SIMT堆栈的设计

• 隐式Mask：
  • SIMT堆栈通过设置一个隐式的mask来控制哪些线程是活跃的。这个mask在warp执行过程中持续生效，直到遇到新的分支指令或特定的管理指令来修改它。
  • 隐式mask可以与RISC-V向量扩展（RVV）中的软件形式的mask叠加，进一步控制线程的执行。
• 自定义扩展指令集：
  • 为了管理分支，GPU架构可能定义了一组扩展指令集，其中包括分支指令。

分支指令示例（vbeq）

以vbeq（向量等于比较）指令为例，该指令执行以下步骤：

• 取源操作数vs2和vs1，这两个操作数通常是向量寄存器。
• valu模块（向量算术逻辑单元）比较这两个向量寄存器中的元素。
• 对于每个元素，如果vs2[i] == vs1[i]，则输出结果out[i]为1，否则为0。
• 最终，valu的输出结果out是一个掩码，表示分支指令的else路径。
• 译码模块会将分支发生标记和else路径的PC（程序计数器）起始值PC_branch发送给分支管理模块。

分支管理过程

• 当执行vbeq指令时，如果所有线程的比较结果相同，则整个warp可以继续执行而不需要分支。
• 如果有分歧，SIMT堆栈会保存当前状态，包括当前PC和活跃线程的mask。
• 分支管理模块根据valu的输出结果决定哪些线程应该跳转到else路径。
• 在分支结束后，SIMT堆栈用于恢复之前保存的状态，确保所有线程都能从正确的位置继续执行。
• 通过这种方式，SIMT堆栈确保了在SIMT执行模型中，即使面对复杂的分支逻辑，也能高效且正确地处理大量的线程。

双发射

增加额外的线程束调度器

为了进一步发掘并行度，“乘影”GPGPU增加了额外的线程束调度器。这种设计允许不同线程束的标量指令和向量指令之间并行调度，理想情况下每周期发射一条标量指令和一条向量指令。这种设计在《基于Chisel HDL的硬件开发进展》文件中有详细描述，并在《乘影GPGPU架构文档手册v2.02》第23页有关双发射的部分也有说明。

双发射的设计动机：

传统的GPU设计中，标量指令和向量指令通常是顺序执行的，这限制了并行度的发掘。为了提高并行度，“乘影”GPGPU引入了双发射机制，允许在同一周期内发射一条标量指令和一条向量指令。这种设计通过增加额外的线程束调度器，实现了更高效的指令调度和执行。

具体实现：

1. 增加额外的线程束调度器：每个SM（流多处理器单元）包含两个线程束调度器（Sche 1 和 Sche 2），分别负责调度标量指令和向量指令。
2. 并行调度：在每个周期内，Sche 1 和 Sche 2 分别从各自的指令缓冲区（Ibuffer）中取出一条标量指令和一条向量指令，并将其发送到相应的执行单元（如sALU和vALU）。
3. 理想情况下的并行发射：在理想情况下，每个周期可以发射一条标量指令和一条向量指令，从而最大化并行度。

示例说明：

假设有一个包含标量指令和向量指令的代码段，如下所示：

float a = b + c;  // 标量指令
vector_add(vd, vs1, vs2);  // 向量指令

在“乘影”GPGPU中，这段代码会被编译成以下汇编指令：

add.s s0, s1, s2  # 标量加法
vadd.vv vd, vs1, vs2  # 向量加法

通过双发射机制，这两个指令可以在同一周期内被发射到不同的执行单元（sALU和vALU），从而实现了并行执行，提高了整体性能。

功能单元数目可配置

“乘影”GPGPU中的vALU（向量算术逻辑单元）、vFPU（向量浮点单元）、vSFU（向量特殊功能单元）、vMUL（向量乘法单元）均为可折叠、全流水配置。这些单元可通过num_lane参数配置硬件单元数目，结合num_thread参数可自由控制每类向量指令的执行cycle数。这种设计在《基于Chisel HDL的硬件开发进展》文件中有详细描述，并在《乘影GPGPU架构文档手册v2.02》第24页有关功能单元配置的部分也有说明。

功能单元的可配置性：

1. 可折叠设计：vALU、vFPU、vSFU、vMUL单元均为可折叠设计，这意味着它们可以根据需要动态调整硬件单元的数量。例如，vALU可以配置为多个独立的执行单元，每个执行单元负责处理一个lane的计算任务。
2. 全流水配置：这些功能单元采用了全流水线设计，确保了每个指令的各个阶段（取指、译码、执行、写回）都可以在不同的周期内并行进行，从而提高了吞吐量。
3. num_lane参数：通过num_lane参数可以配置硬件单元的数量，结合num_thread参数可以自由控制每类向量指令的执行cycle数。例如，如果num_lane设置为8，那么每个vALU单元可以同时处理8个lane的计算任务，从而加速向量指令的执行。

典型latency：

• alu：1 cycle
• mul：2 cycles
• fmul：3 cycles
• fadd：3 cycles
• fmacc：5 cycles

这些latency值反映了不同指令类型的执行周期数，确保了“乘影”GPGPU在处理各种计算任务时的高效性和灵活性。

示例说明：

假设有一个矩阵乘法运算，涉及大量的向量乘法和累加操作。通过配置num_lane参数为8，vMUL单元可以同时处理8个lane的乘法操作，从而加速整个矩阵乘法的执行。结合num_thread参数，可以根据具体的计算需求灵活调整硬件资源的分配，确保在不同场景下都能获得最佳性能。

总结

“乘影”GPGPU通过SIMT-stack分支管理和双发射设计，实现了高效的任务管理和并行计算能力。SIMT-stack通过硬件管理分支，不占用通用寄存器，减少了流水线停顿，支持嵌套分支和快速分支合并。双发射机制允许在同一周期内发射一条标量指令和一条向量指令，进一步发掘并行度。功能单元的可配置性使得“乘影”GPGPU可以根据具体计算需求灵活调整硬件资源，确保在处理各种计算任务时的高效性和灵活性。

缓存设计

乱序机制与指令级并行

什么是乱序机制？

乱序机制（Out-of-Order Execution, OoOE） 是一种优化技术，允许处理器在执行指令时打破原有的程序顺序，以提高指令的并行度和执行效率。具体来说，乱序机制允许处理器在遇到依赖关系之前提前执行不相关的指令，从而减少了流水线停顿的时间。通过这种方式，处理器可以在一个周期内完成更多的工作，提高了吞吐量。

什么是指令级并行？

指令级并行（Instruction-Level Parallelism, ILP） 是指在同一时间执行多条独立的指令的能力。ILP可以通过多种方式实现，包括超标量（Superscalar）、超流水线（Superpipelining）和乱序执行（OoOE）。通过这些技术，处理器可以在一个时钟周期内执行多个指令，从而提高了性能。

为什么乱序机制对性能至关重要？

在现代高性能处理器中，乱序机制和多级缓存是提高性能的关键技术。乱序机制通过打破指令的顺序执行，减少了流水线中的等待时间，尤其是在处理复杂的分支预测、内存访问等操作时。多级缓存则通过减少内存访问延迟，进一步提升了指令的执行速度。然而，这些技术也会带来一些挑战，特别是在线程间同步和缓存一致性方面。

线程间同步与连贯性和一致性问题

什么是线程间同步？

线程间同步（Thread-Level Parallelism, TLP） 是指多个线程之间的协调和协作，确保它们能够正确地共享资源并协同工作。在GPGPU中，线程间同步非常重要，因为多个线程可能会同时访问相同的内存地址或寄存器，如果不加以控制，可能会导致数据竞争和不一致的问题。

什么是连贯性和一致性问题？

1. 连贯性问题：在多核或多线程系统中，不同的处理器核心或线程可能会看到不同的缓存内容。例如，一个核心写入了某个缓存行，但另一个核心还没有看到这个更新。这种现象称为缓存不连贯（Cache Incoherence）。连贯性问题通常发生在流水线中的乱序因素导致不同线程看到不同的缓存状态。

2. 一致性问题：缓存一致性（Cache Coherence） 是指确保所有处理器核心或线程看到相同的内存视图。如果一个核心修改了某个内存地址，其他核心应该能够立即看到这个修改。否则，就会出现缓存不一致（Cache Inconsistency） 的问题，导致程序行为异常。

缓存一致性（Cache Coherence）

定义

缓存一致性是指在一个多处理器或多核心系统中，确保所有处理器或核心看到的内存状态是一致的。具体来说，当一个处理器修改了某个内存位置的数据时，其他处理器必须能够看到这个修改后的最新值，而不是旧值。

重要性

缓存一致性在多处理器系统中非常重要，因为每个处理器都有自己的缓存，这些缓存中可能保存了相同的内存位置的数据副本。如果不确保一致性，可能会导致以下问题：

• 数据不一致：不同处理器看到的数据不一致，导致错误的计算结果。
• 死锁和竞态条件：多个处理器同时访问和修改同一数据，可能导致死锁或竞态条件。

弱存储模型（Weak Memory Model）

定义

弱存储模型是一种内存模型，它允许某些内存操作的顺序与程序中的顺序不同。这种模型通常用于提高并发性能，因为它允许编译器和处理器对内存操作进行重新排序，从而更好地利用硬件资源。

特点

• 放宽顺序性：允许内存操作在一定范围内重新排序，只要最终结果符合程序的逻辑。
• 显式同步：需要显式的同步指令（如FENCE）来确保特定的内存操作顺序。

连贯性指导的缓存一致性（RCC）

什么是RCC？

连贯性指导的缓存一致性（RCC, Release Consistency-directed Cache Coherence） 是“乘影”GPGPU中采用的一种缓存一致性协议。相比传统的硬件一致性协议（如MESI），RCC通过放松一致性要求，减少了L1-L2缓存之间的带宽开销和硬件复杂度。同时，相比GPU的软件一致性方案，RCC减少了编程框架的额外负担，使得开发者更容易编写高效的并行程序。

硬件一致性协议（如MESI）

什么是MESI协议？

MESI协议 是一种常见的缓存一致性协议，广泛应用于多核处理器中。MESI是“Modified, Exclusive, Shared, Invalid”的缩写，表示缓存行的四种状态：

1. Modified（修改）：该缓存行在当前核心中被修改，并且其他核心没有该缓存行的副本。这意味着只有当前核心拥有最新的数据，必须将修改后的数据写回主内存或L2缓存。

2. Exclusive（独占）：该缓存行在当前核心中是唯一的，并且其他核心没有该缓存行的副本。这意味着当前核心可以安全地读取或写入该缓存行，而不需要与其他核心同步。

3. Shared（共享）：该缓存行在多个核心中都有副本。这意味着所有核心都可以读取该缓存行，但不能写入。如果某个核心需要写入该缓存行，则必须先使其他核心的副本无效（即转换为Invalid状态）。

4. Invalid（无效）：该缓存行无效，表示当前核心没有该缓存行的有效副本。如果需要访问该缓存行，必须从主内存或其他核心中获取最新数据。

MESI的工作原理

当多个核心同时访问相同的内存地址时，MESI协议通过消息传递机制来确保所有核心看到一致的内存视图。具体来说：

• 读请求：当一个核心读取某个缓存行时，它会检查本地缓存。如果命中且状态为Shared或Exclusive，则可以直接使用；如果未命中或状态为Invalid，则会向其他核心发送读请求，获取最新的数据。

• 写请求：当一个核心写入某个缓存行时，它会检查本地缓存。如果状态为Modified或Exclusive，则可以直接写入；如果状态为Shared，则必须先使其他核心的副本无效，然后才能写入；如果状态为Invalid，则必须从主内存或其他核心中获取最新数据，再进行写入。

MESI的优缺点

优点：

• 强一致性：MESI协议确保所有核心看到一致的内存视图，避免了数据竞争和不一致的问题。
• 简单易实现：MESI协议相对简单，易于实现和调试。

缺点：

• 带宽开销大：由于每个核心之间的通信频繁，尤其是在多线程程序中，可能会产生大量的消息传递，占用宝贵的带宽资源。
• 硬件复杂度高：为了实现MESI协议，硬件需要支持复杂的控制逻辑和消息传递机制，增加了设计难度和功耗。

GPU的软件一致性方案

什么是GPU的软件一致性方案？

GPU的软件一致性方案 是指通过编程框架（如CUDA、OpenCL）中的显式同步指令（如__syncthreads()、fence等）来管理缓存一致性。与硬件一致性协议不同，软件一致性方案依赖于开发者手动插入同步点，确保多个线程之间的内存操作顺序正确。

软件一致性的工作原理

在GPU编程中，开发者通常需要手动管理缓存一致性，以确保多个线程之间的数据同步。例如，在CUDA中，开发者可以使用__syncthreads()指令来确保同一块内的所有线程都完成了某个操作，然后再继续执行后续代码。类似地，fence指令用于确保在fence之前的所有内存操作都完成之后，才会执行fence之后的操作。

软件一致性的优缺点

优点：

• 灵活性高：开发者可以根据具体的应用需求，灵活地插入同步点，优化性能。
• 减少硬件复杂度：由于一致性管理由软件负责，硬件设计可以简化，降低了设计难度和功耗。

缺点：

• 编程复杂度高：开发者需要手动管理缓存一致性，增加了编程的复杂性。如果不小心插入了过多的同步点，可能会导致性能下降。
• 容易出错：如果开发者没有正确插入同步点，可能会导致数据竞争和不一致的问题，难以调试。

RCC为什么可以减少L1-L2缓存之间的带宽开销和硬件复杂度？

RCC的核心思想

RCC（Relaxed Cache Coherence） 是一种放松的一致性协议，旨在通过放宽一致性要求，减少L1-L2缓存之间的带宽开销和硬件复杂度。具体来说，RCC允许某些类型的内存操作在不同的线程之间看到不同的视图，只要这些操作不会影响程序的正确性。通过这种方式，RCC减少了不必要的消息传递和缓存同步操作，提高了性能。

如何减少带宽开销？

1. 减少消息传递：传统的MESI协议需要频繁地在L1和L2缓存之间传递消息，以确保一致性。RCC通过放宽一致性要求，减少了不必要的消息传递。例如，RCC允许某些读操作直接从L2缓存读取数据，而不必等待其他核心的响应。这显著减少了L1-L2缓存之间的带宽开销。

2. 延迟写回：在MESI协议中，当一个核心修改了某个缓存行时，必须立即将修改后的数据写回L2缓存，以确保其他核心能够看到最新的数据。RCC允许延迟写回，只有在必要时（如释放操作）才将脏数据写回L2缓存。这减少了写操作的频率，进一步降低了带宽开销。

如何减少硬件复杂度？

1. 简化控制逻辑：MESI协议需要复杂的控制逻辑来管理缓存状态和消息传递。RCC通过放宽一致性要求，简化了控制逻辑。例如，RCC不需要跟踪每个缓存行的状态（如Modified、Exclusive、Shared、Invalid），而是只关注某些关键操作（如acquire、release）。这使得硬件设计更加简单，降低了功耗。

2. 减少消息传递机制：MESI协议需要复杂的消息传递机制来确保一致性。RCC通过减少不必要的消息传递，简化了硬件设计。例如，RCC允许某些读操作直接从L2缓存读取数据，而不必等待其他核心的响应。这减少了硬件中的消息传递单元，降低了设计复杂度。

RCC为什么可以减少编程框架的额外负担，使得开发者更容易编写高效的并行程序？

简化一致性管理

1. 自动处理一致性：在传统的GPU编程中，开发者需要手动插入同步点（如__syncthreads()、fence等）来管理缓存一致性。RCC通过硬件支持，自动处理大部分一致性问题，减少了开发者的工作量。例如，RCC在acquire和release操作时自动执行必要的缓存同步操作，开发者无需手动管理。

2. 减少同步开销：在传统的GPU编程中，开发者插入过多的同步点可能会导致性能下降。RCC通过放宽一致性要求，减少了不必要的同步操作，使得程序运行更加高效。例如，RCC允许某些读操作直接从L2缓存读取数据，而不必等待其他核心的响应。这减少了同步开销，提升了性能。

提高编程灵活性

1. 灵活的内存模型：RCC基于弱内存模型（如RVWMO），允许开发者根据具体的应用需求选择合适的内存顺序。例如，开发者可以选择在某些关键操作（如acquire、release）处插入同步点，而在其他地方保持宽松的内存顺序。这种灵活性使得开发者可以更好地优化程序性能。

2. 减少编程复杂度：RCC通过硬件支持，简化了缓存一致性管理，减少了编程的复杂性。开发者无需担心复杂的同步问题，可以专注于算法的设计和优化。这使得开发者更容易编写高效的并行程序，尤其是对于初学者来说，学习曲线更加平缓。

RCC的具体优点

1. 降低L1-L2带宽开销：传统的硬件一致性协议（如MESI）需要频繁地在L1和L2缓存之间传递消息，以确保一致性。这会占用大量的带宽资源，影响性能。RCC通过放松一致性要求，减少了不必要的消息传递，降低了带宽开销。
2. 降低硬件复杂度：传统的一致性协议需要复杂的硬件支持，增加了芯片的设计难度和功耗。RCC通过简化一致性协议，减少了硬件复杂度，使得设计更加高效。
3. 减少编程框架的额外负担：在GPU中，软件一致性通常需要开发者手动管理缓存一致性，增加了编程的复杂性。RCC通过硬件支持，减少了开发者的工作量，使得编程更加简单。

总结

相比传统的硬件一致性协议（如MESI），RCC通过放宽一致性要求，减少了L1-L2缓存之间的带宽开销和硬件复杂度。具体来说，RCC允许某些类型的内存操作在不同的线程之间看到不同的视图，减少了不必要的消息传递和缓存同步操作。相比GPU的软件一致性方案，RCC通过硬件支持，自动处理大部分一致性问题，减少了开发者的工作量，使得开发者更容易编写高效的并行程序。RCC不仅简化了硬件设计，还提高了编程灵活性，使得开发者可以更好地优化程序性能。

RISC-V通过显式指令支持RVWMO

RVWMO是什么？

RVWMO（RISC-V Weak Memory Ordering） 是RISC-V架构中的一种弱内存模型，允许处理器在执行内存操作时进行一定的重排序，以提高性能。为了确保程序的正确性，RISC-V提供了几种显式指令来管理内存顺序：

1. FENCE指令：

   • FENCE R,R：确保所有先前的读操作在所有后续的读操作之前完成。
   • FENCE R,W：确保所有先前的读操作在所有后续的写操作之前完成。
   • FENCE W,R：确保所有先前的写操作在所有后续的读操作之前完成。
   • FENCE W,W：确保所有先前的写操作在所有后续的写操作之前完成。
   • 确保顺序性：FENCE指令确保在该指令之前的内存操作在该指令之后的内存操作之前完成。
   • 防止重排序：编译器和处理器可能会为了优化性能而重新排序内存操作，FENCE指令可以防止这种重排序。

2. .aq（acquire）和.rl（release）标识符：

   • .aq：用于标记内存操作为获取操作，确保该操作之前的读写操作在该操作之前完成。
   • .rl：用于标记内存操作为释放操作，确保该操作之后的读写操作在该操作之后完成。
   • 使用场景
     1. 同步数据访问：在多线程程序中，当一个线程需要确保在获取某个数据之前，所有之前的写操作已经完成时，可以使用.aq标识符。
     2. 同步数据更新：当一个线程需要确保在释放某个数据之后，所有之后的读操作能够看到最新的数据时，可以使用.rl标识符。

3. AMO原子指令：用于执行原子操作，确保多个线程在访问共享资源时不会发生冲突。

缓存微架构设计

为了支持RCC一致性协议，“乘影”GPGPU的缓存微架构设计了一系列操作来确保缓存的一致性和连贯性。具体来说，这些操作包括：

1. Fence操作：FENCE指令用于确保在FENCE指令之前的所有内存操作都完成之后，才会执行FENCE指令之后的内存操作。FENCE指令通过清空MSHR（Miss Status Handling Register）、全局冲刷（Global Flush）和全局无效化（Global Invalidate）等微架构操作来实现。MSHR、全局冲刷、全局无效化的解释看下一节。

2. Acquire/Release操作：

   • Acquire操作：当acquire发生时，无论是否命中，直接从L2缓存读取数据，并无效化缓存中该line之外的所有line。这样可以确保在acquire之前的操作不会被重排序到acquire之后。
   • Release操作：当release发生时，L1缓存中所有脏缓存行会被写回L2缓存（即flush），然后release对应的写入再写回L2。这样可以确保在release之后的操作不会被重排序到release之前。
   • 有关acquire和release的操作也会在后面的章节中进行解释。

示例说明

假设有一个多线程程序，其中两个线程分别执行以下操作：

// 线程A
data[0] = 1;  // 写操作
__release();   // release语义

// 线程B
__acquire();   // acquire语义
if (data[0] == 1) {
    data[1] = 2;  // 写操作
}

在这个例子中，__release() 和 __acquire() 分别表示release和acquire语义。根据RCC一致性协议，线程A的data[0] = 1操作会在__release()之前完成，而线程B的__acquire()操作会确保在读取data[0]之前，所有之前的写操作都已经完成。因此，线程B可以安全地读取data[0]，并根据其值执行后续操作。

FENCE指令的微架构操作

什么是FENCE指令？

FENCE指令 是RISC-V架构中用于确保内存操作顺序的一条同步指令。它确保在FENCE指令之前的所有内存操作都完成之后，才会执行FENCE指令之后的内存操作。FENCE指令可以用于防止编译器或硬件对内存访问进行重排序，确保程序的正确性。

为了实现FENCE指令的功能，处理器需要执行一系列微架构操作，以确保所有未完成的内存操作都得到处理，并且缓存状态一致。这些操作包括清空MSHR（Miss Status Handling Register）、全局冲刷（Global Flush）和全局无效化（Global Invalidate）。接下来我们将详细解释这三个操作。

清空MSHR（Miss Status Handling Register）

什么是MSHR？

MSHR（Miss Status Handling Register） 是一种特殊的寄存器，用于跟踪缓存未命中（cache miss）的状态。当一个核心尝试访问某个缓存行时，如果该缓存行不在L1缓存中（即发生缓存未命中），则MSHR会记录这次未命中的请求，并等待从L2缓存或主内存中获取数据。MSHR的作用是确保多个未命中的请求不会相互干扰，并且能够按正确的顺序处理。

为什么需要清空MSHR？

在执行FENCE指令时，清空MSHR是为了确保所有未完成的缓存未命中请求都得到处理。具体来说，FENCE指令要求在FENCE之前的所有内存操作都必须完成，因此任何未命中的请求都需要等待数据从L2缓存或主内存中返回，然后再继续执行FENCE之后的操作。通过清空MSHR，处理器可以确保所有未命中的请求都已完成，避免了内存操作的重排序。

示例说明

假设有一个多线程程序，其中两个线程分别执行以下操作：

// 线程A
data[0] = 1;  // 写操作
__fence();    // FENCE指令

// 线程B
if (data[0] == 1) {
    data[1] = 2;  // 写操作
}

在这个例子中，线程A的data[0] = 1操作可能会导致缓存未命中，因此MSHR会记录这个请求并等待数据从L2缓存或主内存中返回。FENCE指令确保在__fence()之前的所有内存操作（包括data[0] = 1）都已完成，然后才会继续执行FENCE之后的操作。通过清空MSHR，处理器可以确保data[0] = 1已经写入L1缓存或L2缓存，线程B可以安全地读取data[0]。

全局冲刷（Global Flush）

什么是全局冲刷？

全局冲刷（Global Flush） 是指将所有脏数据（即已经被修改但尚未写回L2缓存或主内存的数据）从L1缓存中写回L2缓存或主内存。全局冲刷确保所有缓存中的脏数据都被刷新，使得其他核心或线程能够看到最新的数据。

为什么需要全局冲刷？

在执行FENCE指令时，全局冲刷是为了确保在FENCE之前的所有写操作都已完成，并且所有脏数据都已经被写回L2缓存或主内存。这样可以避免其他核心或线程读取到过时的数据，确保内存一致性。全局冲刷还可以防止FENCE之后的读操作读取到FENCE之前的未刷新数据，从而保证内存操作的顺序性。

示例说明

假设有一个多线程程序，其中两个线程分别执行以下操作：

// 线程A
data[0] = 1;  // 写操作
__fence();    // FENCE指令

// 线程B
if (data[0] == 1) {
    data[1] = 2;  // 写操作
}

在这个例子中，线程A的data[0] = 1操作可能会导致L1缓存中的数据被标记为脏数据（即已经被修改但尚未写回L2缓存）。FENCE指令通过全局冲刷，确保data[0] = 1已经被写回L2缓存或主内存，线程B可以安全地读取data[0]，并根据其值执行后续操作。

全局无效化（Global Invalidate）

什么是全局无效化？

全局无效化（Global Invalidate） 是指使所有缓存行（无论是干净的还是脏的）都变为无效状态。这意味着所有缓存中的数据都将被丢弃，未来的读操作必须从L2缓存或主内存中重新获取最新数据。全局无效化通常用于确保所有核心或线程看到一致的内存视图，尤其是在执行某些关键操作（如acquire、release）时。

为什么需要全局无效化？

在执行FENCE指令时，全局无效化是为了确保在FENCE之后的读操作不会读取到FENCE之前的旧数据。通过使所有缓存行无效，处理器可以强制未来的读操作从L2缓存或主内存中重新获取最新数据，从而保证内存一致性。此外，全局无效化还可以防止不同核心之间的缓存不一致问题，确保所有核心看到相同的内存视图。

示例说明

假设有一个多线程程序，其中两个线程分别执行以下操作：

// 线程A
data[0] = 1;  // 写操作
__fence();    // FENCE指令

// 线程B
if (data[0] == 1) {
    data[1] = 2;  // 写操作
}

在这个例子中，线程A的data[0] = 1操作可能会导致L1缓存中的数据被标记为脏数据。FENCE指令通过全局无效化，确保线程B在读取data[0]时不会读取到L1缓存中的旧数据，而是从L2缓存或主内存中重新获取最新数据。这确保了线程B能够看到data[0] = 1的最新值，并根据其值执行后续操作。

总结

FENCE指令通过清空MSHR、全局冲刷和全局无效化等微架构操作来确保内存操作的顺序性和一致性。具体来说：

• 清空MSHR：确保所有未完成的缓存未命中请求都得到处理，避免内存操作的重排序。
• 全局冲刷：将所有脏数据从L1缓存中写回L2缓存或主内存，确保其他核心或线程能够看到最新的数据。
• 全局无效化：使所有缓存行无效，确保未来的读操作从L2缓存或主内存中重新获取最新数据，避免读取到旧数据。

这些操作共同作用，确保FENCE指令能够正确地同步内存操作，避免数据竞争和不一致的问题，从而保证程序的正确性和性能。

详细解释Acquire和Release操作

在多线程编程中，Acquire 和 Release 是两种重要的内存屏障（Memory Barrier）操作，用于确保不同线程之间的内存操作顺序正确。它们通过强制某些内存操作的顺序，防止编译器或硬件对这些操作进行重排序，从而保证程序的正确性和一致性。

Acquire操作

什么是Acquire操作？

Acquire操作 是一种内存屏障，用于确保在acquire之前的所有读写操作都完成之后，才会执行acquire之后的操作。具体来说，acquire操作确保在acquire之前的所有内存操作不会被重排序到acquire之后，从而保证了内存操作的顺序性。

Acquire操作的具体实现

当acquire发生时，处理器会执行以下操作：

1. 直接从L2缓存读取数据：无论当前L1缓存中是否命中（即是否存在该缓存行），acquire操作都会直接从L2缓存中读取最新的数据。这确保了acquire操作能够获取到最新的共享数据，避免了读取到过时的数据。

2. 无效化缓存中该line之外的所有line：为了确保其他线程看到一致的内存视图，acquire操作会无效化L1缓存中除当前访问的缓存行之外的所有缓存行。这意味着未来的读操作必须从L2缓存或主内存中重新获取最新数据，而不是使用L1缓存中的旧数据。

为什么需要Acquire操作？

Acquire操作的主要目的是确保在acquire之前的所有内存操作都已完成，并且所有共享数据都已经被刷新到L2缓存或主内存。这样可以防止其他线程读取到过时的数据，确保内存一致性。此外，acquire操作还可以防止编译器或硬件对内存操作进行重排序，从而保证程序的正确性。

示例说明

假设有一个多线程程序，其中两个线程分别执行以下操作：

// 线程A
data[0] = 1;  // 写操作
__release(&flag);  // release语义

// 线程B
__acquire(&flag);  // acquire语义
if (flag) {
    print(data[0]);  // 读操作
}

在这个例子中，线程A首先将data[0]设置为1，然后执行__release(&flag)操作，表示释放锁并确保所有之前的写操作已经完成。线程B随后执行__acquire(&flag)操作，表示获取锁并确保在acquire之前的所有内存操作都已完成。

• 线程A的行为：

  • data[0] = 1：将data[0]设置为1。
  • __release(&flag)：将flag设置为1，并确保所有之前的写操作（如data[0] = 1）已经完成并且写回L2缓存或主内存。

• 线程B的行为：

  • __acquire(&flag)：读取flag，并确保在acquire之前的所有内存操作都已完成。即使flag已经在L1缓存中命中，线程B也会从L2缓存中重新读取flag，以确保获取到最新的值。
  • print(data[0])：只有当flag为1时，线程B才会读取data[0]。由于__acquire(&flag)确保了flag的最新值已经被读取，线程B可以安全地读取data[0]，并打印出正确的值1。

通过Acquire操作，线程B可以确保在flag为1之前，data[0]已经被正确写入并刷新到L2缓存或主内存，从而避免了读取到过时的数据。

Release操作

什么是Release操作？

Release操作 是另一种内存屏障，用于确保在release之后的所有读写操作都不会被重排序到release之前。具体来说，release操作确保在release之后的所有内存操作不会影响release之前的内存操作，从而保证了内存操作的顺序性。

Release操作的具体实现

当release发生时，处理器会执行以下操作：

1. 将L1缓存中所有脏缓存行写回L2缓存：release操作会将L1缓存中所有已经被修改但尚未写回L2缓存的脏数据（即脏缓存行）写回L2缓存。这确保了所有之前的写操作都已经完成，并且最新的数据已经被刷新到L2缓存或主内存。

2. 将release对应的写入再写回L2：release操作本身通常涉及一个写操作（如设置某个标志位），这个写操作也会被立即写回L2缓存或主内存。这确保了其他线程能够立即看到release操作的结果。

为什么需要Release操作？

Release操作的主要目的是确保在release之前的写操作都已经完成，并且最新的数据已经被刷新到L2缓存或主内存。这样可以防止其他线程读取到过时的数据，确保内存一致性。此外，release操作还可以防止编译器或硬件对内存操作进行重排序，从而保证程序的正确性。

示例说明

继续上面的例子：

// 线程A
data[0] = 1;  // 写操作
__release(&flag);  // release语义

// 线程B
__acquire(&flag);  // acquire语义
if (flag) {
    print(data[0]);  // 读操作
}

• 线程A的行为：

  • data[0] = 1：将data[0]设置为1。
  • __release(&flag)：将flag设置为1，并确保所有之前的写操作（如data[0] = 1）已经完成并且写回L2缓存或主内存。release操作还会将flag的写入也写回L2缓存或主内存，确保其他线程能够立即看到flag的变化。

• 线程B的行为：

  • __acquire(&flag)：读取flag，并确保在acquire之前的所有内存操作都已完成。即使flag已经在L1缓存中命中，线程B也会从L2缓存中重新读取flag，以确保获取到最新的值。
  • print(data[0])：只有当flag为1时，线程B才会读取data[0]。由于__acquire(&flag)确保了flag的最新值已经被读取，线程B可以安全地读取data[0]，并打印出正确的值1。

通过Release操作，线程A可以确保在flag被设置为1之前，data[0]已经被正确写入并刷新到L2缓存或主内存，从而避免了线程B读取到过时的数据。

总结

• Acquire操作：确保在acquire之前的所有内存操作都已完成，并且所有共享数据都已经被刷新到L2缓存或主内存。它通过直接从L2缓存读取数据并无效化缓存中其他line，防止其他线程读取到过时的数据，确保内存一致性。

• Release操作：确保在release之前的写操作都已经完成，并且最新的数据已经被刷新到L2缓存或主内存。它通过将L1缓存中所有脏缓存行写回L2缓存，并将release对应的写入也写回L2缓存，确保其他线程能够立即看到最新的数据，防止内存操作的重排序。

通过Acquire和Release操作，开发者可以确保多线程程序中的内存操作顺序正确，避免数据竞争和不一致的问题，从而保证程序的正确性和性能。

不同地址空间的缓存策略

缓存策略

为了优化缓存性能，“乘影”GPGPU针对不同类型的内存访问采用了不同的缓存策略。具体来说：

• 私有内存（Private Memory）：

  • 写命中（Write Hit）：使用write-back策略，即只有当缓存行被替换时才将脏数据写回L2缓存。
  • 写缺失（Write Miss）：使用write non-allocate策略，即不在L1缓存中分配新的缓存行，而是直接将数据写入L2缓存。

• 全局内存（Global Memory）：

  • 写命中（Write Hit）：使用write-through策略，即每次写操作都会立即写入L2缓存，而不保留脏数据。
  • 写缺失（Write Miss）：使用write non-allocate策略，即不在L1缓存中分配新的缓存行，而是直接将数据写入L2缓存。

为什么这么设计？

1. 私有内存：私有内存通常是每个线程独占的，因此可以使用write-back策略来减少不必要的写回操作，提高缓存命中率。write non-allocate策略则避免了不必要的缓存分配，减少了缓存污染。

2. 全局内存：全局内存是多个线程共享的，因此使用write-through策略可以确保所有线程都能立即看到最新的写入结果，避免了一致性问题。write non-allocate策略则避免了全局内存数据在L1缓存中的驻留，减少了缓存污染。

写命中、写缺失、写回、写直通、写分配、写不分配

在缓存系统中，写操作 是指将数据写入缓存或内存的过程。根据缓存的状态和策略，写操作可以分为不同的类型，如写命中、写缺失、写回、写直通、写分配和写不分配等。这些概念描述了处理器在处理写操作时的行为，影响着缓存性能和一致性。

写命中（Write Hit）

写命中 是指当处理器尝试写入某个地址时，该地址的数据已经在缓存中存在（即缓存命中）。此时，处理器可以直接将数据写入缓存，而不需要访问更高级别的缓存或主内存。

• 行为：直接将数据写入L1缓存。
• 优点：写命中可以减少访问主内存的延迟，提高写操作的速度。
• 缺点：如果缓存行是共享的（即多个核心都有副本），则需要确保其他核心的副本也得到更新，以保持一致性。

写缺失（Write Miss）

写缺失 是指当处理器尝试写入某个地址时，该地址的数据不在缓存中（即缓存未命中）。此时，处理器需要从更高级别的缓存或主内存中获取该地址的数据，或者直接将数据写入更高级别的缓存或主内存。

• 行为：处理器需要决定如何处理写缺失。常见的策略包括写分配（Write Allocate）和写不分配（Write No-allocate）。
• 优点：写缺失可以通过写分配将数据加载到缓存中，从而提高后续读取的命中率。
• 缺点：写缺失会增加访问主内存的延迟，降低写操作的速度。

写回（Write Back）

写回 是一种缓存写策略，当缓存中的数据被修改后，只有在该缓存行被替换出缓存时，才会将修改后的数据写回更高级别的缓存或主内存。这种方式称为写回策略（Write-Back Policy）。

• 行为：当缓存行被替换时，如果该缓存行是脏的（即已经被修改），则将其写回L2缓存或主内存；否则，直接丢弃。
• 优点：写回策略减少了写操作的频率，因为只有在缓存行被替换时才需要写回。这可以显著降低带宽开销，提高性能。
• 缺点：写回策略可能会导致缓存污染，尤其是在多线程环境中，多个线程频繁修改同一缓存行时，可能会引发一致性问题。

写直通（Write Through）

写直通 是另一种缓存写策略，当处理器写入缓存时，数据不仅会被写入L1缓存，还会立即写入更高级别的缓存或主内存。这种方式称为写直通策略（Write-Through Policy）。

• 行为：每次写操作都会同时更新L1缓存和L2缓存或主内存。
• 优点：写直通策略确保所有写操作立即可见，避免了缓存一致性问题。它特别适用于全局共享数据或频繁写入的场景。
• 缺点：写直通策略增加了写操作的频率，因为它每次写入都需要访问L2缓存或主内存，导致更高的带宽开销和延迟。

写分配（Write Allocate）

写分配 是一种处理写缺失的策略，当发生写缺失时，处理器不仅会将数据写入更高级别的缓存或主内存，还会将该地址的数据加载到L1缓存中。这种方式称为写分配（Write Allocate）。

• 行为：当发生写缺失时，处理器会从L2缓存或主内存中获取该地址的数据，并将其加载到L1缓存中，然后执行写操作。
• 优点：写分配可以提高后续读取的命中率，因为该地址的数据已经存在于L1缓存中。
• 缺点：写分配会增加缓存污染的风险，尤其是在多线程环境中，多个线程频繁写入同一地址时，可能会导致缓存行频繁替换。

写不分配（Write No-allocate）

写不分配 是另一种处理写缺失的策略，当发生写缺失时，处理器不会将该地址的数据加载到L1缓存中，而是直接将数据写入更高级别的缓存或主内存。这种方式称为写不分配（Write No-allocate）。

• 行为：当发生写缺失时，处理器直接将数据写入L2缓存或主内存，而不将该地址的数据加载到L1缓存中。
• 优点：写不分配减少了缓存污染的风险，因为它不会将不必要的数据加载到L1缓存中。
• 缺点：写不分配可能会降低后续读取的命中率，因为该地址的数据不在L1缓存中，导致更多的缓存未命中。

不同地址空间的缓存策略

在“乘影”GPGPU中，针对不同类型的内存访问采用了不同的缓存策略，具体如下：

• 私有内存（Private Memory）：

  • 写命中：使用写回策略，即只有当缓存行被替换时才将脏数据写回L2缓存。
  • 写缺失：使用写不分配策略，即不在L1缓存中分配新的缓存行，而是直接将数据写入L2缓存。

• 全局内存（Global Memory）：

  • 写命中：使用写直通策略，即每次写操作都会立即写入L2缓存，而不保留脏数据。
  • 写缺失：使用写不分配策略，即不在L1缓存中分配新的缓存行，而是直接将数据写入L2缓存。

为什么这样设计？

• 私有内存：私有内存通常是每个线程独占的，因此可以使用写回策略来减少不必要的写回操作，提高缓存命中率。写不分配策略则避免了不必要的缓存分配，减少了缓存污染。

• 全局内存：全局内存是多个线程共享的，因此使用写直通策略可以确保所有线程都能立即看到最新的写入结果，避免了一致性问题。写不分配策略则避免了全局内存数据在L1缓存中的驻留，减少了缓存污染。

总结

• 写命中：缓存命中时直接将数据写入L1缓存。
• 写缺失：缓存未命中时，处理器需要决定是否将数据加载到L1缓存中（写分配 vs. 写不分配）。
• 写回：只有在缓存行被替换时才将脏数据写回L2缓存或主内存。
• 写直通：每次写操作都会同时更新L1缓存和L2缓存或主内存。
• 写分配：发生写缺失时，将数据加载到L1缓存中。
• 写不分配：发生写缺失时，直接将数据写入L2缓存或主内存，不将数据加载到L1缓存中。

总结

“乘影”GPGPU的缓存设计通过引入连贯性指导的缓存一致性（RCC）协议，解决了乱序机制和多级缓存带来的连贯性和一致性问题。RCC通过放松一致性要求，降低了L1-L2缓存之间的带宽开销和硬件复杂度，同时减少了编程框架的额外负担。此外，RISC-V通过显式指令（如FENCE、AMO、acquire/release）支持RVWMO内存模型，确保了程序的正确性和性能。最后，针对不同类型的内存访问，缓存设计采用了不同的策略，优化了缓存性能，减少了缓存污染。

“乘影”GPGPU v2.1.0微架构

前端（Front-End）

• Fetch：指令获取阶段，从指令缓存中读取指令。指令缓存（Instruction Cache）存储最近使用的指令，以减少内存访问延迟。
• Warp Scheduler：线程束调度器，Warp Scheduler负责在单个流多处理器（SM）内调度多个warp。每个线程束包含32个线程，这些线程并行执行相同的指令。
• Decode：指令解码阶段，将获取的指令解码为微操作（Micro-operations）。
• Ibuffer：指令缓冲区，存储解码后的微操作，等待执行。
• Operand Collector / Issue：操作数收集和分发阶段，收集操作数并将微操作分发到相应的功能单元（Function Unit）。
• Scoreboard：指令发布窗口，用于跟踪指令的依赖关系，确保指令按正确的顺序执行。
• RegFile：寄存器文件，存储线程的寄存器值。

后端（Back-End）

• SIMT-stack：SIMT（Single Instruction, Multiple Threads）栈，用于管理线程束的状态。
• sALU：标量算术逻辑单元，执行标量运算。
• vALU：向量算术逻辑单元，执行向量运算。
• vFPU：向量浮点运算单元，执行向量浮点运算。
• SFU：特殊功能单元，执行特殊功能运算，如三角函数、指数函数等。
• CSR：控制和状态寄存器，用于控制和状态管理。
• TensorCore：张量核心，用于加速矩阵运算，特别适用于深度学习应用。
• LSU：加载/存储单元，处理内存访问操作。
• Writeback：写回阶段，将计算结果写回寄存器文件或内存。

共享内存和数据缓存

• Shared Memory：共享内存，用于线程束之间的快速数据交换。
• Data Cache：数据缓存，存储最近使用的数据，以减少内存访问延迟。

二级缓存和全局内存

• L2 Cache：二级缓存，存储最近使用的数据和指令，减少内存访问延迟。
• Global Memory：全局内存，存储程序和数据，由所有流处理器共享。

示例说明

假设有一个深度学习应用，需要执行大量的矩阵运算。该应用首先将矩阵数据加载到全局内存中。然后，主机设备（Host Device）将计算任务发送到GPGPU。GPGPU的前端部分负责获取和解码指令，将线程束分配到流处理器上执行。后端部分的张量核心（TensorCore）负责加速矩阵运算。计算结果最终被写回共享内存或全局内存。

取指、译码、发射，执行，写回

在“乘影”GPU中，每个warp被视为一段RVV（RISC-V向量扩展）程序，其执行过程遵循传统的指令处理流水线模型。这个模型包括取指（Fetch）、译码（Decode）、发射（Issue）、执行（Execute）和写回（Writeback）几个阶段。以下是每个阶段的详细解释，结合工作台中的文件内容：

取指（Fetch）

功能：取指阶段是处理器从程序存储器中获取指令的阶段。在GPU中，这通常涉及到从全局内存中加载指令，因为GPU程序（如着色器程序）通常存储在全局内存中。

操作：Warp Scheduler负责调度warp，并从指令缓存（Instruction Cache）中请求指令。如果指令不在缓存中，可能需要从更高层的内存（如L2缓存或DDR内存）中获取。

译码（Decode）

功能：译码阶段将取指阶段获取的机器码指令转换为处理器可以理解的微指令或操作码。

操作：译码器（Decoder）解析指令，确定所需的操作类型、操作数以及目标寄存器。这个阶段还会检查指令的有效性和任何潜在的错误。

发射（Issue）

功能：发射阶段是将解码后的指令发送到执行单元之前，将其放入指令缓冲区（Instruction Buffer，IBuffer）或重排序缓冲区（Reservation Station）。

操作：指令从译码阶段进入发射阶段，在这里它们等待必要的寄存器值变得可用。如果所有依赖的寄存器值都已准备好，指令就可以被发射到执行单元。

执行（Execute）

功能：执行阶段是实际执行指令的地方。

操作：在这个阶段，指令被送到执行单元（如ALU、FPU或向量执行单元）。对于RVV程序，这可能涉及到向量运算，其中单个指令操作多个数据元素。执行单元计算指令的结果，并将结果存储在寄存器中。

写回（Writeback）

功能：写回阶段是将执行阶段得到的结果写回到寄存器文件中。

操作：执行完成后，结果被写回到寄存器，这样其他指令就可以使用这些结果。对于向量指令，这可能涉及到将多个结果写回到多个寄存器。

在整个过程中，Warp Scheduler负责管理warp的执行，确保指令流水线的高效运行，并处理任何由于资源冲突或数据依赖性导致的问题。这种周期级的调度和执行模型使得GPU能够有效地处理大量的并行线程，并在单个物理核心上实现高性能的并行计算。

IBuffer

指令缓冲区（Instruction Buffer，IBuffer）是处理器流水线中的一个重要组件，尤其在超标量和乱序执行的处理器中起着关键作用。以下是指令缓冲区的主要作用和详细解释：

存储解码后的指令

指令缓冲区的主要作用之一是存储从译码阶段过来的解码后的指令。这些指令已经从原始的机器码形式转换为处理器可以理解的微操作或操作码，使得后续的发射和执行阶段可以更高效地进行。

指令发射的准备

在发射阶段，指令缓冲区为指令的发射做好准备。它暂存解码后的指令，等待所有必要的源操作数变得可用。一旦所有依赖的操作数就绪，指令就可以从缓冲区中取出并发射到执行单元。

乱序执行的支持

在支持乱序执行的处理器中，指令缓冲区允许处理器发射和执行指令的顺序与它们被解码的顺序不同。这有助于提高执行效率，因为处理器可以跳过等待数据依赖解决的指令，转而执行其他准备好的指令。

减少流水线停顿

指令缓冲区有助于减少流水线的停顿。如果执行单元暂时空闲，而译码阶段仍在持续产生解码后的指令，缓冲区可以确保有足够的指令供给发射和执行阶段，从而保持流水线的连续运行。

提高指令吞吐量

通过暂存多个指令，指令缓冲区可以提高处理器的指令吞吐量。它允许处理器在单个时钟周期内发射多个指令到执行单元，这是提高处理器性能的关键因素。

指令重排序

在某些架构中，指令缓冲区还负责指令的重排序。如果后续指令依赖于前面指令的结果，缓冲区可以暂时保存这些指令，直到所需的数据变得可用，从而维护程序的正确执行顺序。

处理分支和跳转

在遇到分支或跳转指令时，指令缓冲区可以快速丢弃不再需要的指令，并将控制权转移到新的指令流。这有助于快速响应程序中的控制流变化。

总结

指令缓冲区是处理器流水线中的一个关键组件，它通过存储、管理和调度指令来提高处理器的效率和性能。它支持乱序执行、减少流水线停顿、提高指令吞吐量，并帮助处理分支和跳转指令，从而在现代处理器设计中发挥着至关重要的作用。

总结

“乘影”GPGPU v2.1.0微架构通过前端和后端的协同工作，实现了高效的并行计算。前端负责指令的获取、解码和调度，后端负责指令的执行和结果的写回。共享内存和数据缓存用于减少内存访问延迟，二级缓存用于存储最近使用的数据和指令。通过这些机制，GPGPU能够高效地处理大规模并行计算任务，特别适用于深度学习等应用。

基于Verilator的仿真框架

为将RTL模型实现为OpenCL设备做准备

为了使“乘影”GPGPU的硬件设计能够被OpenCL工具链直接调用，开发团队引入了基于Verilator的仿真框架。通过这种方式，可以将RTL（寄存器传输级）模型转换为可以在主机上运行的仿真程序，进而作为OpenCL设备进行测试和验证。这不仅简化了开发流程，还使得硬件设计能够在早期阶段与软件工具链无缝集成。

在原有基于ChiselTest的仿真框架外，额外增加了基于Verilator的仿真测试程序

原有的仿真框架基于ChiselTest，主要用于验证硬件设计的功能正确性。然而，ChiselTest的仿真速度较慢，且难以支持大规模并行仿真。因此，开发团队决定引入基于Verilator的仿真测试程序，以提高仿真的效率和灵活性。Verilator是一个高性能的HDL（硬件描述语言）仿真器，能够将Verilog或SystemVerilog代码编译为C++代码，从而在CPU上高效运行。这种组合不仅保留了ChiselTest的优势，还显著提升了仿真的性能和规模。

精细化控制编译与仿真流程

为了进一步优化仿真过程，开发团队实现了精细化的编译与仿真控制机制。具体来说，该框架能够智能地识别哪些部分需要重新编译和综合，避免不必要的重复工作。例如，当硬件设计中的某些模块没有发生变化时，仿真框架会跳过这些模块的重新编译，直接使用已有的中间结果。这大大缩短了编译时间，提高了开发效率。

避免不必要的编译与综合

通过引入增量编译和缓存机制，仿真框架能够有效地避免不必要的编译与综合步骤。每次修改硬件设计后，框架只会重新编译受影响的部分，而不会对整个设计进行全面编译。此外，框架还会缓存已经编译过的模块，以便在后续仿真中复用。这种优化不仅加快了仿真速度，还减少了资源消耗。

更快的仿真速度：多线程并行仿真

为了进一步提升仿真速度，开发团队利用了多线程并行仿真的技术。通过将仿真任务分配到多个CPU核心上并行执行，仿真框架能够显著减少仿真时间。特别是在处理大规模硬件设计时，多线程并行仿真可以充分发挥现代多核处理器的性能优势。此外，框架还支持分布式仿真，允许用户在多台机器上并行运行仿真任务，进一步加速验证过程。

快速验证，提高调试效率

更快的仿真速度不仅意味着更短的验证周期，还极大地提高了调试效率。开发人员可以在短时间内完成多次仿真，快速定位和修复问题。此外，基于Verilator的仿真框架提供了丰富的调试工具和日志功能，帮助开发人员更好地理解硬件行为。例如，用户可以通过设置断点、查看寄存器状态、跟踪信号变化等方式，深入分析硬件设计的运行情况，从而更快地发现潜在问题。

更大规模的测试用例

基于Verilator的仿真框架不仅支持传统的功能测试，还能处理更大规模的测试用例。这对于验证复杂硬件设计的正确性和稳定性至关重要。开发团队设计了一系列全面的测试用例，涵盖了从基本功能到复杂场景的各种情况。这些测试用例不仅包括单个模块的功能验证，还包括多个模块之间的交互测试，确保整个系统在各种条件下都能正常工作。此外，框架还支持随机测试和压力测试，帮助开发人员发现潜在的边界条件和异常情况。

verilator仿真流程

• Chisel Module：使用Chisel语言编写的硬件模块。这些模块是硬件设计的核心部分，定义了各个功能单元的行为。
• Verilog/SystemVerilog：由Chisel模块生成的Verilog或SystemVerilog代码。这些代码是硬件设计的具体实现，可以直接用于仿真和综合。
• Verilator C++ DUT Module：通过Verilator将Verilog/SystemVerilog代码转换为C++代码生成的DUT（Design Under Test）模块。这些C++代码可以在主机上高效运行，模拟硬件行为。
• Shared Object DUT Module：将C++代码编译为共享对象模块（.so文件）。这些共享对象模块可以在OpenCL应用程序中动态加载和调用，实现硬件仿真的无缝集成。
• C++ Simulation Device Wrapper：用于封装共享对象模块的C++代码。这个包装器提供了一个简单的接口，使得OpenCL应用程序可以方便地调用硬件仿真模块。
• OpenCL Application：使用OpenCL编写的主机应用程序。这个应用程序负责配置和启动仿真任务，并与硬件仿真模块进行交互。它可以通过OpenCL API调用硬件仿真模块，执行各种计算任务。
• Host Executable：编译后的主机应用程序可执行文件。用户可以直接运行这个可执行文件，启动仿真任务并查看结果。
• Call Ventus OpenCL Driver：调用Ventus OpenCL驱动程序的代码。这个驱动程序负责管理硬件仿真模块的加载、初始化和执行，并将仿真结果返回给主机应用程序。
• Application Output：仿真结果输出。仿真结束后，应用程序会生成详细的日志文件和报告，帮助用户分析仿真结果。

总结

基于Verilator的仿真框架为“乘影”GPGPU的开发提供了强大的支持。通过将RTL模型实现为OpenCL设备，开发团队可以更早地进行硬件和软件的协同验证，缩短开发周期。同时，精细化的编译与仿真控制机制、多线程并行仿真技术以及大规模测试用例的支持，使得仿真过程更加高效和可靠。这一框架不仅提高了开发效率，还为复杂硬件设计的验证提供了有力保障。

分页内存管理单元（MMU）设计

目的与特点

“乘影”GPGPU的分页内存管理单元（MMU）旨在支持在GPU上同时运行多个任务，并通过主机端的驱动程序控制内存管理。这一设计使得系统能够高效地管理和分配内存资源，确保不同任务之间的隔离和安全。具体来说，MMU的主要特点包括：

• 多任务支持：通过分页机制，MMU允许多个任务共享同一物理内存，每个任务拥有独立的虚拟地址空间，从而避免了任务之间的干扰。

• 主机端驱动程序控制：内存管理由主机端的驱动程序负责，这意味着所有的内存分配、页面映射和地址转换操作都在主机侧完成，减少了GPU的负担。这种设计简化了硬件实现，并提高了系统的灵活性。

• NVIDIA Unified Memory：类似于NVIDIA的Unified Memory，乘影的MMU也支持自动的页面数据迁移，允许多个GPU和CPU共用相同的地址空间。这不仅简化了编程模型，还提高了内存利用率和访问效率。

• 学术界相关研究：尽管学术界已经对地址转换效率改进和统一内存支持进行了大量研究，但大多数工作仍然集中在仿真器上，如GPGPU Sim和MGPUSim。这些仿真器为研究提供了平台，但实际硬件设计相对较少。乘影的MMU设计填补了这一空白，提供了一个实际的硬件实现。

分页内存管理单元（MMU）是什么，干什么用的？

分页内存

分页内存是指将物理内存划分为固定大小的页面（Page），并通过页表（Page Table）将虚拟地址映射到物理地址。分页机制的主要优点包括：

• 简化内存管理：分页机制使得内存管理更加简单和灵活。操作系统可以将不同的页面分配给不同的进程，而不需要连续的物理内存块。

• 提高内存利用率：通过分页机制，操作系统可以更有效地利用物理内存，避免碎片化问题。未使用的页面可以被标记为空闲，供后续分配使用。

• 支持虚拟内存：分页机制是实现虚拟内存的基础。操作系统可以通过交换技术（Swapping）将不常用的页面移到磁盘上，从而扩大可用的内存空间。

分页内存的申请与释放

分页内存的申请与释放是操作系统和硬件协同工作的结果。具体来说：

• 内存申请：当应用程序需要分配内存时，操作系统会从空闲的物理内存中选择合适的页面，并将它们映射到虚拟地址空间。这个过程涉及到创建新的页表项，并将它们插入到相应的页表中。操作系统还会设置页表项中的权限标志，确保应用程序只能访问自己分配的内存区域。

• 内存释放：当应用程序不再需要某些内存页面时，它可以通过系统调用请求释放这些页面。操作系统会从页表中移除对应的页表项，并将物理页面标记为空闲，供后续分配使用。此外，操作系统还会处理页面的数据回写和同步操作，确保数据的一致性。

MMU的定义与作用

分页内存管理单元（MMU）是计算机系统中的一个重要组件，负责将虚拟地址转换为物理地址。在现代操作系统中，应用程序通常使用虚拟地址来访问内存，而MMU的作用就是将这些虚拟地址映射到实际的物理地址上。MMU的核心功能包括：

• 地址翻译：将虚拟地址转换为物理地址。这是MMU最基本的功能，确保每个进程只能访问自己的内存区域，而不会干扰其他进程。

• 内存保护：通过设置页表项中的权限标志，MMU可以限制某个进程对特定内存区域的访问权限（如只读、可读写、不可访问等）。这有助于防止恶意程序或错误代码破坏系统稳定性。

• 内存隔离：每个进程都有独立的虚拟地址空间，MMU确保不同进程之间不会相互干扰。即使两个进程使用相同的虚拟地址，它们也会被映射到不同的物理地址上。

• 内存分页：将内存划分为固定大小的页面（通常是4KB），并为每个页面分配一个唯一的标识符（页号）。这种方式可以更灵活地管理内存，避免碎片化问题。

NVIDIA Unified Memory是什么，干什么用的？

定义与作用

NVIDIA Unified Memory（统一内存）是一种内存管理技术，允许CPU和GPU共享同一块物理内存。在传统的异构计算系统中，CPU和GPU各自拥有独立的内存空间，数据在两者之间传输需要显式的拷贝操作。而Unified Memory通过自动的页面数据迁移，使得CPU和GPU可以在同一块内存中读取和写入数据，无需手动管理数据传输。这不仅简化了编程模型，还提高了内存利用率和访问效率。

工作原理

Unified Memory的核心思想是将物理内存划分为多个页面，并根据需要动态地将页面迁移到最适合的设备上。具体来说：

• 自动迁移：当CPU或GPU访问某个页面时，如果该页面当前不在本地内存中，系统会自动将其从另一个设备的内存中迁移到本地。这个过程是透明的，应用程序无需关心具体的迁移细节。

• 按需分配：Unified Memory允许系统根据需要动态地分配和回收内存页面，避免了预先分配大量内存的浪费。只有当页面被实际使用时，才会为其分配物理内存。

• 一致性维护：为了确保CPU和GPU看到一致的内存内容，Unified Memory会自动处理页面的脏数据回写和同步操作。当一个设备修改了某个页面的内容时，系统会确保另一个设备也能看到最新的数据。

优势

Unified Memory的主要优势包括：

• 简化编程模型：开发者无需手动管理CPU和GPU之间的数据传输，只需编写一次代码即可在两种设备上运行。这大大降低了编程复杂度，提高了开发效率。

• 提高性能：通过自动的页面数据迁移，Unified Memory可以减少不必要的数据拷贝操作，降低内存带宽占用，提高整体性能。

• 更好的内存利用率：Unified Memory允许CPU和GPU共享同一块物理内存，避免了内存资源的浪费。特别是在多GPU系统中，Unified Memory可以显著提高内存利用率。

地址转换加速是干什么的，乘影是如何进行地址转换加速的？

地址转换加速的定义与作用

地址转换加速是指通过硬件或软件手段，加快虚拟地址到物理地址的转换速度。在现代计算机系统中，地址转换是一个频繁发生的操作，尤其是在多任务环境下，每个进程都需要将自己的虚拟地址转换为物理地址。如果每次地址转换都需要遍历页表，将会带来较大的性能开销。因此，地址转换加速技术的目的是减少页表遍历的次数，提高地址转换的效率。

乘影的地址转换加速设计

乘影的MMU通过两级TLB（Translation Lookaside Buffer）构成的地址缓存系统，实现了高效的地址转换加速。具体来说：

• 一级TLB：位于每个SM（流多处理器）内部，用于快速查找最近使用的地址映射。一级TLB采用全相联结构，尺寸较小但命中率较高。当SM需要访问内存时，首先会在一级TLB中查找虚拟地址到物理地址的映射。如果命中，则直接使用缓存的映射结果，避免了页表遍历的开销。

• 二级TLB：由所有SM共享，用于存储更多的地址映射。二级TLB采用多路组相联结构，相联度和尺寸均可调，以适应不同的应用场景。为了进一步提升效率，二级TLB采用了分块（Banking）和辅助缓存（Auxiliary TLB）等技术。分块将二级TLB划分为多个独立的存储单元，减少了冲突和竞争；辅助缓存则用于存储多级地址转换过程中的中间层页表项（PPN），当主TLB缺失时，辅助TLB可以帮助快速找到正确的页表项，减少页表遍历的时间。

• 页表遍历优化：当TLB缺失时，MMU会启动页表遍历操作，查找虚拟地址对应的物理地址。为了加快页表遍历的速度，乘影的MMU设计了专门的页表遍历单元（PTW），负责逐级遍历页表，并将最终的转换结果回填到TLB中。PTW可以从ASID映射表中获取根页表地址，并每次访存获取连续的多个页表项，减少访存次数。

• 地址转换中间项缓存：在多级页表体系中，地址转换过程中会产生多个中间层的页表项（PPN）。乘影的MMU设计了辅助TLB，用于缓存这些中间层的页表项。当主TLB缺失但辅助TLB命中时，PTW可以直接从中间层开始遍历页表，跳过前面的层级，进一步加快地址转换的速度。

示例说明

假设我们有一个GPGPU应用程序，需要频繁访问内存中的某个数据结构。每当应用程序访问该数据结构时，SM都会发起地址转换请求。如果没有TLB，每次地址转换都需要遍历页表，这将带来较大的性能开销。然而，由于乘影的MMU设计了两级TLB，大部分地址转换请求都可以在TLB中命中，避免了页表遍历的开销。即使出现TLB缺失，PTW也可以快速找到正确的页表项，并将结果回填到TLB中，确保后续的地址转换请求能够命中。通过这种方式，乘影的MMU显著提高了地址转换的效率，减少了内存访问的延迟。

实现RISC-V Sv32多级页表体系

乘影的MMU实现了支持RISC-V Sv32多级页表体系的硬件，具备将虚拟地址转换为物理地址的功能。Sv32是一种32位虚拟地址体系，使用两级页表进行地址转换。每级页表项（PTE）包含一个指向下一級页表或物理页面的指针，以及一些控制标志。Sv32的设计可以扩展为Sv39，以支持64位版本的GPGPU，从而满足未来更高性能的需求。

地址转换加速

为了提高地址转换的效率，乘影的MMU设计了两级TLB（Translation Lookaside Buffer）构成的地址缓存系统。TLB是一种高速缓存，用于存储最近使用的虚拟地址到物理地址的映射，从而减少页表遍历的次数。

内存申请与释放

乘影的MMU为GPGPU设计了支持分页内存申请与释放的模拟内存系统。这个系统允许用户动态地分配和释放内存页面，确保内存资源的高效利用。具体来说：

• 内存申请：当应用程序需要分配内存时，驱动程序会向MMU发送请求，指定所需的页面数量和属性（如读写权限）。MMU会在空闲的物理内存中选择合适的页面，并将其映射到虚拟地址空间。这个过程涉及到创建新的页表项，并将它们插入到相应的页表中。

• 内存释放：当应用程序不再需要某些内存页面时，它可以请求释放这些页面。MMU会从页表中移除对应的页表项，并将物理页面标记为空闲，供后续分配使用。此外，MMU还会处理页面的数据回写和同步操作，确保数据的一致性。

测试框架

为了验证MMU的正确性和性能，开发团队设计了一套完整的测试框架。该框架不仅支持基本的功能测试，还可以模拟复杂的内存分配和释放场景，帮助发现潜在的问题。例如，测试框架可以生成随机的内存请求，模拟多任务并发执行的情况，评估MMU在高负载下的表现。

位置

乘影的MMU设计位于dev-mmu-v2分支，用户可以通过克隆该分支并编译代码来测试和验证MMU的功能。

总结

乘影的分页内存管理单元（MMU）设计旨在支持多任务环境下的高效内存管理，通过主机端的驱动程序控制内存分配和地址转换。MMU实现了RISC-V Sv32多级页表体系，并通过两级TLB和辅助缓存等技术实现了高效的地址转换加速。此外，MMU还支持分页内存的动态申请与释放，确保内存资源的高效利用。NVIDIA的Unified Memory技术通过自动的页面数据迁移，允许CPU和GPU共享同一块物理内存，简化了编程模型并提高了性能。地址转换加速技术则是通过硬件缓存和优化的页表遍历机制，减少了地址转换的开销，提升了系统的整体性能。

虚拟地址和Cache

虚拟内存：Sv32/Sv39内存规范

这两种规范定义了虚拟地址的长度以及相应的页表层次结构，用于将虚拟地址映射到物理地址。以下是具体内容的详细解释：

Sv32 内存规范

• 虚拟地址长度：32位
• 物理地址长度：34位
• 页表层次：两级页表
  • 每一级页表使用10位来索引。
  • 因此，页表的层次结构为10+10。

页表条目大小

• 每个页表条目（PTE）通常占用4字节（32位），这使得每个页表可以包含1024个条目（2^10 = 1024）。
• 第一级页表（PML4）有1024个条目，每个条目指向一个第二级页表。
• 第二级页表（Page Table）也有1024个条目，每个条目指向一个4KB的页面。

虚拟地址分解

• 32位虚拟地址被分为三部分：
  • 高12位：页目录索引（PD Index）
  • 中间10位：页表索引（PT Index）
  • 低10位：页内偏移（Page Offset）
  • 第一级页表（PTE1）：使用虚拟地址的高10位（bit[31:22]）作为索引。
  • 第二级页表（PTE2）：使用虚拟地址的中间10位（bit[21:12]）作为索引。
  • 页内偏移：使用虚拟地址的低12位（bit[11:0]）作为页内偏移。

Sv39 内存规范

• 虚拟地址长度：39位
• 物理地址长度：56位
• 页表层次：三级页表
  • 每一级页表使用9位来索引。
  • 因此，页表的层次结构为9+9+9。

页表条目大小

• 每个页表条目（PTE）通常占用8字节（64位），这使得每个页表可以包含512个条目（2^9 = 512）。
• 第一级页表（PML4）有512个条目，每个条目指向一个第二级页表。
• 第二级页表（Page Directory）也有512个条目，每个条目指向一个第三级页表。
• 第三级页表（Page Table）也有512个条目，每个条目指向一个4KB的页面。

虚拟地址分解

• 39位虚拟地址被分为四部分：
  • 高9位：顶级页表索引（PML4 Index）
  • 中间9位：页目录索引（PD Index）
  • 再中间9位：页表索引（PT Index）
  • 低12位：页内偏移（Page Offset）
  • 第一级页表（PTE1）：使用虚拟地址的高9位（bit[38:30]）作为索引。
  • 第二级页表（PTE2）：使用虚拟地址的中间9位（bit[29:21]）作为索引。
  • 第三级页表（PTE3）：使用虚拟地址的低9位（bit[20:12]）作为索引。
  • 页内偏移：使用虚拟地址的低12位（bit[11:0]）作为页内偏移。

页表项（PTE）格式

每个页表项（Page Table Entry, PTE）通常包含以下信息：

• 物理页面号（PPN, Physical Page Number）：指向物理内存中的页面。
• 权限位：指示该页面的访问权限（如有效位V, 权限位R/W/X等）。
• 缓存属性：指示该页面是否可以缓存。
• 其他标志位：如全局位、脏位、访问位等，用于优化和管理页表。

多级页表映射例子

假设我们有一个虚拟地址0x0000_7fff_f000，我们来看看它是如何通过Sv39页表进行映射的：

1. 第一级页表索引：最高9位为0x000，因此我们访问第一级页表的第0个条目。
2. 第二级页表索引：接下来的9位为0x1ff，所以我们从第一级页表的第0个条目找到对应的第二级页表，并访问它的第511（0x1ff为十进制511）个条目。
3. 第三级页表索引：再接下来的9位为0x3ff，接着我们从第二级页表的第511个条目找到对应的第三级页表，并访问它的第1023（0x3ff为十进制1023）个条目。
4. 页内偏移量：最低12位为0xf000，这告诉我们一旦找到了正确的物理页面，我们还需要在这个页面内的0xf000位置开始读取数据。

总结

• Sv32：适用于32位虚拟地址空间，使用两级页表，每级页表使用10位索引。
• Sv39：适用于39位虚拟地址空间，使用三级页表，每级页表使用9位索引。

这两种内存规范允许RISC-V系统支持不同大小的虚拟地址空间，并通过多级页表有效地管理虚拟地址到物理地址的映射。这对于操作系统和应用程序的内存管理非常重要，尤其是在需要支持大内存或多个进程的系统中。

ITLB和DTLB是什么？

定义

ITLB（Instruction Translation Lookaside Buffer）：用于缓存指令地址的翻译结果。当处理器需要从内存中读取指令时，ITLB会快速提供虚拟地址到物理地址的翻译结果。

DTLB（Data Translation Lookaside Buffer）：用于缓存数据地址的翻译结果。当处理器需要从内存中读取或写入数据时，DTLB会快速提供虚拟地址到物理地址的翻译结果。

工作原理

地址翻译：在虚拟内存系统中，处理器使用虚拟地址访问内存，而实际的物理地址需要通过页表进行翻译。TLB的作用是缓存这些翻译结果，以减少页表查找的开销。

命中与未命中：如果TLB中有对应的翻译结果（命中），可以直接使用；如果TLB中没有对应的翻译结果（未命中），则需要访问页表进行翻译，并将结果存入TLB。

示例

假设有一个程序需要频繁访问某个指令地址和数据地址：

初始状态：

• ITLB和DTLB为空。
• 虚拟地址 0x1000 对应的物理地址是 0x2000。
• 虚拟地址 0x3000 对应的物理地址是 0x4000。

第一次访问：

• 处理器尝试读取指令地址 0x1000：

  • ITLB未命中，访问页表，找到翻译结果 0x1000 -> 0x2000。
  • 将翻译结果存入ITLB。

• 处理器尝试读取数据地址 0x3000：

  • DTLB未命中，访问页表，找到翻译结果 0x3000 -> 0x4000。
  • 将翻译结果存入DTLB。

后续访问：

• 处理器再次尝试读取指令地址 0x1000：ITLB命中，直接使用翻译结果 0x2000。
• 处理器再次尝试读取数据地址 0x3000：DTLB命中，直接使用翻译结果 0x4000。

L2缓存：所有SM共享L2缓存，通过专门的查找表获取ASID对应的根页表地址

查找表工作原理

ASID（Address Space Identifier）：地址空间标识符，用于区分不同进程的虚拟地址空间。每个进程都有一个唯一的ASID。

根页表地址：页表的根节点地址，用于开始地址翻译过程。

工作流程

地址翻译请求：

• 当一个SM需要访问内存时，首先使用虚拟地址和ASID查询L2 TLB。

TLB查找：

• 如果L2 TLB中有对应的翻译结果（命中），直接使用。
• 如果L2 TLB中没有对应的翻译结果（未命中），则需要访问页表进行翻译。

页表访问：

• 使用ASID查找根页表地址。
• 根据根页表地址和虚拟地址进行地址翻译。
• 将翻译结果存入L2 TLB。

示例

假设有一个多线程程序，线程A和线程B分别位于不同的SM上，它们需要访问同一个全局变量data：

初始状态：

• ASID表：

  • 线程A的ASID为1，根页表地址为 0x1000。
  • 线程B的ASID为2，根页表地址为 0x2000。

• L2 TLB为空。

线程A访问data：

• 使用虚拟地址 0x3000 和ASID 1 查询L2 TLB。
• L2 TLB未命中，访问页表 0x1000，找到翻译结果 0x3000 -> 0x4000。
• 将翻译结果存入L2 TLB。

线程B访问data：

• 使用虚拟地址 0x3000 和ASID 2 查询L2 TLB。
• L2 TLB未命中，访问页表 0x2000，找到翻译结果 0x3000 -> 0x5000。
• 将翻译结果存入L2 TLB。

两级TLB构成地址缓存系统

工作原理

地址翻译请求：

• 当处理器需要访问内存时，首先使用虚拟地址（VA）查询L1 TLB。

L1 TLB查找：

• 如果L1 TLB中有对应的翻译结果（命中），直接使用。
• 如果L1 TLB中没有对应的翻译结果（未命中），则查询L2 TLB。

L2 TLB查找：

• 如果L2 TLB中有对应的翻译结果（命中），直接使用，并将结果存入L1 TLB。
• 如果L2 TLB中也没有对应的翻译结果（未命中），则需要访问页表进行翻译。

页表访问：

• 使用根页表地址和虚拟地址进行地址翻译。
• 将翻译结果存入L2 TLB，并将结果存入L1 TLB。

示例

假设有一个程序需要频繁访问某个虚拟地址 0x1000：

初始状态：

• L1 TLB和L2 TLB为空。
• 虚拟地址 0x1000 对应的物理地址是 0x2000。

第一次访问：

处理器尝试访问虚拟地址 0x1000：

• L1 TLB未命中，查询L2 TLB。
• L2 TLB未命中，访问页表，找到翻译结果 0x1000 -> 0x2000。
• 将翻译结果存入L2 TLB，并存入L1 TLB。

后续访问：

• 处理器再次尝试访问虚拟地址 0x1000：
• L1 TLB命中，直接使用翻译结果 0x2000。

全相连是什么结构，尺寸怎么调节的？ASID和VPN是什么？如何并行查找？

全相连结构

定义：全相连（Fully Associative）结构的TLB中，每个条目都可以存储任何虚拟地址的翻译结果。

尺寸调节：可以通过配置TLB的条目数来调节尺寸。例如，可以配置L1 TLB有16个条目，L2 TLB有128个条目。

ASID和VPN

ASID（Address Space Identifier）：地址空间标识符，用于区分不同进程的虚拟地址空间。每个进程都有一个唯一的ASID。

VPN（Virtual Page Number）：虚拟页号，用于标识虚拟地址所在的页。

并行查找

并行查找：在全相连结构的TLB中，可以同时比对所有条目，找到匹配的翻译结果。

查找过程：

• 使用虚拟地址的VPN和ASID同时比对所有条目。
• 如果找到匹配的条目，返回对应的物理地址。
• 如果没有找到匹配的条目，返回未命中。

示例

假设有一个全相连结构的L1 TLB，有4个条目：

ASID VPN PPN V

1 10 20 1

2 11 21 1

1 12 22 1

0 13 23 0

查找虚拟地址 0x1000（ASID=1, VPN=10）：

同时比对所有条目，找到匹配的条目（ASID=1, VPN=10）。

返回对应的物理地址 0x2000。

如何分块和辅助缓存？硬件怎么实现的？如何提升地址转换效率？为什么能提升效率？

分块

定义：将L2 TLB分成多个块（Bank），每个块负责处理一部分虚拟地址的翻译请求。

实现：根据虚拟地址的某些位（如VPN的一部分）决定访问哪个块。

好处：

• 减少冲突：多个请求可以并行处理，减少冲突。
• 提高命中率：每个块可以独立存储翻译结果，提高命中率。

辅助缓存

定义：辅助缓存用于存储多级地址转换过程中的中间层PPN（Physical Page Number，物理页号，用于标识物理地址所在的页。作用：在地址翻译过程中，PPN与页内偏移（Page Offset）组合成完整的物理地址。）。

实现：辅助缓存可以是全相连结构，存储Sv32的一个中间层PPN或Sv39的两个中间层PPN。

好处：

• 减少页表访问次数：如果L2 TLB未命中，但辅助缓存命中，可以直接从中间层PPN开始遍历页表，减少页表访问次数。
• 提高翻译速度：减少页表访问次数，提高地址翻译的速度。

示例

假设有一个L2 TLB，分成4个块，每个块负责处理一部分虚拟地址的翻译请求：

Block ASID VPN PPN V

0 1 10 20 1

1 2 11 21 1

2 1 12 22 1

3 0 13 23 0

查找虚拟地址 0x1000（ASID=1, VPN=10）：

根据VPN的一部分（如高10位）决定访问Block 0。

在Block 0中找到匹配的条目（ASID=1, VPN=10）。

返回对应的物理地址 0x2000。

提升效率的原因

• 减少冲突：分块使得多个请求可以并行处理，减少冲突。
• 提高命中率：每个块可以独立存储翻译结果，提高命中率。
• 减少页表访问次数：辅助缓存存储中间层PPN，减少页表访问次数，提高翻译速度。

总结

• 两级TLB：通过L1 TLB和L2 TLB的多层次缓存，减少页表查找的开销。
• 全相连结构：支持ASID和VPN的并行查找，提高查找速度。
• 分块和辅助缓存：通过分块减少冲突，通过辅助缓存减少页表访问次数，提高地址翻译的效率。

获取数据流程：L1 Cache，L1 TLB，L2 Cache，L2 TLB，页表等之间的工作和协同

初始化状态

假设系统初始状态下，L1 Cache、L1 TLB、L2 Cache和L2 TLB都是空的。

地址翻译请求

当一个SM（Streaming Multiprocessor）需要访问内存中的某个数据时，会发起一个地址翻译请求。假设虚拟地址（VA）为 0x1000，ASID（Address Space Identifier）为 1。

访问L1 TLB

查询L1 TLB：

• 使用虚拟地址 0x1000 和ASID 1 查询L1 TLB。
• 如果L1 TLB中有对应的翻译结果（命中），直接使用翻译结果进行地址转换，跳转到步骤7。
• 如果L1 TLB中没有对应的翻译结果（未命中），继续查询L2 TLB。

访问L2 TLB

查询L2 TLB：

• 使用虚拟地址 0x1000 和ASID 1 查询L2 TLB。
• 如果L2 TLB中有对应的翻译结果（命中），直接使用翻译结果进行地址转换，并将结果存入L1 TLB，跳转到步骤7。
• 如果L2 TLB中没有对应的翻译结果（未命中），继续访问页表。

访问页表

访问页表：

• 使用ASID 1 查找根页表地址。
• 根据根页表地址和虚拟地址 0x1000 进行地址翻译。
• 找到对应的物理地址（PA），假设为 0x2000。
• 将翻译结果 0x1000 -> 0x2000 存入L2 TLB，并将结果存入L1 TLB。

访问L1 Cache

查询L1 Cache：

• 使用物理地址 0x2000 查询L1 Cache。
• 如果L1 Cache中有对应的数据（命中），直接使用数据。
• 如果L1 Cache中没有对应的数据（未命中），继续查询L2 Cache。

访问L2 Cache

查询L2 Cache：

• 使用物理地址 0x2000 查询L2 Cache。
• 如果L2 Cache中有对应的数据（命中），将数据加载到L1 Cache，并返回给处理器。
• 如果L2 Cache中没有对应的数据（未命中），继续访问内存。

访问内存

访问内存：

• 使用物理地址 0x2000 访问内存，读取数据。
• 将读取的数据加载到L2 Cache，并将数据加载到L1 Cache。
• 返回数据给处理器。

流程总结

• 地址翻译请求：处理器发起地址翻译请求，使用虚拟地址 0x1000 和ASID 1。
• 查询L1 TLB：查询L1 TLB，如果没有命中，继续查询L2 TLB。
• 查询L2 TLB：查询L2 TLB，如果没有命中，访问页表进行地址翻译。
• 访问页表：使用ASID查找根页表地址，进行地址翻译，将结果存入L2 TLB和L1 TLB。
• 查询L1 Cache：使用物理地址查询L1 Cache，如果没有命中，继续查询L2 Cache。
• 查询L2 Cache：使用物理地址查询L2 Cache，如果没有命中，访问内存读取数据。
• 访问内存：使用物理地址访问内存，读取数据，将数据加载到L2 Cache和L1 Cache。
• 返回数据：将数据返回给处理器。

通过这种多层次的缓存和地址翻译机制，可以显著减少内存访问的延迟，提高系统的整体性能。

乱序机制：支持指令级并行（ILP），但通过RCC机制减少对线程间同步（TLP）的影响

指令级并行（ILP）

定义：指令级并行是指处理器在同一时钟周期内并行执行多条指令的能力。通过乱序执行（Out-of-Order Execution）和分支预测等技术，可以提高指令的并行度。

实现：

乱序执行：处理器可以重新排序指令，以便更有效地利用硬件资源。

分支预测：预测分支指令的结果，提前执行后续指令，减少分支延迟。

通过RCC机制减少对线程间同步（TLP）的影响

定义：线程级并行（TLP）是指多个线程并行执行的能力。在GPGPU中，线程间的同步操作（如内存屏障）会影响性能。

RCC机制：释放连贯性指导的缓存一致性（RCC）通过特定的微架构操作（如FENCE指令和.acq/.rl标识符）确保内存操作的顺序性，同时减少对线程间同步的影响。

示例

假设有一个多线程程序，线程A和线程B分别执行以下操作：

// 线程A

int x = 0;

int y = 0;

void threadA() \{

    x = 1;  // 写操作1
    
    asm volatile("fence w,w" ::: "memory");  // 确保写操作1在写操作2之前完成
    
    y = 1;  // 写操作2

\}

// 线程B

int r1, r2;

void threadB() \{

    r1 = y;  // 读操作1
    
    asm volatile("fence r,r" ::: "memory");  // 确保读操作1在读操作2之前完成
    
    r2 = x;  // 读操作2

\}

解释

线程A：

• 修改全局变量x。
• 使用fence w,w指令确保所有之前的写操作在设置标志位y之前完成。
• 修改全局变量y。

线程B：

• 读取全局变量y。
• 使用fence r,r指令确保所有之前的读操作在读取标志位x之前完成。
• 读取全局变量x。

乱序执行和RCC机制

乱序执行：

处理器可以重新排序线程A中的写操作，但fence w,w指令确保x = 1在y = 1之前完成。

处理器可以重新排序线程B中的读操作，但fence r,r指令确保r1 = y在r2 = x之前完成。

RCC机制：

通过.aq和.rl标识符，确保线程A的写操作在写回L2缓存之前完成，线程B的读操作在读取L2缓存之后完成。

这样可以减少线程间的同步开销，提高性能。

总结

ITLB和DTLB：用于缓存指令和数据地址的翻译结果，减少页表查找的开销。

L2缓存：所有SM共享L2缓存，通过ASID查找根页表地址，确保地址翻译的高效性。

乱序机制和RCC：通过乱序执行和特定的同步指令，提高指令级并行度，同时减少线程间同步的开销。

虚拟内存设计

常见的缓存组织方式

在计算机体系结构中，缓存（Cache）的组织方式对性能和硬件复杂度有着重要影响。常见的缓存组织方式包括直接映射（Direct-Mapped）、全相连（Fully Associative）和组相连（Set-Associative）。

直接映射缓存（Direct-Mapped Cache）

定义

直接映射缓存是一种最简单的缓存组织方式。在这种结构中，内存中的每个块只能映射到缓存中的一个固定位置。具体来说，内存地址通过哈希函数（通常是地址的低位部分）确定其映射到缓存中的哪个行（Line）。这意味着每个内存块只能存储在缓存的一个特定位置，无法选择其他位置。

工作原理

• 分块映射：内存被划分为多个固定大小的块（Block），每个块的大小通常与缓存行的大小相同。内存地址通过一定的哈希函数（通常是地址的低位部分）确定其映射到缓存中的哪个行。
• 查找过程：当处理器需要访问某个内存地址时，缓存控制器首先根据地址的低位部分确定该地址应该映射到缓存中的哪个行，然后将地址的标签（Tag）与该行的标签进行比较。如果标签匹配，则表示缓存命中（Hit），可以直接从缓存中读取数据；如果标签不匹配，则表示缓存未命中（Miss），需要从主存中加载数据到缓存中。
• 替换策略：由于每个内存块只能映射到缓存中的一个固定位置，因此当缓存未命中且该位置已经被占用时，必须替换掉原有的缓存行。直接映射缓存的替换策略非常简单，因为每次只有一个行可以被替换。

优点

• 硬件实现简单：直接映射缓存的硬件实现最为简单，因为它只需要一个标签比较器和少量的控制逻辑。这使得它适合用于对硬件复杂度要求较低的场景。
• 查找时间短：由于每次访问只需要比较一个缓存行的标签，直接映射缓存的查找时间非常短，尤其是在缓存容量较大时，查找延迟几乎可以忽略不计。
• 功耗低：由于硬件复杂度较低，直接映射缓存的功耗也相对较低，适合用于对功耗敏感的应用场景。

缺点

• 命中率低：直接映射缓存的命中率通常较低，因为在某些情况下，多个内存块可能会映射到同一个缓存行，导致冲突。这种冲突会导致频繁的缓存未命中，进而影响性能。
• 容易发生冲突：由于每个内存块只能映射到缓存中的一个固定位置，因此不同程序或不同时间段的内存访问可能会导致相同的缓存行被频繁替换，进一步降低命中率。
• 不适合随机访问模式：直接映射缓存对内存访问模式较为敏感，尤其是对于随机访问模式，其性能表现较差。相比之下，全相连和组相连缓存能够更好地适应随机访问模式。

应用场景

直接映射缓存通常用于对硬件复杂度和功耗要求较高的场景，例如嵌入式系统、物联网设备等。此外，在某些对缓存容量要求不高且内存访问模式较为规律的场景中，直接映射缓存也可以提供较好的性能。

---

全相连（Fully Associative）

定义

全相连缓存是指每个缓存行可以存储来自内存中的任意一个块（Block）。换句话说，内存中的任何一个块都可以被映射到缓存中的任何一个位置。这种设计的最大特点是灵活性极高，因为没有固定的映射规则，任何内存块都可以放在缓存的任何位置。

工作原理

• 查找过程：当处理器需要访问某个内存地址时，缓存控制器会将该地址的标签（Tag）与缓存中所有行的标签进行比较。如果找到匹配的标签，则表示缓存命中（Hit），可以直接从缓存中读取数据；如果没有找到匹配的标签，则表示缓存未命中（Miss），需要从主存中加载数据到缓存中。
• 替换策略：由于全相连缓存中的每个行都可以存储任意一个内存块，因此当缓存满时，必须选择一个行进行替换。常用的替换策略包括最近最少使用（LRU, Least Recently Used）、**随机替换（Random Replacement）**等。

优点

• 高命中率：由于每个内存块可以映射到缓存中的任何位置，全相连缓存的命中率通常比其他类型的缓存更高。特别是在内存访问模式较为随机的情况下，全相连缓存的表现尤为出色。
• 灵活性强：全相连缓存的设计非常灵活，适用于各种不同的内存访问模式，尤其是在多任务环境下，能够更好地适应不同进程的内存需求。

缺点

• 硬件复杂度高：全相连缓存的缺点是硬件实现较为复杂，因为每次访问都需要同时比较所有缓存行的标签，这需要大量的比较器和逻辑电路，增加了硬件成本和功耗。
• 查找时间长：由于需要同时比较所有缓存行的标签，全相连缓存的查找时间相对较长，尤其是在缓存容量较大时，查找延迟会显著增加。
• 不适合大规模缓存：由于硬件复杂度和查找时间的问题，全相连缓存通常只适用于小规模缓存（如TLB或L1缓存），而不适合用于大容量的L2或L3缓存。

应用场景

全相连缓存通常用于对性能要求极高的场景，例如TLB（Translation Lookaside Buffer）和L1指令缓存。在这些场景中，缓存的容量较小，且对命中率的要求较高，因此全相连结构的优势得以体现。

---

组相连（Set-Associative）

定义

组相连缓存是介于直接映射和全相连之间的一种折衷方案。它将缓存分为多个组（Set），每个组内部有多个行（Way）。内存中的每个块只能映射到特定的一个组，但可以在该组内的任意一行中存储。组相连缓存的灵活性介于直接映射和全相连之间，既保持了一定的灵活性，又避免了全相连缓存的高硬件复杂度。

工作原理

• 分组映射：内存地址通过一定的哈希函数（通常是地址的低位部分）确定其映射到哪个组。每个组内有多个行，内存块可以在该组内的任意一行中存储。
• 查找过程：当处理器需要访问某个内存地址时，缓存控制器首先根据地址的低位部分确定该地址应该映射到哪个组，然后将地址的标签与该组内所有行的标签进行比较。如果找到匹配的标签，则表示缓存命中；如果没有找到匹配的标签，则表示缓存未命中。
• 替换策略：当某个组内的所有行都被占用时，必须选择一个行进行替换。常用的替换策略包括LRU（Least Recently Used）、**FIFO（First In, First Out）**等。

优点

• 平衡命中率和硬件复杂度：组相连缓存的命中率通常高于直接映射缓存，因为它允许每个组内的多个行存储同一个内存块。同时，它的硬件复杂度低于全相连缓存，因为只需要在每个组内进行标签比较，而不是在整个缓存中进行全局比较。
• 适合大规模缓存：由于组相连缓存的硬件实现相对简单，且查找时间较短，因此它非常适合用于大容量的L2或L3缓存。通过合理选择组的数量和每组的行数，可以在命中率和硬件复杂度之间取得良好的平衡。
• 减少冲突：相比于直接映射缓存，组相连缓存减少了内存块之间的冲突，因为即使两个内存块映射到同一个组，它们仍然可以在该组内的不同行中存储。

缺点

• 命中率不如全相连缓存：虽然组相连缓存的命中率高于直接映射缓存，但它仍然无法达到全相连缓存的水平，因为在某些情况下，多个内存块可能会映射到同一个组，导致冲突。
• 硬件复杂度高于直接映射缓存：相比于直接映射缓存，组相连缓存的硬件实现更为复杂，因为每个组内的多个行都需要进行标签比较。随着组的数量和每组的行数增加，硬件复杂度也会相应增加。

应用场景

组相连缓存广泛应用于现代处理器的各级缓存中，尤其是L2缓存和L3缓存。通过合理选择组的数量和每组的行数，可以在命中率和硬件复杂度之间取得良好的平衡。例如，常见的4路组相连（4-Way Set-Associative）缓存意味着每个组内有4个行，而8路组相连（8-Way Set-Associative）缓存则意味着每个组内有8个行。

---

直接映射 vs. 全相连 vs. 组相连

比较维度

特性	直接映射缓存	全相连缓存	组相连缓存
命中率	较低，容易发生冲突	高，接近最优	较高，取决于组的数量和每组的行数
硬件复杂度	最低，只需一个标签比较器	最高，需要全局标签比较	中等，只需在组内进行标签比较
查找时间	最短，每次只比较一个缓存行的标签	较长，尤其在大容量缓存时	较短，适合大规模缓存
适用场景	对硬件复杂度和功耗要求较高的场景	小规模缓存（如TLB、L1缓存）	大规模缓存（如L2、L3缓存）
灵活性	最低，内存块只能映射到固定位置	极高，内存块可以映射到任意位置	中等，内存块只能映射到特定组
冲突概率	较高，容易发生冲突	低，几乎不会发生冲突	较低，但仍可能发生冲突
替换策略	简单，每次只有一个行可以被替换	复杂，需要选择多个行中的一个进行替换	中等，每个组内的多个行可以选择替换

详细分析

1. 命中率：

   • 直接映射缓存：由于每个内存块只能映射到缓存中的一个固定位置，容易发生冲突，导致命中率较低。特别是在内存访问模式较为随机的情况下，直接映射缓存的表现较差。
   • 全相连缓存：全相连缓存的命中率最高，因为它允许每个内存块映射到缓存中的任何位置，几乎不会发生冲突。然而，随着缓存容量的增加，全相连缓存的硬件复杂度和查找时间也会显著增加，限制了其应用范围。
   • 组相连缓存：组相连缓存的命中率介于直接映射和全相连之间。通过合理选择组的数量和每组的行数，可以在命中率和硬件复杂度之间取得良好的平衡。组相连缓存能够在减少冲突的同时，保持较低的硬件复杂度和查找时间。

2. 硬件复杂度：

   • 直接映射缓存：硬件实现最为简单，只需要一个标签比较器和少量的控制逻辑。这使得直接映射缓存适合用于对硬件复杂度和功耗要求较高的场景。
   • 全相连缓存：硬件实现最为复杂，因为每次访问都需要同时比较所有缓存行的标签。这需要大量的比较器和逻辑电路，增加了硬件成本和功耗。因此，全相连缓存通常只适用于小规模缓存。
   • 组相连缓存：硬件复杂度介于直接映射和全相连之间。每个组内的多个行需要进行标签比较，但不需要全局比较，因此硬件实现相对简单。组相连缓存适合用于大容量的L2或L3缓存。

3. 查找时间：

   • 直接映射缓存：查找时间最短，因为每次访问只需要比较一个缓存行的标签。这使得直接映射缓存适合用于对查找时间要求较高的场景。
   • 全相连缓存：查找时间较长，尤其是在缓存容量较大时，需要同时比较所有缓存行的标签。这限制了全相连缓存的应用范围。
   • 组相连缓存：查找时间介于直接映射和全相连之间。通过合理选择组的数量和每组的行数，可以在查找时间和硬件复杂度之间取得良好的平衡。

4. 适用场景：

   • 直接映射缓存：适合用于对硬件复杂度和功耗要求较高的场景，例如嵌入式系统、物联网设备等。此外，在某些对缓存容量要求不高且内存访问模式较为规律的场景中，直接映射缓存也可以提供较好的性能。
   • 全相连缓存：适合用于对命中率要求极高的小规模缓存，例如TLB（Translation Lookaside Buffer）和L1指令缓存。在这些场景中，缓存的容量较小，且对命中率的要求较高，因此全相连结构的优势得以体现。
   • 组相连缓存：广泛应用于现代处理器的各级缓存中，尤其是L2缓存和L3缓存。通过合理选择组的数量和每组的行数，可以在命中率和硬件复杂度之间取得良好的平衡。

5. 灵活性：

   • 直接映射缓存：灵活性最低，因为每个内存块只能映射到缓存中的一个固定位置。这使得直接映射缓存在处理随机访问模式时表现较差。
   • 全相连缓存：灵活性最高，因为每个内存块可以映射到缓存中的任何位置。这使得全相连缓存能够适应各种不同的内存访问模式，尤其是在多任务环境下表现出色。
   • 组相连缓存：灵活性介于直接映射和全相连之间。虽然每个内存块只能映射到特定的组，但在组内的多个行中仍然有一定的灵活性。

6. 冲突概率：

   • 直接映射缓存：冲突概率较高，因为在某些情况下，多个内存块可能会映射到同一个缓存行，导致频繁的缓存未命中。
   • 全相连缓存：冲突概率最低，因为每个内存块可以映射到缓存中的任何位置，几乎不会发生冲突。
   • 组相连缓存：冲突概率介于直接映射和全相连之间。虽然每个内存块只能映射到特定的组，但在组内的多个行中仍然有一定的灵活性，减少了冲突的发生。

7. 替换策略：

   • 直接映射缓存：替换策略非常简单，因为每次只有一个行可以被替换。这使得直接映射缓存的硬件实现更为简单。
   • 全相连缓存：替换策略较为复杂，因为需要选择多个行中的一个进行替换。常用的替换策略包括LRU（Least Recently Used）、Random Replacement等。
   • 组相连缓存：替换策略介于直接映射和全相连之间。每个组内的多个行可以选择替换，常用的替换策略包括LRU、FIFO（First In, First Out）等。

---

一次内存访问示意图

注意事项

• TLB采用组相联
• 页表采用两级页表
• cache采用组相联
• cache仅考虑L1 d-cache，不考虑L1 i-cache、L2 cache和L3 cache
• 未考虑页表缺页
• 简化了cache未命中情况

实际例子

下面展示了现代Intel处理器的CPU cache是如何组织的。

L1 cache – 32KB，8路组相联，64字节缓存线。

由索引拣选缓存组（行）

在cache中的数据是以缓存线（line）为单位组织的，一条缓存线对应于内存中一个连续的字节块。这个cache使用了64字节的缓存线。这些线被保存在cache bank中，也叫路（way）。每一路都有一个专门的目录（directory）用来保存一些登记信息。你可以把每一路连同它的目录想象成电子表格中的一列，而表的一行构成了cache的一组（set）。列中的每一个单元（cell）都含有一条缓存线，由与之对应的目录单元跟踪管理。图中的cache有64 组、每组8路，因此有512个含有缓存线的单元，合计32KB的存储空间。

在cache眼中，物理内存被分割成了许多4KB大小的物理内存页（page）（每一路是一个page？）。每一页都含有4KB / 64 bytes == 64条缓存线。在一个4KB的页中，第0到63字节是第一条缓存线，第64到127字节是第二条缓存线，以此类推。每一页都重复着这种划分，所以第0页第3条缓存线与第1页第3条缓存线是不同的。

**在全相联缓存（fully associative cache）中，内存中的任意一条缓存线都可以被存储到任意的缓存单元中。**这种存储方式十分灵活，但也使得要访问它们时，检索缓存单元的工作变得复杂、昂贵。由于L1和L2 cache工作在很强的约束之下，包括功耗，芯片物理空间，存取速度等，所以在多数情况下，使用全相联缓存并不是一个很好的折中。

取而代之的是图中的组相联缓存（set associative cache）。意思是，内存中一条给定的缓存线只能被保存在一个特定的组（或行）中。所以，任意物理内存页的第0条缓存线（页内第0到63字节）必须存储到第0组，第1条缓存线存储到第1组，以此类推。每一组有8个单元可用于存储它所关联的缓存线，从而形成一个8路关联的组（8-way associative set）。当访问一个内存地址时，地址的第6到11位（译注：组索引，因为有64个组，所以6bit索引）指出了在4KB内存页中缓存线的编号，从而决定了即将使用的缓存组。举例来说，物理地址0x800010a0的组索引是000010，所以此地址的内容一定是在第2组中缓存的。

但是还有一个问题，就是要找出一组中哪个单元包含了想要的信息，如果有的话。这就到了缓存目录登场的时刻。每一个缓存线都被其对应的目录单元做了标记（tag）；这个标记就是一个简单的内存页编号，指出缓存线来自于哪一页。由于处理器可以寻址64GB的物理RAM，所以总共有64GB / 4KB == 224个内存页，需要24位来保存标记。前例中的物理地址0x800010a0对应的页号为524,289。下面是故事的后一半：

​

在组中搜索匹配标记

由于我们只需要去查看某一组中的8路，所以查找匹配标记是非常迅速的；事实上，从电学角度讲，所有的标记是同时进行比对的，我用箭头来表示这一点。如果此时正好有一条具有匹配标签的有效缓存线，我们就获得一次缓存命中（cache hit）。否则，这个请求就会被转发的L2 cache，如果还没匹配上就再转发给主系统内存。通过应用各种调节尺寸和容量的技术，Intel给CPU配置了较大的L2 cache，但其基本的设计都是相同的。比如，你可以将原先的缓存增加8路而获得一个64KB的缓存；再将组数增加到4096，每路可以存储256KB。经过这两次修改，就得到了一个4MB的L2 cache。在此情况下，需要18位来保存标记，12位保存组索引；缓存所使用的物理内存页的大小与其一路的大小相等。（译注：有4096组，就需要lg(4096)12位的组索引，缓存线依然是64字节，所以一路有4096*64B256KB字节；在L2 cache眼中，内存被分割为许多256KB的块，所以需要lg(64GB/256KB)==18位来保存标记。）

如果有一组已经被放满了，那么在另一条缓存线被存储进来之前，已有的某一条则必须被腾空（evict）。为了避免这种情况，对运算速度要求较高的程序就要尝试仔细组织它的数据，使得内存访问均匀的分布在已有的缓存线上。举例来说，假设程序中有一个数组，元素的大小是512字节，其中一些对象在内存中相距4KB。这些对象的各个字段都落在同一缓存线上，并竞争同一缓存组。如果程序频繁的访问一个给定的字段（比如，通过虚函数表vtable调用虚函数），那么这个组看起来就好像一直是被填满的，缓存开始变得毫无意义，因为缓存线一直在重复着腾空与重新载入的步骤。在我们的例子中，由于组数的限制，L1 cache仅能保存8个这类对象的虚函数表。这就是组相联策略的折中所付出的代价：即使在整体缓存的使用率并不高的情况下，由于组冲突，我们还是会遇到缓存缺失的情况。然而，鉴于计算机中各个存储层次的相对速度，不管怎么说，大部分的应用程序并不必为此而担心。

一个内存访问经常由一个线性（或虚拟）地址发起，所以L1 cache需要依赖分页单元（paging unit）来求出物理内存页的地址，以便用于缓存标记。与此相反，组索引来自于线性地址的低位，所以不需要转换就可以使用了（在我们的例子中为第6到11位）。因此L1 cache是物理标记但虚拟索引的（physically tagged but virtually indexed），从而帮助CPU进行并行的查找操作。因为L1 cache的一路绝不会比MMU的一页还大，所以可以保证一个给定的物理地址位置总是关联到同一组，即使组索引是虚拟的。在另一方面L2 cache必须是物理标记和物理索引的，因为它的一路比MMU的一页要大。但是，当一个请求到达L2 cache时，物理地址已经被L1 cache准备（resolved）完毕了，所以L2 cache会工作得很好。

最后，目录单元还存储了对应缓存线的状态（state）。在L1代码缓存中的一条缓存线要么是无效的（invalid）要么是共享的（shared，意思是有效的，真的J）。在L1数据缓存和L2缓存中，一条缓存线可以为4个MESI状态之一：被修改的（modified），独占的（exclusive），共享的（shared），无效的（invalid）。Intel缓存是包容式的（inclusive）：L1缓存的内容会被复制到L2缓存中。

总结

1. 内存层次结构的意义在于利用引用的空间局部性和时间局部性原理，将经常被访问的数据放到快速的存储器中，而将不经常访问的数据留在较慢的存储器中。
2. 一般情况下，除了寄存器和L1缓存可以操作指定字长的数据，下层的内存子系统就不会再使用这么小的单位了，而是直接移动数据块，比如以缓存线为单位访问数据。
3. 对于组冲突，可以这么理解：与上文相似，假设一个缓存，由512条缓存线组成，每条线64字节，容量32KB。
   1. 假如它是直接映射缓存，由于它往往使用地址的低位直接映射缓存线编号，所以所有的32K倍数的地址（32K，64K，96K等）都会映射到同一条线上（即第0线）。假如程序的内存组织不当，交替的去访问布置在这些地址的数据，则会导致冲突。从外表看来就好像缓存只有1条线了，尽管其他缓存线一直是空闲着的。
   2. 如果是全相联缓存，那么每条缓存线都是独立的，可以对应于内存中的任意缓存线。只有当所有的512条缓存线都被占满后才会出现冲突。
   3. 组相联是前两者的折中，每一路中的缓存线采用直接映射方式，而在路与路之间，缓存控制器使用全相联映射算法，决定选择一组中的哪一条线。
   4. 如果是2路组相联缓存，那么这512条缓存线就被分为了2路，每路256条线，一路16KB。此时所有为16K整数倍的地址（16K，32K，48K等）都会映射到第0线，但由于2路是关联的，所以可以同时有2个这种地址的内容被缓存，不会发生冲突。当然了，如果要访问第三个这种地址，还是要先腾空已有的一条才行。所以极端情况下，从外表看来就好像缓存只有2条线了，尽管其他缓存线一直是空闲着的。
   5. 如果是8路组相联缓存（与文中示例相同），那么这512条缓存线就被分为了8路，每路64条线，一路4KB。所以如果数组中元素地址是4K对齐的，并且程序交替的访问这些元素，就会出现组冲突。从外表看来就好像缓存只有8条线了，尽管其他缓存线一直是空闲着的。

总结

• 直接映射缓存：硬件实现最简单，查找时间最短，但命中率较低，容易发生冲突。适合用于对硬件复杂度和功耗要求较高的场景，例如嵌入式系统和物联网设备。
• 全相连缓存：具有最高的命中率和最大的灵活性，但硬件复杂度和查找时间较高，适合用于对命中率要求极高的小规模缓存，例如TLB和L1指令缓存。
• 组相连缓存：是全相连和直接映射之间的一种折衷方案，能够在命中率和硬件复杂度之间取得良好的平衡，广泛应用于现代处理器的各级缓存中，尤其是L2和L3缓存。

在实际应用中，选择哪种缓存组织方式取决于具体的应用场景和性能需求。对于对硬件复杂度和功耗要求较高的场景，直接映射缓存可能是更好的选择；而对于对命中率要求较高的小规模缓存，全相连缓存则更为合适；而对于大容量的L2或L3缓存，组相连缓存能够在命中率和硬件复杂度之间取得良好的平衡。

标签（Tag）和索引（Index）

在缓存（Cache）系统中，**标签（Tag）和索引（Index）**是用于定位和识别缓存行（Cache Line）的两个关键部分。它们是从内存地址中提取出来的，帮助缓存控制器确定某个内存地址是否已经存在于缓存中，以及如果存在的话，具体位于哪个缓存行。下面将详细解释这两个概念，并说明它们在不同类型的缓存组织方式中的作用。

我们一共有8行cache line，cache line大小是8 Bytes。所以我们可以利用地址低3 bits（如上图地址蓝色部分）用来寻址8 bytes中某一字节，我们称这部分bit组合为offset。同理，8行cache line，为了覆盖所有行。我们需要3 bits（如上图地址黄色部分）查找某一行，这部分地址部分称之为index。现在我们知道，如果两个不同的地址，其地址的bit3-bit5如果完全一样的话，那么这两个地址经过硬件散列之后都会找到同一个cache line。所以，当我们找到cache line之后，只代表我们访问的地址对应的数据可能存在这个cache line中，但是也有可能是其他地址对应的数据。所以，我们又引入tag array区域，tag array和data array一一对应。每一个cache line都对应唯一一个tag，tag中保存的是整个地址位宽去除index和offset使用的bit剩余部分（如上图地址绿色部分）。tag、index和offset三者组合就可以唯一确定一个地址了。因此，当我们根据地址中index位找到cache line后，取出当前cache line对应的tag，然后和地址中的tag进行比较，如果相等，这说明cache命中。如果不相等，说明当前cache line存储的是其他地址的数据，这就是cache缺失。在上述图中，我们看到tag的值是0x19，和地址中的tag部分相等，因此在本次访问会命中。由于tag的引入，因此解答了我们之前的一个疑问“为什么硬件cache line不做成一个字节？”。这样会导致硬件成本的上升，因为原本8个字节对应一个tag，现在需要8个tag，占用了很多内存。tag也是cache的一部分，但是我们谈到cache size的时候并不考虑tag占用的内存部分。

Cache的基本原理 - 知乎。

标签（Tag）

定义

**标签（Tag）**是内存地址的一部分，用于唯一标识缓存行中的数据来源。它表示该缓存行中的数据来自内存中的哪个块（Block）。当处理器访问某个内存地址时，缓存控制器会将该地址的标签与缓存行中的标签进行比较，以确定是否存在匹配的缓存行。如果标签匹配，则表示缓存命中（Hit），可以直接从缓存中读取数据；如果标签不匹配，则表示缓存未命中（Miss），需要从主存中加载数据到缓存中。

作用

• 唯一标识：标签用于区分不同的内存块。即使两个内存块映射到同一个缓存行，它们的标签也会不同，因此可以通过标签来区分它们。
• 缓存命中判断：当处理器访问某个内存地址时，缓存控制器会将该地址的标签与缓存行中的标签进行比较，以判断是否发生了缓存命中。

计算方法

标签是从内存地址的高位部分提取出来的。具体来说，内存地址可以分为三个部分：

• 标签（Tag）：用于标识内存块的唯一性。
• 索引（Index）：用于确定缓存行的位置。
• 偏移（Offset）：用于确定块内的具体字节位置。

假设内存地址为32位，缓存行大小为64字节（即6位偏移），缓存有1024行（即10位索引），那么标签的长度为：

标签长度 = 总地址位数 - 索引位数 - 偏移位数
         = 32 - 10 - 6
         = 16位

索引（Index）

定义

**索引（Index）**是内存地址的另一部分，用于确定缓存行的具体位置。它告诉缓存控制器应该查找哪个缓存组或缓存行。在不同的缓存组织方式中，索引的作用略有不同：

• 直接映射缓存：索引用于直接确定缓存行的位置。每个内存块只能映射到缓存中的一个固定位置，因此索引直接决定了该块应该存储在哪一行。
• 组相连缓存：索引用于确定缓存组的位置。每个内存块可以映射到特定的一个组，但可以在该组内的多个行中存储。索引决定了该块应该映射到哪个组，而组内的具体行则由标签来进一步确定。
• 全相连缓存：在全相连缓存中，索引的概念不存在，因为每个内存块可以映射到缓存中的任何位置。所有缓存行都需要进行标签比较，以确定是否存在匹配的缓存行。

作用

• 定位缓存行：索引用于确定缓存行或缓存组的具体位置。它告诉缓存控制器应该查找哪个缓存行或缓存组。
• 减少查找范围：通过使用索引，缓存控制器可以缩小查找范围，减少标签比较的数量，从而提高查找效率。

计算方法

索引是从内存地址的中间部分提取出来的。它的长度取决于缓存的容量和组织方式。假设缓存有1024行，那么索引的长度为10位（因为2^10 = 1024）。具体来说，内存地址可以分为三个部分：

• 标签（Tag）：用于标识内存块的唯一性。
• 索引（Index）：用于确定缓存行的位置。
• 偏移（Offset）：用于确定块内的具体字节位置。

假设内存地址为32位，缓存行大小为64字节（即6位偏移），缓存有1024行（即10位索引），那么索引的长度为10位。

偏移（Offset）

为了完整性，这里也简单介绍一下偏移（Offset）。偏移是内存地址的最低位部分，用于确定缓存行内具体字节的位置。它表示在缓存行中，数据的具体字节偏移量。偏移的长度取决于缓存行的大小。例如，如果缓存行大小为64字节，那么偏移的长度为6位（因为2^6 = 64）。

作用

• 定位字节位置：偏移用于确定缓存行内具体字节的位置。它告诉缓存控制器应该从缓存行的哪个位置读取或写入数据。

总结：标签、索引和偏移的关系

假设我们有一个32位的内存地址，缓存行大小为64字节，缓存有1024行，那么内存地址可以分为三部分：

• 标签（Tag）：16位，用于唯一标识内存块。
• 索引（Index）：10位，用于确定缓存行或缓存组的位置。
• 偏移（Offset）：6位，用于确定缓存行内具体字节的位置。

不同缓存组织方式中的标签和索引

直接映射缓存

• 索引：用于直接确定缓存行的位置。
• 标签：用于验证缓存行中的数据是否来自正确的内存块。
• 查找过程：根据索引找到对应的缓存行，然后将地址的标签与该行的标签进行比较。如果标签匹配，则表示缓存命中；否则，表示缓存未命中。

组相连缓存

• 索引：用于确定缓存组的位置。
• 标签：用于验证缓存组内的多个行中的哪一个行包含正确的数据。
• 查找过程：根据索引找到对应的缓存组，然后将地址的标签与该组内所有行的标签进行比较。如果找到匹配的标签，则表示缓存命中；否则，表示缓存未命中。

全相连缓存

• 索引：不存在，因为每个内存块可以映射到缓存中的任何位置。
• 标签：用于验证所有缓存行中的哪一个行包含正确的数据。
• 查找过程：将地址的标签与所有缓存行的标签进行比较。如果找到匹配的标签，则表示缓存命中；否则，表示缓存未命中。

总结

• 标签（Tag）：用于唯一标识内存块，确保缓存行中的数据来自正确的内存地址。它在缓存命中判断中起着关键作用。
• 索引（Index）：用于确定缓存行或缓存组的具体位置，帮助缓存控制器快速定位目标缓存行。
• 偏移（Offset）：用于确定缓存行内具体字节的位置，帮助缓存控制器读取或写入正确的数据。

在不同的缓存组织方式中，标签和索引的作用有所不同，但它们都是缓存系统中不可或缺的部分，共同决定了缓存的命中率、查找时间和硬件复杂度。

虚拟内存：架构设计

L1缓存改为VIVT，L2保持PIPT

在“乘影”GPGPU的虚拟内存系统中，L1缓存采用了VIVT（Virtual-Index, Virtual-Tag）的设计，而L2缓存则继续保持PIPT（Physically-Indexed, Physically-Tagged）的结构。这种设计的选择是基于性能和复杂度之间的权衡。

• VIVT（L1缓存）：VIVT缓存使用虚拟地址作为索引和标签。由于L1缓存位于每个SM（流多处理器）内部，访问频率极高，使用虚拟地址可以避免每次访问都需要进行地址转换，从而减少延迟。然而，VIVT缓存的一个问题是可能会导致多个虚拟地址映射到同一个缓存行，产生冲突。为了解决这个问题，“乘影”通过逐级Crossbar互联接入的方式，确保不同SM之间的缓存访问不会相互干扰。

• PIPT（L2缓存）：L2缓存使用物理地址作为索引和标签。由于L2缓存由所有SM共享，访问频率相对较低，但容量较大。使用物理地址可以避免虚拟地址映射带来的冲突问题，并且可以更好地支持多任务环境下的内存隔离。此外，PIPT缓存还可以与TLB（Translation Lookaside Buffer）协同工作，进一步提高地址转换的效率。

逐级Crossbar互联接入

为了确保不同SM之间的缓存访问不会相互干扰，“乘影”采用了逐级Crossbar互联接入的方式。具体来说，每个SM内部的L1缓存通过Crossbar连接到L2缓存，而L2缓存又通过更高一级的Crossbar连接到其他系统组件。这种方式不仅减少了冲突，还提高了系统的可扩展性。例如，当多个SM同时访问L2缓存时，Crossbar可以动态地分配带宽，确保每个SM都能获得足够的访问权限。

虚拟内存：L1 TLB

ASID、VPN和PPN

ASID（Address Space Identifier）：地址空间标识符，用于区分不同进程的虚拟地址空间。每个进程都有一个唯一的ASID。

VPN（Virtual Page Number）：虚拟页号，用于标识虚拟地址所在的页。

PPN（Physical Page Number）：物理页号，用于标识物理地址所在的页。作用：在地址翻译过程中，PPN与页内偏移（Page Offset）组合成完整的物理地址。

每个SM内各有ITLB、DTLB服务于L1 Icache和L1 DCache

在每个SM内部，“乘影”设计了两个独立的TLB（Translation Lookaside Buffer），分别为ITLB（Instruction TLB）和DTLB（Data TLB）。这两个TLB分别服务于L1 Icache（指令缓存）和L1 DCache（数据缓存），确保指令和数据的地址转换可以并行进行，减少等待时间。

• 全相联结构，尺寸可调：L1 TLB采用全相联结构，意味着每个TLB项都可以存储任意一个虚拟地址到物理地址的映射。这种设计虽然增加了硬件复杂度，但显著提高了命中率。此外，L1 TLB的尺寸可以根据实际需求进行调整，以平衡性能和资源消耗。

• 同时比对ASID与VPN查找：为了支持多任务环境下的内存隔离，L1 TLB在查找时会同时比对ASID（Address Space Identifier）和VPN（Virtual Page Number）。ASID用于标识不同的进程，确保每个进程只能访问自己的虚拟地址空间。VPN则是虚拟地址的一部分，用于确定具体的页面。通过同时比对ASID和VPN，L1 TLB可以确保地址转换的正确性和安全性。

• 回填替换策略：pLRU：当L1 TLB发生缺失时，系统需要从L2 TLB或页表中获取新的映射，并将其回填到L1 TLB中。为了决定哪个TLB项应该被替换，“乘影”采用了pLRU（Pseudo Least Recently Used）算法。pLRU是一种近似的LRU（Least Recently Used）算法，能够在硬件实现中提供较好的性能，同时保持较低的复杂度。pLRU通过维护一个位图来记录每个TLB项的使用情况，选择最近最少使用的项进行替换。

• 也可配置为不使用L1 TLB：在某些应用场景下，用户可以选择关闭L1 TLB，将所有的地址转换请求直接转发至L2 TLB。这种配置适用于那些对延迟敏感的应用，或者那些不需要频繁进行地址转换的任务。通过灵活配置L1 TLB的使用，用户可以根据实际需求优化系统的性能。

回填替换策略：pLRU

这张图展示的是一个L1 TLB（Translation Lookaside Buffer，翻译后备缓冲器）的结构，特别是其组相连（Set-Associative）缓存的组织方式。红框内的图具体说明了组相连缓存的替换策略，通常采用LRU（Least Recently Used，最近最少使用）策略或其变种。

图解说明

1. 缓存行（Way）：

   • 图中显示了四个缓存行（Way#0, Way#1, Way#2, Way#3），每个缓存行对应一个可能的缓存位置。
   • 每个缓存行包含一个数值（图中显示为0或1），这些数值用于表示缓存行的使用状态。

2. LRU链表：

   • 图中的箭头表示LRU链表的顺序。
   • 数字1和0表示缓存行的使用状态，1表示最近使用，0表示较早使用。
   • 例如，Way#0和Way#1的状态为1，表示它们是最近使用的缓存行；Way#2和Way#3的状态为0，表示它们是较早使用的缓存行。

3. 替换策略：

   • 当需要替换缓存行时，LRU策略会选择最不常用的缓存行进行替换。
   • 在图中，如果需要替换缓存行，会选择Way#2或Way#3，因为它们的状态为0，表示它们是最近最少使用的。

具体步骤

1. 查找：

   • 当处理器访问一个虚拟地址时，TLB会根据虚拟地址的索引部分找到对应的缓存组。
   • 在缓存组中，通过标签部分与缓存行中的标签进行比较，找到匹配的缓存行。

2. 命中：

   • 如果找到匹配的缓存行，则返回对应的物理地址（PPN）。
   • 同时，更新LRU链表，将该缓存行标记为最近使用。

3. 未命中：

   • 如果没有找到匹配的缓存行，则需要从主存中加载新的页表项到TLB中。
   • 选择LRU链表中最不常用的缓存行进行替换，并更新LRU链表。

示例

假设处理器访问一个虚拟地址，TLB查找过程如下：

1. 查找缓存组：

   • 根据虚拟地址的索引部分找到对应的缓存组（例如Way#0, Way#1, Way#2, Way#3）。

2. 标签比较：

   • 比较虚拟地址的标签部分与缓存组中每个缓存行的标签。
   • 假设Way#1的标签匹配。

3. 命中：

   • 返回Way#1中的物理地址（PPN）。
   • 更新LRU链表，将Way#1标记为最近使用。

4. 未命中：

   • 如果没有找到匹配的缓存行，则选择Way#2或Way#3进行替换。
   • 加载新的页表项到替换的缓存行，并更新LRU链表。

总结

红框内的图展示了组相连缓存的LRU替换策略，通过维护一个LRU链表来决定哪些缓存行是最不常用的，从而在需要替换时选择这些缓存行。这种方式可以有效地提高缓存的命中率，减少主存访问次数。

L1 TLB工作流程

在“乘影”GPGPU中，L1 TLB（Translation Lookaside Buffer）位于每个SM（Streaming Multiprocessor）内部，分别服务于L1 Icache（指令缓存）和L1 DCache（数据缓存）。L1 TLB采用全相联结构，尺寸可调，并且能够同时比对ASID（Address Space Identifier）与VPN（Virtual Page Number）进行查找。如果L1 TLB命中，则可以直接返回PPN（Physical Page Number）；如果没有命中，则会将请求转发到L2 TLB。

L1 TLB中的每一项包含ASID、VPN、PPN以及有效位V。当一个虚拟地址vaddr被提交给TLB时，它会被解析为ASID、VPN和其他部分。例如，对于RISC-V Sv32架构，虚拟地址的格式如下：

• ASID：9位（不确定）
• VPN：20位
• 偏移：12位

假设我们有一个虚拟地址vaddr = 0x80004ccc，其中：

• ASID = 0x000（假设）
• VPN = 0x80004
• 偏移 = 0xccc

L1 TLB会检查其内部的所有条目，寻找与当前ASID和VPN匹配的项。如果找到匹配项并且有效位V为1，则直接返回对应的PPN，并使用该PPN和原始偏移构造物理地址paddr。例如，如果L1 TLB中有一项包含ASID=0x000, VPN=0x80004, PPN=0xabcde，那么最终的物理地址将是paddr = 0xabcdeccc。

如果L1 TLB未命中，则请求会被转发到L2 TLB。此时，可以配置L1 TLB不使用，所有查找均转发至L2 TLB。

虚拟内存：L2 TLB系统

L2 TLB由所有SM共享

L2 TLB是一个全局共享的TLB，由所有SM共同使用。与L1 TLB相比，L2 TLB的容量更大，能够存储更多的地址映射，从而提高整体的命中率。L2 TLB的设计考虑了多任务环境下的高效管理和快速查找。

• 从专门的查找表获取ASID对应的根页表地址：L2 TLB通过一个专门的查找表（ASID映射表）来获取每个ASID对应的根页表地址。这个查找表存储了所有活动进程的ASID及其对应的根页表地址，确保L2 TLB可以从正确的页表开始进行地址转换。

• 主TLB存储：L2 TLB的主要部分是主TLB，用于存储虚拟地址到物理地址的映射。主TLB采用多路组相联结构，相联度和尺寸均可调，以适应不同的应用场景。为了提高查找效率，主TLB还进行了分块（Bank）设计，根据来访虚拟页号划分进入不同的L2 TLB端口。这种方式可以减少冲突，提高并发性能。

• TLB表项存放多个PPN（Sector）：每个TLB表项可以存放多个物理页号（PPN），形成一个Sector。这种方式使得L2 TLB的访存宽度与L2 Cache适配，减少了访存次数。具体来说，Sector中的每个PPN可以通过SectorIndex段进行选择，确保每次访问都能准确地找到所需的物理地址。

• 辅助TLB存储：除了主TLB，L2 TLB还包含一个辅助TLB，用于存储多级地址转换过程中的中间层PPN。对于Sv32体系，辅助TLB存储一个中间层PPN；对于Sv39体系，辅助TLB存储两个中间层PPN。辅助TLB采用全相联结构，不分Bank，但分为多个Sector。辅助TLB与主TLB同步查询，当主TLB缺失但辅助TLB命中时，PTW可以直接从中间层PPN开始遍历页表，加快地址转换速度。

页表遍历单元（PTW）

主TLB每Bank拥有一个

每个L2 TLB Bank都配备了一个页表遍历单元（PTW），负责处理TLB缺失时的页表遍历操作。PTW的核心任务是从根页表开始，逐级遍历页表，直到找到目标虚拟地址对应的物理地址。为了提高遍历效率，PTW每次访存可以获取连续的多个页表项（一个Sector），减少了访存次数。

• 根页表地址从ASID映射表获取：PTW从ASID映射表中获取当前进程的根页表地址，作为页表遍历的起点。这种方式确保了每次遍历时都能从正确的页表开始，避免了错误的地址转换。

• 地址转换中间项回填至辅助TLB：在页表遍历过程中，PTW会将遇到的中间层PPN回填到辅助TLB中，供后续查询使用。这样，当主TLB再次缺失时，如果辅助TLB命中，PTW可以直接从中间层开始遍历，跳过前面的层级，进一步加快地址转换速度。

• 最终转换结果回填至主TLB：当PTW完成页表遍历后，它会将最终的转换结果（即虚拟地址对应的物理地址）回填到主TLB中，供后续的地址转换请求使用。这种方式不仅提高了命中率，还减少了TLB缺失的发生频率。

L2 TLB工作流程

L2 TLB是所有SM共享的，负责处理来自L1 TLB的未命中请求。L2 TLB的工作流程更为复杂，涉及到根页表地址的获取、索引分割、多路组相联结构以及Sector机制等。

获取根页表地址

L2 TLB首先需要从专门的查找表中获取ASID对应的根页表地址（PTBR, Page Table Base Register）。这个查找表存储了每个ASID对应的根页表基地址。例如，假设我们有一个ASID=0x000，那么L2 TLB会查询查找表，得到根页表地址ptbr = 0x00001000。

索引分割

接下来，L2 TLB会根据来访的虚拟页号VPN进行索引分割。L2 TLB采用多路组相联结构，相联度和尺寸均可调。为了提高并行度，L2 TLB被划分为多个Bank，每个Bank负责处理一部分虚拟页号。具体的索引分割方式如下：

• SetIdx：用于确定L2 TLB中的Set编号
• BankIdx：用于确定L2 TLB中的Bank编号
• SectorIdx：用于选择正确的Sector

以Sv32为例，虚拟页号的格式如下：

• SetIdx：4位
• BankIdx：2位
• SectorIdx：3位
• Tag：11位

假设我们有虚拟页号VPN = 0x80004，则：

• SetIdx = 0x0
• BankIdx = 0x0
• SectorIdx = 0x4
• Tag = 0x400

L2 TLB会根据SetIdx和BankIdx选择相应的Bank和Set，然后在该Set中查找与Tag匹配的项。如果找到匹配项，则根据SectorIdx选择正确的PPN。

TLB Entry解析

每个L2 TLB条目包含多个PPN（Sector），这些PPN与L2 Cache访存宽度适配。例如，假设L2 TLB条目包含两个PPN（Sector=2），则条目的格式如下：

vpn	ppn0	ppn1	A	B0	B1
0x80004	0xabcde	0xfedcba	0	0	0

根据SectorIdx选择正确的PPN。例如，如果SectorIdx=0，则选择ppn0=0xabcde；如果SectorIdx=1，则选择ppn1=0xfedcba。

回填策略

如果L2 TLB也未命中，则需要通过页表遍历单元（PTW, Page Table Walker）进行逐级查找。PTW会从根页表开始，逐级访问页表项（PTE, Page Table Entry），直到找到对应的物理页号PPN。每次访存获取连续的Sector个页表项，地址转换中间项回填至辅助TLB，最终转换结果回填至主TLB。

辅助TLB

辅助TLB存储多级地址转换过程的中间层PPN。对于Sv32，辅助TLB存储一级页表项；对于Sv39，存储两级页表项。辅助TLB与主TLB同步查询，主TLB缺失但辅助TLB命中时，提示PTW直接从中间层PPN开始遍历页表。辅助TLB采用全相联结构，不分Bank，但分Sector。

示例说明

假设我们有一个虚拟地址vaddr = 0x80004ccc，经过L1 TLB未命中后，请求被转发到L2 TLB。L2 TLB首先从查找表中获取ASID=0x000对应的根页表地址ptbr = 0x00001000。然后，根据虚拟页号VPN = 0x80004进行索引分割，得到SetIdx=0x0, BankIdx=0x0, SectorIdx=0x4, Tag=0x400。L2 TLB在Bank#0的Set#0中查找与Tag=0x400匹配的项，假设找到了一条包含ppn0 = 0xabcde的条目。根据SectorIdx=4，选择ppn4=0xabcde作为最终的物理页号。最后，结合原始偏移0xccc，构造出物理地址paddr = 0xabcdeccc。

如果L2 TLB也未命中，则PTW会从根页表地址ptbr = 0x00001000开始逐级查找页表项，直到找到对应的物理页号PPN，并将其回填到L2 TLB和L1 TLB中，以便后续访问。

通过这种方式，“乘影”GPGPU的L1和L2 TLB系统能够高效地完成虚拟地址到物理地址的转换，确保内存访问的快速性和准确性。

虚拟内存软件支持

使用Scala编写内存空间分配与页表管理系统

“乘影”GPGPU的虚拟内存管理软件使用Scala编写，提供了灵活的内存空间分配和页表管理功能。该系统可以配置为实地址模式或Sv32/Sv39虚拟地址模式，满足不同应用场景的需求。

• 实地址模式：在这种模式下，所有内存访问都使用物理地址，无需进行地址转换。这种模式适用于那些不需要虚拟内存支持的应用，或者在调试和测试阶段使用。

• Sv32/Sv39虚拟地址模式：在这种模式下，系统使用虚拟地址进行内存访问，并通过MMU进行地址转换。Sv32支持32位虚拟地址，使用两级页表；Sv39支持64位虚拟地址，使用三级页表。开发者可以根据实际需求选择合适的模式。

提供GPU侧和主机侧的操作函数

为了方便开发者使用，“乘影”提供了GPU侧和主机侧的操作函数，用于接入仿真框架。这些函数允许开发者在GPU上执行内存分配、页表创建和更新等操作，同时也支持主机侧的控制和管理。通过这种方式，开发者可以在仿真环境中模拟真实的硬件行为，验证虚拟内存管理系统的正确性和性能。

接入内存管理单元的整体测试

额外执行周期开销

引入虚拟内存管理单元（MMU）后，系统会带来一定的额外执行周期开销。根据测试结果，额外开销大约在+15%到+25%之间。这部分开销主要来自于地址转换和缓存未命中等情况。尽管如此，通过优化TLB设计和页表遍历机制，“乘影”仍然能够保持较高的性能水平。

L1 DTLB与L2 TLB命中率统计

为了评估虚拟内存管理单元的性能，“乘影”进行了详细的命中率统计。结果显示，L1 DTLB的命中率较高，尤其是在高频访问的数据区域。L2 TLB的命中率也表现良好，特别是在多任务环境下，能够有效减少页表遍历的次数。通过这些统计数据，开发团队可以进一步优化TLB设计，提高系统的整体性能。

---

VIVT和PIPT分别是什么，为什么这么设计？

VIVT（Virtual-Index, Virtual-Tag）

VIVT缓存使用虚拟地址作为索引和标签。这种设计的优点是访问速度快，因为不需要进行地址转换，可以直接使用虚拟地址进行缓存查找。然而，VIVT缓存的一个缺点是可能会导致多个虚拟地址映射到同一个缓存行，产生冲突。为了避免冲突，“乘影”采用了逐级Crossbar互联接入的方式，确保不同SM之间的缓存访问不会相互干扰。

PIPT（Physically-Indexed, Physically-Tagged）

PIPT缓存使用物理地址作为索引和标签。这种设计的优点是可以避免虚拟地址映射带来的冲突问题，并且可以更好地支持多任务环境下的内存隔离。此外，PIPT缓存还可以与TLB协同工作，进一步提高地址转换的效率。然而，PIPT缓存的缺点是访问速度相对较慢，因为每次访问都需要进行地址转换。因此，“乘影”将PIPT设计应用于L2缓存，以平衡性能和复杂度。

回填替换策略是干什么的？是怎么实现的？

回填替换策略的作用

回填替换策略决定了当TLB发生缺失时，应该用新的映射替换哪个TLB项。合理的替换策略可以提高TLB的命中率，减少地址转换的开销。常见的替换策略包括LRU（Least Recently Used）、FIFO（First In, First Out）等。

pLRU（Pseudo Least Recently Used）的实现

“乘影”采用了pLRU（Pseudo Least Recently Used）算法作为L1 TLB的回填替换策略。pLRU是一种近似的LRU算法，能够在硬件实现中提供较好的性能，同时保持较低的复杂度。pLRU通过维护一个位图来记录每个TLB项的使用情况，选择最近最少使用的项进行替换。具体来说，pLRU将TLB项分为多个层次，每个层次记录一个位，表示该层是否被访问过。当某个TLB项被访问时，相应的位会被置为1；当需要替换时，选择最低层次的位为0的项进行替换。

SectorIndex段是干什么的？为什么用这个选择PPN？

SectorIndex段的作用

SectorIndex段用于选择TLB表项中的具体PPN（Physical Page Number）。在L2 TLB中，每个表项可以存放多个PPN，形成一个Sector。这种方式使得L2 TLB的访存宽度与L2 Cache适配，减少了访存次数。具体来说，Sector中的每个PPN可以通过SectorIndex段进行选择，确保每次访问都能准确地找到所需的物理地址。

为什么用SectorIndex选择PPN

使用SectorIndex选择PPN的原因是为了提高访存效率。传统的单个PPN设计需要每次访问一个单独的页表项，而Sector设计可以在一次访存中获取多个PPN，减少了访存次数。通过使用SectorIndex段，系统可以在Sector中选择正确的PPN，确保每次访问都能准确地找到所需的物理地址。这种方式不仅提高了访存效率，还减少了TLB缺失的发生频率。

页表遍历单元（PTW）干什么的？Sector是什么？

页表遍历单元（PTW）的作用

页表遍历单元（PTW）负责处理TLB缺失时的页表遍历操作。当TLB发生缺失时，PTW会从根页表开始，逐级遍历页表，直到找到目标虚拟地址对应的物理地址。PTW的核心任务是加速地址转换过程，减少TLB缺失带来的性能损失。

Sector的定义

Sector是指L2 TLB表项中存放的多个PPN（Physical Page Number）。在多级页表体系中，每个TLB表项可以存放多个PPN，形成一个Sector。这种方式使得L2 TLB的访存宽度与L2 Cache适配，减少了访存次数。具体来说，Sector中的每个PPN可以通过SectorIndex段进行选择，确保每次访问都能准确地找到所需的物理地址。

PTW的工作流程

当PTW接收到TLB缺失请求时，它会从ASID映射表中获取当前进程的根页表地址，作为页表遍历的起点。然后，PTW会逐级遍历页表，直到找到目标虚拟地址对应的物理地址。在这个过程中，PTW会将遇到的中间层PPN回填到辅助TLB中，供后续查询使用。当PTW完成页表遍历后，它会将最终的转换结果（即虚拟地址对应的物理地址）回填到主TLB中，供后续的地址转换请求使用。

通过这种方式，PTW不仅加速了地址转换过程，还减少了TLB缺失的发生频率，提高了系统的整体性能。

Tensor Core

Tensor Core设计概况

在大模型的浪潮下，神经网络的高效计算成为研究热点。为了加速神经网络的计算，“乘影”GPGPU引入了张量计算单元（Tensor Core），专门用于加速矩阵乘法和矩阵乘加混合操作。这些优化显著提升了神经网络的计算速度。

计算精度支持

“乘影”Tensor Core最初支持FP32（单精度浮点）计算，但为了适应不同精度量化的需求，它最终修改为对FP16（半精度浮点）、INT8（整数）以及混合精度计算的支持。这种多精度支持使得“乘影”能够在不同的应用场景中提供最优的性能和能效比。例如，在推理阶段，使用FP16或INT8可以显著减少计算资源和功耗，而在训练阶段，FP32则提供了更高的数值精度。

微架构细节

在微架构层面，“乘影”通过复用不同精度的运算单元来增加芯片的算力密度。具体来说，为了支持INT8精度的计算，Tensor Core复用了现有的浮点运算乘法器。具体来说，INT8乘法可以通过将整数转换为浮点数的方式，利用现有的FP16乘法器进行计算。这种方式不仅减少了硬件资源的占用，还提高了INT8计算的灵活性。例如，在进行INT8乘法时，操作数收集器会将两个8位整数扩展为16位浮点数，然后通过FP16乘法器进行计算。计算结果再被截断为8位整数，存储到目标寄存器中。这种方式使得Tensor Core能够在不增加额外硬件的情况下，支持多种精度的矩阵运算。

此外，为了支持更大尺寸的矩阵乘加混合操作，“乘影”增加了Tensor Core控制单元，并通过流水线技术使8×4×8的计算阵列能够高效工作。这意味着每个Tensor Core内部有一个计算阵列，可以完成FP16、混合精度和INT8的计算任务。

控制单元与流水线

Tensor Core控制单元负责管理和调度来自不同warp的矩阵计算指令。每条指令在Operand Collector（操作数收集器）中被分解为携带不同寄存器数据的微指令，然后通过多个slot和crossbar发送到计算阵列。多个slot的状态机可以复用同一个计算阵列，从而实现高效的并行计算。例如，一个状态机可以调用4次8×4×8的阵列来完成一次16×8×16的矩阵乘加累计运算。

Tensor数据在寄存器中的分布

为了支持高效的矩阵计算，“乘影”设计了一种特殊的寄存器映射机制，以确保数据在寄存器中的分布能够充分利用硬件资源。

以矩阵A、B为FP16，C为FP16或FP32为例，矩阵元素与寄存器位置的映射关系如下：

• 每个32位寄存器从高位至低位分为两个16位数据，分别表示为A1和A0。
• 例如，寄存器v0可以存储两个FP16数据：v0 = [A1, A0]。
• T0-T31表示第32个线程，每个线程可以访问一组寄存器，用于存储矩阵A、B和C的元素。

图中最下面两个分别为：左图是C为FP32时的情况，右图是C为FP16时的情况。左图中C[0, 0]元素直接存储到T0对应的寄存器中即，C0-C1，为32bit。右图中C[0, 0]元素存到T0对应的寄存器的低16bit，即C0，C[0, 1]元素存到T0对应的寄存器的高16bit，即C1，从而实现寄存器的充分利用。

Tensor Core 硬件结构

Tensor Core硬件结构概述

“乘影”GPGPU的Tensor Core是专门设计用于加速神经网络中矩阵乘法和矩阵乘加混合操作的硬件单元。通过引入张量计算单元，GPGPU能够在大规模矩阵运算中提供显著的性能提升。具体来说，Tensor Core支持FP16、INT8以及混合精度计算，为不同的神经网络应用场景提供了灵活的硬件支持。

计算阵列与指令分解

每个Tensor Core内部包含一个计算阵列，能够执行FP16精度、混合精度以及INT8精度的矩阵乘加操作。计算阵列的设计采用了8×4×8的尺寸。8×4×8的尺寸指的是Tensor Core计算单元共有8行4列点积计算单元，所以是8*4;每个点积计算单元内可以执行操作矩阵A的某行8个数据与操作矩阵B的某列8个数据相乘，然后相加，最后累加上另一个操作矩阵C的某个元素的操作，得到结果矩阵D的一个元素，因此综合起来才是8*4*8。

因此，第一个8*4可以看作是结果矩阵的大小，第二个8是每次操作的元素。

为了实现高效的并行计算，每条Tensor指令在Operand Collector（操作数收集器）中被分解为多个微指令，这些微指令携带不同的寄存器数据，分别对应矩阵A、B和C的不同部分。

例如，假设我们有一个m16n8k16的矩阵乘加累计运算，即矩阵A的维度为16×16，矩阵B的维度为16×8，矩阵C的维度为16×8。计算阵列会分四次调用8×4×8的子阵列来完成整个运算。每次调用时，Operand Collector会从寄存器中收集相应的操作数，并将其传递给计算阵列进行处理。

状态机与流水线

为了提高计算效率，Tensor Core内部的状态机被设计为流水线结构。以m16n8k16的矩阵乘加累计运算为例，状态机会依次调用四个子状态：set1、set2、set3和set4。每个子状态负责处理一部分矩阵运算。

• 状态机计算过程与阵列调用：状态机计算 8×8×16 的矩阵运算，其中 4 个状态调用 4 次 8×4×8 阵列。在计算过程中，状态机分步骤协调数据的处理。例如，对于矩阵乘法运算 A×B = C（这里 A 为 16×16 矩阵，B 为 16×8 矩阵，C 为 16×8 矩阵），状态机首先会组织数据，将 A 矩阵按行、B 矩阵按列进行分组，然后控制计算阵列进行点积计算。在计算 8×8×16 的部分时，每次调用 8×4×8 阵列，通过合理安排数据在阵列中的流动，完成部分点积计算任务。

• FP16 情况：完成 4 个 set 时，FP16 时需要 2 条微指令（4 个 set 整体取 2 次数），每条微指令对应 2 个 set。这意味着在 FP16 精度下，通过合理安排微指令的功能和执行顺序，可以高效地完成计算任务。例如，一条微指令可以负责获取 A、B 矩阵数据并进行部分点积计算，另一条微指令则处理中间结果与 C 矩阵（如果需要）的计算以及最终结果的更新。

• 混合精度计算情况：混合精度计算时需要 4 条微指令（4 个 set 整体取 4 次数，但是可复用 A、B），每条微指令对应 1 个 set。在混合精度计算中，由于不同精度的数据处理方式和计算步骤可能不同，所以需要更多的微指令来精细控制计算过程。例如，对于某些涉及 FP16 和 FP32 混合精度的矩阵计算，可能需要一条微指令专门处理 FP16 数据的计算，另一条微指令处理 FP32 数据的转换和计算，并且在不同的 set 阶段根据精度要求合理安排微指令的执行。

数据路径与复用

为了进一步提高资源利用率，Tensor Core设计了多个slot和crossbar，允许来自不同warp的矩阵计算指令同时进入计算阵列。每个slot的状态机可以复用同一个计算阵列，从而实现多任务并行处理。例如，假设我们有四个不同的warp，每个warp都有一个矩阵乘加指令需要处理。通过crossbar，这四个指令可以被分配到不同的slot中，每个slot的状态机独立工作，但共享同一个计算阵列。这种方式不仅提高了计算效率，还减少了硬件资源的浪费。

Slot

Slot 是指Tensor Core内部用于接收和处理矩阵计算指令的逻辑单元。每个slot可以理解为一个独立的任务处理单元，它负责执行来自不同warp（线程束）的矩阵乘加指令。具体来说，每个slot包含一个状态机，该状态机负责控制指令的执行流程，包括从寄存器中读取操作数、调用计算阵列进行矩阵运算、以及将结果写回到目标寄存器。

例如，假设我们有四个不同的warp，每个warp都有一个矩阵乘加指令需要处理。通过多个slot的设计，这四个指令可以被分配到不同的slot中，每个slot的状态机独立工作，互不干扰。这种方式允许多个矩阵运算指令并行执行，从而提高了整体的计算效率。

Crossbar

Crossbar 是一种多端口交换网络，它连接了操作数收集器（Operand Collector）和计算阵列。Crossbar的作用是将来自不同warp的操作数灵活地路由到计算阵列的不同部分，使得多个slot可以共享同一个计算阵列。具体来说，crossbar可以根据当前正在执行的指令，动态地选择哪些操作数应该被传递给计算阵列的哪些部分，从而实现高效的资源复用。

例如，在一个多任务并行处理的场景中，假设我们有四个不同的warp，每个warp都需要执行一个矩阵乘加指令。通过crossbar，这些指令的操作数可以从不同的寄存器中读取，并被路由到计算阵列的不同部分。这样，即使多个warp同时请求使用计算阵列，crossbar也可以确保每个slot的状态机能够独立地获取所需的操作数，而不会发生冲突。

操作数收集器

操作数收集器（Operand Collector）是Tensor Core中的一个重要组件，负责从寄存器中收集操作数，并将其传递给计算阵列。为了支持高效的矩阵运算，操作数收集器需要根据不同的精度和矩阵大小，灵活地从寄存器中读取数据。

操作数收集器会根据指令的要求，从相应的寄存器中读取数据，并将其打包成微指令传递给计算阵列。这种方式不仅简化了指令的执行流程，还提高了数据传输的效率。

数据异步拷贝机制

背景

H100 的流多处理器架构中的显著变化之一是支持异步数据拷贝功能的张量内存加速器（Tensor Memory Accelerator），这对乘影 GPGPU 的数据异步拷贝机制设计产生了影响。

• 扩展异步传输与地址空间覆盖：扩展了前一代 GPGPU 全局内存到共享内存异步的传输，并且覆盖所有地址空间。这意味着在乘影 GPGPU 中，数据可以在不同内存空间之间更高效地进行异步移动。例如，在处理大规模图像数据时，图像数据可能存储在全局内存中，而在进行图像处理计算时，需要将部分数据快速传输到共享内存供计算单元使用，数据异步拷贝机制能够实现这种高效的传输，且不局限于特定的地址范围，任何地址空间的数据都可以进行异步传输。
• 支持张量内存访问模式：增加了对张量内存访问模式的支持，使得应用程序能够构建端到端的异步流水线。在深度学习模型训练或推理过程中，经常涉及大量张量数据的处理。以卷积神经网络（CNN）为例，在卷积层计算时，需要按照特定的张量内存访问模式读取输入数据和权重数据进行计算，乘影 GPGPU 的数据异步拷贝机制可以更好地适应这种访问模式，提高数据传输与计算的重叠效率。
• 重叠数据移动与计算的优势：通过异步拷贝，应用程序能够将数据移入和移出芯片，同时完全重叠并隐藏数据移动与计算。这极大地提高了系统的整体性能，因为在传统的同步模式下，数据传输和计算是串行进行的，计算单元在等待数据传输完成时处于空闲状态，浪费了计算资源。而采用异步拷贝机制，例如在进行矩阵乘法计算的同时，可以将下一批数据从全局内存异步传输到共享内存，当计算单元完成当前计算任务后，立即可以获取到新的数据进行下一步计算，实现计算资源的充分利用，减少整体计算时间。

DMA engine 硬件设计

• 内部结构与连接关系：DMA engine 在 SM（流多处理器）中的位置涉及多个组件的协同工作。L1tlb 和 L2tlb 用于地址转换，TLBReq 和 TLBRsp 负责处理与 TLB 相关的请求和响应。ISSUE instinfo 模块处理指令信息，InputFIFO instinfo 用于指令信息的缓冲，AddrCalc_2ache 和 AddrCalc_shared 分别负责计算与 L2cache 和共享内存相关的地址，crossbar 用于数据交叉传输，L2Cache 用于缓存数据，TempMem 用于临时存储数据，Arbiter/Crossbar 负责仲裁和交叉开关控制，Warp Scheduler 用于线程束调度。这些组件共同协作，实现数据的异步拷贝。例如，当一个线程束需要从全局内存拷贝数据到共享内存时，Warp Scheduler 首先调度该任务，然后通过 AddrCalc_2ache 计算数据在 L2cache 中的地址，通过 crossbar 将数据传输到 TempMem 进行临时存储，最后再通过 AddrCalc_shared 计算共享内存中的目标地址，将数据从 TempMem 传输到共享内存。
• 工作流程示例：以一个简单的矩阵数据传输为例，假设要将存储在全局内存中的一个 100×100 的矩阵数据传输到共享内存。首先，线程束发出数据传输请求，ISSUE instinfo 模块接收并解析指令，确定这是一个从全局内存到共享内存的拷贝操作。然后，AddrCalc_2ache 根据全局内存地址计算数据在 L2cache 中的位置，通过 L2cache 获取数据并传输到 TempMem。在 TempMem 中，数据可能会根据一定的规则进行整理和缓存，等待后续传输到共享内存。同时，AddrCalc_shared 计算共享内存中的目标地址，当一切准备就绪，Arbiter/Crossbar 控制数据从 TempMem 通过 crossbar 传输到共享内存中指定的位置，完成整个异步拷贝过程。

DMA 硬件设计 Bounding Box, Tempmem

• 数据信息与功能
  • BoxAddress：指定 Box 的起始地址。在处理数据时，例如处理一个图像数据块，BoxAddress 就确定了这个图像数据块在内存中的起始位置，就像在一幅地图上确定一个区域的起始坐标一样。假设一个图像数据存储在内存中，BoxAddress 可以指向图像左上角像素数据的内存地址，后续的数据读取和处理就可以从这个起始地址开始。
  • boxDim array：指定沿每个维度遍历的元素数，必须为非零且小于或等于 256。对于一个三维数据结构（如 RGB 图像数据，可看作三维数组，宽度、高度和颜色通道为三个维度），boxDim array 可以定义在每个维度上需要处理的数据元素数量。例如，对于一个 64×64 像素的 RGB 图像，boxDim array 可能为 [64, 64, 3]，表示在宽度方向遍历 64 个像素，高度方向遍历 64 个像素，颜色通道方向遍历 3 个元素（R、G、B）。
  • elementStrides array：指定沿每个维度的迭代步长，必须不为零且小于或等于 8。在处理一些特殊数据布局时，elementStrides array 非常有用。比如在处理一个稀疏矩阵数据时，元素可能不是连续存储的，elementStrides array 可以定义在每个维度上跳过多少个元素来获取下一个有效数据。如果稀疏矩阵在某一维度上每隔 2 个元素才有一个有效数据，那么 elementStrides array 在该维度上的值可以设置为 2。
  • obbfill：指定使用 0 还是 NaN 来填充越界元素。在数据处理过程中，当访问超出预先定义的 Box 范围时，就需要根据 obbfill 的设置来填充数据。例如，在进行图像边缘处理时，如果算法需要对图像边缘外的虚拟像素进行计算，那么可以根据 obbfill 的设置将这些虚拟像素填充为 0 或 NaN，以便计算能够正常进行，并且在结果中能够正确表示这些越界元素的处理方式。
• 线程束调度与异步栅栏指令执行：线程束调度器执行异步栅栏指令，确保数据拷贝和计算的正确性和同步性。在多个线程束同时进行数据拷贝和计算时，异步栅栏指令就像一个同步点。例如，在一个并行计算任务中，多个线程束分别从不同的数据源拷贝数据到共享内存，当所有线程束都完成数据拷贝后，异步栅栏指令会确保在进行下一步计算之前，所有数据都已经正确地到达共享内存。如果没有这个同步机制，可能会出现部分线程束使用了未完全拷贝的数据进行计算，导致计算结果错误。假设一个深度学习模型的前向传播计算中，多个线程束负责拷贝不同层的权重数据到共享内存，在所有权重数据都拷贝完成后，异步栅栏指令确保所有数据准备就绪，然后计算单元才开始进行前向传播计算，保证计算结果的准确性。

CTA 调度器

重新设计的目标

• 面积削减与性能提升
  • 新设计致力于削减 CTA 调度器的面积，通过优化内部结构来减少不必要的硬件开销。例如，通过对 Local memory、sGPR、vGPR 资源的碎片化进行抑制，能够更高效地利用这些资源，减少资源浪费，进而提升整体性能。
• 更好地兼容各类线程块调度策略
  • 随着应用场景的多样化，CTA 调度器需要能够适应不同类型的线程块调度策略。重新设计的调度器能够更好地处理各种线程块的调度需求，无论是简单的线性调度还是复杂的并行调度策略。

重新设计的结构特点

• 结构更清晰：换用分布式控制结构
  • 新的 CTA 调度器采用分布式控制结构，使得内部结构更加清晰。这种结构有助于将复杂的调度任务分解为多个相对独立的子任务，由不同的模块分别处理，提高了调度器的可维护性和扩展性。
  • 例如，在处理大规模数据并行任务时，分布式控制结构可以让不同的控制单元分别负责不同数据块的调度，避免了集中式控制可能出现的性能瓶颈。
• 数据信号分组与自动连线
  • 通过对数据信号进行分组和自动连线，能够优化数据传输路径，减少信号干扰和传输延迟。在调度器内部，不同类型的数据（如线程块信息、资源分配信息等）可以按照其特性进行分组传输。
  • 例如，在调度多个线程块进行并行计算时，线程块的相关数据可以通过专门的分组和连线快速准确地传输到对应的计算单元，确保计算的高效进行。

重新设计的测试与文档

• 测试更完善：增加中大规模单体随机测试与 GPGPU 整体通过测试用例
  • 为了确保重新设计的 CTA 调度器的可靠性，增加了中大规模单体随机测试。这种测试方法可以模拟各种实际应用场景下的调度需求，发现潜在的漏洞和性能瓶颈。
  • 同时，进行 GPGPU 整体通过测试用例，能够从系统层面验证调度器与其他组件（如计算单元、存储单元等）的协同工作能力。例如，在模拟一个大型深度学习模型训练的场景下，通过整体测试可以检验调度器是否能够合理地分配线程块，保证训练过程的顺利进行。
• 文档更详细
  • 重新设计的过程中，生成了更详细的文档，包括调度器的内部结构、工作原理、接口定义等内容。这些详细的文档有助于开发人员更好地理解和维护调度器，也方便了后续的升级和优化工作。

重新设计的具体模块

• 资源表（Resource table）与资源分配器（Resource allocator）
  • 资源表用于记录和管理可用的资源信息，包括计算单元的空闲情况、存储单元的剩余容量等。资源分配器则根据资源表和当前的调度需求，合理地将资源分配给不同的线程块。
  • 例如，在处理一个多线程的并行任务时，资源分配器会先查询资源表，确定哪些计算单元和存储单元是空闲的，然后将线程块分配到这些空闲资源上，确保资源的充分利用。
• 线程块缓冲区（WG buffer）与 Read & Clear
  • 线程块缓冲区用于临时存储待调度的线程块信息，Read & Clear 模块则负责从缓冲区中读取线程块信息并进行处理，处理完成后清除相应的缓冲区记录。
  • 比如在处理一批新的线程块调度任务时，这些线程块的信息会先进入 WG buffer，然后 Read & Clear 模块会按照一定的顺序从缓冲区中取出线程块信息，进行调度操作，操作完成后将该线程块的记录从缓冲区中清除，为下一批线程块腾出空间。
• 计算单元阵列（CU array）、计算单元接口（CU interface）与待执行 / 已完成线程束
  • 计算单元阵列是实际执行计算任务的硬件单元，计算单元接口负责连接调度器和计算单元阵列，确保数据和指令的正确传输。待执行和已完成线程束则分别表示处于等待调度和已经完成计算的线程束状态。
  • 例如，在执行一个并行计算任务时，调度器通过 CU interface 将线程块分配到 CU array 中的各个计算单元，待执行线程束在 CU array 中等待执行，执行完成后，这些线程束变为已完成线程束状态，调度器可以根据这些状态信息进行后续的调度操作。

Hyper-Threading

Hyper-Threading技术，也称为超线程技术，是英特尔（Intel）公司开发的一种处理器技术。它允许一个物理处理器在操作系统中被识别为多个逻辑处理器（或线程），从而提高处理器的效率和性能。

硬件实现机制

架构状态复制：

每个逻辑处理器都有自己的一套架构状态，包括通用寄存器和控制寄存器。这意味着操作系统可以像管理两个独立的物理处理器一样管理这两个逻辑处理器。

资源共享与分区：

在Hyper-Threading技术中，某些资源是被两个逻辑处理器共享的，如缓存、执行单元和总线。而其他一些资源则被复制或分区，以确保每个逻辑处理器都能独立操作。例如，每个线程都有自己的返回栈缓冲区（RSB）和指令队列。

执行单元和调度：

处理器的执行管线可以同时处理来自两个逻辑处理器的指令。这意味着，如果一个逻辑处理器因为等待缓存数据或其他原因而停滞，另一个逻辑处理器可以利用这些空闲的执行资源继续执行指令。

缓存一致性和资源共享：

L1、L2和L3缓存是被两个逻辑处理器共享的资源。这些共享资源的性能取决于线程的动态行为和对资源的竞争。

硬件开销：

实现Hyper-Threading技术的硬件开销相对较小，因为它主要涉及复制架构状态，而大部分执行资源是共享的。这种设计使得Hyper-Threading技术能够在增加很少的芯片面积和功耗的情况下提供性能提升。

在单个物理核心上虚拟化出两个逻辑核心

即实现Hyper-Threading技术，涉及到硬件和软件的多个方面。以下是实现这一过程的关键步骤和组件：

硬件支持

架构状态复制：

每个逻辑核心都需要有自己的一套架构状态，包括寄存器、程序计数器等。这意味着物理核心需要有能力维护两套这些状态，每套状态对应一个逻辑核心。

执行单元和调度器：

物理核心的执行单元（如ALU、FPU）需要能够同时处理来自两个逻辑核心的指令。这通常通过动态调度器实现，它能够在两个逻辑核心之间动态分配执行资源。

资源共享与分区：

某些资源（如缓存）被两个逻辑核心共享，而其他资源（如寄存器文件）则被分区或复制。共享资源需要有机制来处理资源竞争和一致性问题。

上下文切换：

硬件需要能够快速在两个逻辑核心之间切换上下文。这通常通过特殊的硬件机制实现，以减少上下文切换的延迟。

软件支持

操作系统支持：

操作系统需要能够识别和调度这两个逻辑核心。这包括在调度算法中考虑这两个核心，并正确地分配线程到这些核心上。

线程库和API：

编程库和API需要支持多线程编程，使得开发者能够编写能够在这两个逻辑核心上运行的并发代码。

编译器优化：

编译器需要能够优化代码，以充分利用两个逻辑核心的并行能力。这可能包括线程级别的并行化、指令调度和缓存优化。

缓存和内存管理

缓存一致性：

由于两个逻辑核心共享某些缓存，需要有机制来维护缓存一致性，确保所有核心看到的数据是最新的。

内存访问管理：

内存访问需要被正确地路由到两个逻辑核心，同时确保内存操作的顺序性和一致性。

性能监控和优化

性能计数器：

硬件需要有性能计数器来监控两个逻辑核心的性能，包括指令吞吐量、缓存命中率等。

动态调整：

根据性能反馈，硬件和软件可以动态调整资源分配和调度策略，以优化性能。

通过这些硬件和软件的协同工作，Hyper-Threading技术能够在单个物理核心上有效地虚拟化出两个逻辑核心，提高了处理器的利用率和性能。这种技术特别适用于那些能够并行执行多个线程的应用程序，可以显著提高吞吐量和响应速度。

总的来说，逻辑核心有自己的寄存器、程序计数器，但是执行单元，例如ALU和FPU等是共享的，通过动态调度器来实现。此外，缓存也是共享的，还需要切换上下文。

warp切换是类似hyper-threading的周期级调度

周期级调度（cycle-level scheduling）是指调度器在每个时钟周期都会进行调度决策，选择下一个要执行的线程或线程束（warp）。这种调度方式允许在单个物理处理器核心上模拟出多个逻辑处理器的效果，类似于Intel的Hyper-Threading技术。在GPU环境中，这种调度特别适用于掩盖内存访问延迟和提高执行单元的利用率。

在“乘影”GPU中，warp切换是类似Hyper-Threading的周期级调度，意味着：

Warp的概念：在GPU中，一个warp通常包含32个线程，它们被分组在一起执行相同的指令（SIMT，单指令多线程）。

Warp Scheduler的作用：Warp Scheduler负责在单个流多处理器（SM）内调度多个warp。当一个warp因等待内存操作或其他原因而被阻塞时，Warp Scheduler会迅速切换到另一个准备好的warp，继续执行指令。

掩盖延迟：通过在多个warp之间快速切换，GPU可以掩盖内存访问的高延迟，因为当一个warp等待数据时，其他warp可以继续执行，从而保持处理器的执行单元忙碌。

资源利用率：周期级调度允许GPU在每个时钟周期都尽可能地保持执行单元的高利用率，即使在面对线程执行的不一致性（如分支发散）时也能保持性能。

上下文切换开销：由于线程的寄存器状态是私有的，warp之间的上下文切换开销非常小，这使得在warp被阻塞时可以迅速切换到其他warp，而不需要保存和恢复大量状态信息。

总结来说，周期级调度是一种在每个时钟周期都进行的调度机制，它使得单个物理核心能够通过快速切换执行多个逻辑线程（warp），从而提高处理器的资源利用率和掩盖延迟，类似于CPU中的Hyper-Threading技术。在“乘影”GPU中，这种调度策略允许在单个物理核心上高效地模拟出多个逻辑执行流，提升了GPU的性能和响应能力。

SIMT和SIMD

什么是 SIMT 和 SIMD？

• SIMT（单指令多线程）：

  • 在逻辑上，每个线程独立运行一个执行路径。线程可能分支，硬件需要支持分支的掩码管理和同步。

  • SIMT 模型更适合程序员，因为它隐藏了底层硬件的 SIMD 性质，提供了“线程独立”的编程体验。

• SIMD（单指令多数据）：

  • 硬件上执行同一条指令，并对一组数据进行并行操作。所有执行单元（PE）严格同步，不能分支。

  • SIMD 的硬件实现简单，但对程序员要求更高，因为程序员需要显式管理数据并行。

乘影的特殊性：融合 SIMT 和 SIMD

• SIMT 的表象：

  • 程序员编写的是以线程为单位的代码，每个线程逻辑上可以独立。

  • 硬件通过 SIMT 调度器隐藏线程间的同步和分支差异。

• SIMD 的底层实现：

  • Warp 内的线程被组织成一个向量，使用 RVV 向量指令执行。

  • 每个 PE 同步运行相同的指令，只有在分支时通过掩码处理例外。

CTA Scheduler 和 Warp Scheduler 的区别与联系

Warp Scheduler

Warp Scheduler在GPU架构中的作用是在更宏观的层面上管理和调度warp（线程束）的执行，而不是严格局限在指令处理的五个经典阶段（取指、译码、发射、执行、写回）中的某一个。Warp Scheduler的工作跨越了这些阶段，其主要职责包括：

任务分配：Warp Scheduler负责接收来自上层（如操作系统或运行时环境）的任务，并根据GPU的执行单元和资源情况分配这些任务给不同的流多处理器（Streaming Multiprocessor, SM）。

资源管理：它管理SM内的资源，包括寄存器、共享内存（Shared Memory）和执行单元，确保warp的有效执行。

调度warp执行：Warp Scheduler决定哪些warp应该被发射到执行单元。当一个warp因等待内存操作或其他原因被阻塞时，Warp Scheduler会切换到另一个准备好的warp，以保持执行单元的忙碌和资源的充分利用。

处理分支和同步：在SIMT（单指令多线程）架构中，Warp Scheduler还负责处理线程执行中的分支和同步操作，确保warp中的所有线程在需要时能够正确地汇合（join）或分叉（branch）。

掩盖延迟：通过在多个warp之间快速切换执行，Warp Scheduler帮助掩盖内存访问延迟，提高执行效率。

因此，Warp Scheduler的功能是贯穿整个指令执行流程的，它在指令被取指、译码、发射到执行单元以及执行过程中都在发挥作用。Warp Scheduler的主要作用是在宏观层面上优化warp的执行顺序和资源分配，以提高GPU的整体性能和响应能力。

定义和职责

CTA Scheduler 和 Warp Scheduler 是 GPGPU 中两层调度机制的关键部分，分别负责不同粒度的任务管理：

调度器	CTA Scheduler	Warp Scheduler
调度粒度	线程块（CTA, Cooperative Thread Array），即多个线程组成的任务单元	线程束（Warp），通常是固定数量（如32）的线程集合
职责	分配线程块到不同的流多处理器（SM）。	在单个 SM 内部调度线程束的执行顺序，管理线程束间的切换。
资源管理	分配共享存储、寄存器等硬件资源给线程块。	分配和切换执行硬件单元（如 ALU、寄存器）给不同的线程束。
调度策略	基于线程块的优先级、资源需求等分配到 SM。采用轮询（Round-Robin）的方式，将WorkGroup分配到可用的SM（流多处理器单元）上。	基于指令依赖性、访存延迟等动态调整 warp 的执行顺序。采用贪婪策略（Greedy）和轮询策略（Round-Robin）相结合的方式，选择当前最适合执行的warp。当ICache发生miss或指令缓冲区已满时，会将当前warp的PC回退并切换到下一个warp。

联系

工作流程的层次性：

CTA Scheduler 负责宏观调度，将线程块分配给 SM。

Warp Scheduler 负责微观调度，在 SM 内具体管理线程束的执行。

资源管理的依赖性：

CTA Scheduler 通过为线程块分配 SM 的资源（如共享内存、寄存器），为 Warp Scheduler 提供运行环境。

Warp Scheduler 则根据 CTA Scheduler 分配的资源在内部调度 warp 的执行。

数据共享与同步：

Warp 在 SM 内共享线程块的资源（如共享存储）。CTA Scheduler 确保这些资源分配合理，以支持多个 Warp 的正常运行。

线程块间同步由 CTA Scheduler 通过全局控制指令实现，线程束间同步由 Warp Scheduler 利用栅栏（barrier）实现。

二者的具体工作与例子

3.1 CTA Scheduler 的工作流程

假设需要执行一个矩阵乘法内核，线程块数量为 64，每个线程块有 256 个线程：

主机提交任务：

主机通过驱动程序提交内核任务，并提供元数据，包括线程块数量、线程数量、共享存储需求、寄存器需求等。

线程块分配：

CTA Scheduler 评估每个 SM 的可用资源，将线程块分配到多个 SM。

每个线程块被分解为 256/32 = 8 个线程束。

初始化资源：

CTA Scheduler 为每个线程块分配共享存储（Shared Memory）和私有寄存器（vGPR、sGPR）。

发出执行请求：

分配完成后，将线程块和 warp 列表发送给目标 SM，等待 Warp Scheduler 接管。

3.2 Warp Scheduler 的工作流程

在某个分配了 8 个 Warp 的 SM 内：

线程束初始化：

Warp Scheduler 从 CTA Scheduler 接收到的任务列表中激活线程束，并初始化线程束的寄存器和程序计数器（PC）。

动态调度：

Warp Scheduler 持续监控每个线程束的状态，优先调度已准备好执行的线程束。

如果某个线程束因访存延迟或依赖其他线程的数据而被阻塞，Warp Scheduler 会快速切换到另一个线程束执行（类似超线程）。

执行与切换：

每个周期，Warp Scheduler 向硬件功能单元（如 ALU、FPU）发射指令。

调度过程中 Warp 间的切换不需要重新分配硬件资源，确保高效利用计算单元。

总结