工业界真实版「算子开发工程师技能树」：从 CUDA Kernel 到 AI 编译器与硬件映射

原创已于 2026-08-05 17:14:52 修改 · 456 阅读

本内容遵循CC 4.0 BY-SA版权协议

于 2026-02-10 20:45:37 首次发布

如果你想进入 NVIDIA、AMD、Meta Kernel Team、ByteDance AI Infra、DeepSeek Infra、OpenAI Infra、华为昇腾、寒武纪、燧原、摩尔线程这类团队，那么你需要先建立一个非常重要的认知：

算子开发工程师，本质上不是“写模型的人”，也不是“调 PyTorch 的人”。
算子开发工程师的核心工作是：
为 AI 计算设计高性能 Kernel、Schedule、Memory Layout 与 Hardware Mapping。

这是一个典型的交叉岗位：

AI × HPC Kernel × 编译器 × 体系结构

换句话说，你要懂算法语义，也要懂硬件执行；你要会写 CUDA/Triton，也要理解 loop schedule、tensorize、memory hierarchy、dataflow mapping 和 compiler IR。

这篇文章不是培训机构版本的“CUDA 入门指南”，而是一份更接近工业界真实要求的 算子开发工程师技能树。

0. 一张图理解算子开发工程师的能力结构

如果把算子开发能力从底层到顶层展开，大致可以分成下面几层：

                    LLM 时代核心算子
                         ↑
                 Kernel Fusion / Dispatch
                         ↑
                  算子多实现与优化策略
                         ↑
              Kernel 编程 / Schedule / Layout
                         ↑
              并行计算模型：GPU / AI 加速器
                         ↑
          体系结构：Memory / Cache / Pipeline
                         ↑
              线性代数：GEMM / Reduce / Einsum

再抽象一点：

算子开发工程师 =
    理解计算语义
  + 理解硬件瓶颈
  + 设计并行映射
  + 优化访存与计算
  + 构建可复用的 Kernel Family
  + 接入编译器 / Runtime / Dispatch 系统

下面逐层展开。

第一层：硬核基础

这一层决定你的上限。
如果基础不扎实，后面写再多 CUDA 代码，也很难真正做出高性能算子。

1. 计算机体系结构：算子优化的第一性原理

算子优化的本质，不是背 API，而是理解：

计算为什么慢？
数据从哪里来？
数据搬到了哪里？
哪个环节成为瓶颈？

你需要建立完整的 memory hierarchy 视角：

Register File
    ↑
Shared Memory / L1 Cache
    ↑
L2 Cache
    ↑
Global Memory / HBM

核心知识点包括：

memory hierarchy
cache / shared memory / register file
SIMD / SIMT / warp
instruction pipeline
ILP / TLP
latency hiding
roofline model
memory bound vs compute bound

工业界真正要求的不是“知道这些名词”，而是你能不能看到一个 kernel 后快速判断：

这个 kernel 的瓶颈在哪里？

例如：

是不是 global memory 带宽打满了？
是不是 shared memory bank conflict 严重？
是不是 register pressure 太高导致 occupancy 下降？
是不是访存延迟没有被隐藏？
是不是 warp divergence 导致执行效率下降？
是不是 arithmetic intensity 太低，导致算力完全跑不满？

一个合格的算子工程师，应该具备这样的直觉：

看代码 → 看访存模式 → 看并行结构 → 判断瓶颈 → 给出优化方案

而不是：

看代码 → 凭感觉改 → 再跑一遍 → 不知道为什么变快或变慢

2. 并行计算模型：GPU 与 AI 加速器

GPU 方向

如果你做 NVIDIA / AMD / 通用 GPU kernel，必须深入理解 CUDA execution model：

thread
warp
block
grid
warp scheduling
SIMT divergence
memory coalescing
shared memory bank conflict
occupancy
async copy
warp-level primitives
tensor core execution model

你需要能回答这些问题：

