企业官网建设流程全解析

1. 从“刘胜的告别”说起：一个时代的隐喻

“刘胜的告别”这个说法，最近在一些技术圈子里被反复提起。它并不是指某个具体的人离职，而是一个隐喻——那些曾经靠手写CUDA kernel、手调PTX指令、在寄存器分配和shared memory bank conflict之间反复横跳的底层算子工程师，正在被AI逐步“替代”。这个告别，不是悲情的，而是技术演进到一定阶段的必然。

我做了十多年底层性能优化，从最早的OpenCL到后来的CUDA，再到近几年的各种NPU和AI编译器，亲眼见证了算子开发从“手工艺”走向“工业化”。以前写一个高性能的Matmul，要花几周时间反复调优，现在用TVM或者Ansor，几个小时就能搜出一个接近cuBLAS性能的实现。而最近一两年，大模型开始直接生成CUDA代码，甚至能根据自然语言描述生成完整的算子实现，这个冲击是实实在在的。

但我想说的是，AI学会写算子，并不意味着算子工程师的终结，而是计算机抽象阶梯又往上爬了一层。每一层抽象的出现，都会让上一层的“手艺人”减少，但会让整个系统的生产力大幅提升。汇编语言出现时，没人再手写机器码；C语言出现时，没人再手写汇编；CUDA出现时，没人再直接写GPU机器码。现在AI来写CUDA，逻辑是一样的。

这篇文章，我想从“AI写算子”这个现象切入，拆解背后的技术栈——从CUDA、PTX到LLVM IR，从算子融合到自动调优，从大模型生成代码到实际部署的坑。我会尽量把每个环节讲透，包括我自己的实操经验和踩过的坑。无论你是刚入门的算子开发新手，还是已经写了几年kernel的老手，都能从中找到对自己有用的东西。

2. 计算机抽象阶梯：从晶体管到AI生成算子

2.1 什么是抽象阶梯，为什么它决定了你的饭碗

计算机抽象阶梯这个概念，最早可以追溯到计算机体系结构里的“层次化设计”。简单说，就是每一层都向上层隐藏下层的细节，让上层开发者不用关心底层实现。比如你写Python的时候，不需要知道内存是怎么分配的；你写CUDA的时候，不需要知道SM里的warp scheduler是怎么调度的。

这个阶梯大致是这样的：

每一层抽象的出现，都会让上一层的“手艺人”数量减少，但整个生态的产出会指数级增长。刘胜的告别，本质上就是L3的算子工程师在向L4/L5迁移。

我自己的经历就很典型。2015年左右，我写一个SGEMM，要手动做tiling、向量化、double buffering，调一个参数就要重新编译跑一遍。后来有了cutlass，模板元编程帮你把大部分优化都做了。再后来有了TVM的auto-scheduler，你只需要写计算表达式，调度器自动搜。现在，你甚至可以直接跟大模型说“帮我写一个fp16的matmul，带bias和relu，目标架构是sm_80”，它就能给你生成一份能跑的CUDA代码。

但这里有个关键点：抽象阶梯往上爬，不代表底层不重要。恰恰相反，当AI生成的代码性能不达标时，你还是得回到L3甚至L2去排查问题。所以，理解底层原理，反而变得更加重要——因为你要有能力判断AI生成的代码到底哪里有问题。

2.2 算子开发的三次抽象跃迁

算子开发这个领域，我观察到的抽象跃迁大概经历了三个阶段。

第一阶段：手写CUDA C。这是最原始的方式，直接写.cu文件，用__global__声明kernel，手动管理thread block和shared memory。这个阶段的核心技能是理解GPU的SIMT架构、warp divergence、memory coalescing、bank conflict。我当年为了调一个reduce kernel，把NVIDIA的《CUDA C Programming Guide》翻了不下十遍。

