卓驭 AI Infra 校招 一面

发表于 更新于 分类于 ，

Q: 算子优化从Profiling到落地的完整流程？如何判断一个算子是memory-bound还是compute-bound？

完整流程（工业实践中的闭环）：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
1. Profiling（定位热点）
   ├── Nsight Systems: 系统级时间线，找到耗时最长的kernel
   └── 排序: 按累计GPU时间找top-5热点kernel

2. 深度分析（理解瓶颈）
   ├── Nsight Compute: 单kernel的详细指标
   ├── Roofline分析: 确定memory-bound还是compute-bound
   └── 对比峰值: 实际带宽/FLOPS vs 硬件峰值(达峰率)

3. 优化实施
   ├── Memory-bound: 融合、向量化、shared memory缓存、减少读写次数
   ├── Compute-bound: Tensor Core、循环展开、增加ILP
   └── Latency-bound: 增加活跃warp、prefetch隐藏延迟

4. 验证
   ├── 性能验证: speedup测量（控制变量，多次warmup取稳定值）
   ├── 精度验证: vs FP64参考实现，各dtype容差检查
   └── 边界测试: 各种shape/dtype/device组合

5. 落地集成
   ├── 注册到框架dispatch表（torch.library或C++ TORCH_LIBRARY）
   ├── 实现shape推导（meta function）
   ├── 实现autograd（如需反向）
   └── CI集成（自动化测试+性能回归检查）

判断Bound类型——Roofline Model：

1
2
3
4
5
6
7
8
9
10
Performance (FLOPS)
     ▲
     │         ╱─── Peak FLOPS (Compute ceiling)
     │       ╱
     │     ╱  ← Roofline
     │   ╱
     │ ╱──────── Peak Bandwidth slope
     │╱
     └────────────────── Arithmetic Intensity (FLOPs/Byte) ►
           拐点 = Peak FLOPS / Peak BW

具体计算方法：

以A100为例：

  • Peak FP16 FLOPS = 312 TFLOPS
  • Peak HBM BW = 2 TB/s
  • 拐点 = 312T / 2T = 156 FLOPs/Byte

对于给定kernel：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
# GEMM [M, K] × [K, N] 的算术强度估算
FLOPs = 2 * M * N * K  # 乘加各算一次
Bytes = (M*K + K*N + M*N) * sizeof(dtype)  # 读A,B + 写C
AI = FLOPs / Bytes

# 示例: M=N=K=4096, FP16
FLOPs = 2 * 4096^3 = 137.4 GFLOPs
Bytes = (4096^2 * 3) * 2 = 100.7 MB
AI = 137.4G / 100.7M ≈ 1365 → 远大于156 → Compute-bound

# 示例: Element-wise ReLU, 4096×4096, FP16
FLOPs = 4096^2 = 16.8M  (每元素1次比较)
Bytes = 4096^2 * 2 * 2 = 67.1 MB  (读+写)
AI = 16.8M / 67.1M ≈ 0.25 → 远小于156 → Memory-bound

Nsight Compute直接判断：

  • Memory Throughput / Peak Memory BW > 70%且Compute Throughput / Peak Compute低 → Memory-bound。
  • 反之 → Compute-bound。
  • Nsight Compute的Speed of Light面板直接显示两者的达峰百分比。

Q: 用过哪些Profiling工具？Nsight Systems能看到指令级流水吗？

工具层次对比：

工具 粒度 功能 适用阶段
Nsight Systems 系统级 时间线、kernel/API调用序列 定位热点
Nsight Compute Kernel级 详细性能计数器、Roofline 优化单kernel
cuPTI API级 可编程性能采集 自建监控
nvprof/nvvp 已弃用 旧版profile 兼容旧环境
torch.profiler 框架级 Op级耗时+内存 Python层快速定位

Nsight Systems能做什么：

  • CPU和GPU的联合时间线（看到Python调用→CUDA API→kernel执行的完整链路）。
  • 多GPU时间线对齐（看通信和计算的overlap）。
  • NVTX标注区域的耗时统计。
  • API调用频率和开销（如cudaMalloc/cudaLaunchKernel的频率）。

Nsight Systems不能做什么：

  • 不能看指令级流水（每条指令的issue/execute/retire时序）。
  • 不能看单个warp的执行历史。
  • 不能看寄存器分配和spill详情。

