小米 AI Infra 实习 一二面

发表于 更新于 分类于 ，

Q: Python深拷贝和浅拷贝的区别？

浅拷贝（copy.copy）只复制对象本身（最外层容器），不复制其引用的子对象，新旧对象共享内部可变对象；深拷贝（copy.deepcopy）递归复制所有层级的对象，新旧对象完全独立互不影响。

核心机制对比：

特性 浅拷贝 copy.copy 深拷贝 copy.deepcopy
复制层级 仅最外层 递归所有层级
内部可变对象 共享引用 独立副本
性能开销 低 高（需遍历对象图）
循环引用处理 不涉及 自动处理（memo字典）

关键行为细节：

  • 对于不可变对象（如int、str、tuple），浅拷贝通常不会创建新对象，而是直接复用引用（因为不可变，共享无害）。但如果tuple内部包含可变对象，浅拷贝后修改内部可变对象仍会相互影响。
  • 列表的切片操作 a[:]、list(a)、dict.copy() 都是浅拷贝。
  • deepcopy 内部维护一个 memo 字典追踪已复制对象，避免循环引用导致无限递归。自定义类可通过实现 __copy__ 和 __deepcopy__ 方法控制拷贝行为。
1
2
3
4
5
import copy
a = [[1, 2], [3, 4]]
b = copy.copy(a)       # 浅拷贝：b[0] is a[0] -> True
c = copy.deepcopy(a)   # 深拷贝：c[0] is a[0] -> False
a[0].append(5)         # b[0] = [1,2,5]，c[0] = [1,2] 不受影响

常见陷阱：NumPy数组的切片是视图（view），不是拷贝；需要 .copy() 方法获得独立副本。

Q: C++的三种智能指针及其用法？

C++11引入的三种智能指针通过RAII机制自动管理堆内存生命周期，消除手动new/delete的内存泄漏风险。

1. unique_ptr —— 独占所有权

独占所有权语义，同一时刻只有一个unique_ptr指向该对象。不可复制，只能通过 std::move 转移所有权。析构时自动释放资源。

  • 实现原理：内部仅存储一个裸指针（+可选删除器），禁用拷贝构造和拷贝赋值运算符。
  • 零开销抽象：sizeof(unique_ptr) == sizeof(raw pointer)（无自定义删除器时），无引用计数开销。
  • 适用场景：工厂函数返回值、独占资源（文件句柄、GPU buffer）。

2. shared_ptr —— 共享所有权

多个shared_ptr可指向同一对象，内部通过引用计数管理生命周期。最后一个shared_ptr销毁时自动释放资源。

  • 控制块结构：堆上分配的控制块包含——强引用计数（atomic）、弱引用计数（atomic）、删除器（deleter）、分配器（allocator）。
  • 线程安全性：引用计数的增减是原子操作（线程安全），但所指向对象的读写不是线程安全的。
  • make_shared 优势：一次内存分配同时创建对象和控制块（减少一次heap allocation），提高cache友好性，避免裸new异常安全问题。
  • 开销：每个shared_ptr 16字节（指针+控制块指针），控制块约32-40字节。

3. weak_ptr —— 弱引用观察者

不增加强引用计数，用于解决shared_ptr循环引用问题。使用前需 lock() 提升为shared_ptr，检查对象是否仍然存活。

  • 典型场景：观察者模式中观察者持有目标的weak_ptr、缓存系统中cache条目、树结构中子节点指向父节点。
1
2
3
// make_shared vs new的区别
auto p1 = std::make_shared<Widget>();  // 1次内存分配（对象+控制块连续）
std::shared_ptr<Widget> p2(new Widget()); // 2次分配（对象和控制块分开）

常见陷阱：避免从this创建shared_ptr（应继承 enable_shared_from_this）；避免循环引用（A持有B的shared_ptr，B持有A的shared_ptr）。

Q: 写时拷贝（COW）原理？

写时拷贝（Copy-On-Write）是一种延迟复制的优化策略，核心思想：只有在真正需要修改时才付出拷贝代价。

实现机制：

  1. 多个对象初始共享同一块内存数据。
  2. 通过引用计数追踪共享者数量。
  3. 读操作直接访问共享内存，无任何开销。
  4. 写操作时检查引用计数——若引用计数>1，先复制一份独立副本，再在副本上修改；若引用计数==1，说明独占，直接修改。

