Two and One

Posted Feb 04, 2026Updated Feb 04, 2026blog37 minutes read (About 5620 words)

HPCGames 题解 A B C 题

A 题

B 题

这里又来到了经典的小北问答环节，结合一些理论知识和具体论文的查阅，我们可以对题目进行详细的分析和解答。

1. Amdahl & Gustafson

某程序的代码中 10% 必须串行执行， 90% 可完美并行。

根据 Amdahl’s Law ，无论核心数如何增加，该程序的理论最大加速比极限是 ____ 倍；
若在 10 核系统中通过扩大问题规模来保持每核计算负载不变，根据 Gustafson’s Law ，该系统的加速比将达到 ____ 倍。

首先，根据 Amdahl 定律，加速比 S 可以通过以下公式计算：

$S = \frac{1}{( 1 - P ) + \frac{P}{N}}$

其中 P 是可并行部分的比例， N 是处理器的数量。对于该问题， P = 0.9 （ 90% 可并行），串行部分为 0.1 （ 10% 必须串行）。当 N 趋近于无穷大时，公式简化为：

$S_{ma x} = \frac{1}{( 1 - P )} = \frac{1}{0.1} = 10$

因此，该程序的理论最大加速比极限是 10 倍。

接下来，根据 Gustafson 定律，加速比 S 可以通过以下公式计算：

$S = N - (1 - P) \times (N - 1)$

在 10 核系统中， N = 10 ， P = 0.9 ，因此：

$S = 10 - (1 - 0.9) \times (10 - 1) = 10 - 0.1 \times 9 = 10 - 0.9 = 9.1$

因此，在 10 核系统中通过扩大问题规模来保持每核计算负载不变，该系统的加速比将达到 9.1 倍。

2. OpenMP

以下代码使用 OpenMP 并行执行循环：

int sum = 0; 
#pragma omp parallel for
for (int i = 0; i < 100; i++) { 
 sum += i; 
} 
printf(" sum = %d\n" , sum);

关于该代码，请问以下说法中正确的是 ____ 。

选项	描述
A	代码一定能正确计算出 0 到 99 的和（ 4950 ）
B	代码存在数据竞争，结果不确定。
C	`sum`变量默认为`private`，每个线程有自己的副本。
D	OpenMP 会自动为`sum`变量添加原子操作，保证结果正确。

正确答案是 B 。

解释如下：

选项 A 是错误的，因为代码中存在数据竞争，多个线程同时修改共享变量 sum，导致结果不确定。
选项 B 是正确的，因为在并行执行时，多个线程可能同时读取和写入 sum，导致数据竞争，从而使得最终结果不确定。
选项 C 是错误的，因为 sum 变量在默认情况下是共享的（ shared ），而不是私有的（ private ）。因此，所有线程都访问同一个 sum 变量。
选项 D 是错误的，因为 OpenMP 不会自动为共享变量添加原子操作。要确保结果正确，需要显式地使用 #pragma omp atomic 或其他同步机制来保护对 sum 的访问。

3. 低精度

已知 IEEE 754 标准的 FP32 拥有 8 位指数位。请问：

BF16 拥有 ____ 位指数位，____ 位尾数位
NVFP4 拥有 ____ 位指数位，____ 位尾数位

提示：可以查阅资料，了解 NVFP4 如何在低精度下保持较高的数值范围和动态范围。

经查阅资料可知： - BF16 拥有 8 位指数位， 7 位尾数位 - NVFP4 拥有 2 位指数位， 1 位尾数位

NVFP4 比较特殊，他只有 4 个 bit ，显然如果直接使用的话，其范围会很小，精度也不理想，但经过查阅资料可知：
NVFP4 首先将一组数视为了一个块，一个块中会共享一个高精度的 scale factor ，确定大致的数量级，然后 NVFP4 只存储每个数相对于这个 scale factor 的偏移量，这样就能在保持较大数值范围的同时，使用更少的位数来表示每个数，从而提高了存储效率和计算速度。

4. MPI 通信

4.1 基本原语

