GPGPGPGPGPU——《计算机系统结构》期末 CheatSheet 兼总结

\documentclass[UTF8,10pt]{ctexart} % 使用 CTeX 中文支持
\usepackage[
    paperwidth=8.5in,
    paperheight=11in,
    margin=10pt            % 去除所有页边距
]{geometry}               % 控制页面布局
\usepackage{multicol}     % 分栏支持
\usepackage{lipsum}       % 示例文本（可删除）
\usepackage{xcolor}
\usepackage{tabularx}
\usepackage{amsmath}
\usepackage{graphicx}
\usepackage{multirow}
\usepackage{geometry}
% 可选：调整栏间距（默认10pt）
\setlength{\columnsep}{1cm} 
\begin{document}
\pagestyle{empty} % 移除页眉页脚

% 自动分栏环境（内容填满第一栏后自动进入第二栏）
\begin{multicols}{2}   
\section*{Caches}

各种存储设备在延时和容量间存在 tradeoff。可以使用 locality（时间和空间上的）来建造 memory hierarchy。

使用 cache。不管是 inst 还是 data 都存储于内存之中，都需要 cache，出现 miss 了就要 stall。但是 OoO 的场合，假如 load 和 store 的延时不同，那么 dependency 的处理会更加复杂。

\textbf{Miss Rate 对 CPI 的影响并非决定性的。}当 miss rate 低或访存指令数目少时，performance 不会受内存延时主导。我们更喜欢 \textcolor{red}{MPKI (Misses Per Kilo Instructions)}，它与实际程序更相关。\texttt{AMAT = hit latency + miss rate × miss penalty}。miss 对 CPI 的影响是 \texttt{CPI penalty = miss rate × miss penalty}。

在 32b 机器上，假设有一个总空间为 $2^n$-byte 的 cache，每个 cacheblock 的大小是 $2^m$-byte，且是 $2^w$-\textcolor{red}{ways} associative 的，则地址的后 $m$ 位会是 offset，中间的 $n-m-w$ 位会是 index，然后高 $32-n+w$ 位是 tag。一个 cache entry 中会包含数据、tag、valid bit 以及可能的更多元数据（常见的包括 write-back 时的 dirty bit、replacement policy 的辅助 bit 等）。此时其构成 $2^{n-m-w}$ 个 \textcolor{red}{set}，也即，\#set 与 index bit 的关系是对应的。

产生 cache miss 的 3C 原因：\textcolor{red}{compulsory/cold}，第一次访问的必要开销，解决方案是 prefetch；\textcolor{red}{capacity}，因为 cache 不够大导致的，开大空间即可；\textcolor{red}{conflict}，因为冲突覆盖导致的。

cache 的大小要与代码匹配。如果 cache 太小就会出现 \textcolor{red}{thrashing}（这就是 conflict 的体现），此时后来的数据会直接消除已有的数据，cache 完全没有效果。解决方案包括开大 associativity、调整数组大小（让差值不要是 2 的整数幂次）或使用复杂 hash 函数。

更大的 cache block 会增大 miss penalty—因为要传输更大的数据块。但是它可以减少 miss rate 和 thrashing。但是在固定 cache 总容量时，它反倒会增大 miss rate。

\textcolor{red}{write-through}：将数据写入内存。可以加一个 \textcolor{red}{write-buffer}，让内存慢慢写，满了再 stall，这样可以抹平突然的峰值—但是出现 read miss 时还要从 buffer 里面找。\textcolor{red}{write-back}：将数据只写入 cache，需要一个 dirty bit 表示 cache 是否与内存同步，如果是 dirty 的，则在把这个 block 覆盖时，需要写回内存。它减少了与内存的交互，但是产生了数据的不一致。

\textcolor{red}{write-allocate} 则是另一维度，它认为写入内存的元素将来可能会被再度访问，所以要同步到 cache；\textcolor{red}{no-write-allocate} 则认为写完的东西不会再次被访问，所以不同步到 cache（\textcolor{blue}{优势在于不会污染 \$}）。\textcolor{orange}{一般 write-back 默认搭配有 write-allocate，而 write-through 两种策略都有}。

当 write-allocate 出现 miss（即写入的目标不在 cache 中）后，有 \textcolor{red}{fetch-on-miss} 策略：将这个位置所属整个 cacheblock 一起取过来；另一种策略即为 \textcolor{red}{no-fetch-on-miss}，即仅存储写入的位置。这种策略破坏了 cache 的整体性，于是需要更精细的 valid bit，为每个 word 分开存储它是否合法。

而当 no-write-allocate 出现 hit 时（即被写入的位置已经在 cache 中有备份），为保持一致性，又有两种策略：\textcolor{red}{Write-around} 同步到 cache；\textcolor{red}{write-invalidate} 直接 invalidate cacheline。

% \vspace{-5ex}
\begin{center}
    % 使用 resizebox 强制将表格缩放到当前栏宽
    \resizebox{\linewidth}{!}{%
        \begin{tabular}{|c|c|c|c|c|c|}
        \hline
        \multicolumn{4}{|c|}{\textbf{Write-through}} & \multicolumn{2}{c|}{\textbf{Write-back}} 
        \\\hline\multicolumn{2}{|c|}{\textbf{Write-allocate}} & \multicolumn{2}{c|}{\textbf{No-write-allocate}} & \multicolumn{2}{c|}{\textbf{Write-allocate}} \\\hline\textbf{Fetch} & \textbf{No-fetch} & \textbf{Around} & \textbf{Invalidate} & \textbf{Fetch} & \textbf{No-fetch} \\ \hline
        \multicolumn{2}{|c|}{寻找要覆盖的块} &&& \multicolumn{2}{c|}{寻找要覆盖的块} \\ \hline
        & & & hit 时标记不合法 & \multicolumn{2}{c|}{将脏块同步内存} \\ \hline
        取整块 && & & 取整块 & \\ \hline
        写整块 & 写单点 & hit 时写入\$ & & 写整块 & 写单点 \\ \hline
        \multicolumn{4}{|c|}{将修改写入内存}& & \\ \hline
        \end{tabular}%
    }
\end{center}

一般最常见的组合是 write-back + write-allocate + fetch-on-miss，适合多次读写连续内存的场景；write-through + no-write-allocate + write-around，不需要管理 dirty bit 且易于维护一致性，适合如嵌入式系统等场合。

L1\$：每个 CPU core 都有数据 cache（\textcolor{red}{D\$}）和指令 cache（\textcolor{red}{I\$}）各一个。小但是快，优先优化 hit time 而不是 miss rate，32-64kB，4cycle。L2\$：每个 CPU core 也都有一个，256kB-1MB，10cycle。L3\$：所有的 core 共享，关注 miss rate，8MB，30cycle。如果再 miss 就交由内存处理。

\textcolor{red}{inclusive\$}，底层 \$ 中的数据总会在上层 \$ 有备份，优势是只要检查最上层 \$ 即可得知是否被存储，同时多层 \$ 中，下层脏块在被覆盖时不是直接写入内存而是先写入上层 \$，而 inclusive 架构决定了写入上层时必然会 hit，而其缺点在于空间开销浪费，以及最大痛点 \textcolor{red}{recall}：如果上层 \$ 满了决定覆盖一个块，那么为了保证一致性，该块下层的出现也要一块被干掉，这就要求上层的容量和 associativity 必须显著大于下层，且在多个 core 共用最上层 \$ 时较难以维护；\textcolor{red}{exclusive\$}，所有 \$ 中的数据总是不交；\textcolor{red}{non-inclusive\$}，底层数据可能在上层有（也可能没有）。实现困难但省空间。

有 \textcolor{red}{local miss rate}，即 \# 本层 miss / \# 对本层的访问；以及 \textcolor{red}{global miss rate}，即 \# 本层 miss / \# 对最底层的访问。算多层 \$ 的 AMAT 时，如果按照 local 那就要递归计算，按照 global 就直接带权和即可。

