海康威视 AI Infra

发表于 更新于 分类于 ，

Q: CUDA中什么是Bank Conflict？如何避免？

Bank Conflict 的本质：

Shared Memory 被划分为 32 个等宽的 bank（每个 bank 宽 4 字节），连续的 4 字节地址依次分配到不同 bank（bank 0 -> bank 1 -> … -> bank 31 -> bank 0 循环）。

当同一个 warp（32 线程）中多个线程同时访问同一个 bank 的不同地址时，这些访问必须串行化，称为 bank conflict。N-way conflict 意味着该访问被拆分为 N 次串行操作，延迟变为 N 倍。

特殊情况：

  • 多线程访问同一 bank 的同一地址：触发 broadcast（广播），无冲突，一次完成
  • 32 线程各访问不同 bank：理想情况，一次完成（如连续地址访问）
  • 全部 32 线程访问同一 bank 不同地址：32-way conflict，最坏情况

常见冲突模式：

  • 按 stride=32 访问：smem[threadIdx.x * 32]，所有线程访问同一 bank（32-way conflict）
  • 按 stride=2 访问：smem[threadIdx.x * 2]，每 2 线程访问同一 bank（2-way conflict）
  • 规律：stride 为 32 因子的倍数时产生冲突

避免方法：

  1. Stride=1 连续访问：smem[threadIdx.x]，32 个线程自然访问 32 个不同 bank，零冲突。这是最理想的访问模式。

  2. Padding（填充）：在 shared memory 声明时每行多加一个元素

    • 原始：__shared__ float smem[32][32]（列访问有 32-way conflict）
    • Padding：__shared__ float smem[32][33]（列访问时 stride 从 32 变为 33，打破冲突模式）
    • 代价：浪费 3% 的 shared memory 空间

  3. 数据布局调整：针对算法的访问模式设计 swizzle 布局（如 CUTLASS 中的 SmemSwizzle），将逻辑连续但物理冲突的地址映射为无冲突的物理位置。

  4. 访问模式重组：改变算法使相邻线程访问相邻地址。如矩阵转置中，先连续读（无冲突）到 shared memory，padding 后再连续写。

诊断工具：Nsight Compute 的 Shared Memory 面板直接显示 bank conflict 的次数和 wavefronts，可定量评估优化效果。

Q: CUDA实现Softmax的要点？

CUDA Softmax 实现需要处理数值稳定性和高效 reduction 两个核心问题：

算法步骤（Safe Softmax）：

  1. 求 max（reduction over hidden_dim）
  2. 计算 exp(x_i - max) 并求 sum（reduction over hidden_dim）
  3. 归一化：result_i = exp(x_i - max) / sum

实现要点：

1. 并行策略选择（根据 hidden_dim 大小）：

  • hidden_dim ≤ 32：一个 warp 处理一行，利用 warp shuffle（__shfl_down_sync）做 reduction，零 shared memory 使用
  • 32 < hidden_dim ≤ 1024：一个 block 处理一行，block 内 warp shuffle + shared memory 跨 warp reduction
  • hidden_dim > 1024：一个 block 处理一行但需要多次迭代加载（每个线程处理多个元素后归约）

2. Warp-level Reduction（最高效的小规模归约）：

1
2
3
// 求 warp 内 max
for (int offset = 16; offset > 0; offset >>= 1)
    val = max(val, __shfl_down_sync(0xffffffff, val, offset));

无需 shared memory 和 __syncthreads，延迟最低（~5 周期完成 32 线程归约）。

3. 数值精度：

  • FP16 输入时，reduction（累加 exp）使用 FP32 累加器防止精度损失
  • 结果写出时再转回 FP16
  • Max 求取用 FP16 即可（无精度问题）

4. 向量化加载：

  • 使用 float4/half2 128-bit 加载指令，一次加载 4 个 float32 或 8 个 float16
  • 对 hidden_dim=4096，使用 float4 时线程数可减少到 4096/4=1024

5. 融合优化：

  • Softmax 通常与前序的 QK^T scaling 或后序的 dropout 融合
  • 融合后中间结果（exp 值）保留在寄存器中，避免写出到 Global Memory 再读回

性能参考：对 hidden_dim=4096、batch*seq=4096 的 softmax，优化后 kernel 可达 HBM 带宽的 80%+（因为 softmax 是典型的 memory-bound 操作）。

Q: Online Softmax与FlashAttention的关系？

Online Softmax 是 FlashAttention 能够实现”精确 attention + 分块计算”的数学基础：

Online Softmax 的核心贡献：