为什么 global memory access 要 coalesced？
为什么 shared memory 会出现 bank conflict？
为什么同一个 kernel，不同 block size 性能差异巨大？
为什么 warp divergence 会影响性能？
为什么 tensor core 需要特定的 fragment layout？
为什么高 occupancy 不一定代表高性能？
为什么有些 kernel occupancy 不高，但性能仍然很好？

AI 芯片 / NPU 方向

如果你去华为昇腾、寒武纪、燧原、壁仞、摩尔线程、自研 AI 芯片团队，还需要理解另一类架构：

DMA orchestration
on-chip SRAM scheduling
dataflow mapping
weight stationary
output stationary
row stationary
tile-based execution
NoC / inter-core communication
vector unit / matrix unit / scalar unit 协同

在 AI 芯片公司，算子开发往往不只是“写 kernel”，而是：

如何把一个 tensor 计算切分成 tile，并映射到多个计算单元上，同时最小化片外访存和片上通信。

这比单纯写 CUDA kernel 更接近架构设计。

3. 线性代数：不是数学推导，而是计算结构

算子开发工程师需要懂线性代数，但重点不是手推 SVD，也不是证明矩阵分解。

重点是：

你能不能从算子语义中看出底层计算结构。

你需要重点掌握：

GEMM blocking
tensor contraction
reduce pattern
broadcast pattern
elementwise pattern
einsum mapping
batched GEMM
grouped GEMM
sparse / dense 结构

你需要具备这种能力：

看到一个 operator：
马上判断它是：
    GEMM-like
    Reduce-like
    Elementwise-like
    Memory-bound fusion
    还是复合 pattern

例如：

算子	本质结构
MatMul	GEMM
Batched MatMul	Batched GEMM
Attention	Batched GEMM + Softmax + Mask + Reduce
LayerNorm	Reduce + Elementwise
RMSNorm	Reduce + Elementwise
Softmax	Reduce + Elementwise
BiasAdd	Broadcast Elementwise
GELU / SiLU	Elementwise
Conv	Direct / Implicit GEMM / Winograd / FFT
MoE Routing	TopK + Permute + GEMM + Scatter

工业界非常看重这种“抽象能力”。

因为大多数高性能算子优化，最后都会归结为：

如何把这个计算映射到硬件最高效的执行单元上。

第二层：Kernel 实现能力

这是岗位的核心硬技能。

如果说第一层是内功，第二层就是招式。

1. GPU Kernel 编程：至少精通一个工业主流路径

工业界常见的技术栈包括：

CUDA

最硬核，也最接近底层。

适合：

NVIDIA GPU kernel
高性能算子库
自定义 CUDA extension
推理引擎
训练 infra
底层性能优化

你需要掌握：

thread/block/grid 组织
shared memory 使用
register tile
warp-level programming
async copy
double buffering
memory coalescing
bank conflict 避免
occupancy 调优
CUDA stream/event
PTX / SASS 基本概念

Triton

近年来在 LLM infra 和 AI compiler 领域使用非常多。

适合：

attention kernel
fused kernel
LLM operator
quick prototyping
block-level programming model
GPU kernel DSL

Triton 的优势是：

比 CUDA 更容易表达 tile-level schedule，比手写 CUDA 更适合快速做 fusion 和 block 级优化。

但要注意：

会写 Triton 不等于不需要懂底层。
真正写好 Triton，依然需要理解 shared memory、register pressure、memory coalescing、occupancy 和 GPU pipeline。

CUTLASS

NVIDIA 生态中非常重要的模板库。

适合：

GEMM kernel family
tensor core kernel
convolution kernel
epilogue fusion
mixed precision GEMM
高性能算子库开发

你需要理解：

tile iterator
mainloop
epilogue
mma pipeline
tensor core fragment
layout transformation
fused epilogue

工业界不是要求你只会调用 CUTLASS，而是你要能看懂它的抽象，并基于它构建自己的 kernel family。

2. 你必须能写出的核心 Kernel

至少应该完整写过下面这些内容：

基础版

naive matrix multiplication

Shared Memory 版

tiled matmul with shared memory

Register Tile 版

register blocking + vectorized load/store