4 个进程执行以下代码，每个进程有局部值 local_val ，操作后每个进程都有所有进程的值。

int rank; 
MPI_Comm_rank(MPI_COMM_WORLD, & rank); 
int local_val = rank; // rank 为进程编号， 0~3
int recv_buf[4]; 
/* 填这一行代码 */

4.2 通信器

创建一个 2 维笛卡尔拓扑，尺寸为 2×2 ，行优先排列，允许环绕连接。

MPI_Comm comm_cart; 
int dims[2] = { 2, 2} ;
int periods[2] = { 1, 1} ; // 环绕连接
/* 填这一行代码 */

4.1 基本原语

可以使用 MPI_Allgather 来实现该功能。代码如下：

1	MPI_Allgather(& local_val, 1, MPI_INT, recv_buf, 1, MPI_INT, MPI_COMM_WORLD);

这行代码会将每个进程的 local_val 收集到所有进程的 recv_buf 中。

4.2 通信器

可以使用 MPI_Cart_create 来创建一个二维笛卡尔拓扑。代码如下：

1	MPI_Cart_create(MPI_COMM_WORLD, 2, dims, periods, 1, & comm_cart);

这行代码会创建一个 2x2 的笛卡尔拓扑通信器 comm_cart，并允许环绕连接。

5.NCCL 延迟

在深度学习的并行推理与训练中，进程之间会频繁进行集合通信操作。 NVIDIA 的开源集合通信库 NCCL ，提供了在 GPU 之间进行集合通信的高性能解决方案。

当在异构的硬件上进行大规模集合通信时，如何选择通信的算法将很影响集合通信操作的效率。为了解决这个问题， NCCL 的解决方案是：基于一套硬编码的调优常数，估算不同集合通信算法下的集合通信完成时间，由此选择最优的算法。

问题：在 NCCL 2.28 的默认调优常量中，用 NVLink 连接的两 GPU 、采用 Tree 算法和 LL 协议时，在估算时每跳（单步）的硬件延迟取值为 ______ µs 。

根据 NCCL 2.28 的默认调优常量，当使用 NVLink 连接的两 GPU ，采用 Tree 算法和 LL 协议时，每跳（单步）的硬件延迟取值为 0.6 µs 。

具体参考资料可见 NCCL Github src/graph/tunning.cc中的 151 行。

6. 高性能网络

Rail-optimized networking 与 Clos 都是高性能网络设计方案。以下说法正确的有：

选项	描述
A	在 Rail-optimized 网络中，来自不同 HB 域（ High-Bandwidth Domain ）但具有相同 local rank 的 GPU 会被连接到同一个 rail switch 上，以减少跨域通信的延迟
B	常见部署模式下， Rail-optimized 网络相比传统 Clos 网络的主要优势是完全不需要 Spine 层交换机，因此可以大大节省网络设备成本
C	Clos 网络因其使用 Spanning Tree Protocol (STP) 而在大规模部署时存在扩展性问题，这是 Rail-optimized 网络要解决的核心问题之一
D	Rail-optimized 网络保证了任何情况下集群内任意两个 GPU 之间都能以网络线速（如 400 Gbps InfiniBand ）进行通信，无论它们是否在同一个 rail 中
E	NCCL 2.12 引入的 PXN 特性可以结合 NVLink 和 PCI 通信来优化网络流量，这个优化对于 Rail-optimized 网络尤为重要
F	对于 LLM 训练工作负载，最优的通信策略会将大部分网络流量集中在相同 local rank 的 NIC 之间，并且会多用 NVLink 等高速互联进行跨 rail 交换，这使得 Rail-optimized 架构特别适合此类场景

正确答案是 A 、 E 、 F 。

解释如下：