第二阶段：模板库和DSL。代表就是cutlass、TVM、Halide。你不再直接写CUDA，而是用更高层的抽象描述计算，由编译器或模板库生成底层代码。比如cutlass里，你定义一个GemmTraits，指定tile大小、warp数量、指令类型，剩下的它帮你生成。TVM里，你用TIR描述计算，用schedule描述调度，auto-scheduler自动搜最优配置。

第三阶段：AI生成算子。这是最近两年才真正落地的。大模型通过学习海量CUDA代码，能够根据自然语言描述生成kernel实现。比如你输入“写一个sobel算子，3x3卷积，边缘检测，输入是uint8的灰度图”，它就能生成一份完整的CUDA代码，包括host端的内存分配和kernel launch。

这三个阶段，抽象层级越来越高，但底层原理的掌握程度决定了你能走多远。我见过太多人，只会调库，一旦遇到性能瓶颈就束手无策。也见过一些人，虽然用TVM，但能通过分析生成的CUDA代码找到调度问题，手动改schedule，性能提升30%以上。

2.3 为什么AI写算子现在才成为可能

AI写算子这件事，其实不是新概念。早在2018年，就有用强化学习做schedule搜索的工作（比如Google的AutoTVM）。但为什么最近才火起来？我觉得有三个原因。

第一，大模型代码能力的提升。GPT-4、Claude这些模型，在代码生成任务上已经达到了相当高的水平。CUDA代码虽然相对小众，但在GitHub上有大量高质量的开源项目（cutlass、cub、thrust等），模型有足够的学习材料。

第二，算子开发的标准化程度提高。以前每个公司的算子接口都不一样，现在有了ONNX、TorchScript这些中间表示，算子的输入输出格式相对统一，模型更容易生成可移植的代码。

第三，硬件架构的收敛。虽然NPU百花齐放，但GPU上的编程模型基本收敛到CUDA这一套。模型只需要学好CUDA，就能覆盖大部分场景。

但这里有个坑：AI生成的CUDA代码，性能往往不是最优的。我实测过，让GPT-4生成一个fp16的matmul，它能写出来，但性能只有cuBLAS的30%左右。原因很简单，它不知道你的具体硬件配置，不知道L2 cache大小，不知道shared memory容量，只能生成“通用”的代码。要让它生成高性能代码，你需要给它非常详细的提示，包括目标架构、tile大小、向量化宽度等。

3. 拆解AI写算子的技术栈：从CUDA到PTX再到LLVM

3.1 CUDA、PTX、SASS：三层代码的转换关系

要理解AI写算子，首先得搞清楚CUDA代码是怎么变成GPU能执行的指令的。这个过程大致是：

CUDA C/C++ → PTX → SASS

CUDA C是你写的代码，PTX是NVIDIA定义的中间表示（类似汇编），SASS是GPU实际执行的机器码。编译器（nvcc）负责把CUDA C编译成PTX，然后ptxas把PTX编译成SASS。

为什么要有PTX这一层？因为GPU架构迭代很快，每一代SM的指令集都有变化。PTX作为中间表示，可以保证向前兼容——你编译的PTX可以在未来的GPU上重新编译成SASS。这也是为什么CUDA程序能在不同代GPU上运行的原因。

AI生成算子，通常生成的是CUDA C代码。但有时候，为了更精细的控制，也会直接生成PTX。比如在一些对性能要求极高的场景，直接写PTX可以避免编译器的一些“保守”优化。

我举个例子。假设你要写一个warp-level的reduce，用CUDA C写是这样的：

__inline__ __device__ float warpReduceSum(float val) { for (int offset = 16; offset > 0; offset /= 2) val += __shfl_down_sync(0xffffffff, val, offset); return val; }

这段代码编译成PTX后，会变成一系列shfl.sync.down指令。如果你直接写PTX，可以控制指令的调度顺序，避免一些不必要的依赖。但大多数情况下，CUDA C生成的PTX已经足够好了。

3.2 LLVM在算子生成中的角色

