寒武纪高性能计算库开发一面：项目深挖 + 手撕 transpose

轮次: 一面
时间: 2026-09
来源: 牛客网

寒武纪高性能计算 C++矩阵乘 OpenMP HIP 手撕代码

《面试题目》

自我介绍。
介绍一个你认为花费精力最多、最有意义的项目。
如何定位出具体问题？
你说你减少了内核寄存器的使用，如何实现的？
对于 dgemm，你们有考虑去优化它吗？
HIP 重构怎么做的？只是用 HIP 的接口重写代码吗？
遇到过的状态同步缺陷具体是什么样的？你有试过通过改写状态语句和逻辑的方式，依然写同一个文件来解决吗？
CPU 矩阵乘，分块你是怎么确定的？
如何基于 OpenMP 分块并行并动态调整分块粒度？

手撕：用 C++ 写一个 transpose，要求用 tiling 分块来避免 cache miss，写完之后讲一下边界处理。

反问：① 业务方向；② 后续流程（至少还有一个技术面和一个 HR/主管面）。

《参考解析》

手撕 transpose：tiling 怎么写，边界怎么处理

朴素二维转置的问题是写端跨步、读端连续（或反过来），每次写都落在不同的 cache line 上，缓存命中率极差。Tiling 的思路是：把矩阵切成 BS × BS 的小块，块内数据能装进 L1，块内先转置到临时缓冲再按行写出，这样读和写都尽量按行顺序访问。

constexpr int BS = 32;  // 32*32*4B = 4KB，正好贴合常见 32KB L1
void transpose_tiled(const float* A, float* B, int M, int N) {
    for (int i0 = 0; i0 < M; i0 += BS)
        for (int j0 = 0; j0 < N; j0 += BS) {
            int imax = std::min(i0 + BS, M);
            int jmax = std::min(j0 + BS, N);
            float tile[BS][BS];
            for (int i = i0; i < imax; ++i)
                for (int j = j0; j < jmax; ++j)
                    tile[i - i0][j - j0] = A[i * N + j];   // 读 A：连续
            for (int j = j0; j < jmax; ++j)
                for (int i = i0; i < imax; ++i)
                    B[j * M + i] = tile[i - i0][j - j0];   // 写 B：连续
        }
}

块大小怎么定：让一个块的工作集落在 L1（常见 32KB）内，BS² × sizeof(T) ≤ L1/2 比较安全，float 取 32、double 取 16 是常用值；再往上还有 L2/L3 级的分块（多级 blocking）和寄存器级分块（如 4×4 展开）。

边界处理是这题的真正考点，必须主动讲：M、N 不整除 BS 时，尾部块用 std::min 截断（上面代码就是这么写的），只遍历真实存在的行列，绝不能越界读；临时 tile 里超出部分不用初始化，因为不会被读到。另一个常见坑是非方阵且维度很大时 B 的索引要用 j * M + i（转置后行数是 N、列数是 M），写错就成了未定义行为。还可以提一句优化方向：用 4×4 寄存器分块减少 tile 缓冲的读写次数，或者用 _mm256_permute/vpermd 之类做 SIMD 块内转置，写完先保证正确性再谈性能。

矩阵乘的分块怎么定

不要只答「试出来的」，要给出受硬件约束的推导 + 调优两条腿：

受约束的部分：目标是让被复用的数据留在正确的层级。经典的三级分块对应三级缓存 —— 寄存器块（MR × NR，如 8×6，由寄存器数量和 FMA 吞吐决定）→ L1/L2 块（MC × KC、KC × NC，由 L1/L2 容量决定，BLIS/GotoBLAS 用参数化模型算出来）→ L3 块（MC × NC，由 L3 容量和 TLB 覆盖范围决定）。推导时的经验值：L1 分块工作集占 L1 的 50%~70%，L2 分块占 L2 的 50% 左右，再留出空间给其他数据。
受布局的部分：分块之后通常还要 packing——把 A 的 MC × KC 和 B 的 KC × NC 拷进连续缓冲，使内层核（micro-kernel）的访存变成纯顺序，同时消掉 A 的转置需求。dgemm 优化与否，一半功夫在 packing 和 micro-kernel 上。
调优的部分：真实机器上 cache 关联度、TLB 大小、NUMA、频率都会影响最优值，所以要靠 autotuning 在小范围枚举（BS 在 16~~256 之间、KC 在 64~~512 之间）跑基准选参数，并把参数按 CPU 型号固化。被追问 dgemm 优化时，可以直接说「我们当前只做了 blocking + packing，没做 micro-kernel 汇编和多线程 NUMA 亲和，这是已知的优化空间」——承认边界比硬吹强。

怎么减少内核寄存器的使用

面官问的是具体手段，按收益从高到低列：

降低寄存器分块与展开因子：寄存器压力主要来自 micro-kernel 的 MR × NR 累加器（8×6 的 double 就要 48 个寄存器，加上 A/B 的操作数寄存器很容易爆）。把 MR × NR 调小是最直接的减负。
消除寄存器溢出（spilling）：用编译器报告确认（CUDA/HIP 用 -Xptxas -v 或 --ptxas-options=-v 看 registers/spills，C++ 用 -Rpass-analysis 或看汇编）。溢出说明某个循环变量/指针没被放进寄存器，常见修法是缩小作用域、把地址计算提到循环外、用 __restrict__/const 帮助编译器、避免别名分析失败。
把常量与地址挪出寄存器：编译期常量用模板参数或宏；循环不变量提前算好；CUDA/HIP 上把只读常量放 __constant__ 或直接传立即数，减少常驻寄存器。
减少活跃变量数：拆循环、分阶段计算、把复用的中间值写回 shared memory（代价是访存变多，要和 occupancy 权衡）。
注意 occupancy 的反向关系：寄存器少了 occupancy 高，但可能因为并行度下降而变慢。这个话题的正确说法是在寄存器和占用率之间找平衡，用 profiling 数据说话，而不是「寄存器越少越好」。