选项 A 是正确的，这是 Rail-optimized 网络的核心定义。在这种架构中，网络拓扑是根据 GPU 的 rank 进行物理隔离的。例如，所有服务器上的 0 号 GPU 都连接到同一组交换机（ Rail 0 ）， 1 号 GPU 连接到另一组（ Rail 1 ）。这使得在进行数据并行（ Data Parallelism ）训练时， AllReduce 等操作只需在同一个 Rail 内进行，无需跨越复杂的交换层级，大大降低了拥塞和延迟。
选项 B 是错误的， Rail-optimized 仍然基于 Clos 架构，通常需要 Spine 交换机。 Rail-optimized 描述的是 Leaf 层交换机与 GPU 的连接方式以及流量的导向方式，而不是一种去除了 Spine 的新型拓扑。对于大规模集群（超过一个 Leaf 交换机的容量）， Rail 0 的 Leaf 交换机之间仍然需要通过 Spine 交换机互联，以构成一个完整的 Rail 0 网络平面。
选项 C 是错误的， Clos 网络并不使用 STP ，且 STP 是传统以太网的痛点。传统二层以太网使用 STP (Spanning Tree Protocol) 防止环路，这会导致大量链路被阻塞，带宽利用率低。而现代 Data Center Clos 网络（无论是基于 IP 路由的 ECMP 还是 InfiniBand ）的设计初衷就是利用所有链路进行负载均衡，完全摒弃了 STP 。因此， C 选项描述的前提本身就是错误的。
选项 D 是错误的， Rail-optimized 并不保证“跨 Rail”通信的效率等同于“同 Rail”。 Rail-optimized 的设计哲学是“专路专用”。虽然物理上可以通过 Spine 进行跨 Rail 通信（例如 Node A 的 GPU 0 发给 Node B 的 GPU 1 ），但这通常不是最优路径，且可能面临 oversubscription （收敛比）的问题。实际上，这种架构倾向于利用 E 选项和 F 选项提到的技术来避免在网络层面上进行跨 Rail 数据传输。
选项 E 是正确的， PXN (PCIe/NVLink Cross-NIC) 是解决 Rail 架构灵活性的关键。在 Rail-optimized 网络中，如果 GPU 0 需要向网络中的 Rail 1 发送数据，传统的路径非常低效（走 PCIe -> CPU -> NIC -> Switch ->…）。 NCCL 的 PXN 特性允许 GPU 0 通过 NVLink 直接把数据传给同机的 GPU 1 ，然后由 GPU 1 的 NIC （连接着 Rail 1 ）发送出去。这相当于在节点内部利用 NVLink 完成了“变轨”，从而充分利用 Rail 网络的优势。

选项 F 是正确的，这准确描述了 LLM 训练中的混合通信模式。在 LLM 训练中，通常结合了数据并行（ DP ）和模型并行（ TP/PP ）。

+ TP (Tensor Parallelism) 流量极大，通常限制在单机内部，完全走 NVLink 。

+ DP (Data Parallelism) 需要跨机同步梯度，流量发生在相同 rank 的 GPU 之间，这完美契合 Rail-optimized 的网络路径。

+ 如果需要跨 rank 的操作（如 Pipeline Parallelism 的某些阶段或特定的 All-to-All ），结合 NVLink （节点内）+ Rail （节点间）是目前最优的策略。

7. GPU

NVIDIA 的 Hopper 架构引入了 TMA （ Tensor Memory Accelerator ）以提升 GPU 内存访问效率。以下说法正确的有：

选项	描述
A	相比 cp.async ， TMA 可以直接将数据从全局内存加载到共享内存，无需经过寄存器中转，从而能节省寄存器
B	在 cutlass 的异步流水线抽象中， Producer 调用 producer_acquire 获取空闲的 buffer stage ，完成数据加载后调用 producer_commit 通知 Consumer ； Consumer 则通过 consumer_wait 等待数据就绪，使用完毕后调用 consumer_release 释放 buffer
C	在使用 TMA 进行数据传输时，所有参与的线程都需要执行相同的 TMA 指令， TMA 硬件会自动处理线程间的协调
D	Cutlass Pipeline 使用多级缓冲（ multi-stage buffering ），通过 PipelineState 追踪当前读写的 stage index 和 phase ，实现 Producer 和 Consumer 之间的流水线重叠
E	TMA 的 multicast 功能允许一次 TMA 操作将同一块数据广播到 Cluster 内的多个 Thread Block 的共享内存中，减少了重复的全局内存访问
F	TMA 描述符（ TMA Descriptor ）需要在 kernel 启动前在 host 端创建，描述符中包含了张量的形状、步长和 swizzle 模式等信息， kernel 执行时通过预取描述符（ prefetch_tma_descriptor ）来减少首次 TMA 操作的延迟

