B. 小北问答 -- 超速版¶
分数:100 分
1. Amdahl & Gustafson¶
📝 答题格式:填数字
某程序的代码中 10% 必须串行执行,90% 可完美并行。
- 根据 Amdahl's Law,无论核心数如何增加,该程序的理论最大加速比极限是
____倍; - 若在 10 核系统中通过扩大问题规模来保持每核计算负载不变,根据 Gustafson's Law,该系统的加速比将达到
____倍。
2. OpenMP¶
📝 答题格式:填大写字母(A/B/C/D)
以下代码使用 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 变量添加原子操作,保证结果正确 |
3. 低精度¶
📝 答题格式:填数字
已知 IEEE 754 标准的 FP32 拥有 8 位指数位。请问:
- BF16 拥有
____位指数位,____位尾数位 - NVFP4 拥有
____位指数位,____位尾数位
💡 提示:可以查阅资料,了解 NVFP4 如何在低精度下保持较高的数值范围和动态范围。
4. MPI 通信¶
📝 答题格式:填一行代码(分号结尾,区分大小写,空白会被忽略)
💡 MPI 常量可以使用,如
MPI_COMM_WORLD、MPI_INT等。
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}; // 环绕连接
/* 填这一行代码 */
5. NCCL 延迟¶
📝 答题格式:填数字
在深度学习的并行推理与训练中,进程之间会频繁进行集合通信操作。 NVIDIA 的开源集合通信库 NCCL,提供了在 GPU 之间进行集合通信的高性能解决方案。
当在异构的硬件上进行大规模集合通信时,如何选择通信的算法将很影响集合通信操作的效率。为了解决这个问题,NCCL 的解决方案是:基于一套硬编码的调优常数,估算不同集合通信算法下的集合通信完成时间,由此选择最优的算法。
问题:在 NCCL 2.28 的默认调优常量中,用 NVLink 连接的两 GPU 、采用 Tree 算法和 LL 协议时,在估算时每跳(单步)的硬件延迟取值为 ______ µs 。
💡 NCCL 用这类静态常数估算不同算法/协议的完成时间并做选择,这种静态建模也成为后续优化工作的切入点。
6. 高性能网络¶
📝 答题格式:不定项选择,多选少选错选都不得分
填大写字母,中间不加空格等符号,如
ABCDEFG
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 架构特别适合此类场景 |
7. GPU¶
📝 答题格式:不定项选择,多选少选错选都不得分
填大写字母,中间不加空格等符号,如
ABCDEFG
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 操作的延迟 |
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,第一个线性层按照输出维度进行拆分,第二个线性层按照输入维度进行拆分
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 的通信带宽
10. Cache 行为分析¶
📝 答题格式:填大写字母(A/B/C/D)
假设我们需要进行一个矩阵乘法 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 \times 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 总共有 S = 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 的跨度。
问题 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 \times 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 \times 8 的分块没有任何意义 |