通用及图形处理器架构与系统

计算机抽象与技术

• 冯诺伊曼架构要素
• ISA：指令集架构，定义了机器语言指令的集合，包括指令格式，寻址模式，数据类型等，计算机硬件和软件之间的接口
  • RISC：精简指令集计算机
    • 指令数量少，格式和寻址模式比较简单，可以通过硬布线控制逻辑
    • 处理器能够以更高的频率运行，提升整体性能
    • 对编译器优化要求较高
    • 功耗较CISC较低
  • CISC：复杂指令集计算机
    • 丰富的指令集，复杂的指令，完成复杂操作，通常需要使用微程序控制实现复杂指令的解释和执行
    • 软件编程更方便，减少完成任务所需的指令数量
    • 每条指令要多周期
    • 对编译器的依赖相对较小
• 性能计算
  • CPI(Cycles Per Instruction)
  • $CPU\ time=\frac{{CPU\ Clock \ Cycles}}{Clock\ rate}=CPU\ Clock\ Cycles \times Clock\ Cycle\ Time$
  • $CPU\ Clock\ Cycles=Instruction\ Count \times CPI$
  • $Instruction\ Count=\frac{Instructions}{Program}$
  • $CPI=\frac{{Clock\ cycles}}{Instruction}$
  • $Clock\ Cycle\ Time = \frac{Seconds}{Clock\ cycle}$
  • 加权平均CPI：每个指令占总指令比例之后乘以对应的CPI并加和
• 影响性能的因素
  • 算法：影响 Instruction Count 和 CPI
  • 编程语言：影响 IC 和 CPI
  • 编译器：影响 IC 和 CPI
  • ISA: 影响 IC ，CPI，$Clock\ Cycle\ Time$
• 改变上面的量可以提升性能
• 并行
  • 指令级并行（ILP）：处理器执行程序时，同时执行多条指令
  • 数据级并行（DLP）：对大量数据进行相同操作时，可以同时对多个数据项进行处理
  • 任务级并行（TLP）：将一个大任务分解为多个子任务，在不同处理器/核心上同时执行

指令

• 机器语言：
  • 指令以二进制形式编码，称为机器码
  • 用于计算机内部通信
• 汇编语言：
  • 用特定符号表示及其指令
• 指令编码的方法与思路
  • 简单源于规整：固定操作数数量的硬件比操作数数量可变的硬件更复杂
  • 更少则更快：寄存器的个数需要32个而不是更多
  • 加快常见的情况：因为比较小的常数在指令中比较常见，因此加入立即数指令比较合理
• 大端序与小端序
  • 大端序：最低地址处存放的是最高有效字节
  • 小端序：最低地址存放的是最低有效字节
  • 数字的书写是从最高有效位（MSB）开始，而在内存中最高有效位却位于较高的地址处
• 寻址方式
• 跳转方式
  • PC存储目前指令的地址
  • 条件分支：bne、beq（比较近）
    • 相对PC寻址：目标地址为PC+offset
    • $offset=\{ Imm[12:1],0 \}=2^{13}$ ，13 位有符号数 ，因此 offset 范围是 $[-2^{12},2^{12}-1]$ ，$\pm4KiB$
    • 4 字节对齐（指每条指令占 4 个字节，因此合法指令地址只能是 4 的倍数）
  • 无条件跳转：jal（跳转到很远的地方，如果超了，汇编器会修改代码）
    • 相对PC寻址：目标地址为PC+offset
    • $offset=\{ Imm[20:1],0 \}=2^{21}$ ，21 位有符号数 ，因此 offset 范围是 $[-2^{20},2^{20}-1]$ ，$\pm1MB$
    • 在 rd 寄存器中记录 PC+4
    • 对于过程调用，jal起到跳转+链接的作用
  • 绝对 / 寄存器间接无条件跳转：jalr
    • 直接寻址：目标地址为 rs1+offset
    • 设置offset的方法：12bit 立即数/或者使用 lui 指令将值放入立即数的高20位
    • 在 rd 寄存器中记录 PC+4
  • 写入具体值：相对 PC 寻址不要忘记×4，以及正负号的问题
• 函数调用
  • 栈指针 $x2
  • Caller：给 callee 提供参数
    • 调用前：跳转+链接
      • jal x1, ProcedureLabel
      • 把调用函数的下一条指令地址放在 $x1
    • 参数放入 x10-x17
    • 如果临时寄存器里有值，需要将其先保存到栈中
    • 调用后：返回到对应的地址
      • jalr x0, 0(x1)
      • 把 PC 置为之前存入 $x1 的地址（也就是调用函数的下面的指令）
  • Callee：被调用的程序，根据参数进行对应的运算
    • 参数寄存器 x10~x17 ，传递参数或返回值
    • 如果不是叶子Callee，需要在函数的开头保存其返回值（因为如果在这个函数里调用它自己，这个值会发生变化，出现错误）
    • 如果用到了保存计数器，也要在函数开头保存
      • 过程调用中保存的对象，若软件依赖PC，也要保存
  • 函数调用时的栈
    • FP：帧指针（指向给定过程的局部变量和保存的寄存器地址的值，主要作用是为局部变量引用提供一个稳定的基址寄存器）/SP：栈指针
• 内存布局：冯诺依曼架构
  • RISC-V分配
    • 栈：局部数据，过程帧
    • 动态数据：堆，动态内存管理，内存泄漏，空指针
    • 静态数据：全局变量，全局指针在程序开始时被初始化到这个段的开头，方便对这里面的数据进行访问
  • 指令和数据一样被保存为二进制数字
  • 程序可以在程序上运行
  • 二进制兼容性允许编译后的程序在不同的计算机上运行
• 编译器
  • 传统编译器
    • 将源代码翻译成机器代码，并生成可执行文件。
    • 包括词法分析、语法分析、语义分析、中间代码生成、代码优化以及目标代码生成等
    • 执行机器代码，执行效率较高。
  • 即时编译器（JIT）
    • 常用于解释型语言
    • 在程序运行时，将字节码编译成机器码来执行
    • 即时编译， 提高程序的执行效率。