正确答案是 B D E F.

解释：

- 选项 A 是错误的。 Ampere 架构引入的 cp.async 指令同样也是绕过寄存器（ Register File ），直接将数据从全局内存（ GMEM ）搬运到共享内存（ SMEM ）。
- 选项 B 是正确的。这描述了 Cutlass 异步流水线中 Producer 和 Consumer 之间的交互方式，符合 Cutlass 的设计理念。
- 选项 C 是错误的。 TMA 操作允许线程组内的线程根据需要选择性地执行 TMA 指令，而不是所有线程都必须执行相同的指令。如果所有线程都执行，会导致重复发射多个拷贝操作（除非有特殊的掩码处理）。这一点与 Ampere 的 cp.async （通常每个线程负责一部分）不同。
- 选项 D 是正确的。 Cutlass Pipeline 确实使用多级缓冲，通过 PipelineState 来追踪当前读写的 stage index 和 phase ，从而实现 Producer 和 Consumer 之间的流水线重叠。
- 选项 E 是正确的。 TMA 的 multicast 功能允许一次 TMA 操作将同一块数据广播到 Cluster 内的多个 Thread Block 的共享内存中，减少了重复的全局内存访问。
- 选项 F 是正确的。 TMA 描述符需要在 kernel 启动前在 host 端创建，包含张量的形状、步长和 swizzle 模式等信息， kernel 执行时通过预取描述符来减少首次 TMA 操作的延迟。

8. LLM

对于参数如下的一个标准的 Transformer-Decoder 模型，所有的 all reduce 操作都使用 ring all reduce 。假设一共有 4 张卡。

模型参数

参数	值
层数	32 层
隐藏层维度 (h)	4096
FFN 结构	两层线性层，中间层维度为 4h
序列长度	2048
Batch Size	32
优化器	Adam + 混合精度训练
精度设置	参数和梯度使用 fp16 ， Adam 优化器状态使用 fp32 （包括 momentum 、 variance 和 master weights ）

问题

请计算在以下三种并行方式下，进行一个 batch 的前向传播和反向传播，每张卡需要的发送量（以 GB 为单位）：

数据并行：每张卡上存放完整的模型，把 batch 均匀拆分到每张卡上，分别计算完成后对梯度进行 All-Reduce 操作
流水并行：按层拆分模型放到不同卡上，只需要前向传播的时候发送 activation ，反向传播的时候发送 gradient 。（计算通信量时只考虑中间的卡）
张量并行：对于 MHA 操作，按照 head 拆分到不同卡上。对于 FFN ，第一个线性层按照输出维度进行拆分，第二个线性层按照输入维度进行拆分
计算过程如下：
- 数据并行：
  - 模型参数量：
    
    $P = 32 \times 12 \times (409 6^{2})$
  - 梯度大小：
    
    $G = P \times 2 bytes (fp16) = 12 GB$
  - 通信量：
    
    $通信量 = 2 \times \frac{N - 1}{N} \times G = 18 GB$
- 流水并行：
  - 每层激活大小：
    
    $A = 32 \times 2048 \times 4096 \times 2 bytes (fp16) = 0.5 GB$
  - 通信量：
    
    $通信量 = 2 \times A = 1 GB$
- 张量并行：
  - 单次 All-Reduce 大小：
    
    $A R = 32 \times 2048 \times 4096 \times 2 bytes (fp16) = 0.5 GB$
  - 单次 Ring All-Reduce 通信量：
    
    $单次通信量 = 2 \times \frac{N - 1}{N} \times A R = 0.75 GB$
  - 总通信量：
    
    $通信量 = 32 \times 4 \times 0.75 GB = 96 GB$