典型应用：

  • **Linux fork()**：子进程创建时不复制父进程的物理页面，只复制页表（标记为只读）。当任一进程写入某页时触发page fault，内核此时才复制该页。对于大量fork()+exec()的场景极大减少了无用拷贝。
  • 早期 std::string（GCC < 5）：多个string对象共享同一缓冲区，修改时才复制。但在多线程环境下引用计数需原子操作，性能反而下降，因此C++11后std::string标准禁止COW实现。
  • QEMU/虚拟化快照：qcow2镜像格式基于COW实现快照链。

性能权衡：COW在读多写少的场景收益大；如果写操作频繁，COW的引用计数检查和潜在的延迟拷贝反而增加开销。此外COW在多线程环境下需要原子操作保护引用计数，可能产生cache line争用。

Q: 零拷贝原理？

零拷贝（Zero-Copy）指数据传输过程中避免CPU在内核空间和用户空间之间进行冗余数据拷贝，从而降低CPU开销和内存带宽消耗。

传统I/O路径（4次拷贝+4次上下文切换）：

1
磁盘 → DMA → 内核页缓存 → CPU → 用户缓冲区 → CPU → Socket缓冲区 → DMA → 网卡

零拷贝技术演进：

技术 拷贝次数 CPU参与拷贝 适用场景
传统read+write 4次 2次 -
mmap+write 3次 1次 需要应用层处理数据
sendfile 3次 1次 静态文件传输
sendfile+SG-DMA 2次 0次 网卡支持scatter-gather
splice 2次 0次 管道中转

各技术详解：

  • mmap：将内核页缓存映射到用户空间地址，应用可直接读取（少一次内核→用户拷贝）。但仍需一次从页缓存到socket缓冲区的CPU拷贝。
  • sendfile：数据直接在内核空间从页缓存传到socket缓冲区，不经过用户空间。减少两次上下文切换。
  • sendfile + DMA scatter-gather：网卡支持scatter-gather DMA时，只需将文件描述符（地址+长度）传给socket缓冲区，DMA引擎直接从页缓存搬数据到网卡，实现真正的零CPU拷贝。
  • splice：通过内核管道作为中介，利用页面引用计数在两个fd间移动数据，无实际数据拷贝。

实际应用：Kafka使用sendfile实现高吞吐消息投递；Nginx的sendfile指令；RDMA在GPU通信中实现真正的零拷贝（绕过CPU和OS内核）。

Q: 大模型分布式训练的流程和并行策略如何选择？

分布式训练完整流程：

  1. 数据准备：数据分片、tokenization、构建DataLoader（每卡加载不同数据分片）。
  2. 模型初始化：加载预训练权重、按并行策略分片模型（TP切权重、PP分层、ZeRO分片参数）。
  3. 前向计算：各卡并行执行前向，TP需AllReduce/AllGather通信，PP需跨stage传输激活。
  4. 梯度计算：反向传播计算局部梯度。
  5. 梯度同步：数据并行维度做AllReduce聚合梯度（可与反向重叠）。
  6. 参数更新：优化器更新参数（ZeRO下各卡更新自己负责的分片）。

并行策略选择决策树：

场景 推荐策略 原因
模型放得进单卡 纯数据并行（DDP） 最简单、扩展性最好、通信开销最小
单层参数超单卡显存 张量并行（TP） 切分权重矩阵到多卡
模型层数多、跨节点 流水线并行（PP） 通信量小（仅传激活），适合高延迟网络
超长序列（>128K） 序列并行（SP/CP） 切分序列维度减少激活内存
训练规模极大 3D并行组合 TP节点内 + PP跨节点 + DP扩batch

实践中的选择逻辑：

  • 节点内（NVLink 600GB/s）：TP = 4/8，利用高带宽做频繁通信。
  • 节点间（IB 400Gb/s ~ 50GB/s）：PP和DP，通信量相对小。
  • 典型配置举例：训练Llama-70B，TP=8（节点内），PP=4（4个节点），DP=4（32卡总共），global batch通过梯度累积扩大。

Q: 各种并行策略介绍？

1. 数据并行（DP/DDP/ZeRO）

每卡持有完整模型副本，将训练数据的不同mini-batch分到各卡。前向和反向独立计算，梯度通过AllReduce同步后统一更新。通信量为 2 x 模型参数量（reduce-scatter + all-gather）。

  • DDP相比DP的改进：多进程（避免GIL）、梯度桶通信与计算overlap。
  • ZeRO-1/2/3逐步将优化器状态/梯度/参数分片，显存从16Φ降至16Φ/N。

2. 张量并行（Tensor Parallelism, TP）

将单个算子的权重矩阵按行或列切分到多卡，每卡计算部分结果后通过AllReduce或AllGather通信汇总。通信频率极高（每层2次），因此必须使用节点内NVLink高带宽互联。

  • 通信量：每层前向一次AllReduce（2 x hidden_size x seq_len x batch / TP），反向同等。