Double Buffer 版

shared memory double buffering

Async Copy 版

cp.async pipeline / asynchronous data movement

Tensor Core 版

warp-level mma / mma.sync / CUTLASS-style GEMM

Fusion 版

matmul + bias + activation

这些不是“练习题”，而是工业界面试和实际工作中最常见的基本功。

尤其是 tiled matmul，几乎是算子工程师的试金石。

因为它同时包含：

global memory access
shared memory staging
register blocking
compute reuse
memory bound / compute bound 判断
pipeline overlap
tensor core mapping
layout design

一个 matmul 写得好不好，基本能看出这个人的底层功力。

3. Loop Schedule 设计：DL Compiler 的核心能力

如果你只写 CUDA，不碰编译器，也建议理解 schedule。
如果你想往 AI compiler、TVM、MLIR、Triton compiler、自研编译器方向走，那 schedule 是核心中的核心。

你需要理解：

loop tiling
loop reorder
loop fusion
loop splitting
loop unrolling
vectorization
parallelization
tensorization
schedule primitive
compute at / inline
cache read / cache write

常见工具体系包括：

TVM TensorIR
MLIR Linalg
Halide
Triton IR
XLA HLO
IREE
自研 polyhedral / schedule system

本质能力是：

把一个高层 tensor 计算，转换成一个适合硬件执行的最优 loop nest。

以 GEMM 为例：

原始形式：

for m in range(M):
    for n in range(N):
        for k in range(K):
            C[m, n] += A[m, k] * B[k, n]

这显然不能直接高性能执行。

你需要逐步变成：

block tiling
    ↓
thread tiling
    ↓
register tiling
    ↓
shared memory staging
    ↓
vectorized load
    ↓
unrolled compute
    ↓
tensorization / mma

最终形成一个适合 GPU 或 AI 加速器的执行结构。

这就是 schedule 的价值。

4. Memory Layout 设计：新人最容易忽视的部分

很多新人写 kernel，只关心：

index 怎么算？

但工业界更关心：

这个 layout 对硬件友好吗？

常见的 layout 问题包括：

NCHW vs NHWC
row-major vs column-major
interleaved layout
blocked layout
COL32
tensor core fragment layout
swizzle layout
padded layout
split-k layout
KV-cache layout
quantized tensor layout

Memory layout 的目标非常明确：

让数据访问更连续，让 shared memory 无冲突，让 tensor core 更容易消费数据，让 L2 cache 利用率更高。

你需要能判断：

当前 layout 是否导致 uncoalesced access？
是否造成 shared memory bank conflict？
是否导致 L2 thrashing？
是否增加了不必要的 transpose？
是否让 tensor core fragment 转换开销过大？
是否影响 epilogue fusion？

一个真实工业经验是：

很多时候，不是你的计算写得不好，而是你的 layout 从开始就选错了。

第三层：算子级优化能力

这一层决定你能不能独立负责一个算子模块。

工业界不会只问你“会不会写 kernel”，更常见的问题是：

给你一个算子，你会怎么设计它的实现策略？

1. 算子多实现策略

一个真实算子往往不是一个 kernel 就够了。

以 Conv 为例，可能有：

direct convolution
implicit GEMM
Winograd
FFT convolution
tensor core convolution
depthwise convolution
grouped convolution
split-k convolution
NHWC / NCHW 特化版本

不同 shape、不同 batch、不同 kernel size、不同 stride，最优实现可能完全不同。

例如：

small kernel + large batch   → implicit GEMM / tensor core
3x3 stride=1               → Winograd 可能更优
depthwise conv             → 专用 kernel 更优
odd shape                  → direct 或 padded GEMM
inference low latency      → 选择 launch overhead 更小的实现
training high throughput     → 选择 tensor core utilization 更高的实现

所以工业界真正需要的是：

不只是实现一个 kernel，而是设计一个 kernel family，并根据 workload 自动选择最佳实现。

2. Kernel Fusion：LLM 时代的核心技能

在大模型时代，fusion 的重要性被进一步放大。