9. UB 互联

在高性能计算系统中，集合通信（ Collective Communication ）的性能主要受带宽（ Bandwidth ）与延迟（ Latency ）两个因素制约。

NVIDIA 通过 NVLink 与 NVSwitch 构建 GPU 间的高速 Scale-up 互联网络，而华为则提出了 Unified Bus （ UB ）协议，作为面向 NPU 的统一互联与内存访问机制。 UB 协议基于华为自研的 UB Switch 交换芯片，并通过高带宽物理链路 HCCS （ High-Capacity Coherent System ）进行连接。

传统 AI 集群通常以 8 卡服务器为基本单元进行 Scale-out 扩展，而华为在 CloudMatrix 384 （ CM384 ）架构中，通过两级 UB Switch 组网，将 384 颗昇腾 910C NPU 构建为一个统一的超节点（ SuperPod ）。在该超节点范围内，所有 NPU 均处于同一个低延迟的轨道优化网络中，实现全对等 Scale-up 互联。

CM384 进一步将 UB 网络划分为 7 个相互独立的物理平面。每颗 NPU 的 7 个 HCCS 接口分别接入不同的交换平面，从而保证大规模并行通信过程中，数据流在物理路径上完全隔离、无链路冲突。

问题

在 CloudMatrix 384 的标准满配部署方案中，为了支撑 384 颗昇腾 910C NPU 实现无收敛、全对等的 Scale-up 互联，系统采用两级交换架构。在该超节点的物理拓扑中，分别使用了：

____ 个 Level 1 UB Switch
____ 个 Level 2 UB Switch
最终实现了理论上 ____ GB/s 的系统级聚合带宽

假设 switch chip 提供的单个 Port 可以提供 28GB/s 的通信带宽

该问题直接查阅华为 CloudMatrix 384 的白皮书即可得到答案。

10. Cache 行为分析

假设我们需要进行一个矩阵乘法 $C = A \times B$ 。

测试环境

为了简化分析，假设：

参数类型	配置
数据类型	double (8 Bytes)
L1 Cache 大小	4KB (4096 Bytes)
相联度	直接映射 (Direct Mapped, E=1)
块大小	64 Bytes （ 1 个 Cache Line 可存 8 个 double ）
矩阵规模	A, B, C 均为 64×64 的方阵 (N=64)
存储方式	数组按行优先存储
内存对齐	A, B, C 的起始地址均对齐到 Cache 的起始 Set

代码实现

// 假设变量 sum 已优化到寄存器中，忽略 C 的访存影响
// 仅考虑内层循环中 A 和 B 的读取
for (int j = 0; j < 64; ++j) { // Loop 1
 for (int i = 0; i < 64; ++i) { // Loop 2
 double sum = 0.0; 
 for (int k = 0; k < 64; ++k) { // Loop 3
 sum += A[i][k] * B[k][j]; 
 } 
 C[i][j] = sum; 
 } 
}

问题 10.1

我们试图分析上述代码中最内层循环 Loop 3 对矩阵 $B$ 的访存行为。

已知 Cache 总共有 $4096/64 = 64$ 个 Set 。

在计算
$C [i] [j]$ 的过程中（即一次完整的 Loop 3 ），关于
$B [k] [j]$ 的 Cache Miss Rate （不命中率），下列说法正确的是：

选项	描述
A	12.5% - 这里有良好的空间局部性，每 8 个 double 只有 1 次 Miss
B	25% - 虽然是列优先访问，但 Cache 够大，只有冷不命中
C	约 50% - A 和 B 互相打架（冲突），导致一半的数据被驱逐
D	100% - 发生了严重的 Cache Thrashing （抖动），每次读取都是 Miss

💡 提示：计算一下访问 $B [k] [j]$ 和 $B [k + 1] [j]$ 时的内存地址差值（ Stride ），以及它们映射到的 Set Index 的跨度。

正确答案是 D 。