3. 流水线并行（Pipeline Parallelism, PP）

将模型按层切分为多个stage，各stage分布在不同设备上流水执行。每个micro-batch依次经过各stage。

  • 调度策略：GPipe（先全前向再全反向，bubble大）、1F1B（交替执行减少bubble到 (pp-1)/micro_batches）。
  • bubble占比 ≈ (PP-1) / (PP-1 + num_microbatches)，micro-batch越多bubble越小。

4. 序列并行（Sequence Parallelism, SP）

将序列维度切分到多卡，主要处理LayerNorm和Dropout等非张量并行区域。Megatron-SP将AllReduce拆为Reduce-Scatter + AllGather，每卡只处理1/TP的序列片段，激活内存降为1/TP。

5. 专家并行（Expert Parallelism, EP）

MoE模型中将不同expert分配到不同卡。Router决定token路由，通过All-to-All通信将token发送到对应expert所在卡。关键挑战是负载均衡。

Q: TP为什么有按行/列两种切分方式，分别对应什么？

在Transformer的MLP层中，存在两个线性变换：第一层升维（h→4h），第二层降维（4h→h）。两种切分方式配合使用能最小化通信次数。

列并行（Column Parallel）—— 应用于第一个线性层：

对权重矩阵A按列切分：A = [A1, A2]（每卡持有一半列）。计算 Y = XA 时：

  • 每卡计算 Yi = X × Ai，得到输出的不同列。
  • 输出在列维度上自然分割，无需通信即可直接作为下一层（第二个线性层）的输入。
  • 配合GeLU激活：GeLU(XA1) 与 GeLU(X) × A1 不等价，因此GeLU必须在分割后的输出上执行（这也是列切分放在第一层的原因）。

行并行（Row Parallel）—— 应用于第二个线性层：

对权重矩阵B按行切分：B = [B1; B2]。输入已经在各卡上按列分割：

  • 每卡计算 Zi = Yi × Bi，得到部分结果（维度相同但只是部分和）。
  • 需要一次 AllReduce求和 得到最终输出 Z = Z1 + Z2。

通信分析：

  • MLP前向：1次AllReduce。反向：1次AllReduce。
  • Attention层类似：QKV投影列并行，输出投影行并行，前向1次AllReduce。
  • 整体每个Transformer层：前向2次AllReduce（MLP + Attention），反向2次。

Q: Megatron-SP（序列并行）介绍？

Megatron-SP是在张量并行基础上的进一步优化，核心目标是减少非TP区域的激活内存占用。

背景问题：在标准TP中，LayerNorm和Dropout操作需要完整的激活张量（因为它们不在TP切分的维度上操作）。这意味着即使使用TP=8，这些操作的激活仍占完整大小，造成内存浪费。

Megatron-SP解决方案：

  1. 将LayerNorm和Dropout的输入在序列维度上切分到TP组内的各GPU上。
  2. 每卡只处理 1/TP 的序列长度对应的LayerNorm/Dropout操作。
  3. 进入TP区域（如QKV投影）前，通过AllGather收集完整激活。
  4. TP区域输出后，通过Reduce-Scatter同时完成规约和分散。

关键洞察：原来TP中的AllReduce = Reduce-Scatter + AllGather。Megatron-SP只是将这两步分开使用——AllGather在进入TP区域时执行，Reduce-Scatter在离开TP区域时执行。总通信量不变，但激活内存在非TP区域降为1/TP。

具体效果：

  • 激活内存节省：LayerNorm和Dropout的激活占比约30-40%，这部分降为1/TP。
  • 对于TP=8，非TP区域的激活内存减少8倍。
  • 通信量与标准TP完全相同（只是拆分了AllReduce的时机）。

Q: Transformer架构中有哪些层和算子？

Transformer由多个相同结构的层堆叠而成，每层包含以下核心组件和对应的计算算子：

Multi-Head Attention子层：

  • QKV线性投影：3个GEMM操作（或融合为1个），将hidden_state投影为Q/K/V
  • 注意力分数计算：矩阵乘 QK^T，计算量 O(n^2 x d)
  • Scale：逐元素除以 sqrt(d_k)
  • Causal Mask：上三角mask填负无穷（Decoder）
  • Softmax：指数+归一化，数值稳定需先减max
  • Attention加权：矩阵乘 softmax_output x V
  • 输出投影：1个GEMM，将多头拼接结果映射回hidden_size

Feed-Forward Network（MLP/FFN）子层：

  • Gate投影 + Up投影：2个GEMM（SwiGLU架构），将hidden_size映射到intermediate_size
  • 激活函数：SiLU/GeLU（逐元素非线性）
  • Down投影：1个GEMM，将intermediate_size映射回hidden_size