LLVM在算子生成里扮演什么角色？简单说，它是很多AI编译器的后端。比如TVM，它的TIR会lower到LLVM IR，然后LLVM生成目标代码（可以是x86、ARM、CUDA PTX等）。

为什么用LLVM？因为LLVM有成熟的优化pass，比如循环展开、向量化、死代码消除。而且LLVM支持多种后端，一套IR可以生成不同硬件的代码。

但LLVM在GPU上的表现，说实话不如NVIDIA自己的编译器。我做过对比，同样的TIR，走LLVM生成PTX，性能通常比nvcc低10%-20%。原因是LLVM对GPU架构的特定优化（比如shared memory swizzling、warp-level primitive）支持不够好。

所以现在很多AI编译器，在GPU后端上会选择直接生成CUDA C，然后调用nvcc编译。比如TVM的cuda后端，就是生成CUDA C代码，然后交给nvcc。这样虽然多了一步，但性能更有保障。

3.3 AI生成算子的典型流程

AI生成算子，通常不是一步到位的。我总结了一个比较靠谱的流程：

第一步：自然语言描述计算。比如“实现一个3x3的sobel边缘检测，输入是HxW的uint8灰度图，输出是HxW的uint8边缘图”。

第二步：生成计算表达式。模型会先生成一个高层的计算描述，比如用TIR或者Python DSL。这一步主要是确定算子的数学定义。

第三步：生成调度。这一步是关键。模型需要决定怎么tile、怎么向量化、怎么用shared memory。通常需要给模型提供目标硬件的参数，比如SM数量、shared memory大小、寄存器数量。

第四步：生成CUDA代码。根据调度，生成具体的CUDA C代码。

第五步：编译和测试。用nvcc编译，跑正确性测试和性能测试。

第六步：迭代优化。如果性能不达标，把性能数据反馈给模型，让它重新生成。

这个流程里，第三步和第六步是最难的。调度空间巨大，模型很难一次搜到最优。而且性能测试需要真实硬件，反馈周期长。

我自己的做法是，先用AI生成一个baseline，然后手动调优。AI生成的代码通常结构是对的，但细节需要打磨。比如它可能忘了用__restrict__，或者没有做向量化加载，这些手动改一下，性能就能提升不少。

4. 实操：用AI生成一个高性能算子

4.1 环境准备与工具选型

在开始之前，你需要准备好环境。我推荐用WSL2 + Ubuntu 20.04，因为Windows上CUDA的支持不如Linux完善。具体步骤：

1. 安装WSL2和Ubuntu 20.04。
2. 安装CUDA Toolkit 11.8或12.x。注意，CUDA版本要和你的GPU驱动匹配。比如4060Ti需要CUDA 12.x以上。
3. 安装cuDNN（可选，如果你要做深度学习相关的算子）。
4. 安装Python 3.10+，以及PyTorch（用于对比性能）。

工具选型方面，我推荐几个：

• 代码生成：GPT-4或Claude 3.5 Sonnet。实测下来，Claude在CUDA代码生成上更稳，生成的代码更符合NVIDIA的编程规范。
• 编译：nvcc。直接用命令行编译，方便调试。
• 性能分析：nsight compute（ncu）和nsight systems（nsys）。这两个是NVIDIA官方的profiler，能看kernel的occupancy、memory throughput、instruction mix。
• 对比基准：cuBLAS、cuDNN、cutlass。用来对比AI生成的代码性能。

注意：安装CUDA时，不要用apt安装，建议用NVIDIA官方的runfile安装。apt安装的版本往往比较旧，而且容易和系统里的其他库冲突。

4.2 提示词设计：让AI生成可用的CUDA代码

提示词的设计，直接决定了生成代码的质量。我总结了一个模板：

请生成一个CUDA kernel，实现以下计算： - 算子名称：matmul - 输入：A (MxK, fp16), B (KxN, fp16) - 输出：C (MxN, fp16) - 目标架构：sm_80 (A100) - 要求： 1. 使用shared memory tiling，tile大小128x128x32 2. 每个thread block 256个线程 3. 使用float4向量化加载 4. 使用double buffering 5. 输出使用__restrict__指针 6. 包含完整的host端代码，包括内存分配和kernel launch