指令级分析需要什么：

  • Nsight Compute的Source页面：可以看每条SASS指令的采样hit count（类似perf的热点行）。这是统计近似，不是精确的逐指令时序。
  • Nsight Compute的Warp State统计：按stall原因分类（memory dependency/barrier/instruction fetch等），给出warp时间分布。
  • SASS级汇编分析：cuobjdump反汇编kernel的SASS代码，人工分析指令依赖链和吞吐瓶颈。
  • 硬件模拟器（GPGPU-Sim等）：学术工具，可模拟逐cycle的流水线行为，但与实际硬件可能有偏差。

实践中的Profiling工作流：

1
2
3
4
5
第一步: nsys看全局 → 发现matmul_kernel占总时间40%
第二步: ncu分析该kernel → 发现是memory-bound (DRAM 85%利用, Compute仅20%)
第三步: ncu的Memory分析 → 发现L2 hit rate仅30%，大量冷数据访问
第四步: 检查代码 → 发现tiling不够，数据复用不足
第五步: 优化shared memory tiling → 重新profile验证

Q: Warp利用率低怎么归因？负载不均衡怎么解决？

Warp利用率低的归因分析：

在Nsight Compute中查看Warp State Statistics面板，各stall原因对应不同的root cause：

Stall原因 含义 典型root cause 解决方向
Long Scoreboard 等待全局内存/L2返回 数据预取不足、cache miss prefetch、增大tiling
Barrier 等待__syncthreads() Block内warp进度不一致 减少sync次数、减小block
Not Selected 就绪但未被调度 正常（其他warp在执行） 增加ILP
Wait 等待固定延迟指令完成 计算依赖链过长 循环展开、增加独立指令
Stall MIO Throttle shared memory bank conflict 多线程同时访问同bank padding、改访问模式
Dispatch Stall 指令发射受限 指令cache miss或调度器满 简化kernel逻辑

利用率低的常见原因：

  1. Block尾部不满warp：block大小非32倍数，最后一个warp只有部分线程活跃。

    1
    Block=100 → warp 0(32线程), warp 1(32线程), warp 2(32线程), warp 3(仅4线程活跃!)

    解决：block大小始终设为32的倍数。

  2. 分支Divergence：同一warp内线程走不同分支，串行执行两条路径。

    1
    2
    if (tid % 2 == 0) heavy_work();  // 50% divergence!
    else light_work();

    解决：数据重排使同warp线程走相同路径，或用predication替代分支。

  3. 数据依赖导致活跃warp不足：每线程使用大量寄存器→SM能驻留的warp少→延迟隐藏不足。
    解决：减少寄存器用量（__launch_bounds__限制）或使用double buffering流水化。

负载不均衡的解决方案：

1. 动态任务分配（Work Stealing）：

1
2
3
4
5
6
7
8
9
__device__ int global_counter = 0;  // 全局任务计数器

__global__ void dynamic_kernel(float* data, int N, int tasks_per_thread) {
    while (true) {
        int task_id = atomicAdd(&global_counter, 1);
        if (task_id >= N) break;
        process(data, task_id);  // 每次取一个任务
    }
}

每个线程做完当前任务后从全局计数器取下一个，自动均衡。代价：原子操作开销（但任务粒度够大时可忽略）。

2. Grid-stride Loop（最简单有效）：

1
2
3
4
5
6
7
__global__ void kernel(float* data, int N) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; 
         i < N; 
         i += gridDim.x * blockDim.x) {  // stride = grid总线程数
        process(data[i]);
    }
}

每线程处理多个元素，天然均衡（除非process耗时不均）。

3. Persistent Kernel（极致负载均衡）：

1
2
3
4
5
6
7
8
__global__ void persistent_kernel(TaskQueue* queue) {
    while (true) {
        Task t = queue->dequeue();  // 阻塞式获取任务
        if (t.is_done_signal) return;
        process(t);
    }
}
// 只启动刚好填满GPU的block数量，kernel长驻不退出

适合任务粒度差异大的场景（如稀疏矩阵、图算法）。

4. 数据重排/Padding：

1
2
3
// 不均衡：不同行长度不同（如CSR稀疏矩阵的各行非零元数差异大）
// 解决：按行长度排序，相近长度的行分给同一个warp
sort_by_row_length(rows);  // 预处理

Q: 昇腾NPU和NVIDIA GPU架构差异？内存层级设计？

核心架构对比：