归一化与残差：

  • RMSNorm/LayerNorm：逐元素操作（求均值/方差、归一化、缩放）
  • 残差连接：逐元素加法

输入输出层：

  • Token Embedding：查表操作（大词表时可视为GEMM）
  • 位置编码：RoPE（旋转位置编码，在QK上应用旋转矩阵）
  • LM Head：GEMM + Softmax（或直接取argmax）

计算量分布（典型LLM如Llama-70B）：GEMM（MLP+Attention投影）占90%+，Softmax/Norm/激活占<10%。因此GEMM优化是核心。

Q: Encoder和Decoder的介绍与特点？

Encoder（双向注意力）：

  • 注意力机制：每个token可以看到输入序列中的所有token（full attention），没有mask。
  • 适合任务：需要全局理解的任务——文本分类、命名实体识别（NER）、句子相似度、特征提取。
  • 代表模型：BERT、RoBERTa、DeBERTa。
  • 训练方式：Masked Language Modeling（随机mask 15% token预测）。
  • 特点：不能直接做生成，但理解能力强。

Decoder（因果注意力）：

  • 注意力机制：因果mask（causal mask），每个token只能attend到自身及之前的token。训练时通过下三角mask实现并行计算。
  • 适合任务：文本生成、对话、代码生成——所有自回归生成任务。
  • 代表模型：GPT系列、Llama、Qwen、DeepSeek。
  • 训练方式：Next Token Prediction（每个位置预测下一个token）。
  • 特点：天然支持生成，通过in-context learning也能做理解任务。

Encoder-Decoder（交叉注意力）：

  • 结构：Encoder编码输入得到表示，Decoder通过cross-attention关注encoder输出并生成。
  • 适合任务：输入和输出结构不同的seq2seq任务——翻译、摘要、ASR。
  • 代表模型：T5、BART、Whisper。
  • 特点：encoder处理输入不需要因果约束，decoder生成时有因果mask。

为什么Decoder-only成为主流：训练效率高（每个token都贡献loss）、结构简单（易于并行和优化）、scaling law表现最好、统一了理解和生成（通过prompt工程）。

Q: FlashAttention介绍？

FlashAttention是一种IO-aware的精确注意力算法，核心创新在于通过分块计算（tiling）和在线统计量更新（online softmax）避免在HBM中显式存储 N x N 的注意力矩阵。

为什么需要FlashAttention：

  • 标准Attention的内存复杂度O(N^2)，当序列长度N=8192时，注意力矩阵占64M个float16 = 128MB/head。
  • GPU的HBM带宽（如A100 2TB/s）远慢于SRAM（19TB/s），大量中间矩阵的HBM读写成为瓶颈。
  • 实测标准Attention中，HBM访问时间占总时间的60-80%。

核心算法：

  1. 将Q按行分块（block size Br），K和V按行分块（block size Bc）。典型值：Br=Bc=128。
  2. 外层循环遍历Q的每个block。
  3. 内层循环遍历K、V的每个block：
    • 从HBM加载Qi、Kj、Vj到SRAM（片上共享内存，如192KB/SM）。
    • 在SRAM中计算局部注意力分数 Sij = Qi x Kj^T。
    • 使用online softmax增量更新全局最大值和归一化因子。
    • 更新输出结果（无需等所有block计算完）。
  4. 最终输出直接写回HBM，中间不存储N x N矩阵。

效果与性能数据：

  • 内存复杂度：O(N^2) → O(N)，节省的内存可用于增大batch或加长序列。
  • HBM访问量：从O(N^2 x d)降至O(N^2 x d^2 / SRAM_size)。
  • 实测加速：比PyTorch标准Attention快2-4倍（A100上），训练端到端加速15-20%。
  • 精度：数学等价（不是近似），结果bit-wise相同（除浮点重排序误差）。

V2改进：减少非matmul FLOPs、改进并行度（外循环在Q block上）、优化warp分工（4个warp各处理不同KV block而非reduce），速度比V1提升约2倍。

Q: Online Softmax介绍？

Online Softmax允许在不知道全部元素的情况下增量计算softmax，是FlashAttention实现分块计算的数学基础。

传统Softmax的问题（需两次完整遍历）：

  1. 第一遍：求全局max（数值稳定性）：m = max(x1, x2, …, xN)
  2. 第二遍：求归一化因子和输出：d = sum(exp(xi - m)), softmax(xi) = exp(xi - m) / d

这要求所有元素同时可用，不适合分块处理。