• 不同的程序语言
  • 编译型语言
    • 源代码直接翻译成机器码
    • 编译过程复杂，分为多个阶段
    • 执行效率高
    • 可执行文件与机器架构相关，在不同架构的计算机上运行时，需要重新编译源代码
    • C/C++
    • 高效，并且能提前发现错误，但编译耗时长，可移植性差
  • 解释性语言
    • 源代码被编译为中间代码（如 Java 的字节码），这种代码与平台无关
    • JVM解释器执行字节码，并识别频繁执行的代码（热点代码）解释性语言依赖解释器执行
    • JIT编译器将热点代码编译为本地机器码，在后续调用中直接执行
    • Java/Python/Javascript
    • 开发效率高，具有良好的跨平台性，但执行效率相对较低，并且依赖解释器

计算机的算数运算

• 定点数据的表示方法：补码=反码+1
• 浮点数据
  • 分为符号位(1bit)，小数(single:8bits/double:11bits)以及指数位(single:23bits/double:52bits)
  • $x=(-1)^S\times(1+Fraction)\times2^{Exponent-Bias}$
  • 注意，最后呈现出的指数位是无符号数（恒为正数）
  • single: Bias=127/Double:Bias=1203
  • 注意指数位不能是全零或全1（用于表示非规格化浮点数）
• 非规格化浮点数
  • 有时允许非规格化的浮点数，在下溢的时候
• 浮点数据计算
  • 加
    • 对齐：小数与大数对齐
    • 对阶（对齐）时两个阶码只差的绝对值大于等于25，则无需进行后续处理，多出来的一位主要是因为IEEE 754的舍入规则。
  • 乘
    • 需要注意指数偏移量的加减
    • 最后需要对符号位进行处理

处理器

• 指令级并行（ILP）：处理器同时处理多条指令
• 流水线：将指令的执行过程分解为多个阶段，每个阶段完成指令执行过程中的特定任务
• 流水线设计与结构
  • 流水线性能分析
    • 如果流水线的切分是均匀的，那么满足 $Time\ between\ instructions_{pipelined}=\frac{{Time\ between\ instructions_{nonpipelined}}}{number\ of\ stages}$
    • 若不均匀，加速的程度就会减少
    • 加速源于吞吐量的提升，但实际上每条指令化的时间不变
  • 冒险
    • 结构冒险
      • 如果数据内存和指令内存是同一块内存（冯诺依曼架构），那么下图中的IF就需要延迟一个周期
      • 哈佛架构分离数据内存和指令内存，解决了这一问题
    • 数据冒险
      • RAW(Read after Write)写后读：即后条指令要读取的数据依赖于前条指令的写操作。
        • EX 冒险
          • 在EX/MEM级如果要向寄存器堆中写入并且写入的寄存器编号不是x0
          • 且EX/MEM级要写入的寄存器名字与此时ID/EX级的rs1或rs2名字相同
          • 解决方法：前递
        • MEM 冒险
          • 在MEM/WB级如果要向寄存器堆中写入并且写入的寄存器编号不是x0
          • 且MEM/WB级要写入的寄存器名字与此时ID/EX级的rs1或rs2名字相同
          • 解决方法：前递
        • Load_use 冒险
          • ID/EX级要求从内存中读到值并写入寄存器堆
          • 写入寄存器堆的地址正好与IF/ID级的rs1或rs2名字相同
          • 解决方法：插入一个空指令或者对指令的顺序重新排列
      • WAR(Write After Read)读后写冒险：一个指令尝试写入寄存器之前，另一条靠后的指令却先读取了它
      • WAW(Write After Write)写后写冒险：一个指令尝试写入寄存器之前，另一条靠后的指令先写入了同一个寄存器
      • 在 RISC-V 五级流水线中，WAR（Write After Read）和 WAW（Write After Write）冒险实际上不会真正发生，因为RISC体系结构中，寄存器读写明确，在IF阶段读取寄存器值，在WB阶段写回，且指令操作简单
        • 由于WB是流水线的最后一个阶段，读操作总是发生在写之前，因此不会出现WAR
        • 由于五级流水线中指令按照顺序进入流水线，且写回是最后一个阶段，因此前面的指令会先于后面的指令进入写回阶段，从而避免了WAW
    • 控制冒险
      • 在运行跳转指令之前，并不知道要不要跳转
      • 因此可以先插入一个空指令，等待正确的跳转地址计算完了再继续运行
      • 或者进行分支预测：假设不跳转，若跳转的时候对已经进入流水线的指令进行冲刷，之后继续运行。相比直接插入空指令，这提升了跳转的准确率
      • 若流水线太长，可能无法迅速确定分支结果，阻塞浪费的时间变得不可接受，因此要尽早预测分支的结果，并且仅在预测错误的时候刷新流水线
  • 解决冒险的办法
    • 转发/旁路
      • 通过上述可以用转发的情况产生转发指令，通过一个选择器来选择正确的数据，之后输入ALU
      • 分支指令如果遇到了上一条尚未写入，也需要进行转发，前递单元可以在ID或者EX
    • 插入空指令
      • 将ID/EX级的control values都置为0
      • 保持PC以及 IF/ID的寄存器不变
    • 冲刷
      • 只将ID/EX级的control values都置为0
    • 分支预测
      • 静态：预测向后分支被执行，向前分支不被执行
      • 动态：记录每个分支的最近历史，并假设未来行为会继续这个趋势，出错时刷新同时更新历史
        • 使用BHT(branch history table),通过PC索引对应的表项（是否跳转）
          • 1bit：记录上次的跳转，并按照上次的跳转情况进行预测。但当一个条件分支总是跳转，只有一次不跳转的时候，就会造成两次预测错误
          • 2bit：纠正上述缺点
• 中断与异常
  • 中断(Interrupt)：处理器外部事件引发的控制流变化
  • 例外(Exception)：CPU导致的控制流变化
  • 处理中断：向量式中断
    • 用基址寄存器加上例外原因（偏移量）作为目标地址完成控制流转换
    • 目标地址要么直接处理中断，要么跳转到真正的处理部分
  • 异常处理机制
    • 精确异常
      • 异常处理机制能准确定位引发异常的指令，
      • 处理异常时，程序的状态与该指令单独执行时的状态一致（除去后续指令对其的干扰）
      • 过程
        • 保存PC到SEPC
        • 保存出现问题的原因（例外原因）到SCAUSE
        • 将流水线的取值阶段，将指令变为NOP，使用ID.Flush/EX.Flush控制mux，使得ID/EX和EX/MEM中的寄存器为0（WB阶段最好也置零一下）
        • 跳转到预定义的处理程序（某条特定指令）
          • 该程序读取原因，并把请求发送给对应的处理程序（读取SCAUSE寄存器以确定异常类型）
          • 确定之后的操作
            • 可重试
              • 处理程序执行
              • 重新获取并执行指令
              • 使用SEPC回到程序
            • 否则
              • 中止程序
              • 使用SEPC，SCAUSE报告错误
      • 处理流水线中的多处例外：处理最早的一个，并刷新后续置零
    • “不精确异常”
      • 不精确地定位引发异常的指令
      • 采取更简单的硬件设计，让异常处理程序来解决复杂的问题。
      • 过程
        • 停止流水线并保存状态
        • 让处理程序来处理异常（可能要手动完成某些指令的执行）
      • 简化了硬件，但处理程序更复杂了
      • 对复杂的多发射，乱序执行流水线不可行