原因很简单：

大模型里很多操作不是算力打满，而是被 global memory traffic 卡住。

Fusion 的本质目标是：

减少中间结果写回 global memory 的次数。

典型 fused kernel 包括：

matmul + bias + gelu
matmul + bias + silu
add + layernorm
rmsnorm + residual add
softmax + mask + scale
attention + kv-cache update
flash attention
fused rope
fused moe gate / permute / unpermute
fused quantization / dequantization

你需要能设计：

fused attention
fused rmsnorm
fused layernorm
fused kv-cache update
fused matmul epilogue
fused elementwise chain
fused reduce + broadcast

但也要注意：

Fusion 不是越多越好。

过度 fusion 可能导致：

register pressure 过高
occupancy 下降
kernel 复杂度上升
编译时间变长
shape 泛化变差
某些 shape 下反而更慢

所以真正工业级的 fusion 策略是：

根据 shape、硬件、访存特征、寄存器使用率综合判断。

3. Autotuning & Dispatch

真实系统里，不是一个 kernel 打天下。

工业系统通常是：

operator registry
    ↓
multiple kernel implementations
    ↓
shape / dtype / layout dispatch
    ↓
heuristic selection
    ↓
autotuning
    ↓
runtime cache

你需要理解：

kernel registry
dispatch key
shape bucketing
heuristic rules
autotuning space
benchmark cache
offline tuning
online warmup
fallback kernel

一个成熟的算子库往往长这样：

同一个算子：
    有多个实现
    有多个 tile config
    有多个 block size
    有多个 pipeline stage
    有多个 layout 支持
    根据 shape 自动选择

这才是工业界真实的算子系统。

第四层：硬件映射能力

这是从“写 kernel”走向“设计执行系统”的关键一层。

尤其是对于 AI 芯片公司来说，这一层几乎决定岗位核心价值。

1. Loop 到 Dataflow 的映射

你需要理解：

一个 loop nest 如何映射到硬件执行单元？

常见 dataflow 包括：

Weight Stationary

weight 固定在计算单元附近，
activation / input 不断流入。

适合：

weight reuse 高
inference 场景
权重可以常驻片上 SRAM

Output Stationary

output 固定在局部存储，
input / weight 流入并累加。

适合：

reduction 维度较长
GEMM / Conv 累加结构

Row Stationary

按 row 做局部驻留，
平衡 input、weight、output 的 reuse。

适合：

某些 spatial accelerator
需要平衡片上通信和复用的场景

你需要进一步理解：

tile = communication unit
tile = scheduling unit
tile = memory unit
tile = noc packet

在多核 / 多 PE / NoC 架构中，算子开发不是只看单个计算单元，而是要考虑：

数据如何切分？
tile 如何在核间移动？
哪个维度并行？
哪个维度复用？
片上 SRAM 如何分配？
DMA 如何 overlap？
通信和计算能否流水？

这就是 hardware mapping。

2. 对 AI 芯片团队意味着什么？

在 GPU 团队，你可能更多关注：

warp / block / shared memory / tensor core / cache

在 AI 芯片团队，你可能需要关注：

SRAM partition
DMA schedule
core mapping
dataflow reuse
NoC traffic
tile dependency
double buffering
pipeline stage
vector / matrix / scalar 协同

所以如果你想冲顶，不能只停留在 CUDA。
你需要逐渐建立：

从算法到 schedule，从 schedule 到 dataflow，从 dataflow 到硬件资源分配的全链路能力。

第五层：工具链能力

工业界不是靠猜性能，而是靠 profile。

没有 profiling 的优化，基本等于盲改。

1. Profiling 工具

NVIDIA

常用：

Nsight Systems：系统级 timeline
Nsight Compute：kernel 级分析
nvprof：传统 profiler
cuobjdump：查看 PTX / SASS
nvdisasm：反汇编 SASS

你需要能从 Nsight Compute 里看懂：

memory throughput
compute throughput
achieved occupancy
warp stall reasons
shared memory bank conflict
L1/L2 hit rate
instruction mix
tensor core utilization
register per thread
branch divergence