Online Softmax算法：

在一次遍历中维护两个统计量：

  • m：当前已见元素的最大值
  • d：当前累积的 sum(exp(xi - m))

当新block（含新的一组元素）到来时，更新规则：

1
2
3
m_new = max(m_old, m_block)        // 更新全局最大值
d_new = d_old × exp(m_old - m_new) // 修正旧累积和（因为max变了）
      + sum(exp(x_block - m_new))  // 加入新block的贡献

对于FlashAttention中输出O的更新：

1
2
O_new = O_old × (d_old × exp(m_old - m_new) / d_new)  // 修正旧输出的权重
      + P_block × V_block × (exp(m_block - m_new) / d_new)  // 加入新block的贡献

为什么有效：

  • 每个block处理后，O已经是基于当前所有已见block的正确softmax加权结果。
  • 只需存储O、m、d这三个统计量（O(N)空间），无需存储完整的N x N注意力矩阵。
  • 数学上完全等价于标准softmax（不是近似），只是改变了计算顺序。

Q: CUDA中的Block是软件还是硬件概念？

Block（Thread Block）是软件（编程模型）概念，是程序员定义的线程组织单位，用于表达哪些线程需要协作和共享数据。

软件层面：

  • 程序员在kernel launch时指定block大小（如 <<<grid, block>>>）。
  • Block内线程可以通过共享内存通信、通过 __syncthreads() 同步。
  • Block大小通常为32的倍数（如128、256），最大不超过1024个线程。

硬件映射：

  • 一个Block被调度到一个SM（Streaming Multiprocessor）上执行，整个生命周期不会迁移。
  • Block内线程共享该SM的shared memory资源。
  • 一个SM可同时驻留多个Block（受寄存器文件、共享内存、warp槽位限制）。
  • Block内的线程进一步被组织为Warp（32个线程），Warp是硬件的实际执行和调度单位。

层次关系：

1
2
3
4
Grid（软件，一个kernel的所有线程）
  └── Block（软件，线程协作单位） → 映射到 SM（硬件）
        └── Warp（硬件调度单位，32线程）→ 在SM的执行单元上锁步执行
              └── Thread（逻辑最小单位）→ 映射到 CUDA Core/Tensor Core

设计哲学：这种软硬件解耦使得同一CUDA程序可以在不同代GPU（SM数量不同）上运行——运行时由硬件调度器决定Block到SM的映射，程序员无需关心具体硬件配置。

Q: CUDA有哪些优化方法？

CUDA优化是一个系统工程，需要根据Profiling结果确定瓶颈类型（memory-bound vs compute-bound）后针对性优化。

1. 访存优化（解决memory-bound瓶颈）：

  • 合并全局内存访问（Coalesced Access）：确保warp内32个线程访问连续128字节地址，一次事务完成。错位/跳跃访问会拆分为多次事务，带宽浪费可达32倍。
  • 共享内存缓存：将频繁复用的数据从HBM（2TB/s）加载到shared memory（19TB/s/SM），典型如GEMM的tile。
  • 向量化读写（float4/int4）：一条指令搬运128位，减少指令数和事务数。要求16字节对齐。
  • 利用只读缓存：__ldg() 或 const __restrict__ 提示走texture/L1 readonly路径。

2. 计算优化（解决compute-bound瓶颈）：

  • 减少Warp Divergence：将相似条件的线程映射到同一warp，或用无分支代码（bitwise ops）替代if-else。
  • 利用Tensor Core：FP16/BF16/INT8矩阵运算，吞吐比CUDA Core高8-16倍。通过WMMA API或cuBLAS自动使用。
  • 快速数学函数：__expf()、__sinf() 等，牺牲少量精度换取数倍速度。
  • 循环展开：#pragma unroll 减少循环控制开销，增加ILP。

3. 并行度优化：

  • 提高Occupancy：平衡寄存器/共享内存使用，让更多warp同时驻留。用 __launch_bounds__ 限制寄存器。
  • 合理设置Block/Grid大小：Block为32倍数（满warp），Grid block数为SM数的整数倍（避免tail effect）。

4. 延迟隐藏：

  • Double Buffering/Prefetch：计算buffer A的同时异步加载buffer B（cp.async），交替使用。
  • 增加活跃Warp数：当一个warp stall在访存时，调度器切换到另一个就绪warp。

5. 算法/系统级优化：

  • 算子融合：多个kernel合并减少launch开销和中间tensor的HBM读写。
  • 异步执行：多Stream实现计算与通信overlap。
  • 持久化Kernel：kernel不退出，循环处理任务，减少launch开销。

Q: 访存优化方法有哪些？