这个提示词里，我明确指定了tile大小、线程数、向量化宽度、优化技术。如果不指定，模型会生成一个“通用”的版本，性能通常很差。

我实测过，用这个提示词生成的matmul，在A100上能达到cuBLAS的70%左右。如果再手动调一下tile大小和warp数量，能到85%以上。

4.3 从生成到部署：完整流程记录

下面我记录一次完整的实操过程。目标是生成一个fp16的matmul，M=N=K=4096。

第一步：生成代码。用上面的提示词，Claude生成了大约200行CUDA代码。包括一个__global__kernel和一个host端的launch函数。

第二步：编译。用nvcc编译：

nvcc -O3 -arch=sm_80 -o matmul matmul.cu -lcublas

编译通过，没有报错。

第三步：正确性测试。写一个简单的测试，对比cuBLAS的结果。最大相对误差在1e-3以内，符合fp16的精度要求。

第四步：性能测试。用cudaEvent计时，跑100次取平均。结果是12.5ms。cuBLAS是8.2ms。性能比约65%。

第五步：性能分析。用ncu分析，发现两个问题：一是shared memory bank conflict比较严重，二是global memory的加载没有完全coalesced。

第六步：手动优化。我做了两处修改：一是把shared memory的padding从0改成4，消除bank conflict；二是调整了加载顺序，让warp内的线程访问连续的内存地址。

第七步：重新测试。性能提升到9.8ms，达到cuBLAS的84%。

这个流程里，AI负责生成baseline，人负责调优。我觉得这是目前最现实的协作方式。

4.4 性能对比与调优记录

为了更直观地展示AI生成算子的性能，我整理了一个对比表格：

从表格可以看出，AI生成的代码经过调优后，能达到官方库80%左右的性能。这个水平对于很多场景已经够用了，尤其是那些官方库没有覆盖的算子。

但要注意，这个数据是在A100上测的。在消费级GPU（比如4060Ti）上，由于架构差异，AI生成的代码性能可能更差。我实测过，4060Ti上AI生成的matmul只有cuBLAS的60%左右。原因是消费级GPU的L2 cache更小，AI生成的代码没有针对这一点做优化。

5. 常见问题与排查技巧实录

5.1 AI生成代码的典型问题

在用AI生成CUDA代码的过程中，我遇到了不少问题。这里整理一个速查表：

这些问题里，最常见的是性能极差。AI生成的代码往往“能跑”，但“跑不快”。原因通常是它没有做足够的优化，比如没有向量化、没有用shared memory、没有做double buffering。

5.2 性能不达标的排查思路

如果AI生成的代码性能不达标，我通常按以下顺序排查：

第一步：看occupancy。用ncu看achieved occupancy。如果低于50%，说明线程块太小或者寄存器用量太高。可以尝试增大block size，或者减少寄存器使用。

第二步：看memory throughput。如果global memory throughput很低，说明没有coalesced access。检查线程的访问模式，确保warp内的线程访问连续地址。

第三步：看shared memory bank conflict。如果shared memory的throughput很低，说明有bank conflict。可以尝试padding，或者改变数据布局。

第四步：看instruction mix。如果FMA指令占比很低，说明有大量整数运算或者分支。检查循环展开和边界处理。

第五步：看warp divergence。如果branch efficiency很低，说明有warp divergence。检查边界条件，尽量让warp内的线程走同一分支。

这个排查顺序，是我多年调优总结出来的。大部分性能问题，都能通过这五步定位。

5.3 独家避坑技巧

这里分享几个我踩过的坑，希望能帮你省点时间。

坑一：不要相信AI生成的launch配置。AI经常把grid和block的维度搞反，或者算错。比如一个MxN的矩阵，它可能把grid设成(N/32, M/32)，但实际应该是(M/32, N/32)。一定要手动检查。