AMD

常用：

rocprof
rocprofiler
ROCm profiling tools

自研 AI 芯片

通常有自研 profiler，但核心思想类似：

看计算单元利用率
看 DMA 时间
看 SRAM 占用
看核间通信
看 pipeline bubble
看 dependency stall

2. IR / Compiler 工具链

如果你想往 AI compiler 或更高层优化发展，需要理解：

MLIR
TVM
Triton IR
LLVM IR
HLO
Linalg
TensorIR
Polyhedral IR

至少要能看懂：

loop nest 如何表示？
tile 之后 IR 怎么变化？
fusion 发生在哪一层？
vectorization 如何 lowering？
tensorization 如何映射到硬件指令？

3. 反汇编与底层分析

GPU 方向尤其是 NVIDIA，常需要看：

PTX
SASS
register allocation
instruction scheduling
mma 指令
ld/st 指令
cp.async
shared memory address pattern

不要求你一开始就能精通，但至少要能建立意识：

高级语言写出来的每一行代码，最终都会落到具体指令和访存行为上。

第六层：LLM 时代新增核心技能

2023 之后，尤其是 2024、2025 以后，LLM infra 对算子工程师提出了新的强需求。

如果你目标是 LLM infra、推理引擎、训练加速、大模型系统方向，下面这些内容非常重要。

1. FlashAttention 结构

FlashAttention 的本质不是“更快的数学公式”，而是：

通过 tiling 和 recomputation，
避免 materialize 巨大的 attention matrix，
减少 HBM 读写。

你需要理解：

online softmax
attention tiling
Q/K/V block 计算
rescaling trick
memory-efficient backward
causal mask
dropout fusion
register/shared memory usage
occupancy 与 tile size 的权衡

现在 FlashAttention 已经成为 LLM kernel 的基础课。

2. PagedAttention

PagedAttention 是推理侧非常重要的能力。

核心问题：

KV cache 如何高效管理？
如何避免内存碎片？
如何支持 continuous batching？
如何减少 KV cache 的拷贝？

你需要理解：

block table
page table
paged kv cache
gather/scatter access
variable sequence length
batched decoding
speculative decoding 相关 kernel
cache reuse

3. KV-cache Layout

KV-cache 的 layout 会直接影响：

memory bandwidth
attention kernel 效率
batch 调度灵活性
多轮对话性能
tensor core 利用率
量化支持

常见考虑包括：

batch / head / seq / dim 的排列
是否连续存储
是否分页
是否量化
是否支持 MQA / GQA
是否支持 prefix cache
是否支持 cross-attention cache

4. MoE Routing Kernels

MoE 大模型普及后，MoE kernel 成为新的高频需求。

典型操作包括：

router logits
top-k selection
expert assignment
token permute
grouped gemm
expert compute
token unpermute
combine outputs
capacity handling
load balancing
all-to-all overlap

你需要理解：

MoE 不只是 GEMM，
而是 routing + permutation + grouped GEMM + communication 的组合。

5. Quantization Kernel

量化是推理加速的核心方向。

你需要理解：

int8 quantization
fp8 quantization
per-tensor scale
per-channel scale
per-token scale
dynamic quantization
static quantization
dequantization fusion
quantized GEMM
mixed precision kernel
outlier handling
weight-only quantization
activation quantization

尤其在大模型推理中，常见组合包括：

W8A8
W4A16
W8A16
FP8 GEMM
INT8 GEMM
smoothquant
awq/gptq style weight-only kernels

未来算子工程师如果不懂量化，会明显少一块核心竞争力。

第七层：真实面试考察点

下面这些，是算子岗位面试中非常真实、非常高频的问题。

1. 手写题

常见：

手写 tiled matrix multiplication

要求可能逐步升级：

1. naive version
2. shared memory tiling
3. register blocking
4. double buffering
5. vectorized load/store
6. async copy
7. tensor core version

面试官真正看的不是你会不会背代码，而是：