传统 softmax 需要全局 max 和 sum，看似无法分块计算。Online Softmax 证明可以增量式维护这两个统计量：

1
2
3
处理第 j 个元素时：
  m_new = max(m_old, x_j)
  d_new = d_old * exp(m_old - m_new) + exp(x_j - m_new)

关键洞察：当发现新的 max 时，历史累积的 sum 可以通过乘以修正因子 exp(old_max - new_max) 精确校正。

在 FlashAttention 中的应用：

FlashAttention 将 K/V 分为多个块（如每块 128 个 token），逐块计算：

  1. 处理第 i 块 K/V：

    • 计算局部 attention score：S_i = Q * K_i^T / sqrt(d)
    • 计算局部 max: m_i = max(S_i)
    • 更新全局 max: m_new = max(m_prev, m_i)
    • 修正之前块的输出：O_prev *= exp(m_prev - m_new) / exp(m_i - m_new)
    • 计算当前块贡献：O_i = softmax_local(S_i) * V_i
    • 合并：O = O_prev + O_i（加权）

  2. 所有块处理完后：O 即为精确的 attention 输出（非近似）

为什么不需要存储 N*N 矩阵：

  • 每个 Q 块只需要与当前 K 块计算一个小的 attention score 矩阵（如 128*128）
  • 这个小矩阵完全存在于 SRAM（shared memory）中
  • 处理完后用 Online Softmax 修正因子合并到累积输出中
  • 下一块 K/V 覆盖当前 SRAM 内容，无需保留

性能收益链：Online Softmax -> 允许分块 -> 小块在 SRAM 中计算 -> 避免 NN 矩阵写入 HBM -> IO 从 O(N^2) 降为 O(N^2d/M) -> 实际 2-4x 加速

Q: 大模型的端侧部署挑战和方法？

端侧部署（手机/嵌入式/Edge 设备）的核心约束是”在有限资源下运行尽可能大的模型”：

核心挑战：

约束 具体限制 影响
内存 手机 6-12GB RAM（与系统共享） 7B FP16 模型 14GB 无法直接加载
算力 手机 GPU ~2-5 TFLOPS 生成速度慢（<5 tokens/s for 7B）
功耗 手机电池 3-5W TDP 长时间推理导致发热降频
存储 模型下载和存储空间限制 用户不愿下载 >5GB 的应用
延迟 用户期望实时响应 First token latency 需 <1s

解决方法：

1. 激进量化（最直接有效）：

  • W4A16：权重 4-bit，激活 16-bit。7B 模型压缩到 ~3.5GB（可加载到手机内存）
  • W4A8/W3A8：更极端压缩，配合 GPTQ/AWQ 保持精度
  • 1-2 bit 量化（BitNet/AQLM）：研究前沿，7B 模型 <2GB
  • 量化格式选择：Q4_K_M（llama.cpp 的 k-quant）是精度/大小的甜点

2. 模型蒸馏/小模型：

  • 大模型蒸馏到小模型（如 7B -> 1.5B/3B），保留大部分能力
  • 专用小模型：Phi-3-mini（3.8B）、Gemma-2B 等天然适合端侧
  • 领域蒸馏：针对特定任务蒸馏的小模型可以匹敌大模型在该任务的表现

3. 架构优化：

  • 减少层数/head 数（减少 KV Cache 和计算量）
  • 使用 GQA/MQA 减少 KV Cache 大小
  • Sliding Window Attention 限制上下文长度
  • 结构化剪枝移除不重要的 attention head 或 FFN 神经元

4. 专用推理框架：

  • llama.cpp：纯 C/C++ 实现，支持所有主流量化格式，CPU/GPU/Metal 后端。单线程 7B Q4 在 M2 上 ~30 tokens/s
  • MLC-LLM：基于 Apache TVM，编译优化适配不同硬件（手机 GPU/NPU）
  • MediaPipe LLM：Google 的端侧 LLM 框架，利用 GPU delegate
  • ONNX Runtime Mobile：跨平台推理引擎

5. 异构加速：

  • CPU + GPU + NPU 联合推理：GEMM 在 GPU/NPU 上执行，轻量计算在 CPU
  • iOS Metal / Android Vulkan GPU 加速
  • 高通骁龙的 Hexagon NPU、Apple 的 Neural Engine
  • 挑战：数据在 CPU/GPU/NPU 间搬运的开销可能抵消加速收益

实际性能参考：iPhone 15 Pro 运行 LLaMA-2 7B Q4_K_M，约 15-20 tokens/s（Metal GPU 加速）。这是目前端侧大模型的实用水平。