访存优化是CUDA性能优化中最重要的方向，因为大部分kernel都是memory-bound的。

1. Coalesced Access（合并访问）

GPU内存系统以32字节/128字节为粒度的事务（transaction）传输数据。Warp内32个线程同时发起的访存请求，如果地址连续且对齐，可以合并为最少的事务：

  • 理想情况：32个线程访问连续128字节，仅1次128B事务。
  • 最差情况：32个线程各访问不同cache line，需要32次事务，带宽利用率仅1/32。
  • 实践：让thread_id对应连续的内存地址（如 data[threadIdx.x] 而非 data[threadIdx.x * stride]）。

2. 共享内存（Shared Memory）

片上SRAM（如A100每SM 192KB），带宽约19TB/s，延迟约20-30 cycles（vs HBM的~400 cycles）：

  • 使用模式：线程协作加载全局内存tile到shared memory → __syncthreads() → 多次重复访问shared memory中的数据。
  • Bank Conflict：32个bank，每bank宽4字节。同一warp中的不同线程访问同一bank的不同地址会串行化。解决：padding（在shared memory数组每行末尾加1个元素）或调整访问模式。

3. 向量化读写（Vectorized Load/Store）

使用 float4（128位）一次性读写4个float值：

  • 减少load/store指令数量（4倍）。
  • 减少内存事务数。
  • 前提条件：数据起始地址必须16字节对齐。

4. 数据预取（Prefetch/Double Buffering）

在计算当前数据的同时，异步加载下一批数据：

1
2
3
4
5
// 使用cp.async实现异步拷贝
cp.async.cg.shared.global [shared_ptr], [global_ptr], 16;  // DMA搬运，不占计算单元
cp.async.commit_group;
// ... 做计算 ...
cp.async.wait_group<0>;  // 等待搬运完成

5. 只读缓存路径

通过 __ldg() 或声明 const __restrict__ 指针，提示编译器使用L1 read-only cache路径，避免L1/shared memory的容量争用。

6. 减少冗余访存

算子融合避免中间结果写回HBM再读取。例如将LayerNorm的reduce、normalize、scale三步融合为一个kernel，中间结果保留在寄存器中。

7. 寄存器复用

让数据尽量停留在寄存器中：每个线程处理多个元素（增加ILP），循环展开后中间值不溢出到local memory。

Q: 计算和访存如何overlap？

计算和访存overlap的核心思想是让GPU的不同硬件单元同时工作——在计算单元执行运算的同时，内存子系统（DMA/TMA）执行数据搬运。

Warp级别天然延迟隐藏（TLP）：

GPU通过大量Warp切换天然隐藏访存延迟——当一个Warp的load指令发出后进入等待状态，调度器立即切换到另一个数据已就绪的Warp执行计算。这是GPU最基本的延迟隐藏机制，要求有足够多的活跃Warp（即足够的Occupancy）。

Double Buffering（显式软件流水线）：

1
2
3
4
阶段0: Load A0 到 buffer_A
阶段1: Compute buffer_A | Load A1 到 buffer_B  (overlap!)
阶段2: Compute buffer_B | Load A2 到 buffer_A  (overlap!)
...

实现要点：

  1. 分配两份shared memory buffer（或寄存器buffer）。
  2. 使用 cp.async（Ampere+）或 TMA（Hopper）指令发起异步拷贝，数据搬运由DMA硬件执行，不占用CUDA Core。
  3. 通过 cp.async.wait_group 或 mbarrier 等待数据就绪。
  4. 交替使用两个buffer，实现计算和搬运的完美重叠。

Hopper架构的Warp Specialization：

更进一步，将Block内的Warp分为Producer（专职搬运）和Consumer（专职计算）：

  • Producer Warp使用TMA硬件单元持续搬运数据到shared memory。
  • Consumer Warp从shared memory读取数据执行GEMM（WGMMA指令）。
  • 通过异步mbarrier协调，实现零等待流水线。

Stream级别的Compute-Communication Overlap：

在分布式训练中，使用不同的CUDA Stream让AllReduce通信和下一层的前向计算并行执行：

1
2
3
# DDP中的梯度桶通信与反向计算overlap
stream1: backward(layer_N)
stream2: allreduce(grad_layer_N-1)  # 已完成层的梯度立即开始通信

Q: GPU中L1和L2缓存的区别？

特性 L1 Cache L2 Cache
作用域 每个SM私有 全部SM共享
容量（A100） 192KB/SM（与Shared Memory共用） 40MB（全局）
容量（H100） 228KB/SM 50MB
访问延迟 ~30 cycles ~200 cycles
带宽 ~19 TB/s（每SM） ~5 TB/s（总）
管理方式 硬件自动 + 可配置与shared memory的比例 纯硬件自动管理
缓存行大小 128 bytes 128 bytes
功能定位 缓存当前SM的工作集 跨SM数据共享和全局访问加速