维度 NVIDIA GPU (A100) 华为昇腾NPU (910B)
计算核心 CUDA Core(通用) + Tensor Core(矩阵) Cube(矩阵) + Vector(向量) + Scalar(标量)
计算粒度 Warp(32线程SIMT) 向量运算(2048-bit一次处理128个FP16)
矩阵运算 Tensor Core: 4×4 FMA/cycle Cube: 16×16 FMA/cycle
编程模型 CUDA (相对隐式内存管理) Ascend C (显式数据搬运)
Cache策略 L1/L2 Cache硬件自动管理 无自动Cache，靠程序员显式管理Buffer
内存层级数 3级(Registers→L1/Shared→L2→HBM) 4级(L0→L1 Buffer→L2→HBM)
总带宽(HBM) 2 TB/s ~1.6 TB/s
峰值算力(FP16) 312 TFLOPS ~320 TFLOPS

内存层级详细对比：

1
2
3
4
5
6
7
8
9
10
11
NVIDIA GPU:                          华为昇腾NPU:
┌────────────┐                      ┌────────────┐
│ Registers  │ ~256KB/SM, 1 cycle   │ L0 Buffer  │ ~1KB, 寄存器级
├────────────┤                      ├────────────┤
│ Shared Mem │ 192KB/SM, ~20 cy     │ L1 Buffer  │ 512KB/AI Core, ~5 cy
│ / L1 Cache │ 可配置比例            │            │ 程序员显式管理
├────────────┤                      ├────────────┤
│ L2 Cache   │ 40MB, ~200 cy        │ L2 Cache   │ ~96MB, ~100 cy
├────────────┤                      ├────────────┤
│ HBM        │ 80GB, ~400 cy        │ HBM        │ 64GB, ~300 cy
└────────────┘                      └────────────┘

编程模型的核心差异：

GPU（CUDA）：

1
2
3
4
5
6
7
// 程序员无需显式搬运数据，硬件Cache自动管理
__global__ void add(float* a, float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        c[i] = a[i] + b[i];  // 直接访问全局内存，L1/L2 Cache自动缓存
}
// 可选优化：手动使用shared memory提升性能

NPU（Ascend C）：

1
2
3
4
5
6
7
8
9
10
11
12
13
// 必须显式管理数据搬运，否则每次计算都从HBM读取（极慢）
AscendC::GlobalTensor<float16_t> inputGlobal;
AscendC::LocalTensor<float16_t> inputLocal;  // L1 Buffer中的局部tensor

// 1. 显式搬入: HBM → L1 Buffer
DataCopy(inputLocal, inputGlobal, count);
PipeBarrier<PIPE_MTE2>();  // 等待搬运完成

// 2. 计算（数据必须在L1 Buffer中）
Add(outputLocal, inputA_local, inputB_local, count);

// 3. 显式搬出: L1 Buffer → HBM
DataCopy(outputGlobal, outputLocal, count);

关键差异总结：

  • GPU有Cache兜底：即使不优化shared memory，程序也能正确运行（只是慢）。
  • NPU无Cache兜底：不手动搬运=每次从HBM读=极慢。编程者必须理解内存层次。
  • 这使得NPU编程更接近传统DSP/FPGA的思维模式，对程序员要求更高。
  • NPU优势：显式管理带来更可预测的性能，适合固定计算模式的深度学习算子。

Q: 多进程和多线程的性能区别？

系统级对比：

维度 多线程 多进程
地址空间 共享（同一虚拟地址空间） 独立（各自的地址空间）
创建开销 ~10us (clone少量数据结构) ~1ms (fork/exec,复制页表)
上下文切换 ~1-5us (仅保存寄存器) ~10-100us (切换页表+刷TLB)
通信方式 直接读写共享变量 IPC(管道/shm/socket)
同步 mutex/condvar/atomic 信号量/消息/文件锁
故障隔离 无（一线程crash全进程挂） 有（一进程crash不影响其他）
内存开销 仅栈空间（~8MB/thread） 完整进程空间（虚拟内存COW）

性能关键差异：

1. 数据共享效率：

1
2
3
4
5
6
7
# 多线程：直接共享内存，零拷贝
shared_data = np.zeros(1000000)  # 所有线程直接访问同一数组

# 多进程：需要显式IPC
from multiprocessing import shared_memory
shm = shared_memory.SharedMemory(create=True, size=1000000*8)
# 或用 torch.multiprocessing 的share_memory_()