Q: 模型导出时的动态维度问题？

模型导出（如 PyTorch -> ONNX -> TensorRT）时，动态维度是最常遇到的工程问题：

问题本质：推理时 batch_size 和 sequence_length 通常是变化的（不同请求的长度不同），但很多推理引擎需要在编译时确定 tensor shape 来做内存分配和 kernel 选择。

ONNX 导出的动态维度处理：

1
2
3
4
5
6
torch.onnx.export(model, dummy_input, "model.onnx",
    dynamic_axes={
        'input_ids': {0: 'batch', 1: 'seq_len'},
        'attention_mask': {0: 'batch', 1: 'seq_len'},
        'output': {0: 'batch', 1: 'seq_len'}
    })
  • 未指定 dynamic_axes 的维度会被固定为 dummy_input 的 shape
  • 指定后该维度可在推理时变化，但需要后续引擎支持

TensorRT 的动态 Shape 处理：

  • 需要设置 Optimization Profile，指定每个动态维度的 min/opt/max 三个值：
    • min：该维度的最小值（如 batch=1, seq=1）
    • opt：最常见值（用于 kernel auto-tune 的目标 shape）
    • max：最大值（决定预分配的内存上限）
  • TensorRT 会根据 opt shape 选择最优 kernel，对 min-max 范围内的 shape 保证正确性
  • 多 profile 支持：不同的 shape 范围可以设置不同的 profile，各自优化

动态维度对性能的影响：

  • 内存分配：必须按 max shape 预分配内存（可能浪费显存）
  • Kernel 选择：动态 shape kernel 通常比固定 shape kernel 慢 5-15%（需要额外的边界检查/动态调度逻辑）
  • 算子融合受限：某些融合模式依赖 shape 已知（如将 reshape+transpose 折叠为单步）

最佳实践：

  • 对 batch_size 设为动态（请求量变化），seq_len 根据场景选择：
    • Padding 到固定长度（简单但浪费计算）
    • Bucketing（将请求按长度分桶，每个桶一个编译好的 engine）
    • 完全动态（灵活但性能略差）
  • 生产环境常用 bucketing：如 seq_len 分为 [128, 256, 512, 1024, 2048] 五个桶

Q: 量化相关技术？

量化技术涵盖多个维度的设计选择，每个选择都影响精度-性能权衡：

量化类型：

  • 对称量化：zero_point = 0，范围 [-max_abs, +max_abs]。实现简单（无 zp 偏移计算），适合权重（近零均值分布）
  • 非对称量化：zero_point ≠ 0，范围 [min, max]。充分利用量化范围，适合激活（如 ReLU 后全正值）

量化粒度（从粗到细，精度递增，开销递增）：

粒度 Scale 数量 适用场景 精度
Per-tensor 1 快速验证 最低
Per-channel C_out 权重量化标配 中
Per-group (128) C_out * C_in/128 GPTQ/AWQ 高
Per-element C_out * C_in 理论最优（不实用） 最高

PTQ 校准方法对比：

  • MinMax：取绝对值最大值。简单但 outlier 敏感
  • Percentile（99.99%）：截断极端值，多数场景优于 MinMax
  • MSE 最小化：搜索使量化误差 E[(x-x_q)^2] 最小的 scale
  • KL 散度：TensorRT 方法，最小化量化前后分布差异
  • 学习方法（LSQ）：QAT 中将 scale 作为可训练参数

QAT（Quantization-Aware Training）：

  • 训练时在前向传播中插入伪量化节点（quantize -> dequantize），模拟量化误差
  • 反向传播使用 STE（Straight-Through Estimator）：量化函数的梯度不可微，直接将下游梯度传到上游
  • 精度最优（通常比 PTQ 好 0.5-1% accuracy），但需要训练资源和时间

混合精度量化：

  • 对不同层/算子使用不同精度：敏感层 FP16，非敏感层 INT8/INT4
  • 敏感度判定标准：逐层量化后 PPL/accuracy 变化量
  • 典型策略：embedding/LM head 保持 FP16，attention QK 投影 INT8，FFN INT4

先进量化算法：

  • GPTQ：逐列量化权重，用 Hessian 信息补偿已量化列对未量化列的影响。速度快（几分钟完成 7B 模型量化）
  • AWQ：识别对输出贡献大的”重要”权重通道（通过激活值大小判断），保护这些通道的量化精度
  • SmoothQuant：将激活的量化困难（outlier）通过数学等价变换迁移到权重侧，使 W8A8 PTQ 可行