• 提高处理器性能的方式：并行
  • 时间并行：
    • 使用流水线，减少每一级的操作，获得更短的时钟周期（更深的流水线）
    • 但当冒险与中断到来之时会带来很多时间上的延迟
  • 空间并行
    • 使用更多核或者工作单元
    • 可能会遇到资源不能在多个核之间同步的问题
• 可以采用多发射来同时增加空间和时间的并行性，每个时钟周期启动多条指令，CPI<1，可以用IPC衡量
  • 但数据依赖会减少这种并行性
• 多发射
  • 静态多发射
    • 过程
      • 编译器将指令打包在一起进行发射，并检测依赖关系，避免冒险。
      • 将一个发射包看作一个超长指令字
    • 结构修改：Registers加两个输入端口，并复制一份ImmGen和ALU
      • 两个指令包/流水线，一个用于ALU/branch 指令，另一个用原图load/store指令
      • 指令为64bit对齐，在顺序上先是ALU/branch，之后是load/store。（为了简化）用nop填充未使用的指令槽
    • 冒险：
      • EX 数据冒险
        • 前递可以解决包和包之间的问题
        • 若一个指令包中的ALU结果在load/store指令中被使用，需要将指令分成两个包
      • Load_use冒险
        • 仍然是一个周期的使用延迟，但现在受到影响的指令有两个（数据包）
    • 指令调度
      • 可以把Loop展开(Loop Unrolling)
      • 之后根据流水线进行调度
        • 除去重复的跳转指令
        • 注意这里可能出现WAR(读后写的问题，比如前面的load指令还没有从数据内存中取出值，就急吼吼地用那里的不知道啥值进行了计算)因此应该合理安排load和之后指令的顺序
        • 数据争用要以两条指令为单位
        • 遇到修改存储地址的指令时，要对其前后存储有关的指令的偏移量进行修改
      • 之后考虑多发射问题
        • 需要注意冒险之类的问题
  • 动态多发射
    • CPU检查指令流并选择每个周期要发射的指令，编译器可协助，但主要是CPU在运行时解决冒险问题
      • 超标量处理器，CPU决定每个周期发射多少指令以避免结构冒险和数据冒险
      • 允许CPU无序执行指令以避免停顿，但按顺序提交结果到寄存器
    • 使用动态多发射的原因
      • 并非所有的停顿都可预测（cache）
      • 并非所有的分支都可以预测（分支是动态的）
      • 不同ISA实现有不同延迟和风险
    • 需要更多电力，多个简单核可能更好

层次化存储

• 局部性
  • 程序在任何时候都只访问其地址空间的一小部分
  • 时间局部性：程序中某条指令或数据被访问过，在不久的将来会再次被访问
  • 空间局部性：如果程序访问了某个存储位置，在不久的将来可能访问附近的存储元素

Cache

• 概念
  • 一种速度比主存（如 DRAM）快得多的存储芯片（通常采用 SRAM）
  • 位于 CPU 和主存之间
  • 缓解 CPU 和主存之间在速度上的巨大差异，让 CPU 能够更高效地获取数据和指令。
• 内部结构