多线程共享大数据（如模型权重、特征图）时性能远超多进程（后者需要显式共享内存映射或数据序列化）。

2. Python GIL的影响：

1
2
3
4
5
6
7
# Python多线程受GIL限制，无法真正并行CPU计算
# 但IO密集型（网络/文件读写）不受影响

# CPU密集型任务必须用多进程
from concurrent.futures import ProcessPoolExecutor
with ProcessPoolExecutor(max_workers=8) as pool:
    results = pool.map(cpu_heavy_func, data_chunks)

3. GPU训练中的选择：

DDP使用多进程而非多线程的原因：

  • 绕过Python GIL（每个进程独立Python解释器）。
  • 每进程独立管理一块GPU的CUDA context（CUDA context不能跨线程安全共享所有操作）。
  • 进程隔离：一个进程OOM/crash不影响其他（可以重启恢复）。
  • 内存隔离：每进程独立的显存空间，NCCL通信明确定义。

4. 数据加载的选择：

1
2
3
4
5
# PyTorch DataLoader使用多进程
loader = DataLoader(dataset, num_workers=8, pin_memory=True)
# 8个worker进程并行做数据预处理（CPU密集），通过shared memory传给主进程

# 如果用多线程：GIL导致预处理无法并行利用多核

性能选择决策树：

1
2
3
4
5
任务需要共享大量数据且IO密集？ → 多线程
任务是CPU计算密集（Python）？ → 多进程（绕GIL）
任务需要故障隔离？ → 多进程
任务是GPU计算？ → 多进程（每进程一块GPU）
高并发网络IO（10K+连接）？ → 协程/asyncio

Q: KV Cache、算子融合、量化分别如何优化推理？

三种技术从不同层面加速推理，且可叠加使用：

KV Cache（消除冗余计算）：

1
2
3
4
5
无KV Cache (每步计算全序列attention):
  Step n: 计算token 0~n的全部QKV → O(n²) 每步

有KV Cache (增量计算):
  Step n: 只计算token n的QKV，与缓存的K[0:n-1], V[0:n-1]做attention → O(n) 每步
  • 效果：decode从O(n^2)总计算降为O(n)每步（累计仍O(n^2)但节省了n倍重复计算）。
  • 代价：需要O(n×layers×kv_heads×head_dim)的显存存储KV Cache。
  • 进一步优化：PagedAttention(内存管理)→GQA(减少KV头数)→量化KV(减少每token字节)。

数值示例（LLaMA-7B, seq=2K, FP16）：

  • 无KV Cache每步计算量：~2K × 2K × 4096 × 32层 × 2 ≈ 34T FLOPs（不现实！）
  • 有KV Cache每步计算量：1 × 2K × 4096 × 32层 × 2 ≈ 17G FLOPs（合理）

算子融合（减少内存带宽浪费）：

1
2
3
4
5
6
7
8
9
10
11
未融合（3个独立kernel）:
  Kernel 1: bias_add   → 读input(HBM) + 读bias → 写output(HBM)
  Kernel 2: layer_norm → 读input(HBM) → 写output(HBM)
  Kernel 3: relu       → 读input(HBM) → 写output(HBM)
  总HBM访问: 6次读 + 3次写 = 9次tensor大小的HBM I/O

融合为1个kernel:
  Fused Kernel: 读input(HBM) + 读bias → bias_add → layernorm → relu → 写output(HBM)
  中间结果全在寄存器/shared memory中
  总HBM访问: 2次读 + 1次写 = 3次tensor大小的HBM I/O
  I/O减少67%！

对于memory-bound操作（大多数非GEMM操作），I/O减少直接等比加速。

典型融合收益：

  • Add+LayerNorm融合：2x加速。
  • QKV三个GEMM合一：1.3x加速（减少kernel launch + 输入只读一次）。
  • GEMM+bias+activation: 1.1x加速（bias和activation几乎零开销融合到epilogue）。

量化（减少数据传输量）：

1
2
3
FP16权重 (70B模型):  140GB
INT8权重 (70B模型):   70GB  → HBM带宽需求减半 → decode加速~2x
INT4权重 (70B模型):   35GB  → HBM带宽需求降75% → decode加速~2-3x（受反量化计算限制）

Decode阶段是memory-bound（AI≈2），权重读取是瓶颈。量化直接减少需要读取的字节数：