你是否理解为什么这样写；
你是否知道访存模式；
你是否能解释性能提升来源；
你是否能处理边界条件；
你是否知道 occupancy 和 register pressure 的影响。

2. 解释题

常见：

为什么 shared memory 要 double buffer？

参考思路：

如果不做 double buffer：
    load 和 compute 往往串行。

如果做 double buffer：
    可以让当前 buffer 做 compute，
    同时预取下一个 tile 到另一个 buffer，
    从而 overlap memory latency 和 compute。

进一步可以延伸：

cp.async
multi-stage pipeline
register pressure
shared memory capacity
pipeline depth

3. 分析题

常见：

给你一个 kernel，分析它的 memory pattern。

你需要能判断：

global load 是否 coalesced？
global store 是否 coalesced？
shared memory 是否有 bank conflict？
是否有 redundant load？
是否有不必要的写回？
是否可以融合 elementwise？
是否 memory bound？
是否可以提升 arithmetic intensity？

4. 设计题

常见：

设计一个 fused operator。

例如：

matmul + bias + gelu
layernorm + residual add
attention + kv-cache update
rmsnorm + cast + quantization

你需要回答：

输入输出是什么？
哪些中间结果可以避免写回？
哪个维度做并行？
哪个维度做 reduction？
是否需要 shared memory？
是否适合 tensor core？
是否会 register spill？
不同 shape 下是否仍然高效？

5. 判断题

常见：

这个算子是 compute bound 还是 memory bound？

你需要会看：

arithmetic intensity
memory bandwidth
peak FLOPS
tensor core utilization
actual achieved throughput
roofline position

例如：

elementwise / reduce / layernorm / softmax 大概率 memory bound
large GEMM / conv 大概率 compute bound
small batch attention 可能 memory bound
large shape grouped GEMM 可能 compute bound
quantized inference 可能受 dequant 和 memory traffic 影响

第八层：工业界真实能力分级

下面是一个更贴近工业内部判断的分级模型。

Level	能力描述	可交付成果
L1	会写基础 CUDA kernel	naive matmul、elementwise、reduce
L2	会 tiling + shared memory + 基本性能优化	可运行的 tiled SGEMM，能解释 coalescing / bank conflict
L3	会 tensor core、fusion、profiling	高性能 matmul、fused bias/gelu、能使用 NCU/Nsys 定位瓶颈
L4	能设计 kernel family 与 dispatch	支持多 shape、多 dtype、多 layout 的算子实现
L5	能设计 compiler schedule / hardware mapping	参与编译器、算子库、芯片执行引擎设计

更具体一点：

L1：会写 CUDA matmul

特征：

知道 CUDA 基本语法；
能启动 kernel；
能写 naive GEMM；
知道 global/shared memory 的区别；
但性能优化能力较弱。

L2：会 tiling + shared memory

特征：

能写 shared memory tiled matmul；
能解释为什么 tile；
知道 coalescing；
能处理边界条件；
能初步调 block size；
可以进入工业团队做基础算子开发。

L3：会 tensor core + fusion

特征：

能使用或理解 tensor core；
能写 mma / CUTLASS / Triton 风格 kernel；
能做 matmul epilogue fusion；
能使用 profiler；
能分析 memory/compute bound；
是团队里很有战斗力的工程师。

L4：能设计新算子 kernel family

特征：

不只是写单个 kernel；
能设计一组 kernel；
能处理 shape dispatch；
能做 autotuning；
能维护性能回归；
能负责一个算子模块。

L5：能设计 compiler schedule + dataflow

特征：

能从高层 IR 设计 schedule；
能理解 hardware mapping；
能做 tile / dataflow / pipeline 设计；
能影响芯片架构或编译器架构；
是核心系统 / 架构级人才。

第九层：一条现实的学习路径

如果你真的想往这个方向走，不建议一上来就啃编译器和架构论文。

更现实的路径是：

阶段一：打基础

目标：

能写 CUDA，能理解 GPU 执行模型。

重点：