L1 Cache详解：

  • 与Shared Memory共享同一物理SRAM（如A100共192KB），可通过 cudaFuncSetAttribute 配置比例。
  • L1命中时无需访问L2和HBM，是全局内存访问的第一级加速。
  • 对于只读数据（__ldg），L1有单独的只读缓存路径。

L2 Cache详解：

  • 所有SM共用，当多个SM访问相同数据时L2提供缓存（如reduction操作的中间结果）。
  • Ampere+架构支持L2 cache residency control（cudaAccessPolicyWindow），可将热数据持久化在L2中。
  • L2是HBM的前端缓存，HBM bandwidth（2TB/s）是数据从DRAM到L2的带宽。

数据流路径：

1
Register → L1/Shared Memory（~30 cycles）→ L2 Cache（~200 cycles）→ HBM（~400 cycles）

Q: 共享内存和L1缓存的区别？

特性 Shared Memory L1 Cache
管理方式 程序员显式管理（__shared__声明） 硬件自动管理，对程序透明
生命周期 程序员控制（block存在期间有效） 缓存替换策略决定（LRU）
可靠性 数据一定在（程序员保证） 可能miss（需重新从下级加载）
访问模式 32个bank并行访问 Tag查找 + Hit/Miss判断
特殊问题 Bank Conflict（同bank不同地址串行） Cache Miss导致stall
编程复杂度 需要手动load/sync/访问 自动，无需额外代码
适用场景 明确知道数据复用pattern 不确定访问pattern时的通用加速

物理实现：

  • 两者共享同一块物理SRAM（如A100每SM 192KB），通过硬件配置决定各占多少。
  • 可选配置：prefer shared（164KB shared + 28KB L1）或 prefer L1（28KB shared + 164KB L1），默认平分。
  • Hopper架构（H100）最大支持228KB shared memory per SM。

Bank Conflict详解：

  • Shared memory被分为32个bank，每bank宽4字节，地址按 bank_id = (address / 4) % 32 映射。
  • 同一warp中如果两个线程访问同一bank的不同行地址，会产生conflict导致串行化。
  • 特殊情况：所有线程访问同一地址是broadcast（无conflict）。
  • 解决方案：padding（每行末尾加一个元素偏移bank映射）。

选择建议：

  • 数据复用次数多、访问pattern规律 → 用Shared Memory。
  • 临时数据、不确定是否复用 → 依赖L1 Cache。
  • GEMM中的tile → Shared Memory；逐元素kernel → L1自动缓存。

Q: 为什么需要共享智能指针（shared_ptr）？

当多个对象需要共同拥有某个资源的所有权，且无法确定谁最后使用完该资源时，unique_ptr的独占语义不适用，这正是shared_ptr的应用场景。

典型使用场景：

  1. 观察者模式：多个观察者引用同一个Subject对象，Subject应在所有观察者都不再需要时才析构。
  2. 图/DAG结构：多个节点指向同一子节点（如AST中多个表达式共享同一子表达式节点）。
  3. 缓存系统：缓存池和使用者同时持有引用，使用者释放后若缓存也不再持有则真正释放。
  4. 异步回调/Future：多个callback或continuation共享同一个上下文对象，最后一个完成时释放。
  5. 多线程共享数据：多个线程共同读取的只读数据结构。

为什么不用裸指针：

  • 裸指针无法表达所有权语义，不知道谁负责释放。
  • 手动管理多所有者场景极易出bug（double free或memory leak）。

shared_ptr的代价：

  • 控制块堆分配开销（make_shared优化为与对象同一次分配）。
  • 引用计数的原子操作开销（每次拷贝/析构涉及atomic increment/decrement，在高争用下可能成为瓶颈）。
  • 对象大小增加（16字节：对象指针 + 控制块指针）。

设计原则：优先使用unique_ptr（零开销），只在确实需要共享所有权时才使用shared_ptr。用weak_ptr打破循环引用。

Q: AllReduce介绍及实现方式？

AllReduce是分布式训练中最核心的集合通信原语，功能是将所有节点的数据做规约（如求和）后将完整结果广播到所有节点。

语义：假设N个节点各持有向量Vi，AllReduce后每个节点都得到 V1 + V2 + … + VN。

实现方式对比：