坑二：AI生成的代码往往缺少边界处理。如果矩阵大小不是tile的整数倍，AI生成的代码可能会越界。一定要加边界检查，或者把矩阵padding到tile的整数倍。

坑三：AI对__syncthreads()的位置把握不准。有时候它会漏掉同步，导致race condition。有时候它会加多余的同步，导致性能下降。建议手动检查每一处shared memory的读写，确保同步正确。

坑四：AI生成的代码可能用了不支持的指令。比如它可能用了__shfl_sync，但你的目标架构是sm_70以下，不支持这个指令。一定要检查目标架构的指令集。

坑五：不要直接用AI生成的host代码。AI生成的host代码往往很粗糙，比如没有检查CUDA错误，没有释放内存。建议自己写host代码，只把kernel部分交给AI。

6. 算子开发的未来：AI与人的协作模式

6.1 AI会完全替代算子工程师吗

我的判断是：不会完全替代，但会大幅改变工作内容。

AI擅长的是生成“通用”的代码，但高性能算子往往需要针对特定硬件、特定shape做极致优化。比如在A100上，一个M=1024, N=1024, K=1024的matmul，最优的tile大小可能是128x128x32；但在M=64, N=4096, K=4096时，最优tile可能是64x256x16。这种细粒度的调优，AI目前还做不好。

另外，AI不擅长处理“非标准”的算子。比如一些自定义的融合算子，或者带有复杂控制流的算子。这些还是需要人来写。

所以，未来的算子工程师，工作内容会从“写代码”变成“调AI+调优”。你需要会写提示词，会分析性能瓶颈，会手动改代码。纯手写CUDA的需求会减少，但理解底层原理的需求会增加。

6.2 算子工程师的技能树更新

如果你是一个算子工程师，我建议你更新一下技能树：

第一，学会用AI工具。不是简单地让AI生成代码，而是学会设计提示词、分析生成结果、迭代优化。这本身就是一门手艺。

第二，深入理解硬件架构。当AI生成的代码性能不达标时，你需要知道问题出在哪里。这要求你理解SM架构、memory hierarchy、instruction pipeline。

第三，掌握性能分析工具。ncu、nsys这些工具，是你和AI协作的“眼睛”。没有它们，你就是在盲调。

第四，学习编译器原理。理解TVM、MLIR、LLVM这些编译器的原理，能帮你更好地理解AI生成的代码，也能帮你手动优化。

第五，保持对新硬件的敏感度。NPU、TPU、RISC-V这些新硬件，它们的编程模型和GPU不同。AI生成的代码可能不适用，你需要手动适配。

6.3 一个可能的协作流程

最后，我分享一下我目前用的协作流程：

1. 需求分析：确定算子的数学定义、输入输出格式、目标硬件。
2. AI生成baseline：用详细的提示词，让AI生成一个能跑的版本。
3. 正确性验证：写测试用例，对比参考实现。
4. 性能分析：用ncu分析瓶颈。
5. 手动调优：针对瓶颈，手动修改代码。
6. 迭代：重复4-5，直到性能达标。
7. 部署：集成到系统中，做端到端测试。

这个流程里，AI负责“从0到1”，人负责“从1到100”。我觉得这是目前最现实的模式。

我个人在实际操作中的体会是，AI生成的代码，最大的价值不是性能，而是“省时间”。以前写一个算子，从零开始要几天；现在用AI生成baseline，半天就能跑通。省下来的时间，可以用来做更有价值的优化。所以，不要害怕AI，学会用它，你的生产力会大幅提升。

最后再分享一个小技巧：如果你用AI生成的代码性能不达标，不要急着放弃。把ncu的分析结果贴给AI，让它根据性能数据重新生成。我试过几次，效果不错。AI虽然不懂硬件，但它能根据反馈调整代码结构。这就像和一个不懂硬件的聪明人合作，你需要给它足够的“上下文”。