C/C++
CUDA programming
thread/block/grid
shared memory
synchronization
memory coalescing
bank conflict
occupancy

练习：

vector add
reduce sum
softmax
layernorm
naive matmul
tiled matmul

阶段二：把 MatMul 吃透

这是最关键的一步。

目标：

从 naive matmul 一路优化到接近可用的高性能 matmul。

路线：

naive matmul
    ↓
shared memory tiled matmul
    ↓
register tiled matmul
    ↓
vectorized load/store
    ↓
double buffering
    ↓
async copy
    ↓
tensor core matmul

每做一次优化，都要回答：

为什么变快了？
瓶颈从哪里转移到了哪里？
profile 数据怎么解释？

阶段三：做 Fusion 和 LLM Kernel

目标：

能写现代 LLM infra 常见 kernel。

建议实现：

fused matmul + bias + gelu
fused rmsnorm
fused layernorm
softmax
flash attention minimal version
kv-cache update
int8/fp8 quantized matmul

这一步非常适合作品集化。

阶段四：进入编译器 / Schedule / IR

目标：

从写 kernel，升级为生成 kernel。

可以学习：

Triton
TVM
MLIR
Linalg
TensorIR
Halide schedule
loop transformation
autotuning

最终能力：

不是只手写一个 kernel，
而是设计一套 schedule，
自动生成一批 kernel。

第十层：常见误区

误区一：会写 CUDA 语法，就是算子工程师

不是。

会写：

__global__ void kernel(...)

只是入门。

算子工程师真正要解决的是：

为什么慢？
如何变快？
如何稳定地快？
如何在不同 shape 下都快？

误区二：背很多 API 就够了

不是。

API 会变，硬件会变，框架会变。
但底层能力相对稳定：

memory hierarchy
cache
parallelism
data reuse
schedule
layout
roofline
profiling

这些才是长期壁垒。

误区三：只看框架层，不碰底层

如果你只会：

torch.compile
model.half()
torch.cuda.synchronize()

但不知道 kernel 为什么慢，很难进入算子岗位核心。

误区四：只写 kernel，不做 profiling

没有 profile 的优化很容易变成玄学。

工业界更相信：

Nsight Compute 数据
带宽利用率
计算利用率
stall reason
occupancy
register usage

而不是：

我感觉这样应该更快。

误区五：盲目 fusion

Fusion 是好事，但不是所有东西都应该 fuse。

有些 fusion 会：

增加寄存器压力；
降低 occupancy；
增加控制流复杂度；
使 kernel 难以维护；
在小 shape 下反而更慢。

好的工程师会做 cost model，而不是无脑 fuse。

最后总结：算子开发工程师的核心竞争力

如果用一句话总结：

算子开发工程师的核心竞争力，不是会用某个框架，而是能把一个数学计算高效映射到具体硬件上。

你需要同时具备四种能力：

1. 理解算法语义
2. 理解硬件瓶颈
3. 设计并行与访存结构
4. 构建可复用的工程系统

从成长路径看：

入门：会写 CUDA kernel
进阶：会 tiling / shared memory / profiling
高级：会 tensor core / fusion / autotuning
专家：会设计 kernel family / dispatch / schedule
顶层：会设计 compiler schedule / dataflow / hardware mapping

如果你能沿着这条路持续积累，那么无论是 GPU kernel、AI compiler、LLM infra，还是自研 AI 芯片团队，你都会有非常强的竞争力。

附：一个最小作品清单

如果你想证明自己具备算子开发能力，可以做下面这个项目清单：

1. naive SGEMM
2. shared memory tiled SGEMM
3. register tiled SGEMM
4. double buffered SGEMM
5. vectorized SGEMM
6. CUTLASS / Triton GEMM
7. matmul + bias + gelu fused kernel
8. fused rmsnorm
9. fused softmax
10. mini FlashAttention
11. int8 / fp8 quantized matmul
12. 每个 kernel 附带 Nsight Compute / profiling 分析报告

做完这些，你就不再只是“学过 CUDA”，而是真正开始具备工业界算子开发的基本盘。

标签