概念	作用	通俗比喻
缓存块大小 (Block Size)	缓存一次性从主存读取数据的最小单位。	书架上每一格能放多少本书。
缓存块数量 (## of Blocks)	缓存中总共有多少个这样的存储单位（块）。	整个书架总共有多少个格子。
偏移量 (Offset)	在一个缓存块内部，精确定位到所需的字节。	在书架的一格里，找到你要的第几本书。
索引 (Index)	决定一个主存块应该被放置到缓存的哪个位置。	根据书的编号规则，决定它应该放在哪个格子里。
标记 (Tag)	验证当前缓存块中的数据是否是CPU真正想要的那个。	确认格子里放的书，是不是你要找的那一本（因为很多书可能都按规则要放进同一个格子）。

• 块大小的考量
  • 优点：更大的块可以降低Miss Rate
  • 缺点：
    • 块的数量变小，竞争加剧
    • 缓存污染：大块数据里可能有很多程序用不到的数据
    • Miss Penalty更大，一旦发生未命中，CPU等待数据的时间会更长
• 参数计算：当题目给出总缓存大小、块大小和主存地址时，可以遵循以下步骤来拆分地址。
  1. 偏移量位数 = log₂(块大小)
     • 例如：块大小为 64 字节 (B)，则偏移量需要 log₂(64) = 6 位。
  2. 缓存块数量 = 总缓存大小 / 块大小
     • 例如：总缓存为 4KB，块大小为 64B，则缓存块数量为 4096 / 64 = 64 个。
  3. 索引位数 = log₂(缓存块数量)
     • 例如：缓存块数量为 64 个，则索引需要 log₂(64) = 6 位。
  4. 标记位数 = 地址总位数 - 索引位数 - 偏移量位数
     • 例如：在一个32位地址系统中，标记位数等于 32 - 6 - 6 = 20 位。
  • 最右边的 偏移量位数 对应 Offset。（注意这个是以字节为单位的，因为risc指令集中也有按字节取数的）
  • 往左接着的 索引位数 对应 Index。
  • 剩下最高位的部分就是 Tag。 示例：
  • 地址: ( ... 标记 ... ) ( ... 索引 ... ) ( ... 偏移量 ... )
  • 位数: ( 20位 ) ( 6位 ) ( 6位 )
• 基本结构
  • 直接映射
    • 缓存块位置 = (主存块地址) % (缓存中的总块数)
    • 简单快速
    • 冲突率高
  • 全相联
    • 数据随便放
    • 为了找到数据，系统需要搜索缓存中的每个条目
    • 冲突miss rate最低，因为缓存满时，数据块才会发生冲突
    • 构建成本高，复杂，需要每个块都有一个比较器电路
  • 多路组相联
    • 缓存分为多个组，每个组包括多个路，n路组相联每组有n个块
    • 放置规则：一个内存块被映射到一个特定的组，但它可以被放置在该组内的任意路上。
      • 地址的索引位 (Index bits) 决定了数据块应该去往哪个组。公式是：组号 = (块地址) % (组的总数)。
      • 数据块可以被放置在所选组内的任何一个空的“路”中。
    • 查找过程：
      • 索引 (Index) 指向一个特定的组。
      • 同时搜索该组内的n个块，需要n个比较器。
    • 显著减少了冲突未命中率
    • 比直接映射缓存更复杂
• 读/写Cache
  • 读
    • 命中时CPU正常工作
    • 缺失时
      • 阻塞CPU流水线，并从下一级内存中获得对应内存块
      • 如果是指令缓存缺失，重启IF模块
      • 如果是数据缓存缺失，重新完成数据访问
  • 写
    • 命中
      • 写通（Write-Though）：数据同时写入缓冲区和主存
        • 实现简单，保持缓存和主存的数据一致性
        • 速度较慢，写入需要等待速度较慢的主存完成操作，会导致CPU停顿
        • 可以使用写缓冲，CPU写入缓冲后可以继续执行，缓冲区负责在后台将数据写入主存
      • 写回（Write-Back）：数据只写入缓存而不写入主存
        • 写入缓存后标记为脏块
        • 被替换时，脏块的内容被写回到主存中
        • 写入速度快，不受主存限制，对于连续多次写入的场景，可以减少对主存的访问，节省带宽
        • 实现复杂，数据写回之前，主存中的数据过时，在多处理器中需要额外的机制保持一致性
    • 未命中
      • 写分配
        • 向将写入的数据库从主存加载到内存中，之后成为命中的情况
        • 利用了空间的局部性
      • 非写分配（写绕过）
        • 数据直接写入主存（与写通策略搭配）
        • 适合程序需要初始化内存区域时（连续写入，短期不会读取）
• 替换算法
  • 直接映射：直接替换
  • 组相联
    • 若组内路有空，直接补空
    • 没有空
      • LRU算法：替换最长时间没有访问的块
        • 相联度高的时候很难追踪
      • 随机替换：随机选择进行替换
        • 相联度高的时候跟LRU比较接近
• 3C与4C模型
  • 强制性未命中 (Compulsory Misses)/冷启动未命中 (Cold Start Misses)。
    • 数据块第一次被访问时，它肯定不在缓存中，必须从主存加载。这种未命中是不可避免的。
    • 随着程序运行，这类未命中会越来越少。
  • 容量性未命中 (Capacity Misses)
    • 缓存的容量有限，无法容纳程序需要使用的所有数据（即工作集大于缓存容量）。
    • 一个数据块被加载进缓存，但后来因为缓存空间不足被替换出去，之后程序又需要访问它，从而导致未命中。
    • 即使是全相联缓存（最灵活的缓存），只要容量不够，这类未命中依然会发生。唯一的解决方法是增大缓存容量。
  • 冲突性未命中 (Conflict Misses)/碰撞未命中 (Collision Misses)。
    • 在直接映射或组相联缓存中，多个不同的主存块可能会被映射到同一个缓存位置（或同一个组）
    • 即使缓存中还有很多空闲空间，但程序需要交替访问的两个数据块正好映射到了同一个位置，导致它们互相“踢出”对方，从而产生未命中。
    • 这种未命中是由于映射规则的限制造成的。在相同容量的全相联缓存中，这种未命中不会发生。
  • 一致性未命中 (Coherence Miss)
    • 这一类未命中专属于多处理器或多核系统，它是由维持缓存一致性 (Cache Coherence)的协议所引起的。
    • 在一个多核系统中，同一个数据块可能会被多个不同的核心缓存。当一个核心（比如核心A）修改了这个数据块，缓存一致性协议为了保证所有核心看到的数据都是最新的，会强制其他核心（比如核心B）中该数据的旧副本（copy）失效 (invalidate)。
    • 当核心B再去访问那个已被标记为失效的数据时，就会发生未命中，必须重新从内存或其他核心的缓存中获取最新版本。这个为了数据同步而导致的“被动”未命中，就是一致性未命中。
    • 与冲突性未命中不同，一致性未命中是为了保证数据正确性而付出的必要代价。
• 性能计算
  • CPU时间
    • 程序执行周期（包括缓存命中时间）
    • 内存停滞周期（缓存未命中）
  • AMAT（平均内存访问时间）
    • AMAT=Hit Time+Miss rate × Miss panelty
    • $\text{Memory stall cycles} = \frac{\text{Memory accesses}}{\text{Program}} \times \text{Miss rate} \times \text{Miss penalty} = \frac{\text{Instructions}}{\text{Program}} \times \frac{\text{Misses}}{\text{Instruction}} \times \text{Miss penalty}$
  • 有效CPI = 基础CPI + 全局未命中率
    • 全局未命中率=(L1未命中率 × L2命中时间) + (L1未命中率 × L2局部未命中率 × 主存惩罚)
    • Performance ratio: 新CPI/原CPI

主存与虚拟存储器

• 主存
  • 使用DRAM，其中的数据以电荷形式存储在电容器中，需要一个晶体管访问。由于电荷会泄漏，需要周期性地刷新
  • 组织成一个矩形阵列，通过突发模式提供行中的连续字
• 虚拟存储器
  • 虚拟地址：虚拟内存中的地址
  • 物理地址：在主存中的真实地址
  • 虚拟内存（一种数据管理机制）
    • 由MMU（内存管理单元）硬件和OS共同管理
    • 将主存用作辅助存储（磁盘）的缓存
  • 地址转换：CPU生成的虚拟地址分为虚拟页号（→查找物理页号）和页内偏移量（与物理页号一起组合形成物理地址）。
    • 页内偏移量表示页的大小，而物理地址/虚拟地址的总位数代表了对应内存的大小
    • 页也就是一块数据
• 页表（虚拟内存系统用来实现地址翻译的核心数据结构）
  • 存放在物理内存中
    • CPU内部有一个特殊的页表寄存器，存放着当前进程页表的起始物理地址，操作系统切换进程时，会更新这个寄存器
  • 存储虚拟地址到物理地址的映射信息
  • 这些映射信息组成了页表条目PTEs
  • 每个进程有其独立的页表，由OS进行管理，用户进程不能修改
    • 进程包括其PC，寄存器与页表
  • 工作过程
    • 有效位（valid）表示虚拟页是否在内存中
      • 有效位为1时，可以将虚拟页号转换为物理页号
      • 有效位为0时，表示虚拟页不在物理内存（主存）中，访问时会触发页错误，存储的地址不是物理地址，而是页在磁盘上的存放位置（也就是交换空间）
    • 页错误
      • OS执行页错误处理程序
        • 首先找到导致错误的虚拟地址，找到页表中的条目PTE
          • 取指令导致的：虚拟地址保存在SEPC中
          • 读写数据导致的：从SEPC中找出来指令并且计算
        • 磁盘上找到页，在物理内存中用替换算法找一个替换
          • 如果被替换的页是脏的，需要先写回磁盘交换空间
        • 之后加载新页，回复进程，重启指令
      • 如何避免页处理器自身发生页错误？（双重错误）
        • 将页处理器的核心代码和异常堆栈放在不经过地址映射的内存区域中，直接使用物理地址，不会发生页错误
      • 需要尽量最小化页错误率（使用智能替换算法和全相联的策略实现）
      • 需要数百万个时钟周期，一般会在这时切换进程
• TLB
  • 由于对页表的访问也需要引用内存，会降低效率，但对页表的访问有比较好的局部性，因此引入TLB作为页表条目的cache，称为Translation Look-aside Buffer, TLB。
  • 目标：加速从虚拟地址到物理内存地址的转换 ，TLB只会缓存有效位为1的页表条目
  • 位于CPU内部，全相联，随机替换，写回
    • 复制了TLB写访问位，从页表中复制而来，指示对应内存页面是否可写，主要提供保护（若强制写入不能写的内存，会触发写保护异常，之后OS来接管）
  • 未命中
    • TLB未命中，但所需页面在主存中
      • 从内存中加载PTE到TLB，之后重试
      • 可用硬件（复杂）或软件（引发一个特殊异常）进行处理
    • 所需页面不在主存中（page fault）
      • OS执行错误处理程序，并重新执行

数据访问的过程

1. CPU生成一个虚拟地址 。
2. TLB查询: MMU首先检查TLB，看其中是否包含该虚拟页号的映射。
   • TLB命中: 如果命中，则获得物理地址，然后进行缓存查询 。
   • TLB未命中: 如果未命中，则会产生TLB未命中异常。系统会查询页表。若页在内存中，则将PTE加载到TLB后重试；若页不在内存中，则发生页错误，由操作系统处理 。
3. 缓存查询: 使用TLB转换得到的物理地址来访问缓存。
   • 缓存命中: 直接从缓存中读取或写入数据 。
   • 缓存未命中： 从下一级存储（如L2缓存或主存）中获取数据块到当前缓存，然后再将数据提供给CPU 。

TLB 状态	页表 状态	缓存 状态	是否可能？	原因解释
命中	命中	未命中	可能	地址翻译成功，但数据不在缓存中，需要访问主存。这种情况很常见。
未命中	命中	命中	可能	TLB未命中后，从页表加载翻译信息；重试指令后，在缓存中找到数据。
未命中	命中	未命中	可能	TLB未命中后，从页表加载翻译信息；重试指令后，缓存也未命中，需访问主存。
未命中	未命中	未命中	可能	发生页错误。数据从磁盘加载后，在缓存中必然也是未命中的。
命中	未命中	任意	不可能	TLB命中意味着页在内存中，这与页表未命中（页不在内存中）在逻辑上是矛盾的。
未命中	未命中	命中	不可能	页表未命中意味着数据所在的页不在主存中，那么它也绝不可能存在于作为主存子集的缓存中。

并行处理器

• Flynn 分类法
  • SISD：单指令流单数据流
  • SIMD：单指令流多数据流
    • 通过对数据向量进行元素级操作来工作
    • 所有处理器在同一时间执行完全相同的指令，但处理的是不同的数据（称之为空间并行）
    • 简化了处理器之间的同步过程
    • 指令控制相关的硬件减少
    • 利用了数据级并行
  • MISD：多指令流单数据流
  • MIMD：多指令流多数据流
    • 软件层面可以进行任务级并行，通过运行多个独立的任务获得高吞吐量
    • 创新或并发的软件都可以运行在串行或并行的硬件上
    • 难点：任务划分，协作以及通信开销
• 向量处理器
  • 较早的SIMD实现，数据从内存加载到专门的向量寄存器中进行计算
  • 向量代码的效率体现在它可以用一条指令处理整个数据集合（向量）
  • 可以通过向量指令的链接技术使得两个前后相连的指令同时进行：可以对向量中的每个数据进行流水化操作，避免了由于数据相关性造成的流水线停顿
• SIMD扩展指令集
  • MMX (1996)：主要针对整数运算，可以一次性处理八个8位整数或四个16位整数 。（64位）
  • SSE (Streaming SIMD Extensions) (1999)：在MMX的基础上进行了扩展，支持更宽的数据和浮点运算 。它可以处理八个16位整数、四个32位整数/浮点数，或两个64位整数/浮点数 。（128位）
  • AVX (Advanced Vector eXtensions) (2010)：进一步将寄存器宽度扩展到256位，可以一次性处理四个64位整数或浮点数，并行能力翻倍 。（256位）
    • 向量架构
    • 优势：简化编程，硬件简化，内存访问优化，低功耗
• 向量架构与标量架构
  • 向量架构的优势
    • 简化数据，并行编程
    • 显式声明循环体前后的元素没有迭代关系
      • 减少硬件对此的检查
    • 常以连续，规则的方式访问内存，它可以高效地利用交错式和突发式内存传输技术
      • 交错式：物理内存分为多个独立模块。连续内存放在多个模块中。它比较高效，因为单个内存模块完成一次读写之后，需要一小段恢复时间才能接受下一次请求。可以为向量处理器提供平滑，高效的数据流
      • 突发式：CPU只需发送一个起始地址，内存控制器会自动连续地传输接下来几个地址的数据块，省去了后续几次的地址传输开销
    • 避免循环，因此避免了控制冒险
    • 完成同样的工作需要更少的指令，在功耗和能源上更具优势
    • 比专用的多媒体扩展指令集更通用
    • 通用性和规则性使得编译器更容易识别代码中的并行性，并将其转化为向量指令
• Amdahl's Law
  • 有关改变前后的加速
  • $Speedup=\frac{1}{1-F+\frac{F}{S}}$
  • $F$ 可以改变的部分占总时间的比例
  • $\frac{F}{S}$ 改变后时间变成改变前时间 $F$ 的 $\frac{1}{S}$ 倍
  • 相关计算：增加处理器数量仅仅能提高并行部分的性能，当可并行的工作量远大于串行工作量时，系统可以获得更好的加速效果和更高的效率
  • 扩展性模型
    • 强扩展：问题总规模固定不变，通过增加处理器数量来缩短计算时间，以追求更快的速度
    • 弱扩展：在增加处理器数量的同时，按比例增大问题的总规模，使得每个处理器上的计算负载保持不变，以在处理更大的问题的同时维持恒定的计算时间
• 多处理器（属于MIMD）
  • 至少包含两个处理器的计算机系统
  • 多核微处理器：单个芯片上集成多个处理器
    • 单芯片，多核心
    • 多芯片，多核心
• SMP（shared memory multiprocessor）共享内存多处理器架构
  • 硬件为处理器提供一个统一的物理地址空间
  • 处理器通过内存中的共享变量进行通信
  • 用“锁”机制确保数据同步
  • 内存访问时间
    • UMA（统一内存访问）：任何处理器访问任何内存模块的时间都相同
    • NUMA（非统一内存访问）：处理器访问本地内存和远程内存的时间不同
  • 浮点数的加法不一定满足严格的结合律，改变加法的顺序可能因为精度误差累积而导致最终结果有差异。
• Cache一致性
  • 多个处理器核心共享一个物理地址空间时，会出现一个核更新了缓存而另一个核不知道的情况，这时另一个核中的数据已经过时
  • 解决：缓存一致性协议
    • 为确保一致性的操作
      • 数据复制：多个处理器需要读取同一份数据时，给每个处理器的本地缓存复制一份，以实现高效率的并行读取
      • 数据迁移：某个处理器要对一份数据进行频繁写入时，执行迁移操作，将数据的所有权和最新版本转移到该处理器的本地缓存中，于是该处理器可以直接在cache中进行读写，无需与主内存通信，减少对共享内存的范围带宽
    • 监听协议：每个缓存监听总线上的读写活动，判断自己是否需要更新或废弃缓存中的数据
      • 无效化监听协议：缓存对数据块进行写操作的时候，首先获得对该块的独占访问权，在总线上广播一个无效化消息，其他所有拥有该数据块的缓存接到消息，将自己的副本标记为无效
    • 目录协议：使用一个集中的目录来记录每个内存块的共享状态
  • 另一种解决方案：每个处理器拥有自己的私有物理地址空间，不能直接访问对方的内存，处理器之间通过硬件提供的显式发送/接收消息的操作进行通信
• I/O
  • 强调可靠性（而不像处理器/内存关注性能和成本）
  • OS在I/O管理中扮演核心角色，因为有多个程序共享I/O资源，且I/O操作会导致异步中断
    • OS发给I/O命令
    • 设备在I/O发生错误时通知操作系统
    • 数据在I/O和内存之间传输
  • 由I/O controller hardware管理
    • 命令寄存器
    • 状态寄存器
    • 数据寄存器（进行写入和读取的操作）
  • I/O寻址
    • 内存映射I/O
      • 系统划分一部分内存地址空间给I/O，用load/store进行访问
      • 地址解码器负责区分地址指向物理地址还是指向I/O
      • 操作系统会利用地址转换机制将I/O地址标记为仅内核可见
    • I/O指令
      • 一套专门为I/O设计的指令，利用它们访问I/O
      • I/O设备拥有独立于主内存地址空间的I/O空间
      • 处于内核模式下才能被执行
  • 与处理器通信
    • 轮询：处理器周期性检查I/O设备的状态寄存器，如果准备好了就进行操作，如果出现错误就解决错误
      • 常见于小型/低性能的实时嵌入式系统
        • 时序可预测，硬件成本低
      • 最简单的方式，但浪费CPU时间
    • 中断：设备准备就绪或出错时，向CPU发出错误信号
      • 类似于异常，但不与指令执行同步，可以在指令之间调用处理程序
      • 优先级中断：需要更紧急关注的设备获得更高的优先级，可以中断低优先级中断的处理程序
  • I/O数据传输
    • 轮询和中断驱动的I/O
      • CPU将数据在内存和I/O之间来回搬运
      • 对高速设备而言比较耗时
    • 直接内存访问（DMA）
      • 一种机制，赋予设备控制器直接与内存进行数据传输的能力
      • 无需CPU介入
      • 步骤：处理器向DMA控制器提供参数（设备ID，操作类型，内存地址，传输字节数等）→执行传输→传输完成后通过中断通知处理器
      • 与缓存的交互
        • 因为提供了另一条访问内存的路径，且不经过地址转换或缓存层次结构，造成了与cache数据不一致的问题
        • 解决
          • 如果某些内存块要被用于DMA操作，提前将其从缓存中刷出
          • I/O操作使用不可缓存的内存区域
      • 与虚拟内存的交互
        • 连续的虚拟地址在物理内存中可能不是连续的
        • 若DMA使用物理地址（有可能不连续），单次DMA传输的大小被限制在一个内存页内
        • 若DMA使用虚拟地址，需要DMA控制器具备地址转换能力，控制器变得复杂
        • 解决
          • 将大传输任务分解成页大小的块
          • 将多个小的传输链接起来
          • OS为DMA分配物理上连续的页面

GPGPU

• CPU/GPU
  • CPU多核处理器每个核心都是一个完整的处理器
  • GPU多核每个核心都小得多而且更简单，专用于计算密集型处理而不是缓存与控制
• GPU：为执行图形渲染设计专用硬件
  • 采用固定管线架构，硬件逻辑固定，为图形渲染流程服务
• GPGPU：通用图形处理器
  • 并行计算架构
  • 关键：抽象复杂硬件的编程模型
  • 核心是一个统一的，可编程的处理器阵列，处理单元不是专用的，而是可以灵活低执行图形指令或通用的并行计算指令
• CUDA(Compute Unified Device Architecture)：
  • NVIDIA：并行计算平台和编程模型
    • 如何分配工作量
    • 如何在分离的工作之间进行通信
    • 如何同步划分过的工作量
  • 使用类C语言
• OpenCL(Open Computing Language)
  • 开放，免版税的行业标准，用来对各种异构平台进行并行编程
• 异构执行
  • 异构计算任务同时利用两种或多种不种类型的处理器完成
    • 结合CPU处理复杂逻辑和串行任务的优势，以及GPGPU在执行大规模并行计算上的优势
  • 程序员视角的GPGPU程序
    • 主机端代码 __host__ ：在CPU上运行，负责程序流程控制内存管理和任务调度
    • 设备端代码 __device__/__global__ ：称为内核，在GPGPU上并行执行的计算密集型函数
    • 执行流程：CPU将数据从主机内存复制到GPGPU设备内存中→CPU向GPGPU发出指令启动设备上的内核函数→GPGPU并行执行内核代码→计算完成后GPGPU将结果数据从设备内存复制回主机内容，供CPU进行后续处理或输出
  • 架构师视角的异构执行平台
    • CPU和主机端存储器：负责运行主机驱动程序，处理应用程序的逻辑部分
    • GPGPU：大量计算单元阵列（SM/CU）用于并行处理数据
    • 设备端存储器（GPGPU自带的高速显存），存放待处理的数据和计算结果
    • PCI-E总线：连接CPU和GPGPU的告诉通道，负责在主机和设备之间传输命令指令常数以及输入输出数据
• SIMT与SIMD
  • SIMT (Single Instruction, Multiple Threads) 是GPGPU的核心执行模型，从更传统的SIMD (Single Instruction, Multiple Data)模型演变而来
  • SIMD模型：一条指令同时操作多个数据元素，通常使用宽向量寄存器和专门的向量指令（一个线程）
    • 其并行性由指令本身显式管理
    • 要求待操作的数据在空间上是连续的（要借助其他数据进行重组和拆分）
    • 分支时需指令显式地对活跃掩码寄存器进行设置
    • 需要使用SIMD指令，对线程数有限制，造成指令规模的膨胀
  • SIMT模型：一条指令由多个独立的标量线程执行，每个线程处理自己的数据 。
    • 可以通过独立的数据索引降低空间连续性的需求
    • 分支时通过硬件自动对谓词寄存器进行管理，并允许线程进行分支，实现多种执行路径
    • 对线程数没有限制，硬件管理模式使得SIMT不需设计性的指令
• 线程模型：CUDA中有关线程的三级层次结构
  • Grid（线程网格）最高层级：一次内核启动所要执行的全部线程，由多个线程块组成
  • Thread Block（线程块）：Grid下一级，一组可以相互协作的线程，通异构块内的线程可以通过共享数据并进行同步，线程块会被调度到GPU的流式多处理器（SM）上执行
  • Thread（线程）：最基本的执行单元，负责执行内核代码中的指令，由SM中的流式处理器（SP）执行
    • 一个线程束由最多 32 个线程组成，一般来说，执行完全相同的指令。是GPGPU硬件进行调度和执行的基本单位
• 数据索引
  • blockIdx 线程块在Grid中的索引（分xyz）
  • blockDim 线程块的维度，一个线程块中包含的线程总数
  • threadIdx 线程在其线程块内的索引（分xyz）
  • 计算公式int index = threadIdx.x + blockIdx.x * blockDim.x;
• 线程分支
  • 线程束中的线程遭遇条件分支语句，不同线程需要执行的代码路径不一样
  • 谓词执行
    • 硬件为线程束中的每个线程的每个分支都配备了1bit的谓词寄存器，用来存储条件指令的判断结果
    • 根据前面的计算来判断每个线程的各个分支谓词寄存器是否有效
    • 对不同分支设置不同谓词，只有对应谓词急寄存器有效时，线程执行该分支，否则执行的结果会被屏蔽掉
    • 消除了分支跳转，并且没有分支预测失败的惩罚
    • 硬件更复杂，处理嵌套分支困难，存在冗余计算
  • SIMT堆栈
    • 每一项记录一个分支点的信息
    • 图示算法，参考
      • 活跃掩码：当前线程束中哪些线程需要执行下面的路径
      • NPC：分支指令的下一地址
      • RPC：两条或多条分支路径汇合点的地址
      • 算法
        • 跳转到TOS中NPC，并根据“活跃掩码”（Active Mask）来执行指令。（每个基本块的第一条指令地址即代表该基本块的NPC。）
        • 将TOS的NPC修改为其IPDOM的地址。（由于IPDOM可以在编译时通过控制流图分析确定，所以可以在此时而不是在基本块结束时就进行修改。）
        • 持续执行……直到当前基本块的末尾。在运行时，获取最后一条指令并评估其分支条件（如果最后一条指令是条件跳转）。
        • 如果基本块的最后一条指令是条件跳转，则在SIMT堆栈上为所有分支路径分配新的条目，将新条目的RPC设置为原栈顶条目的NPC，并根据刚才评估的分支条件设置各自的“活跃掩码”。否则，不执行任何操作。
        • 检查当前执行的地址（即当前栈顶条目的NPC）是否与RPC相等，如果相等，则从SIMT堆栈中弹出并移除该条目。（注意：步骤4和5是互斥的！）
    • 传统的SIMT堆栈机制通常将重聚点设置在分支发生处的直接后支配节点（Immediate Post-Dominator, IPDOM） 。但在某些复杂的控制流中，这可能不是最早的可以安全重聚的地方，导致线程保持分支状态的时间过长 。
      • 优化思路：通过更先进的静态或动态分析，找到一个比IPDOM更早的实际重聚点 。这样可以让发散的线程尽快恢复同步，从而提高利用率
    • Thread Frontier：采用该技术后，线程在更早的、局部的重聚点就恢复了同步，显著减少了空闲周期，提高了整体执行效率
    • 分支线程重组
      • 当多个线程束发生分支时，它们内部都会有部分线程处于非活跃状态。
      • 将来自不同线程束、但将要执行相同代码路径（即有相同PC）的活跃线程动态地“凑”在一起，组成一个新的、完全活跃的线程束来执行，可以极大地提高计算单元的利用率 。
• 线程束调度
  • 将工作分配给资源，旨在实现高吞吐量、最短延迟和公平性等目标，在GPGPU中，就是将线程束分到对应的执行单元上
  • 一个SM上允许有48或64个线程束。
  • 与可实现的最大线程级并行度（TLP）有关。
  • 通常有很多线程束可供选择，以充分占用执行单元并掩盖内存访问延迟。
  • 调度策略
    • 轮询：按顺序选择下一个
    • GTO：优先选择能够最快完成的指令，优势也会选择等待时间最长的线程束（防止饿死）
    • GPU性能与线程束调度高度相关，尤其是在计算密集型（大部分时间花在计算上，让计算单元满负荷运转）和内存密集型（性能瓶颈在访存速度，访存时间掩藏是主要任务）的应用中。
    • 局部性仍然很重要，包括线程束内部的局部性和线程束之间的局部性。
    • 但所有线程束行为模式相似（先计算后访存）时，轮询调度会让它们内存请求集中爆发时计算核心陷入等待。
    • GTO调度可能比轮询调度更好，但它可能会破坏线程束间的局部性，从而导致平均内存访问时间（AMAT）变长。
  • 利用两级轮询制度对访存时间掩藏
    • 假设一组指令是先计算后访存
    • 将其分为两组，第一组的计算指令执行时，采用轮询方式
    • 第一组遇到访存指令时，切换到第二组，对其进行轮询
    • 第一组的内存访问与第二组的计算执行重叠，第一组的访存时间被掩藏了
    • 问题：如何合理对线程束分组
  • CPU访存时间与GPU访存时间掩藏
    • CPU 重在减少延迟，使用多级缓存层层设防，避免等待/通过乱序执行先执行后续准备好的指令/硬件预取，提前取需要的指令
    • GPU 关注整体吞吐量，在单位时间内完成尽可能多的计算任务，使用海量多线程与零开销切换/简单但专门化的访存/进行访存合并
• 卷积操作与GEMM
  • 卷积参数
    • $N$ 一次并行处理的输入特征图的数量
    • $C_{in}$ 输入数据的通道数量
    • $C_{out}$ 输出数据的通道数量
    • $H_{in}$ 输入图像高度
    • $W_{in}$ 输入图像宽度
    • $K$ 滤波器（正方形）的高度和宽度
    • $H_{out}$ 输出特征图的高度
    • $W_{out}$ 输出特征图的宽度
    • 步长：滤波器每次移动的步长，决定了输出特征图的尺寸
  • 卷积操作：滤波器扔到输入矩阵里从左到右从上到下移动，相乘，相加，加偏置（学习到的值）并输出
  • 卷积本身涉及到很多计算的嵌套，计算量巨大，把它转化为GEMM（标准矩阵乘法好处多多）
  • 转换方法：im2col
    • 把滤波器滑窗覆盖的位置展开成一行，滤波器自己展开为一列，偏置展开为一列
    • 多通道输入
    • 多通道输出
    • 增多并行处理图的数量

GEMM的CUDA实现

• 朴素实现 Host Code

dim3 threadsPerBlock(16, 16);
dim3 numBlocks((N + threadsPerBlock.x - 1)/threadsPerBlock.x,(M+threadsPerBlock.y-1)/threadsPerBlock.y);//定义线程网格中线程块的数量和布局，本来应该是N/threadsPerBlock.x，但是这种算法无法保留余数，因此必须要这样算
Gemm_kernel<<<numBlocks, threadsPerBlock>>>(A, B, C, M, N, K);

Device Code

__global__ void gemm_kernel(float *A, float *B, float *C, int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; ++k) {
            sum += A[row * K + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

• 共享内存优化
  • 线程块不直接从全局内存中找数据计算，而是直接将A,B的一部分数据加载到共享内存中，被块内的所有线程多次使用，增加了数据的复用率，避免了重复缓慢的全局内存访问

global void gemm_kernel_reg_opt(float *A, float *B, float *C, int M, int N, int K) { 
	shared float tile_A[tb_m][tb_k]; 
	shared float tile_B[tb_k][tb_n];
	int bx = blockIdx.x, by = blockIdx.y;
	int tx = threadIdx.x;
	
	int row_start = by * tb_m;
	int col_start = bx * tb_n;
	
	float Csub[8][16] = {0}; // 每个线程负责的C子块
	
	for (int m = 0; m < (K / tb_k); ++m) {
	    // 加载A和B到共享内存
	    for (int i = 0; i < 8; ++i) {
	        int row = row_start + tx * 8 + i;
	        int col_A = m * tb_k + tx % tb_k;
	        tile_A[i][col_A] = (row < M && col_A < K) ? A[row * K + col_A] : 0.0f;
	
	        int col_B = m * tb_k + (tx * 8 + i) % tb_k;
	        tile_B[col_B][tx % tb_n] = (col_B < K && tx % tb_n + col_start < N) ? B[col_B * N + tx % tb_n + col_start] : 0.0f;
	    }
	    __syncthreads();
	
	    // 计算C的子块
	    for (int k = 0; k < tb_k; ++k) {
	        for (int i = 0; i < 8; ++i) {
	            for (int j = 0; j < 16; ++j) {
	                Csub[i][j] += tile_A[tx * 8 + i][k] * tile_B[k][tx % tb_n + j];
	            }
	        }
	    }
	    __syncthreads();
	}
	
	// 将计算结果写回C
	for (int i = 0; i < 8; ++i) {
	    for (int j = 0; j < 16; ++j) {
	        int row = row_start + tx * 8 + i;
	        int col = col_start + tx % tb_n + j;
	        if (row < M && col < N) {
	            C[row * N + col] = Csub[i][j];
	        }
	    }
	}
}

• 数据流优化：数据保持不动（驻留在寄存器/缓存）中，其他数发生变化，从而最大化数据复用
• Tensor Core：用于执行大规模矩阵加法乘法单元，计算速度很快，从共享内存加载数据到寄存器的延迟称为瓶颈
  • 使用软件流水线，将“从全局内存加载到共享内存”、“从共享内存加载到寄存器”和“进行计算”这三个步骤进行重叠。在一个主循环体内，当核心正在对第 i 块数据进行计算时，可以异步地预取第 i+1 块数据到共享内存。这样，计算的延迟就和内存访问的延迟相互掩盖了，使得计算单元能够持续满负荷工作 。