解释如下：

矩阵 B 是按行优先存储的，因此访问 B[k][j] 时， k 的变化会导致访问的内存地址以列为单位跳跃。
计算地址差值（ Stride ）：

$Stride = sizeof(double) \times N = 8 Bytes \times 64 = 512 Bytes$
每次访问 B[k][j] 时，地址增加 512 Bytes ，而每个 Cache Line 大小为 64 Bytes ，因此每次访问都会跨越多个 Cache Line 。
计算 Set Index 的跨度：

$Set Index Span = \frac{Stride}{Cache Line Size} = \frac{512 Bytes}{64 Bytes} = 8$
因为 Cache 有 64 个 Set ，跨度为 8 意味着每次访问都会映射到不同的 Set ，但由于 k 从 0 到 63 ，共有 64 次访问，这些访问会循环映射到同一组 Set 上，导致频繁的冲突和驱逐。
最终结果是每次读取 B[k][j] 都会导致 Cache Miss ，即 Cache Thrashing 。

问题 10.2
为了进一步提升矩阵乘法的效率，我们决定使用分块技术。你将矩阵分成了 $8 \times 8$ 的小块（ Block Size = 8 ）。

// 8x8 分块优化演示
for (int jj = 0; jj < 64; jj += 8) { 
 for (int ii = 0; ii < 64; ii += 8) { 
 for (int kk = 0; kk < 64; kk += 8) { 
 // 在这里处理 8x8 的子块乘法
 for (int j = jj; j < jj + 8; ++j) { 
 for (int i = ii; i < ii + 8; ++i) { 
 double sum = C[i][j]; // 简化写法
 for (int k = kk; k < kk + 8; ++k) { 
 sum += A[i][k] * B[k][j]; 
 } 
 C[i][j] = sum; 
 } 
 } 
 } 
 } 
}

针对一个 $8 \times 8$ 的 $B$ 矩阵子块（假设该子块已预加载），在处理该子块内部的计算时，关于其在 L1 Cache 中的状态，下列分析正确的是：

选项	描述
A	一个 8×8 的子块大小为 512 Bytes ，远小于 Cache 大小，因此完全没有冲突，所有数据都能驻留在 Cache 中
B	尽管子块很小，但由于 B 的原始列宽（ Stride ）很大，导致子块内的 8 行数据全部映射到了同一个 Set 中，依然存在严重的冲突
C	子块内的 8 行数据分别映射到了 8 个不同的 Set 中（ Set 索引间隔为 8 ），且在子块计算期间不会发生自我冲突（ Self-Conflict ）
D	分块主要是为了利用 L2/L3 Cache ，对这么小的 L1 Cache (4KB) 来说， 8×8 的分块没有任何意义

正确答案是 C 。

解释如下：

一个 8×8 的子块大小为：

$Block Size = 8 \times 8 \times sizeof(double) = 64 \times 8 Bytes = 512 Bytes$
该子块的大小（ 512 Bytes ）确实远小于 L1 Cache 大小（ 4KB ），但关键在于访问模式。
在处理该子块时，访问 B[k][j] 时， k 的变化会导致访问的内存地址以列为单位跳跃。
计算 Set Index 的跨度：

$Set Index Span = \frac{Stride}{Cache Line Size} = \frac{512 Bytes}{64 Bytes} = 8$
因此，子块内的 8 行数据分别映射到了 8 个不同的 Set 中，且在子块计算期间不会发生自我冲突（ Self-Conflict ）。
分块技术有效地利用了 Cache 的空间局部性，减少了冲突，提高了数据的命中率。

D. Hyperlane Hopper

A 题

B 题

1. Amdahl & Gustafson

2. OpenMP

3. 低精度

4. MPI 通信

4.1 基本原语

4.2 通信器

4.1 基本原语

4.2 通信器

5.NCCL 延迟

6. 高性能网络

7. GPU

8. LLM

模型参数

问题

9. UB 互联

问题

10. Cache 行为分析

C 题

Categories

Recents

Archives

Tags