Q: 多线程、OpenMP、MPI的区别？

三种并行编程模型针对不同的硬件架构和应用场景：

多线程（Pthreads/std::thread）：

  • 模型：共享内存并行，程序员手动创建/管理线程
  • 通信：共享变量（需要 mutex/condition variable/atomic 同步）
  • 粒度：细粒度控制，可精确管理每个线程的行为
  • 适用：单机多核，需要精细控制的场景（自定义线程池/生产者消费者等）
  • 缺点：手动管理复杂（死锁/竞态/数据竞争），代码可读性差

OpenMP：

  • 模型：共享内存并行，编译器指令驱动（#pragma omp parallel for）
  • 通信：自动处理共享变量同步，支持 reduction/private/shared 属性
  • 粒度：粗粒度，适合循环并行化
  • 适用：科学计算中的 for 循环并行、矩阵运算、CPU 上的算子并行
  • 优点：改动小（加 pragma 即可并行化）、可增量并行化
  • 示例：#pragma omp parallel for reduction(+:sum) 自动并行化累加循环

MPI（Message Passing Interface）：

  • 模型：分布式内存并行，进程间通过消息传递通信
  • 通信：MPI_Send/Recv（点对点）、MPI_Allreduce/Bcast/Gather（集合通信）
  • 粒度：进程级（每个进程有独立内存空间）
  • 适用：多节点集群、超算、分布式训练（NCCL 底层类似 MPI 语义）
  • 优点：可扩展到数千节点，无共享内存限制
  • 缺点：编程复杂（需显式 pack/unpack/send/recv），调试困难

对比总结：

特性 Pthreads OpenMP MPI
内存模型 共享 共享 分布式
编程难度 高 低 中-高
扩展性 单机 单机 多节点
典型节点数 1 1 1-10000+
AI Infra 应用 推理线程池 CPU 算子 分布式训练（via NCCL）

混合使用：大规模并行程序常用 MPI+OpenMP 混合模式——节点间 MPI 通信，节点内 OpenMP 多线程共享内存。类比分布式训练：节点间 NCCL AllReduce，节点内 NVLink TP。

Q: NCNN交叉编译？

NCNN 是腾讯开源的高性能移动端推理框架，交叉编译是将其从开发机（x86）编译为目标平台（ARM）可执行的过程：

Android 交叉编译：

1
2
3
4
5
6
7
mkdir build-android && cd build-android
cmake .. \
  -DCMAKE_TOOLCHAIN_FILE=$ANDROID_NDK/build/cmake/android.toolchain.cmake \
  -DANDROID_ABI=arm64-v8a \
  -DANDROID_PLATFORM=android-24 \
  -DNCNN_VULKAN=ON
make -j$(nproc)

关键配置项：

  • 目标架构：arm64-v8a（64位，现代手机标配）、armeabi-v7a（32位，旧设备兼容）
  • Toolchain：Android NDK 提供的 clang 交叉编译工具链
  • API Level：android-24+ 支持更多系统调用
  • Vulkan 支持：开启 GPU 加速（需要目标设备支持 Vulkan）

iOS 交叉编译：

1
2
3
4
cmake .. \
  -DCMAKE_TOOLCHAIN_FILE=../toolchains/ios.toolchain.cmake \
  -DIOS_PLATFORM=OS64 \
  -DENABLE_BITCODE=OFF

需要注意的关键问题：

  1. NEON SIMD 支持：arm64-v8a 默认支持 NEON，armeabi-v7a 需要确认目标设备支持。NCNN 大量使用 NEON intrinsics 加速计算，未启用 NEON 会导致性能剧降。

  2. 线程模型：Android NDK 默认使用 libc++ 的 pthread，需要确保链接正确的 C++ 标准库（-DANDROID_STL=c++_shared 或 c++_static）。

  3. C++ 标准库选择：

    • c++_static：静态链接到每个 .so 中，APK 更大但兼容性好
    • c++_shared：动态链接，APK 小但需要打包 libc++_shared.so

  4. 浮点 ABI：armeabi-v7a 有 softfp（软件浮点传参）和 hard（硬件浮点传参）两种，NDK 默认 softfp。

  5. 性能验证：交叉编译后需要在真机上验证正确性和性能，adb push 到设备执行 benchmark。不同 SoC（骁龙/天玑/麒麟）的 NEON/GPU 性能差异大。

部署流程：开发机编译 -> adb push 到手机 -> 运行 benchncnn 验证性能 -> 集成到 Android/iOS 应用中（通过 JNI/Objective-C++ 桥接）。