replacement policy。其包括 \textcolor{orange}{要不要替换}（\textcolor{blue}{可以选择 normal 也即替换以及 bypass 也即不替换两种}）以及 \textcolor{orange}{替换哪一项} 两个部分。\textcolor{red}{random} 随机替换一个。\textcolor{red}{LRU}，覆盖上次访问时间最早的 entry，通过维护 timestamp 实现；\textcolor{red}{LFU}，覆盖访问最不频繁的 entry，通过 aging counter 实现，hit 了增加、定期减半；\textcolor{red}{Belady's optimal}，覆盖下一次访问最晚的项，但非常不切实际，只有理论价值。

所有 policy 的本质都是维护一个 rank，eviction 时干掉 rank 最低的 entry。描述一个 rank 需要在插入新项时的 insertion policy 和 hit 时的 promotion policy 两部分。典型例子就是 \textcolor{orange}{Recency Stack} 也就是 LRU，分为记录完整 rank 但空间开销大的 precise 以及记录少量 bit 的 approximate 两种。

LRU 的 insert 和 promote 都是直接设为 MRU，但在 thrashing 和 scanning 这种流式数据的场合效果不好。\textcolor{red}{LIP} 在 insert 时设为 LRU、promote 时设为 MRU，其对 scan 是友好的，但在其它访问模式下效果很差，适应性不佳。\textcolor{red}{BIP} 是一种结合方法，在 insert 时大概率设为 LRU、小概率设为 MRU，这样持续访问的 block 会有更大概率在某次被设为 MRU 然后被记住，而仅访问一次的 block 则是大概率设为 LRU 然后立刻被覆盖。

不同的 access pattern 有不同的最优策略。因此有 \textcolor{red}{DIP}，动态在 LRU 和 BIP 之间切换。评估切换的方式是在主 cache 旁维护两个只存 tag 的 \textcolor{red}{shadow tag arrays}，一个模拟 LRU，另一个模拟 BIP，竞争上岗。另一种开销更小的解决方案是 \textcolor{red}{set dueling}：考虑到实际的 cache 已经被 hash 打乱了，有理由相信所有 set 的访问模式差不多，于是可以取一个 \textcolor{orange}{SDM-P0} 集合，强制它们使用策略 P0，取另一个 \textcolor{orange}{SDM-P1} 强制使用 P1，使用一个 \textcolor{red}{PSEL 饱和计数器} 统计双方的 miss 差，然后所有其它的都是 \textcolor{orange}{follower set}，谁赢听谁。

LFU 使用 counter。但缺点是对 pattern 的改变不是很敏感，较小的 counter 需要时间累积。因此也有修正如 \textcolor{red}{FBR}，如果一个 block 近期被访问则不重复增加，防止短期突发流量的过量影响；\textcolor{red}{LRFU}，就是带权重衰减的 LFU。

不论是 superscalar 的处理器还是多核共用的 cache 都会发生多个 load 和 store 同时进行的场景。cache 需要支持 \textcolor{red}{hit-under-miss}，即在处理 miss 的同时允许其它 hit；支持 \textcolor{red}{hit-to-miss}，即如果一个正在处理 miss 的 entry 又被访问，则应该标记为 hit，避免重复取数据；支持 \textcolor{red}{miss-under-miss}，允许多个 miss 同时被处理（这需要取数据的带宽高）。

专属名词：\textcolor{red}{outstanding miss}，即正在被处理的 miss。\textcolor{red}{MSHR} 会以 block 为单位追踪所有 outstanding miss，并处理对这个 block 的 load 和 store。不同块的不同 miss 由不同 MSHR 完成，相同块的多个 load 和 store 由同一个 MSHR 完成。一个 MSHR 包括 valid bit（其是否正在被使用）、block address（这个 MSHR 正在处理哪个块的请求），以及多个 entry，每个 entry 对应一条正在等待该块的指令，包括 valid bit（是否有效）、type/fmt（其是 load 还是 store、数据粒度—byte、halfword 还是 word）、offset（在缓存块中的位置）、Src/Dst Reg（读/写指令的源寄存器）。每个 cache 一般有几十个 MSHR，有好的 Memory Level Parallelism。

Cache miss 时，首先搜索所有的 MSHR 找到其是否正在被处理，如果是的话就在 MSHR 的 entry 里面追加一项（处理 hit-to-miss），否则分配一个新 MSHR 并追加第一项 entry（处理 miss-under-miss），而如果没有空余 MSHR 或 entry 就 stall。通过将 miss 外包给 MSHR，cache 可以处理其它的请求（处理 hit-under-miss）。MSHR 可以有更细粒度的处理，只要有一部分准备好就可以检查有哪些 load/store 在等待它，从内核中 forward 数据处理掉这些 load/store 后，即可 free 对应的 entry；当整个 MSHR 的数据全都准备好后，其所有 entry 也应该被处理完毕，可以直接释放之。

除了被动地在 miss 时获取数据以外，还可以做 \textcolor{red}{data prefetching}，在处理器请求之前就主动预测并获取数据。可以减少 cold miss，但代价是需要更大带宽，且需要 non-blocking 的 \$（因为预取不能阻塞正常请求）。可以通过软件显式预取（\texttt{GCC \_\_builtin\_prefetch()}，但与微架构强相关，移植性差，且必须需要 OoO（同理，不应该阻塞正常请求）。有软件就有硬件，这通常是对程序更友好的方式，且不会带来额外的指令开销。

衡量预取策略的好坏：\textcolor{red}{accuracy = \#(prefetched and accessed) / \#prefetched}；\textcolor{red}{coverage = \#(prefetched and accessed) / \#accessed}；\textcolor{red}{timeliness}，在足够早的时刻发出预取请求（以掩盖内存访问延时），但又不至于过早（污染 \$）；\textcolor{red}{资源竞争 (resource contention)}，预取消耗带宽和容量，可能与常规请求竞争。

例如，对于顺序访问，可以始终保持预取当前访问的后继 $N$ 个块，不论当前访问是 hit 还是 miss。过大的 $N$ 无法掩盖预取延迟，过小的则产生带宽浪费且容易污染 \$。此外，预取的块也有两种选择：直接像常规 access 一样写入 \$，优势是简单，劣势是污染；另一种是建立独立的 \textcolor{orange}{stream buffer}，优点是无污染，缺点是需要分别检查 \$ 和 buffer，且在命中 buffer 时要移入 \$。为了避免浪费，只在检测到流式访问时开启这种预取。另一种模式是固定步长访问，为了知道步长，可以用 2bit 饱和计数器来维护，只在高置信度时才会以该步长预取。

有一些更高级的预取方式，例如，在遍历链表、树或图时，一个地址经常会接着另一个被访问，因此可以维护查找表来列出后继的备选方案以及置信度。缺点是查找表需要巨大内存。另一种方式是空间相对性：如果在一处以某种特定的 offset 访问了基地址的若干相对位置，则在另一处也可能遵照相同的模式。还有神秘 ML-based 方法！

总结：对于 cache 相关问题，主要关注无非是两点：\textcolor{red}{hit 了要干什么，miss 了又要干什么}，两侧必须都要考虑！

\section*{DRAM}

DRAM 比起 SRAM，容量更大但更慢，高度并行因此带宽很大，和计算逻辑位于不同的片上。DRAM 间以及 DRAM 与 \$ 的通信都需要通过 memory channel。

现代服务器通常包含多个 CPU \textcolor{red}{socket}。每个 CPU 拥有独立的内存控制器。一个 socket 既可以访问自己的内存，也可以通过桥访问别人的内存，但后者的延时势必较大，这就是 \textcolor{red}{NUMA (non-uniform memory access)}。

一个 CPU socket 上可以有多个 \textcolor{red}{channel}，每个 channel 上连接不同的内存条。它们拥有完全独立的物理连线（\textcolor{orange}{channel bus}），\textcolor{blue}{它们的并行是真正的空间并行（完全同步访问）}。channel 数受限于 \textcolor{orange}{CPU 引脚数}。\textcolor{blue}{DDR (Double Data Rate)} 是描述 channel 的特性，即每个时钟周期可以发送两次数据，因此 400MHz 其实对应着 800MT。

\textcolor{red}{DIMM} 是物理概念，即一根内存条；\textcolor{red}{rank} 是逻辑概念，即连接到同一个 \textcolor{blue}{chip select} 的 chip 集合。一个 rank 必须提供填满内存通道位宽（通常为 64b）的数据。例如，8 个 x8 的芯片组成一个 rank。双面 DIMM 通常包含两个 rank。同一个 channel 上的 rank 共享 \textcolor{blue}{address/command 总线} 和 \textcolor{blue}{data 总线}，但独占 \textcolor{blue}{chip select} 信号线。rank 数目受限于 \textcolor{orange}{总线信号完整性 (bus signal integrity)} 即电信号在传输线上保持清晰、未失真的能力：更多的 rank 会有更多的电容负载。因为共享 I/O，不同 rank 不能同时传输数据，只能进行 \textcolor{orange}{时分复用 (interleaving)}。这不能提高峰值带宽，但可以通过 pipeline 掩盖延迟，提高有效带宽。

一个 rank 上放了多个 \textcolor{red}{chip}。同一个 rank 中的所有 chip 完全协同，它们共享 \textcolor{blue}{address/command 信号}（接收的指令同步），但有独立的 \textcolor{blue}{data line}。其不提供独立的访问，必须将多个芯片组合起来填满 64b 的 channel 总线。\textcolor{blue}{同一个 rank 上的芯片之间不存在并行性，只存在协同性。}

一个 chip 上则是有多个 \textcolor{red}{bank}，是芯片内部的并行性。所有 bank 共享 \textcolor{blue}{address/command/data 总线} 也即 \textcolor{blue}{chip I/O channel}，bank 内部的操作彼此独立，但因为共用 I/O，\textcolor{blue}{它们的并行是 pipeline。}限制 bank 数量的是 \textcolor{orange}{chip area efficiency}，因为每个 bank 都需要配套的外围电路，切分太细就会导致外围电路占用过大，因此 DDR3 使用 8bank、DDR4 则是 16bank。

因此，bank 和 rank 具有相似的结构（channel 中的所有 rank、chip 内的所有 bank 均共享 I/O），而 chip 是用来提高带宽的，是纯粹物理结构，逻辑上等效于一个大单元。

一个 bank 内部有 $2^n$ 行，每行称作一个 \textcolor{red}{wordline}；$2^m$ 列，共 $2^m\times w$ 个 \textcolor{red}{bitline}。$w$ 是 I/O 粒度，有 x4、x8、x16 等规格。有一个 \textcolor{red}{row buffer}，其有 $2^m\times w/8$ Bytes，和一行的形状相同，也被称作一个 \textcolor{red}{DRAM page}。行和列共用引脚（以减少总引脚数）。

CPU 请求数据时，步骤为：1. 解码行地址，给对应的 wordline 加电压，所有晶体管打开；2. 电容里的电荷顺着 bitline 流动，但因为电容容量很低，信号很微弱；3. row buffer 上附有 \textcolor{red}{sense-amp}，能感知到微小信号并将其放大，存储到 row buffer 中，于是整整一行的数据都被读取，这个操作即为 \textcolor{blue}{open row}；4. 之后列选择器从 row buffer 中挑选需要部分并传输；5. 最后，为了保证下一次访问可以进行，必须保证 bitline 的电压恢复平衡，这个操作即为 \textcolor{blue}{precharge}。

综合起来，DRAM 的五种常见动作（受内存管理器支配）：\textcolor{red}{ACTIVATE}，即开行，前三步；\textcolor{red}{READ}，读取数据，只要行已经打开（\textcolor{blue}{row hit}）就会很快；\textcolor{red}{WRITE}，写入数据，直接修改 row buffer 并在后台同步回电容；\textcolor{red}{PRECHARGE}，关闭一行，即最后一步；\textcolor{red}{REFRESH}，因为 DRAM 总是在漏电所以要定期把所有数据读出来再写回去，此时内存无法响应 GPU。

如果连续读取的数据位于同一行，则一次 row miss 后多次都会是 row hit；不同行则必须先 PRECHARGE 再 ACTIVATE，就会很慢。如前文所说，一个 chip 里的不同 bank 可能在干不同的事，例如有的 bank 在 ACTIVATE，其它 bank 占用总线在传输。

当 CPU 访问数据，比如说一个 64B cacheline 时，假如 channel bus 的带宽是 8B，则会有 64B/8B=8 个 \textcolor{red}{burst}。每次 burst 的 8B 数据，假设 chip 是 x8 的，则需要由 8 个 chip 各自贡献 8b 的数据出来，最后拼成 8B 的数据在半周期（因为是 DDR）内输出。

然而内部阵列的充放电是很慢的。为了对外有高的带宽（8b per chip），阵列必须一次取出来 64b 的数据，然后在接下来的 8 个 burst 中一点点吐出来。这样，芯片对外可能有 1200MHz = 2400MT 的速率，但每次只输出 8b；对内只需要有 300MT 的速率，每次输出 64b 即可。

总结：channel 并行是真并行，address/command/data 都完全独立；rank/bank 并行的访问可以 overlap，但 address/command/data 是共享的；chip 必须协同访问，有共享的 address/command 和独立的 data。

处理器芯片里面有 \textcolor{red}{memory controller}，其功能包括将 load 和 store 翻译为五种指令；为了保证 DRAM 的充放电正确进行，控制指令发送间隔；将 address 翻译为对对应 channel/rank/bank/row/column 的访问；重排内存访问顺序以最大化吞吐量；管理 row buffer 和 DRAM 刷新；管理电源模式等。

一次内存访问的时间开销包括：CPU 到 memory controller 的延时；controller 内部延时；DRAM bank 的延时，包括 \textcolor{orange}{tCAS = RD/WR}—命中 open row 的延时，\textcolor{orange}{tRCD + tCAS = ACT + RD/WR}—要访问一个 close row 但 row buffer 为空时的延时，\textcolor{orange}{tRP + tRCD + tCAS = PRE + ACT + RD/RW}—要访问一个 close row 且当前 row buffer 非空时的延时；DRAM 数据传输延时，即 channel width × burst length / bandwidth；memory controller 传回 CPU 的延时。

一个物理地址会对应 channel $x$、rank $y$、bank $z$、row $r$、column $c$ 的一个数据。一般这些 ID 都会是物理地址的若干位。一次 cacheline 访问是 64B，因此最低 6 位会是定死的：最低 3 位是 bus，即当前传输的 Byte 位于一次 8B 的 burst 中的哪个 Byte；次低 3 位是 column 的低 3 位，即这次传输属于第几个 burst。高位的排列相对自由，因为核心目标有两个相互冲突的维度：是尽量命中同一个 row buffer，还是尽量让多个 channel 和 bank 同时工作。

粗粒度划分：在读取连续数据时只会激活一个 channel，带宽利用率低。

% \vspace{-7ex}
% \begin{center}
    % 使用 resizebox 强制将表格缩放到当前栏宽
\noindent\resizebox{\linewidth}{!}{%
    \begin{tabular}{|c|c|c|c|c||c|c|}
    \hline
    C(2b)&R(1b)&Row(14b)&Bank(3b)&Column(H8b)&Column(L3b)&Bus(3b)
    \\ \hline
    \end{tabular}%
}
% \end{center}

行交错：连续的地址位于同一个 row 内，但跨行即会切换 channel 或 bank。

% \vspace{-7ex}
% \begin{center}
    % 使用 resizebox 强制将表格缩放到当前栏宽
\noindent\resizebox{\linewidth}{!}{%
        \begin{tabular}{|c|c|c|c|c||c|c|}
        \hline
        Row(14b)&C(2b)&R(1b)&Bank(3b)&Column(H8b)&Column(L3b)&Bus(3b)
        \\ \hline
        \end{tabular}%
    }
% \end{center}
% \vspace{-9ex}
% \begin{center}
    % 使用 resizebox 强制将表格缩放到当前栏宽
\noindent\resizebox{\linewidth}{!}{%
        \begin{tabular}{|c|c|c|c|c||c|c|}
        \hline
        Row(14b)&Bank(3b)&R(1b)&C(2b)&Column(H8b)&Column(L3b)&Bus(3b)
        \\ \hline
        \end{tabular}%
    }
% \end{center}

缓存行交错：跑一小段就会切换 channel 和 bank。

% \vspace{-7ex}
% \begin{center}
    % 使用 resizebox 强制将表格缩放到当前栏宽
\noindent\resizebox{\linewidth}{!}{%
        \begin{tabular}{|c|c|c|c|c||c|c|}
        \hline
        Row(14b)&Column(H8b)&Bank(3b)&R(1b)&C(2b)&Column(L3b)&Bus(3b)
        \\ \hline
        \end{tabular}%
    }
% \end{center}
不论如何，row 都应该在高位以保证 row buffer 被良好利用，而 channel 和 bank 应该在低位以并行工作。

此外对 row buffer 的管理策略也有影响。\textcolor{orange}{open-page policy} 期待下一次 hit，因此在 access 后会保持其 open，不会 PRECHARGE；\textcolor{orange}{close-page policy} 期待下一次 conflict，因此 access 后会立刻 PRECHARGE；\textcolor{orange}{adaptive policy} 则会预测下一次访问是否 hit 并动态决定是否立刻 PRECHARGE。\textcolor{blue}{close-page 并没有减少指令数，而是将 PRECHARGE 从 critical path 中移除}。

DRAM 的 scheduling 策略则有 \textcolor{red}{FCFS}（先到的先处理）、\textcolor{red}{FR-FCFS}（优先处理 row-hit，如果没有再处理先到的，可以最大化 row-hit 数目，最大化吞吐量）。

DRAM 每隔一段时间就要把所有的 row 都 ACT + PRE 一遍。当然这不用手动操作，有专门的指令来自动成块刷新。刷新策略可以统一刷新，而代价是一长段不可用时间；也可以分布式刷新。后者是如今主流。

\section*{Cache Coherence}

当多个 core 共用同一个 cache 时，我们理应期望读取到最新写入的数据—即使数据是其它 core 写入的。然而，如果上级 cache 是公有的而下级 cache 是私有的，就可能出现不同步现象。

最粗暴的解决方案是禁止对共用的数据使用 cache。更好的解决方案是设立缓存一致性协议，即 cache controller 不能只管自己 core 的读写，还要监听其它 core 的写入。一致性的基本原则是 \textcolor{red}{SWMR}：允许多核同时读入，但只允许一个人写，且写之前要让其它私有 \$ 中的副本失效。

\textcolor{red}{MSI} 协议：所有 cacheline 被标记为以下三者之一：\textcolor{orange}{M(odified)}，修改态，当前 core 拥有该数据的唯一有效副本，主存中的数据是过时的，当前 core 可以随意读写，不用通知别人；\textcolor{orange}{S(hared)}，当前 core 拥有数据的有效副本之一，所有 S 态的 core 的数据是互相一致且与主存亦一致的，其可以读但不能写；\textcolor{orange}{I(nvalid)}，当前 core 不拥有数据或已经过时，此时既不能读也不能写，必须与最新状态同步。可以发现，该协议是对 valid bit 的一种扩展。

一个 core 如果要写，就要先把自己提升成 M、把其它所有 S 态的 core 降级到 I，然后就可以独占来写了；同理，如果一个 core 想要读，则必须先把目前处于 M 的 core 降级为 S（同时把修改后的数据写入主存或直接 forward 到当前 core），然后再把自己提升成 S。可以搭一个 M/S/I 三个态在自己读、自己写、别人读、别人写的四种 transition 下的自动姬描述之。

现在考虑如何在多核间通讯以实现上述提升/降级策略。策略一：\textcolor{red}{snooping}，在出现 miss 或需要修改数据时就广播，所有 core 时刻监听广播，如果发现有人要读写自己 cache 中的数据时就依照 MSI 修改。缺点是扩展性差。策略二：\textcolor{red}{directory-based}，系统中有一个专门的目录区，存储了每个 data block 目前在哪些 core 的 \$ 中，修改时查目录并只通知其中记载的 core，扩展性更好但硬件设计更复杂。

但是 MSI 也有一个问题：\textcolor{red}{其强制所有 \$ 以 cacheline 为单位保持一致性}。这意味着如果我们希望开一个 counter 数组并令每个线程独立地维护其中的一个 counter 时，因为有多个 counter 处于同一个 cacheline 中，所以它们的线程会不断争抢这个 cacheline，这就是 \textcolor{red}{false sharing}。解决方案是强行把 counter 中的每一项 padding 到 block size。这说明，\textcolor{blue}{开大 block size 并不一定是好事，它有可能会增大 false sharing 的概率}。

\textcolor{red}{shared \$ 不一定会让单个线程的效率下降，典型例子是使用共享数据时，一个线程可以为其它线程 warmup \$。}

\section*{NoC}

共享的 cache（比如说 L3\$）常常使用 \textcolor{red}{banking} 实现，即划分为很多 bank，每个 bank 类似一个独立的、临近 core 的小 \$，而一个 block 只能位于至多一个 bank 中，这样不同bank 间可以并行，降低 hit latency。

如何决定每个 block 归属于哪个 bank？静态的方法选择地址中某一段作为 bank address（过一个 hash 后得到真实的 bank ID），动态的方法倾向于把 block 放在访问其更多的 core 旁边，且允许 block 在 bank 间迁徙—但任意时刻仍然只能归属于一个 bank！

现在有很多 core、很多私有的 \$ 或公有的 bank，它们之间通过 NoC 连接。

首先是信息在 NoC 上传输的方式。一条 \textcolor{red}{message} 是长度任意的 01 串，其被切割成长度可变但有限的 \textcolor{red}{packet}。packet 是路由的基本单位，其中包含一个 \textcolor{orange}{header} 和一个 \textcolor{orange}{payload}，前者包括控制信息，如路由信息 \textcolor{orange}{route} 或在原始消息中的顺序 \textcolor{orange}{seq\#}，后者则是真实要传输的信息。packet 被继续切成固定大小的 \textcolor{red}{flit}，其是流控制的基本单位，分为三种类型：\textcolor{orange}{head flit}：包含完整的 route，可以用于寻路。在开辟出路径之后，剩余 flit 即跟随之。\textcolor{orange}{body flit}，主要承载数据。\textcolor{orange}{tail flit}：告知 packet 已经传输完毕，并提供校验机制等。此外，如果整个 packet 太短，那也不会硬切，而是直接用一个 \textcolor{orange}{head-tail flit} 全部承载。最后，flit 可能被进一步切成若干 \textcolor{red}{Phit}，是物理链路上单周期可传输的数据量，在 NoC 的场合一般和 flit 大小相同。\textcolor{blue}{每切一次都会带来额外的 header 开销}。

一次传输（发一个 packet）的延时，假设网络是零负载的（\textcolor{red}{zero-load latency}），是 \textcolor{orange}{header latency} 和 \textcolor{orange}{serialization latency} 两部分：前者是 header 开辟道路到达终点的延时，有 $T_h=H\times(t_r+t_l)$，其中 $H$ 是跳转数，$t_r$ 是 router 延时，$t_l$ 是 link 延时；后者是 $T_s=L/B$，其中 $L$ 是连同 header 额外开销的 packet size，$B$ 是带宽。\textcolor{blue}{然而，实际延时会因为排队而大幅延长，甚至排队成为延时的主导！}

\textcolor{red}{路由距离 (routing distance)} 是从路由起点到终点的链路跳数。\textcolor{red}{直径} 是最大路由距离。\textcolor{red}{bisection bandwidth} 是等大小切成两半的最小割。

在路由算法中，\textcolor{red}{同一个 packet 的所有 flit 都沿着同一条路线行进}，但是 \textcolor{red}{就算起讫点相同，不同的 packet 可以也可以不沿着同一条路线行进}。

一些属性：\textcolor{red}{确定性 (deterministic)}：$x\to y$ 总是走同一条路。\textcolor{red}{最小 (minimal)}：总是走最短路。\textcolor{red}{无知觉 (oblivious)}：寻路时不管实际情况（例如，拥塞）；deterministic 是 oblivious 的一种。\textcolor{red}{自适应 (adaptive)}：寻路时管实际情况；实际情况通过 link availability, buffer occupancy, history of channel load 等方式衡量—但是这些信息常常无法完整获取。\textcolor{red}{Source Routing}：在源点处即决定整条路径。\textcolor{red}{Incremental Routing}：路径在每次跳转时动态地决定。

在 mesh/torus 上，可以使用 XY 或 YX 的 \textcolor{red}{dimension order}，它们是 minimal 且 oblivious 的。在 mesh 上是 deterministic 的，在 torus 上可以确定性或随机地进行 tie break。

mesh 上的自适应路由方法则包括 \textcolor{red}{minimal adaptive}—总是走最短路中负载最小的一条；\textcolor{red}{fully adaptive}—可以走任何路，但是会遇到死锁。

\textcolor{red}{flow control} 负责在多个 packet 之间仲裁。其有 \textcolor{red}{bufferless} 的协议，即 router 处不设额外的 buffer，input link 输入的东西如果不在下一周期转发出去就会丢弃。这方面的策略包括 \textcolor{red}{dropping}，竞争失败的 packet 直接扔掉，回传一个 \textcolor{orange}{NACK} 指令或是直接 \textcolor{orange}{time out}；\textcolor{red}{misrouting}，竞争失败的 packet 被随便丢到一条 out link 中，并期望其最终流回去，适用于 incremental 的模式，但可能导致活锁；\textcolor{red}{circuit switching}，传输前发送一个 \textcolor{orange}{setup probe} 建立一整套路径并预留其上所有 link，tail flit 会释放整条路径，优势是简单，劣势是对短 packet 浪费很大，且延迟更长。\textcolor{red}{buffered} 协议则在 router 处设有 buffer，但这又牵涉到 buffer 资源的分配：模式包括 \textcolor{red}{store-and-forward}，以 packet 为单位传输，只有整个 packet 全都到达同一个 router 后 head flit 才继续；\textcolor{red}{virtual cut-through}，只要下一个 router 有空间塞得下整个 packet 就可以继续 forward（避免一个包发了一半后面的卡着动不了了）。这两个模式都简单，但 lose 的包要等的时间更久。\textcolor{red}{wormhole}，其中所有链路按照 packet 为单位分配（一个 packet 传完前，这个 link 都被占用—不管这个 packet 的第一个 flit 已经跑了多远）—带来的后果就是，只要 head flit 被阻塞了，后面的 flit 一个也动不了，且因为一个 package 的所有 flits 可能长长地拖在一众 router 上，所以这可能导致多个 link 同时阻塞 (head-of-line blocking, HOL blocking)。\textcolor{red}{virtual channel flow control}，对每个物理链路设立多个并行 link，缺点是更复杂。

如何感知拥堵？使用 \textcolor{red}{backpressure} 机制，buffer 满了的下游 router 会给上游发信号暂停传输。另一种思路是 \textcolor{red}{credit-based flow control}，即上游记录下游可用的 buffer 数目。下游每发出一个 flit，就给上游传信增加上游的 credit，同时减少自身的 credit（因为下游多了一个包）。

在 circuit switching 模式下，一次传输会同时占用多个 link；在使用 packet buffer 时，一个 packet 会占用一个（若是 store-and-forward）或多个（若是 virtual cut-through）packet buffer；在使用 flit buffer 时，一个 packet 会占用多个 flit buffer。所有这些都会引起 deadlock。

mesh 的场景，有 \textcolor{red}{turn model}：在 mesh 上 ban 掉一些 turn 的方式。例如，XY model 和 YX model 分开用都没问题，但是合在一起就会出现被称为 \textcolor{red}{01 turn} 的经典死锁现象。所以解决方法是把逆时针和顺时针的环上各自 ban 掉一种 turn 的方式。当然，不是随便 ban 都行：有一种叫做 \textcolor{red}{six turn} 的错误案例，可以拼出一个八字形的环来。

另一种方式是 \textcolor{red}{resource ordering}，即为所有 link 编号，一条 path 必须满足编号的单调性。如 XY model 和 turn model 等都是这种方式的子集。（相当于人工编了一个拓扑序出来）VC channel 可以与这种方法结合，强制路径上 VC channel ID 不降，人工造 ordering。

\section*{Virtual Memory}

每个进程都有一份独占的 \textcolor{red}{virtual address space}，受 ISA 决定。系统具有唯一的 \textcolor{red}{physical address space}，受硬件决定。\textcolor{red}{address translation} 就是将虚拟地址翻译为物理地址的过程，以 page 为单位。

操作系统为每个进程维护一个 \textcolor{red}{page table}，其储存在 DRAM 中，基地址保存于一个特殊的寄存器中。虚拟地址刨除最后若干位的 \textcolor{orange}{offset}（4kB pagesize 对应着 12b offset）后，前面的位加上基地址后被用来定位页表中的一个 \textcolor{orange}{PTE}；PTE 中存储的信息包括： \textcolor{orange}{valid bit} 标志其是否合法，\textcolor{orange}{physical page number} a.k.a. \textcolor{orange}{frame number} 指向物理地址（刨除相同长度的 offset）；\textcolor{orange}{access right}（R/W/X 权限），以及 \textcolor{orange}{dirty bit} 或更多 metadata 等。\textcolor{red}{注意，metadata 指非 frame number 的一切东西}。

VM 可以保证进程间的隔离、高效管理物理资源，但更重要的是作为一种特殊的 \$。

进程中一般只有一部分 virtual page 是 \textcolor{red}{allocated} 的；这其中，又只有一部分是 \textcolor{red}{mapped} 到 DRAM 的，剩下已分配但未映射的会被 \textcolor{red}{swapped out} 到外存（SSD/磁盘）中。在进程要求访问一个 page 时，会先定位 PTE 然后检查其 valid bit 以及权限，如果都通过就是 \textcolor{orange}{page hit}，可以直接访问；否则，如果是权限不足就报 \textcolor{orange}{permission fault}，交由 OS 处理；如果不在内存中（valid bit 不合法）就是 \textcolor{orange}{page fault}，交由操作系统中的 \textcolor{orange}{page fault handler} 处理，需要从外存中进行 \textcolor{orange}{page in}。

VM 并非简单的一对一映射，还可以以三种方式实现复杂效果：\textcolor{red}{synonym}，一个进程的多个虚拟地址映到同一个物理地址，方便不同模块访问；\textcolor{red}{page sharing}，不同进程可以 \textcolor{blue}{只读地} 将虚拟地址映到同一个物理地址，可以实现库文件的共享或是 \texttt{fork} 子进程时继承父进程的状态；\textcolor{red}{homonym}，不同进程访问数值上相同的虚拟地址，但因为它们的 page table 是独立的，它们实际上会经由不同的 PTE 访问不同的物理地址。

page table 的一大特征是其上 allocated 的 entry 非常稀疏。所以使用 radix tree 维护：内存中仅仅固定了最上层的页表，而其中的每一项指向次一层页表所处的物理地址，可以是未分配—如果次级页表中所有项均未分配、分配且映射或是分配但被交换出去了。一个典型的例子是，虚拟地址被分为顶级页表 index、次级页表 index 和 offset 三部分，访问时首先根据 index 定位顶级页表中的 PTE，访问之获取次级页表基址所处位置；然后定位次级页表 PTE，访问之获取真实物理地址。\textcolor{blue}{缺点如果遇到多次 page fault 就会导致效率低。}

因为 PTB 和物理地址都在内存中，这样取一次数据就要访问两次（$n$ 级 PTB 则需要 $n+1$ 次）内存，很慢。所以 CPU 中集成了 \textcolor{red}{Translation Look-aside Buffer (TLB)}，存储了最近遇到过的虚拟地址→物理地址的映射，这样把多次内存访问减少到一次。TLB 一般只有 10\textasciitilde100 的项数，但是具有高的 associativity，使用随机或 FIFO 的 policy。和 \$ 的设计一样，其也可以将指令和数据访问分开或者采用多层设计。TLB 中的每一项包括某个 PTE 的完整数据，以及其自身的 metadata（valid bit、tag、LRU info 等）。

TLB 的 miss 被称作 \textcolor{red}{page table walk}，可以被软件或硬件处理。miss 后就要 fallback 回常规 PTE 访问，如果 page hit 了就把 PTE 拷回 TLB，如果 page fault 了就 handle 后把 remapped PTE 拷回 TLB。拷回后，再次询问 TLB（\textcolor{blue}{之所以 retry 而不是绕过 TLB 直接利用获取的 PTE 访问内存，是因为这样让路径更简洁}）

多个进程可能试图访问同一个虚拟地址。为了正确处理 homonym，要么在 context switch 时清空 TLB，要么给 TLB entry 中加一个 PID 区分来源进程。另一个问题是 TLB 实在太小了，而 miss 的代价又很大，因此如前文所述，可以使用多级 TLB，或者简单粗暴增大 page size（代价是 internal fragmentation 加剧，映射更不flexible），或者允许非均匀的 page size（代价是管理困难）。

TLB 和 \$ 的区别：前者负责「从虚拟地址翻译物理地址」，后者负责「获取某个物理地址的数据」。在最好的场合下，也要先访问 TLB 获知物理地址后，再访问 \$ 获得数据，二者是顺序的。

如果要想省去这个顺序访问，解决方法之一使用虚拟地址而非物理地址来建立 \$。这就是 \textcolor{red}{Virtually Indexed, Virtually Tagged Caches}。这样 TLB 和 \$ 可以并行进行，只有 \$ miss 了才需要等待 TLB。为了处理 homonym，和 TLB 的解决方案相同，要么做 context switch flush 要么存储 PID。然而如果有 synonym，就会在单个 \$ 中产生 coherence issue，处理起来是困难的。

所以折衷的方案是 \textcolor{red}{Virtually Indexed, Physically Tagged Caches}。首先，虚拟地址的效果是将一个 virtual page \#（虚拟地址刨除 offset 的部分）映到 physical page \#（物理地址刨除 offset 的部分），offset 本身是不变的。而如果 \$ 的 index 完全由 page offset 决定（换句话说，$|\text{index}|+|\text{cache offset}|\leq|\text{page offset}|$；或者，$\text{cache size}/\text{associativity}\leq\text{page size}$），则可以一边使用虚拟地址的 index 去查对应的 cacheline（\textcolor{red}{里面存储的 tag 是关于物理地址的}），另一边查 TLB 获取 physical page \# 并（可能需要搭上 offset 中的某几位）与 cacheline 中的 tag 比较。如果二者匹配即 cache hit，否则 miss。在满足前述关系时，它等效于 PIPT（也即传统的）\$。这就是为什么 L1 \$ 一般是 32kB 8-way 的，刚好与 4kB page size 匹配。

\section*{Virtual Machine}

操作系统已经把物理内存包装成虚拟内存、把物理 CPU 时间分割成处理器时间、把磁盘存储用文件系统维护，然后提供给每个进程。为何不进一步，直接将整个物理计算机也给虚拟化呢？这就是 \textcolor{red}{Virtual Machine}。虚拟机的三原则：\textcolor{orange}{fidelity}，和真实物理机等价；\textcolor{orange}{performance}，带来最小的额外开销；\textcolor{orange}{safety}，保证 VM 间的隔离，且不能让 VM 绕开管理程序独占硬件。

\textcolor{red}{Host} 指背景的物理系统；\textcolor{red}{Guest} 指模拟出的某个虚拟机；\textcolor{red}{Virtual Machine Monitor (VMM) a.k.a. hypervisor} 是实现虚拟化的软件。\textcolor{orange}{Type 0 VMM}：基于硬件的 VMM，将资源物理地分配给不同的 VM，优势是资源全都专用、没有争抢，劣势是很不灵活。\textcolor{orange}{Type 1 VMM}，VMM 像一个特殊的操作系统运行在硬件最底层，不依赖于 Windows 或 Linux 等通用 OS，直接调动硬件。\textcolor{orange}{Type 2 VMM}，仅仅是一个普通的软件进程。

VM 和 VMM 的关系如同进程和 OS 的关系：VMM 要决定如何为 VM 提供资源。对于处理器，其采用 \textcolor{orange}{时分复用 (Time Multiplexing)}，每个 VM 跑一段时间后就 context switch。对于物理内存，使用 \textcolor{orange}{资源分区 (Resource Partitioning)}，将虚拟内存映射到物理内存的某一段。对于网络、键盘、鼠标等共享 I/O 设备，使用 \textcolor{orange}{调解硬件接口 (Mediating Hardware Interface)}，VMM 保留所有权，在 VM 请求时根据特定的策略转发数据。

\textcolor{red}{Paravirtualization} 是一种模式，其中 Guest OS 知道自己运行在虚拟环境中（但是其上的软件不知道），因此会主动配合 VM，不直接发送硬件指令而是以正确的方式通知 VM，优势是简单，劣势要改 OS 源码，对闭源 OS 不太适合。\textcolor{red}{Full Virtualization} 是另一种模式，连 Guest OS 都不知道，优势是可以不改 OS，劣势是要付出额外代价处理这些问题。

传统的硬件一般仅支持 \textcolor{red}{user} 和 \textcolor{red}{kernel} 两种模式。VMM 显然是属于内核态的，Guest 上的应用仍然是用户态，但是 Guest 的 OS 呢？不能给它内核态的权限，因此物理用户态必须被虚拟地分成虚拟用户态（Guest App）和虚拟内核态（Guest OS)。一种解决方案是 \textcolor{red}{Trap-and-Emulate}，当 Guest OS 试图执行一条 privileged instruction 时，CPU 会捕获该行为，将其移交给 VMM，然后 VMM 模拟该效果后，将控制权返回给 Guest OS。但是这种方法需要 \textcolor{blue}{所有的 sensitive 指令均是 privileged}，其中敏感指令是检查或修改硬件资源配置的指令，而特权指令是会触发 trap 的指令。但是 x86 有一些敏感指令并非特权指令。解决方法是\textcolor{orange}{半虚拟化}；\textcolor{orange}{binary translation}：VMM 检查 Guest OS 的代码，动态将那些敏感但非特权指令翻译为安全的模拟代码，开销较大但可以用 \$ 优化；\textcolor{orange}{hardware support}：引入了给 VMM 的 \textcolor{red}{VT root mode}，给 guest OS 的 \textcolor{red}{VT non-root mode}。此时 CPU 会自动捕获敏感指令并移交权限。\textcolor{red}{VM control structure (VMCS)} 由 VMM 配置，指定了 host 和 guest 的 OS 态，控制虚拟机的相关任务。

虚拟机的地址翻译也是一个问题。首先有 \textcolor{blue}{guest virtual address (gVA)} 到 \textcolor{blue}{guest physical address (gPA)} 的翻译；然后 gPA 在 host 端被看做虚拟地址 \textcolor{blue}{(hVA)}，再被翻译为 \textcolor{blue}{hPA}。因此，guest OS 要维护 gVA → gPA 的页表，host 则要维护 gPA → hPA 的页表。解决方案是让硬件支持 \textcolor{red}{Nest Page Table}，然后 TLB 直接一步到位实现 gVA → hPA。

这个 Nest Page Table 的开销非常恐怖：guest OS 的四级页表查询需要 4 次访问，而每一级都是对 gPA 的访问，Host 界面就要走一套完整的虚拟地址翻译流程。所以总访存次数会非常大！其具体实现方式是，\textcolor{blue}{硬件}维护一个负责 gPA → hPA 的 \textcolor{red}{extended page table (EPT)}，并增加 EPTP 指针指向其基地址。TLB 必须带标签，避免 VM 切换时清空 TLB。

当然也有\textcolor{blue}{软件}的解决方案：VMM 为每个虚拟进程维护一个 \textcolor{red}{shadow page table}，直接一步到位实现 gVA → hPA 的映射，但需要与 guest OS 维护的页表相一致。这需要标记 guest OS 的页表为 write-protected，然后 trap 对其的写入，并依照之更新 shadow page table。好处是 page walk 更快且不需要太多硬件支持，坏处是需要更多的捕获，以及额外存储开销。

\noindent\resizebox{\linewidth}{!}{%
        \begin{tabular}{|c|c|c|c|}
        \hline
        &non-VM & Nested & Shadow
        \\ \hline
        TLB 储存信息& VA→PA & gVA→hPA & gVA→hPA
        \\ \hline
        TLB worst-case miss cost &4&24&4
        \\ \hline
        guest PTE 更新速度&快&快&慢（需捕获）
        \\ \hline
        额外存储 & 无 & per VM & per virtual app
        \\ \hline
        \end{tabular}%
    }

\section*{Security}

软件和硬件上都存在一些安全性漏洞。这些漏洞除了被用来进行恶意行为外，还有积极作用，例如 \textcolor{red}{非侵入性系统监控}，通过设备工作时散发的电磁辐射来监控系统。

一种漏洞是通过检测软件或硬件的非直接效果来获取信息。\textcolor{red}{Covert Channel} 是双方主动利用原本不用于通信的机制进行通信，例如一方故意在 CPU 上运行高耗能程序、另一方监控。\textcolor{red}{Side Channel} 则是不经意间的信息泄露，例如分析加密运算时的电量消耗。

对攻击行为建模，则受害者和攻击方各自处于一个逻辑隔离的 \textcolor{red}{domain}，但信息却经由 \textcolor{red}{信道 (channel)} 传递了。信道是任何可以被受害者方的 \textcolor{red}{transmitter} \textcolor{red}{调制 (modulate)}、被攻击方的 \textcolor{red}{receiver} \textcolor{red}{检测 (detect)} 的信息，一般处于底层，初衷是为了监测而非通讯。

常见信道包括 \textcolor{red}{物理信道} 如能耗、声音或电磁波等，需要能物理接触设备；\textcolor{red}{时序信道} 如运行时间，可以远程处理；\textcolor{red}{微架构信道} 如 \$ 或 TLB 等，需要能访问微架构状态。

例如，计算 RSA 也即快速幂（密文在指数）时，物理侧信道可以通过耗能获取指数上为 1 的位，时序侧信道则是通过计时获取。微架构则可以实现隐蔽信道通讯，方法是通过对共享资源的取用，而如果对密文的处理涉及到对共享资源的不同取用方式则也可以实现侧信道。常见的微架构信道包括核内的 branch predictor 和私有 \$/TLB，核间但 CPU socket 内的公用 \$/TLB 和 NoC，跨 socket 的 cache coherence directory 和 DRAM。

前述建模适用于被动攻击者。主动攻击者可以通过 \textcolor{red}{信道预处理 (preconditioning)} 让信道处于已知状态，并与受害者的动作 \textcolor{red}{同步 (synchronize)}。一个例子是 \textcolor{red}{Prime+Probe}：攻击者首先占满一个特定的 \$ set，然后受害者依据密文不同，决定是否驱逐被占满的 set 中的某个位置，之后攻击者再次访问其占领的 set，如果发现出现了 miss 就认定受害者访问了这个 set。\textcolor{blue}{这个方法适用于一切共享存储场合，但受害者的内存访问必须依赖于密文，同时攻击者需要通过逆向工程知晓 set-associative 场合的地址映射，再额外考虑虚拟/物理地址、同步和噪声等种种问题}。

另一种攻击是 \textcolor{red}{瞬态攻击 (transient attack)}，它依赖于处理器的 OoO：预测错误的指令是 \textcolor{red}{瞬态的}，它们对寄存器的影响会被回滚，但一般不回滚对 \$ 的影响。于是攻击者可以探测 \$ 来获取瞬态指令的功能并窃取数据。例如，对于一个 \textcolor{orange}{访问数组前检查索引} 的数据段，可以使用多个合法值把检查索引的分支预测器 saturate，然后突然输入一个任意值读取内存中\textcolor{orange}{任意位置}的数据。为了让这个数据存到 \$ 里，就以其为下标访问另一个数组，然后做一遍 probe 看另一个数组中什么地方被加载入 \$ 即可。另一个变体是 \textcolor{orange}{分支目标注入}：受害者代码中只要有任意一个片段的操作是 \textcolor{orange}{读取密文，使用密文为下标访问数组}，称之为 \textcolor{red}{gadget}，则攻击者可以利用 BTB 只存储较低位的 aliasing 性质，提前在受害者代码中找到任何一个跳转指令（\texttt{branch/jump}），并在自己的代码中的 aliasing 处用 \textcolor{orange}{跳转到 gadget} 这一操作 saturate BTB，然后调用受害者。受害者就会受被投毒的 BTB 影响跳转到 gadget 并写入 \$，然后再 probe 即可获取密文。

防御侧信道攻击的方法主要有三种：\textcolor{red}{Data-Oblivious Execution}，让程序的行为与密文无关。对于 \texttt{if(secret)a=...;else a=...;} 的指令可以使用 x86 中的 \texttt{cmov cond, src, dst}，依据 \texttt{cond} 来判断是否把 \texttt{src} 赋给 \texttt{dst}，省去了条件判断，外界看来是 \texttt{src} 和 \texttt{dst} 同时被访问了。对于 \texttt{a=buffer[secret]} 的指令，可以扫过整个数组并结合 \texttt{cmov}（但代价是时间开销变大非常多）。另一种做法是 \textcolor{red}{Hardware Resource Partitioning}，可以时序划分，但在切换时清空所有 \$ 等；也可以空间划分，以 ways 或 sets 为单位划分给不同的设备，如 \textcolor{blue}{Intel Cache Allocation Technology}。但缺点是分区导致资源利用率低、扩展性弱（应用数可能多于可提供的划分数），且对于 TLB 等其它共享资源难以划分。还有就是 \textcolor{red}{Randomization/Fuzzing}，在时间测量中引入噪声，或者降低计时器的精度。优点是简单且广谱，缺点是对不依赖时间的攻击无效，对时间差异过大的受害者无效，且妨害某些需要高精度计时器的应用。

除了作为应用程序的攻击者，随着操作系统越来越臃肿，OS 或 hypervisor 也可能存在漏洞，不能完全信任。因此，\textcolor{red}{Trusted Computing Base} 也即为了运转而信任的组件集合必须被缩小：\textcolor{blue}{传统模式}信任硬件、OS、hypervisor 但怀疑其它应用，\textcolor{blue}{ARM TrustZone} 将 OS 和 hypervisor 划分为 trusted 和 normal 两半，\textcolor{blue}{AMD SEV} 不信任 hypervisor 但信任 guest OS，\textcolor{blue}{Intel SGX} 只信任硬件和应用中一块受保护的区域也即 \textcolor{red}{enclave}，这块区域与其它部件彻底隔离，就算是权限最高的 OS 或 hypervisor，若不在 enclave 内部也无法访问。这就是 \textcolor{red}{Trusted Execution Environment}。其设计时考虑的因素包括：软件攻击，此时只有寄存器是安全的，\$ 和片外内存都是易受攻击的；硬件攻击，此时假设片上寄存器与 \$ 的数据是安全的（难以破拆硅片）但片外存储如 DRAM 是可以通过喷液氮来降低 DRAM 在断电时数据消失速率的，也可以通过直接监听总线获取数据。

TEE 的三大目标：\textcolor{red}{Attestation}，确保在真硬件上运行真数据。验证硬件通过 \textcolor{orange}{platform attestation}，在设备中烧录私钥，然后由可信第三方（如 Intel）验证。验证软件则是 \textcolor{orange}{enclave measurement}，使用数字签名，计算程序的 hash 值并发送摘要。\textcolor{red}{Isolation}，确保 OS 不通过修改物理地址来访问 enclave。首先使用 \textcolor{orange}{processor reserved memory}，通过硬件检查只允许 enclave 的访问；但 OS 可以把 enclave 的虚拟地址翻译为 PRM 以外的物理地址，解决方案是维护 \textcolor{orange}{reversed page table}，存储在受保护的区域 (\textcolor{orange}{enclave page cache mapping})，将物理地址映到虚拟地址、enclave ID 等，每次翻译后硬件核对是否匹配。\textcolor{red}{Off-chip data protection}，保护片外数据。目标有 \textcolor{orange}{confidentiality} 即保证不可见，可以加密；\textcolor{orange}{Integrity} 即保证不被篡改，可以 MAC 校验；\textcolor{orange}{Freshness} 即保证没有用旧加密数据覆盖新的，可以使用 MAC+counter 或 Merkel Tree。

替代方案包括 \textcolor{red}{Secure Remote Computing}，用户将加密数据发给不可信的云端，云端在可信的 enclave 中解密、计算、加密再回传，缺点是仍然能被侧信道攻击。还有 \textcolor{red}{FHE}，绝对安全但很慢。

\section*{Multithreading and GPU}

多线程可以每个 core 各维护一个线程，也可以单核多线程。这样做硬件成本更低，且因为 core 有很多时候在 stall 所以可以在某个线程等待时运行其它线程，但只有在 \textcolor{blue}{context switch 的代价远低于 stall 长度}时这样做是有意义的。具体而言，I/O 事件一般会有 >1000 cycle 的 stall；\$ miss 是 10\textasciitilde100 cycle，pipeline stall 则是 2\textasciitilde10。

传统单线程由 OS 主导，两段线程间需要运行 OS 代码，花费数百周期，因此一般几千周期才换一次；粗粒度多线程仅在高延迟事件如 \$ miss 时 switch，切换开销低，几十周期就可以换一次；细粒度多线程每个时钟周期都在交替处理不同线程，切换频率极高；同步多线程也称超线程，同一个时钟周期里的不同部件可以服务于不同线程（因为 OoO 在 ROB 满或依赖太多就会出现闲置），真正做到指令集并行并利用闲置资源。

为了避免在 context switch 时前往片外存储调取数据，有些资源必须有多份 \textcolor{red}{replicate}。这些东西是每个线程私有的状态量，如 PC、寄存器和页表基指针（\textcolor{blue}{虽然同一个进程下的线程共用同一张页表，但内核可能同时运行来自不同进程的线程}）。另外有些资源则可以 \textcolor{red}{share}，它们是不存储状态的元件（如 ALU）和存储元件（如 \$ 和 TLB）。OoO 中的指令窗口究竟是复制还是共享要视具体场合而定。

\textcolor{red}{GPU} 一开始只用来算图形学，用户只能配置 pipeline 但不能在上面编程。中期加上了少量的编程功能，之后有些人试图把输入输出转成图片交由 GPU 处理。直到 nvidia 推出了 CUDA，行业又设置了 \textcolor{red}{OpenCL} 的通用标准，由此转向 \textcolor{red}{General-Purpose GPU}。

GPU 支持 \textcolor{red}{Single Program Multiple Data} a.k.a. \textcolor{red}{Single Instruction Multiple Threads}。此时不同线程上运行同一份代码，通过显式调用 \texttt{threadIdx} 来处理不同数据。在 CUDA 中，每个线程都有其私有存储。并行的线程以 \textcolor{red}{block} 为单位，可以访问 block 共有的存储，并直接通过 \textcolor{orange}{barrier} 同步。block 可以以 1D、2D 或 3D 的形式组织，其中的线程也可以通过 \texttt{threadIdx.x/y/z} 获取其在组织中的标号。block 可以进一步 1D、2D 或 3D 地组成 \textcolor{red}{grid}，同样通过 \texttt{blockIdx.x/y/z} 获取编号，只能访问全局内存，block 之间的同步只能通过分开启动内核实现。

程序员为单个线程编写程序，并确保其适用于所有线程。GPU 会自动启动多个块，再为每个块启动多个线程。硬件将块内线程以 \textcolor{red}{warp} 为单位打包，warp-size 与硬件有关（标准大小为 32），同一个 warp 中的线程执行是完全 lockstep 同步的。\textcolor{blue}{于是 stall 等引起的 context switch 可以以 warp 为单位切换}。为了保证急速切换，状态必须存储在片上，因此需要大的物理寄存器堆和片上存储。

因为 warp 保证同步，所以如果碰到分支结构就会造成低效率，此时处理器不得不串行化，使用 \textcolor{red}{active thread mask} 维护 wrap 中活着的线程集合，硬件需要使用 stack structure 处理串行结果。新的 VOLTA 架构也允许动态 warp 分组。

block 内的并行通过 \texttt{\_\_syncthreads()} 强制块内所有线程均执行完某行代码后才能继续，也有 warp 粒度的 \texttt{\_\_syncwarp()}。硬件实现非常高效，但快线程必须等待慢的，会限制并行度。grid 内的并行则只能依赖 kernel launch，CUDA 保证前一个 kernel 执行完之前后一个不会开始，但是非常昂贵，需要尽量避免。同时对于 L2\$ 和全局存储的操作必须使用原子指令。block 内的通信通过块内 shared 实现，而 grid 内通过 global 存储实现。

\textcolor{red}{Streaming Multiprocessor} 是 GPU 真正的核心，类似 CPU 的 core。\textcolor{blue}{一个 block 必须完整地归属于一个 SM，且一旦被分配就会一直停留直到处理完毕；一个 SM 中可以放入多个 block；一个 grid 并非所有 block 都要进入 SM，可以排队}。SM 会被分成多个 \textcolor{orange}{sub-core}，每个 sub-core 都配有一个 \textcolor{orange}{warp scheduler}。sub-core 的数量就是单周期内能运行的 warp 数目。SM 有较大的 register file 和较小的 L1\$ 与共享存储。之所以需要大的寄存器堆是为了达成 replicate 下的\textcolor{blue}{零上下文切换开销}。因此，SM 能容纳的 wrap 数一般远大于能同步执行的数目。

SM 的前端是 \textcolor{orange}{warp 级别的多线程}，需要调度各个 warp 何时 stall、何时 switch，且在 branch divergence 时参考条件 mask 介入处理；后端是 \textcolor{orange}{warp 内的 SIMD}。

现代 GPU 使用统一的 \textcolor{red}{Unified Virtual Address}，不需要显式指定地址的来源：全局、块内共享和本地存储（乃至 CPU 存储）各自占据了一段地址域，硬件会自动将对地址的访问翻译。SM 内有硬件管理的 L1 \$ 和\textcolor{blue}{软件管理}的共享存储（虽然访问会由硬件翻译，但是将什么数据从片外搬到片上要手动处理），它们共享例如 96kB 的容量，可以依照共享存储的具体开销分割这些容量。最后，片外有全局的 L2\$ 和 DRAM，分成公有和私有的部分。

% \noindent\resizebox{\linewidth}{!}{%
        \begin{tabular}{|c|c|c|c|}
        \hline
        存储&位置&\$&周期\\
        \hline
        寄存器&片上&无&线程\\
        \hline
        块内共享&片上&无&块\\
        \hline
        线程私有&片外&L1/L2&线程\\
        \hline
        全局&片外&L1/L2&程序\\
        \hline
        \end{tabular}%
    % }

线程数除了受限于 SM 容量，还受限于单线程的寄存器与共享内存消耗。也即，单线程可获得的资源量与并行度间存在 tradeoff。此外，GPU 这种多线程环境的 \$ 会比较困难，因此 GPU 依赖于更高的离片带宽，并需要塞满大量的待处理请求以填满带宽，进而需要高效的 warp switch 能力。因此，写代码时最好让一个 warp 的访存连续（\textcolor{red}{coalesced}），这样硬件就会把它们合并为一个 transaction；进一步，起始地址 \textcolor{red}{aligned}，这样就不用拆成两次。

GPU 为了实现高带宽，使用的 \textcolor{red}{GDDR5} 牺牲了容量和通用性，不能通过插内存条简单扩展，而是直接焊死在电路板上。其和 DDR 一样有 8 的 burst，单引脚效率提高数倍但代价是不支持多 rank。rank 内部结构和 DRAM 类似，但为了速度牺牲了容量或增加了延时。\textcolor{red}{HBM} 则是另一种 2.5D 堆叠的技术，带宽极大提升而耗能降低，延迟也略有减少。

GPU 和 CPU 通过高速总线如 \textcolor{red}{PCIe} 连接，连接的 Host 是 CPU 及其主存，Device 是 GPU 及其显存。在计算时，主机要显式地分配设备显存，将主存复制到显存，启动运算，将显存复制回主存，释放显存。CPU 和 GPU 的地址空间在物理和逻辑上均独立，但现代 GPU 也支持统一的访问，此时主存和显存之间的复制自动进行了。因为通讯耗时高，所以必须重叠计算与通信。

CPU 可以把 GPU 当成一个加速器，CPU 负责管理和辅助，GPU 是计算主力；也可以让二者协同处理（\textcolor{red}{heterogeneous computing}），此时应合理分配 CPU 和 GPU 的工作量。

\section*{Customized Hardware}

通用处理器更加灵活，但会带来额外的指令、控制开销；定制硬件或加速器更高效，但是只适用于特定任务。

同时，硬件速率不仅受限于运算本身，还受限于进入和输出的数据速率，特别是如果数据存在 DRAM 中时，而加速器进一步加剧了这一受限。因此，不仅需要对运算设置加速器，还需要定制专用内存层级。

例如，在图像卷积时，可以维护包含卷积核的若干行，这样每扫一格只需取一个像素并弹出一个像素，这就是 \textcolor{red}{line buffer}。另一种更通用的方式是 \textcolor{red}{double buffer} a.k.a. \textcolor{red}{ping-pong buffer}，将 buffer 分成两个 bank，一个与存储通信进行预取，一个与加速器通信进行运算，二者角色交替互换，因此重叠访存和运算。代价是需要两倍的成本，且仅适用于特定访问方式（如流式数据）。

在 CPU 的视角，加速器是一个外围设备，需要由软件提供输入（\textcolor{blue}{通过 DMA}）、启动计算（\textcolor{blue}{写入控制寄存器}）、等待完成并处理报错（\textcolor{blue}{要么 CPU 轮询、要么加速器 interrupt}）、接受输出（\textcolor{blue}{DMA}）。而总时间为传入数据+计算+传回数据的综合，所以加速器一般对计算密集型任务才有需求。

\textcolor{red}{Reconfigurable Architecture} 是由可重构的基本模块组成的结构，包含基础的计算、存储和互联功能，在芯片上构成 1D 或 2D 的阵列。configuration 即为实现一种特定的硬件设计，即对硬件编程。对于计算模块，其配置具体运算；对于互联模块，其配置 multiplexer。通过不同的配置，可以实现不同的定制硬件。

FPGA 是一个 2D 可重构逻辑单元阵列。每个单元包括实现组合逻辑的 \textcolor{red}{Loop-Up Table (LUT)}，是一个真值表，一个 $n$-bit 至 $m$-bit 的 LUT 需要 $m2^n$-bit 的 SRAM，或者 $2^n$ 个 $m$-bit 的寄存器和 $2^n$-way 的 MUX，小 LUT 可以组合实现大 LUT；实现时序逻辑的 \textcolor{red}{flip-flop}，以及其它 \textcolor{red}{dedicated blocks}：如针对加法/乘法特别优化的 \textcolor{red}{DSP block}，用于片上数据缓存的 \textcolor{red}{block RAM (BRAM)}（大小适中，灵活）和 \textcolor{red}{ultra RAM (URAM)}（大小更大，不灵活）。最后，为了实现灵活性，需要有一张可以实现任意互联的 \textcolor{red}{static programmable interconnect}，由提供引脚连接的 \textcolor{red}{connection box} 和提供控制的可编程的 \textcolor{red}{switch box} 组成，功耗和占地面积均大。

FPGA 比定制硬件更灵活、更便宜，比通用处理器能效比更高；但是比定制硬件效率低，比通用处理器编程更困难，且 bit 级别的可重构性有时过于复杂。综合来看，最大的缺点在于面积效率低—为了实现通用 LUT 和通用互联的代价。

\textcolor{red}{Domain-Specific Architecture} 是对特定领域的定制硬件。例如，TPU 就有大量的乘法累加器以实现矩阵乘法，集成大量片上 SRAM，能效远超 GPU。

\textcolor{red}{Systolic Array} 是 TPU 的核心硬件，能高效执行矩阵乘法，每个单元从上方和左侧接受数据并相乘、累加，然后将上方数据下传、左侧数据右传。其并行度高，且仅需相邻通信。

DSA 的核心在于让多种算子使用同种硬件。例如，软件（编译器）可以通过重排把卷积变成矩阵乘法，也可以把小矩阵复制或大矩阵分块来匹配乘法维度。

TPU 的主要指令只有五条：\texttt{Read\_Host\_Memory} / \texttt{Write\_Host\_Memory} 读写数据、\texttt{Read\_Weights} 读取并存储权重、\texttt{MatrixMultiply / Convolve} 应用乘法/卷积、\texttt{Activate(ReLU / Sigmoid / Maxpool / LRN /...)} 应用对应操作，控制逻辑完全由 CPU 负责。

\section*{Conclusion}

\textbf{并行}：子任务必须独立且均衡。例子包括多核、superscalar、DRAM 的多 channel/bank。Data-level 的例如 SIMD、Instruction-level 的例如 OoO、Thread-level 的例如多核以及单核多线程。

\textbf{流水线}：一种特殊的并行方法。为了处理依赖关系，使用 forwarding。为了实现均衡，必须切分 stage。其提高流量但增加了延迟，需要更多的功耗和面积。

\textbf{乱序}：动态的探索并行能力的方式。

\textbf{投机}：存在依赖关系时，预测其输出。包括 branch target 预测、prefetch、load value prediction 等。同时其必须拥有 check 和 flush 的能力。

\textbf{缓存}：利用时域或空域的临近性。临近性意味着特定的 pattern，意味着可预测性。例子包括对指令、数据的 \$、对虚拟地址的翻译（TLB）等。

\textbf{间接访问}：允许拦截访问并引入额外功能。间接访问的能力和代价间存在 tradeoff。例子包括虚拟地址、虚拟机和 OoO 的寄存器重命名等。

\textbf{均摊}：cacheline 和 DRAM row 是一次取多个数据的均摊；FPGA 的编程开销；DSA 的成本开销。

\textbf{冗余}：加速处理并容忍错误。成本与表现/安全的权衡。各个核都有私有的 L1\$ 存储（一份数据存在多个复制）；使用双倍计算资源以减少单个服务器的偶然降速（减少 tail latency）；Error-Correcting Code、数据备份、分布式存储。

\textbf{专门化}：定制组件。

\textbf{Amdahl's Law}。

\$ capacity 增大可以降低 capacity miss，但一般会增加 hit time，是 trade-off。associativity 增大降低 conflict miss，增加 hit time 且边际效用递减。block size 增大可以降低 compulsory miss，但增加 miss penalty，产生 cache pollution 且减少 entry 数目。private 可以减少共享数据的 miss（即 warmup）。

多核随机 DRAM 访问的问题：多核的 false sharing、对 DRAM 带宽或 MSHR 的争抢；\$ miss；prefetch 的有用或污染；TLB 的 page table walk；DRAM 没有 row hit（取来的东西没用；每次取的代价高）；磁盘的高延时。但是，如果使用 close-page policy，或者考虑多 channel/rank/bank 的并行占满带宽，或者不好的 address mapping 把 row 乱映射，顺序和随机的 DRAM 访问可能结果相近。
\end{multicols}
\end{document}