精度 模型大小(70B) 理论decode加速(vs FP16) 实际加速
FP16 140GB 1x baseline
INT8 70GB 2x 1.7-1.9x
INT4 35GB 4x 2.5-3x
FP8 70GB 2x 1.8x

实际加速低于理论值因为：反量化计算开销、kernel效率、量化不能加速attention中的KV Cache读取。

三者叠加效果：

  • KV Cache（基础必备）+ 算子融合（全链路优化）+ W4A16量化：总计可达5-10x加速（相比naive实现）。

Q: 模型输出与预期不符怎么debug？误差累积怎么解决？

系统化Debug方法：

Step 1: 确认问题范围

1
2
3
4
5
6
7
8
9
10
# 首先检查是全部输出错还是部分
with torch.no_grad():
    ref_out = reference_model(input)
    my_out = my_model(input)
    
    # 全局误差统计
    max_err = (ref_out - my_out).abs().max()
    mean_err = (ref_out - my_out).abs().mean()
    cos_sim = F.cosine_similarity(ref_out.flatten(), my_out.flatten(), dim=0)
    print(f"Max: {max_err:.6f}, Mean: {mean_err:.6f}, Cosine: {cos_sim:.6f}")

Step 2: 逐层二分定位

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
# 在模型中间插入hook，逐层对比
hooks = {}
def make_hook(name):
    def hook(module, input, output):
        hooks[name] = output.detach().clone()
    return hook

# 给参考模型和自实现模型的对应层都加hook
for name, layer in ref_model.named_modules():
    layer.register_forward_hook(make_hook(f"ref_{name}"))
for name, layer in my_model.named_modules():
    layer.register_forward_hook(make_hook(f"my_{name}"))

# 前向后对比每层输出，找到第一个误差超标的层
for i in range(num_layers):
    ref = hooks[f"ref_layer_{i}"]
    my = hooks[f"my_layer_{i}"]
    err = (ref - my).abs().max()
    if err > threshold:
        print(f"Layer {i}: first divergence! Max err = {err}")
        break

Step 3: 单层深入分析

  • 确认权重完全一致（bit-exact对比加载的权重）。
  • 确认输入完全一致（包括attention mask、position ids）。
  • 检查数据类型是否匹配（FP32 vs FP16可能在softmax等处产生差异）。
  • 检查是否有NaN/Inf（torch.isnan(output).any()）。

Step 4: 常见bug类型：

症状 可能原因 验证方法
输出全为NaN 除以0/exp溢出 在softmax前检查score范围
输出全为0 ReLU过度kill/权重错位 检查激活分布
输出正确但偏移 bias遗漏/scale错误 检查各层的bias和scale
精度逐层下降 累积误差/精度不够 逐层比较误差增长曲线
随机错误 未初始化内存/竞态 固定seed+deterministic模式

误差累积的解决方案：

问题来源：每层的微小误差（如量化舍入、低精度计算）经过N层传播后可能放大到不可接受。

1
2
3
4
Layer 0: err = 1e-4
Layer 10: err = 1e-3 (10x放大)
Layer 40: err = 1e-2 (100x放大)
Layer 80: err = 1e-1 (可能导致输出质量明显下降)

解决策略：

1. 关键层用高精度：

1
2
3
4
5
# Softmax和LayerNorm对精度最敏感（涉及exp和除法）
class SafeLayerNorm(nn.Module):
    def forward(self, x):
        # 在FP32下计算，结果转回原精度
        return F.layer_norm(x.float(), self.weight.shape).to(x.dtype)

2. 混合精度累加：

  • GEMM：输入FP16/INT8，累加器FP32/INT32。Tensor Core天然支持。
  • Reduce操作（sum/mean）：始终用FP32累加。
  • Softmax的exp和sum：FP32计算后转回FP16。

3. 量化感知训练（QAT）：

1
2
3
# 训练时模拟量化误差，让模型适应
x_quant = fake_quantize(x, scale, zero_point, quant_min, quant_max)
# 前向用量化值，反向用STE(Straight-Through Estimator)传梯度

4. 逐层校准优化：

  • 使用更精细的量化粒度（per-group而非per-tensor）。
  • 选择更好的校准方法（percentile/KL散度而非minmax）。
  • 对误差敏感层使用更高精度（如attention层FP16，其他INT8）。

5. 输出校正：

  • 如果已知系统性偏移（如量化导致的均值偏移），可以加learnable correction layer。
  • AdaRound等方法在量化时考虑层间误差传播进行联合优化。