HIP 重构不只是换接口

原帖面试官专门追问了这句，因为很多人的「HIP 移植」就是跑一遍 hipify 换函数名。真正的重构要处理：

wavefront/warp 宽度：AMD 是 64、NVIDIA 是 32。所有依赖 __shfl、ballot、reduce、以及「一个 warp 处理一行」这类假设的代码，都要按 64 重新推导，否则结果正确但性能腰斩。
内存层级与容量：LDS（shared memory）大小、bank 冲突行为、__syncthreads 语义、L2 大小与缓存策略都不同，块大小和分块参数要重调。
指令与内建函数：__syncwarp、原子操作、__ldg、数学内建（__fmaf_rn 等）在两边语义或性能不同；双精度在消费级卡上的吞吐差异极大，dgemm 的优化策略要据此调整。
编译与工具链：用 __HIP_PLATFORM_AMD__/__HIP_PLATFORM_NVIDIA__ 宏做条件编译，性能剖析工具从 ncu 换成 rocprof，构建脚本和 math library（rocBLAS vs cuBLAS）也要分叉。
可移植性设计：把平台相关部分收敛到一层薄封装（kernel launch、内存分配、同步、reduce 原语），上层算法代码只写一份，这样新增后端时改动可控。这才是「重构」和「重写」的区别。

OpenMP 分块并行与动态调整粒度

CPU 矩阵乘并行化的关键是按块分配任务，而不是按行/单点，否则任务粒度过细、调度开销吃掉并行收益：

#pragma omp parallel for schedule(dynamic, 1) collapse(2) num_threads(T)
for (int i0 = 0; i0 < M; i0 += BM)
  for (int j0 = 0; j0 < N; j0 += BN)
    block_gemm(A, B, C, i0, j0, BM, BN);

调度策略：static 开销最小但负载不均（尾部块小、或 NUMA 分布导致的访存差异）；dynamic, chunk 按需取任务，负载均衡好但有原子操作开销；guided 是两者的折中——任务块从大到小递减，前期开销低、后期均衡好，矩阵分块这种「任务数远多于线程数、单个任务耗时差异大」的场景，guided 通常最合适。
动态调整粒度：任务粒度取「单个块的计算量 ≈ 调度开销 × 10~100 倍」这个量级；实践中可以先按 schedule(dynamic, 1) 上的块尺寸跑一遍，再用 schedule(guided) 配合更大的块。也可以用 #pragma omp taskloop grainsize(k) 做任务化，让运行时自己合并小任务。
别忘掉的细节：collapse(2) 把两层循环合并成一个任务空间；块大小要让「每个线程的工作集」落在私有 L2 内（否则多线程互相踢缓存，出现并行反而更慢的异常）；NUMA 机器上要注意首次触碰（first touch）——C 由主线程分配、其他线程写，会造成跨节点访存，正确做法是在并行区里分配或按节点交错初始化；再用 OMP_PROC_BIND=close/OMP_PLACES=cores 绑核，避免线程迁移。

「如何定位出具体问题」与状态同步缺陷

这两问其实是一类题：性能/正确性问题怎么复现、隔离、归因。可复用的排查框架：

建立基线与可复现用例：把问题固化成最小规模的输入（小矩阵、单卡、固定随机种子），能稳定复现才有后面的分析。随机失败的问题要用循环复现 + 日志把「失败那次」的输入固化下来。
二分与隔离：换硬件/换规模/换线程数/关掉某层优化，逐项二分。比如矩阵乘结果不对，先验证单线程、单块（一个 block 只算一小块）、单精度 vs 双精度，缩小到具体 kernel。
工具定位：用 rocprof/nvprof/ncu 看 kernel 耗时与占用率，用 cache miss、stall reason 判断是访存瓶颈还是计算瓶颈；用 compute-sanitizer / cuda-memcheck、-fsanitize=thread,address 抓越界与数据竞争。
状态同步缺陷：多线程/多卡场景里的典型问题有——missing barrier（某条线程提前读了未写完的 shared memory 或全局缓冲）、原子性与内存序（用了放松的内存序或普通读，读到旧值）、读写同一文件或缓冲区时的竞态（一个 kernel 正在写、另一个已经开始读，需要 stream/event 或双缓冲同步）、以及对「同步点」的错误假设（以为 __syncthreads 会同步整个 grid，实际只同步 block）。「改写状态语句和逻辑但依然写同一个文件」这个追问，本质是在问你有没有想过用双缓冲/乒乓缓冲来解耦读写时序——回答时可以说明：先定位竞争窗口（打印时序、用 sanitizer），再选择加同步原语、换无锁结构、或改成双缓冲/流水线；只改语句不改时序，多数时候只是把概率问题变成了另一个概率问题。

《面试题目》

《参考解析》

相关面经