算法 通信量/节点 延迟（步数） 适用场景
Ring AllReduce 2(N-1)/N × M ≈ 2M 2(N-1) 大消息、节点数适中
Tree AllReduce 2M 2×log(N) 小消息、节点数多
Recursive Halving-Doubling 2M 2×log(N) 小消息、2的幂次节点
Chunked Pipeline ~2M O(N) + O(chunks) 超大消息

1. Ring AllReduce（NCCL大消息默认）：

N个节点组成逻辑环，数据切分为N个chunk，分两阶段：

  • Reduce-Scatter阶段（N-1步）：每步每节点向下一个发送一个chunk，接收并累加，最终每节点持有1/N数据的完整规约结果。
  • AllGather阶段（N-1步）：每步发送已完成的chunk，N-1步后所有节点持有完整结果。

2. Tree AllReduce（NCCL小消息默认）：

构建树形拓扑，先Reduce到根（logN步），再Broadcast从根到叶（logN步）。延迟优势在节点数多时显著，但根节点成为带宽瓶颈。

3. NCCL实际行为：

NCCL根据消息大小和网络拓扑自动选择算法：

  • 小消息（<256KB）：Tree算法（延迟优先）。
  • 大消息（>256KB）：Ring算法（带宽优先）。
  • NVLink拓扑：使用NVLink的全连接特性优化，不严格遵循ring。
  • 跨节点：利用IB RDMA的特性做multi-rail优化。

Q: Ring AllReduce的通信容量分析？

设定：N个节点，每个节点持有大小为M的数据需要AllReduce。

Ring将数据均分为N个chunk，每个chunk大小 M/N：

Reduce-Scatter阶段：

  • 共执行 N-1 步。
  • 每步每节点向下一个邻居发送 M/N 数据，同时从上一个邻居接收 M/N 数据并累加。
  • 单节点总发送量 = (N-1) × M/N。
  • 结束时每节点持有 1/N 数据的完整规约结果。

AllGather阶段：

  • 同样 N-1 步。
  • 每步每节点发送 M/N 已规约完成的chunk给下一个邻居。
  • 单节点总发送量 = (N-1) × M/N。
  • 结束时每节点持有完整的AllReduce结果。

总通信量分析：

  • 单节点总发送量 = 2 × (N-1)/N × M
  • 当N较大时趋近 2M（与节点数无关），这是带宽最优的。
  • 理论下界证明：AllReduce的通信量下界就是2M，Ring达到了这个最优。

延迟分析：

  • 总步数 = 2(N-1)
  • 每步延迟 = alpha（启动延迟）+ M/(N × BW)（传输延迟）
  • 总延迟 = 2(N-1) × [alpha + M/(N × BW)]
  • 延迟随N线性增长——这是Ring的主要缺点，节点数很多时延迟累积严重。

数值示例（A100集群，8节点Ring，梯度1GB）：

  • 每节点发送 ≈ 2GB，带宽200GB/s（IB HDR），传输时间 ≈ 10ms。
  • 14步 × alpha ≈ 14 × 5us = 70us。
  • 总时间 ≈ 10ms + 70us ≈ 10ms（大消息下alpha可忽略）。

Q: Tree AllReduce相比Ring有什么优点？

**核心优势：延迟 O(log N) vs Ring的 O(N)**。

在集合通信的延迟模型 T = alpha × steps + M/BW × steps 中：

  • Ring步数 = 2(N-1)，延迟与节点数线性增长。
  • Tree步数 = 2×log2(N)，延迟与节点数对数增长。

当alpha主导（小消息场景）时：

假设N=256节点：

  • Ring延迟 ≈ 2×255×alpha = 510×alpha
  • Tree延迟 ≈ 2×8×alpha = 16×alpha
  • Tree快 32倍！

Tree的适用场景：

  1. 小消息：梯度较小时启动延迟alpha占主导，Tree的logN步数优势大。
  2. 节点数多：成百上千个节点时，Ring延迟不可接受。
  3. 跨数据中心：高延迟WAN网络中，减少通信步数极为关键。
  4. 层次化拓扑：自然匹配fat-tree网络架构。

Tree的缺点：

  • 带宽利用率低：叶子节点带宽充分利用，但根节点附近成为瓶颈（需要处理所有子树的数据），总有效带宽约为Ring的 1/log(N)。
  • 大消息时根节点带宽饱和，性能远不如Ring。

NCCL的策略：

  • 消息 < 256KB：使用Double Binary Tree（两棵树交错，提高带宽利用率）。
  • 消息 > 256KB：使用Ring。
  • 实际阈值会根据拓扑和NIC数量动态调整。

Q: 手撕：反转矩阵？

（编程题）

Q: 手撕：合并两个有序链表？

（编程题）