算子是什么？深度学习算子从数学定义到GPU执行与开发优化全解析

看到算子是什么这个问题很多人第一反应是去翻书然后被一堆符号劝退。我在做深度学习模型部署和算子开发这几年几乎每隔一段时间就要跟人解释一遍这个概念。前阵子有个跑模型的同事看着profiling报告问我这上面全是kernel和operator到底哪个是算子我写的Python模型里也没见过它。这个问题问得特别好因为它戳中了一个事实算子是深度学习从框架到底层硬件之间最关键的翻译官但几乎没人认真介绍过它。这篇文章我打算彻底讲清楚算子到底是什么。我会把它放在三个语境里拆开——数学、深度学习框架、AI编译器/算子库——然后沿着一条真实的工作链路走一遍GPU上一个算子是怎么被执行的Sobel、Canny这些经典算子为什么到现在还有人用算子开发到底在开发什么以及当一张模型图里堆了几千个算子时性能瓶颈会出在哪里。想聊的内容有点多但每个都是实操中绕不开的东西。1. 算子的三重身份数学工具、计算单元、工程抽象要回答算子是什么绕不开一个事实它在不同场合指的东西不完全一样。很多人的困惑其实就来自这里。我在跟团队新人聊天时会把算子的身份拆成三层来讲每一层都有对应的真实场景。1.1 数学意义上的算子函数空间上的映射在数学里算子的定义很宽。简单说它是把一个函数或者一个向量、一张函数值组成的数组变成另一个函数或数组的规则。比如拉普拉斯算子它把一个空间分布温度场、电势场变成另一个分布——把每个点的数值求成一圈邻居值和中心值之间的差异程度反应的就是这个点在多大程度上是个局部极值。用生活化的类比算子的输入是一桌人的座位分布输出是每个人的人际关系张力。拉普拉斯算子干的事情就是看每个人跟四周朋友之间差多少、往哪个方向差最后输出一张谁会先吵起来的标注图。推而广之sobel算子做的是看哪个方向变化最剧烈canny算子做的是先降噪再找突变处然后连成边界线——它们都是把一张像素图变成另一张特征图的规则。这就是数学语境下算子最朴素的含义一种输入和输出同类的变换规则。1.2 深度学习框架语境下的算子可调用的计算接口进了深度学习框架的世界算子的定义被工程化了。PyTorch里的torch.nn.Conv2dTensorFlow里的tf.nn.conv2d这些API背后都对应一个算子。你调用它时框架给你一个统一接口但真正执行时它要调度到GPU上跑一个预先写好的kernel函数。这里要特别纠正一个常见误解算子不等于kernel但算子的最终归宿是kernel。算子是框架层的逻辑单元它规定了输入两个张量、输出一个张量计算规则是逐元素相乘再累加kernel是硬件层的实现单元它规定了GPU上每个线程负责哪一段数据、用什么样的指令去算。同一个算子在NVIDIA GPU上是一个CUDA kernel在华为昇腾上是一个AscendC kernel在CPU上是一段手写向量化的C代码。这也是为什么你在学习时经常看到一个词叫kernel算子——它指的是已经被编译、准备在具体硬件上执行的算子形态通常出现在分析性能时的CUDA trace、Ascend profiling数据里。1.3 算子库语境下的算子一个家族式的功能清单如果你翻过Halcon的算子中文手册或者用过OpenCV里那些图像处理函数你会看到另一番景象。那里面的算子更像一个功能库Blob分析、边缘提取、模板匹配、畸变校正……动词化、接口化、数量庞大。数学上一个sobel算子到了Halcon里可能对应了好几个算子和一堆参数组合。理解算子这三个字在不同语境下长什么样子比死记定义重要得多。因为实际工作中你基本不会遇到纯数学层面的算子问题你遇到的是为什么我这张图的sobel算子提不出边、为什么框架里这个算子在GPU上这么慢、为什么这个算子和那个算子可以融合——这些都是第二层、第三层语境里的事。2. GPU上的算子调用从主机端到kernel的一次完整旅行我经常爱说一句话不搞清算子怎么在GPU上跑你就不知道为什么程序慢。这一节咱们完整走一遍这个流程。这里以NVIDIA GPU和CUDA为例昇腾/AscendC的逻辑大体相似只是专有名词不同。2.1 主机端的下单过程你写下一行output torch.add(a, b)之后这一步在CPU主机端上做的事情其实很轻但步骤不少检查输入张量的shape、dtype、是否连续。根据算子名称Add在算子表中找到对应的实现句柄。把输入张量的地址、输出张量的地址、形状信息打包成一个调用参数结构体。调用CUDA runtime的接口把这个结构体和kernel函数指针一起交给GPU驱动排进GPU的命令队列。这一步我习惯叫它下单。它本身不干重活但它决定了整个算子执行的姿势。如果形状检查不通过就要先触发广播broadcast逻辑或报错如果张量不连续还可能要插入一个contiguous算子把数据先搬运成紧致排布才能继续做后面的计算。2.2 GPU端的真正干活kernel发射与线程调度命令队列里的任务到了GPU端才是算子的执行核心。GPU上跑算子时不是一个算子函数从头到尾执行一遍而是把一个算子拆分成成百上千个并行线程每个线程只处理数据中的一个或几个元素。拿向量加法来说CUDA kernel的原理是这样一串代码__global__ void vecAddKernel(const float* a, const float* b, float* out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) { out[i] a[i] b[i]; } }看着很简单但背后有一整套调度机制。GPU会按grid、block、thread三级组织线程一个grid处理整个张量block是能共享一块片上内存的线程小组thread是实际计算的最小单位。现代GPU一个SM流式多处理器上能同时驻留上千个线程硬件的线程调度器负责把block分配到各个SM上并且靠切换线程来隐藏访存延迟。如果你做过kernel级别的性能分析最后看到的图通常是这样每个kernel一行起止时间、SM占用率、显存带宽利用率、寄存器使用量一目了然。看这张图的能力基本就是算子开发和优化入门的第一道门槛。2.3 访存开销算子被饿死的那个坑新手看kernel性能最常见的问题是明明所有线程都在满负荷跑为什么时间还是很长答案十有八九在访存上。深度学习算子绝大多数是访存密集型而不是计算密集型。一个1000×1000的矩阵相加总计算量很小但你要把几百万个浮点数从显存搬进SM、算完再搬回去。GPU计算单元的速度远快于显存带宽于是大量时间花在等数据上。这也是为什么在分析GPU算子耗时的时候通常要同时看计算吞吐FLOPS和访存吞吐Bytes大部分人盯着的FLOPS反而不是瓶颈。举一组我自己实测过的数据在一个常见推理模型上某个简单的ElementwiseAdd算子kernel本身计算只花了不到5%的时间剩下95%都在从显存里搬数据。想要优化这种算子光靠把代码写得更高效是没用的真正有效的路子是后面要讲的算子融合减少数据在全局显存和SM之间的往返次数。2.4 一个容易被忽略的环节算子启动开销还有一个经常被忽略的现象叫kernel launch开销。每次向GPU发出一条指令哪怕任务再小也有一个固定的CPU-GPU交互成本大约在几微秒到几十微秒量级。如果你把一个大模型拆成几千个小算子一个个调光启动开销就能吃掉你总延迟的大头。这也是为什么现在主流的推理引擎TensorRT、MindSpore Lite、昇腾应用都会做模型图优化把一个一个的小算子合并成大算子块减少启动次数提高单次kernel里每个线程的工作量。3. Sobel、Canny与拉普拉斯经典图像算子在深度学习时代依然能打的理由聊完GPU执行链路我们把视角拉回到传统图像算子。发热词里有Sobel算子、Canny算子、拉普拉斯算子还有反对称化算子手写体这些都是我平时在图像处理和算子开发中用得非常频繁的东西。我单独开一节聊聊它们的价值因为它们和深度学习算子的关系比大多数人想的更近。3.1 Sobel与Canny的实际选择逻辑Sobel算子的作用是把一张灰度图变成梯度强度图它在水平和垂直方向各做一次卷积分别捕捉横边和竖边然后合成总梯度。原理上就是拿两个固定的3×3卷积核去扫图像本质上是卷积操作只不过卷积核的数值是固定的。Canny算子是个更大的流程先用高斯平滑降噪再用Sobel算梯度然后做非极大值抑制NMS把粗边细化最后用双阈值连接断开的边缘。很多人问既然深度学习模型都能做边缘检测了这些老算子还拿来干嘛我的回答通常是老算子在模型的前处理阶段依然是最稳定的选择。比如文档扫描、工业质检里先用Sobel/Canny检测边缘提取ROI区域再交给深度学习模型做分类稳定可靠还能省下一大截训练成本。数据预处理流程里这些算子在CPU上跑一遍也就几毫秒用好了比dataset增强的各种花活实在得多。3.2 拉普拉斯算子与卷积算子的数学血缘说到拉普拉斯算子它在图像里常用作二阶导数操作的离散近似作用是找图像里灰度突变更灵敏的地方也常用于图像锐化和作为深度学习中一些网络结构的初始化先验。有意思的是拉普拉斯算子离散化之后就是一个3×3的卷积核0 1 0 1 -4 1 0 1 0你拿这个卷积核去图像上扫一遍等于在每个像素上算自己和上下左右邻居的差异之和。它和CNN里的卷积算子本质上用的是同一套卷积数学。我在做网络结构设计时经常说你能把Canny拆成高斯模糊求梯度两遍阈值处理再把每一步都改写成可微的算子那你就基本理解了如何把传统算子和深度学习算子揉在同一个pipeline里。很多工业视觉方案就是这么做的传统算子负责粗定位深度学习算子负责细分类。3.3 反对称化算子和手写体识别中的现身方式反对称化算子这个词看起来吓人但它的内核其实很直观。一个矩阵或一个算子总可以拆成正对称部分和反对称部分。反对称部分满足转置后加负号通常对应着某种旋转或者方向差的特征。在手写体识别、字符对齐、图像配准这些任务里反对称化算子常用于提取局部方向的非对称响应也就是笔迹往哪边偏、扭转程度多大。它的数学形式和sobel算子的梯度方向提取异曲同工只是从的张量运算和旋转不变性的角度重新组织了一遍。我之所以特意提它是因为很多人遇到XX化算子这种词容易发怵以为是什么高深莫测的东西。实际在工程落地时它就是某个特征提取过程里的一组矩阵运算你完全可以把它先放到深度学习框架里用张量运算把这个算子调通再决定是否需要写成高效kernel。先当成数学函数理解再当成性能热点来优化是我处理这类问题的一贯次序。4. 算子开发到底在开发什么从手写kernel到LLVM自发现进入算子开发这个领域你会发现它跟写业务代码完全是两回事。它要同时懂硬件架构、懂编译原理、懂math库、懂并行计算。我把它拆成手写和自发现两条路径来讲。4.1 手写一个GPU算子从零到能跑的最小步骤假设我们现在要开发一个自定义算子对输入张量做exp(x) x * 2在GPU上跑。一般人用PyTorch直接写torch.exp(x) 2 * x就完事了。但如果你想把它改成融合算子减少多次kernel launch和中间张量的显存读写手写一份CUDA kernel是最直接的方案。开发流程大致是写kernel代码核心逻辑尽量简单循环交给线程组织grid-stride loop。用cudaMalloc分配显存用cudaMemcpy把数据拷进去。launch kernel指定grid尺寸和block尺寸。用cudaDeviceSynchronize等待执行完成。拷回结果和PyTorch的参考实现做对比验证。用ncu/nsys做profiling观察访存效率、SM占用率和寄存器溢出情况。手动写kernel最容易翻车的地方有两个一是越界访问二是精度不一致。前者会导致极其难查的报错甚至不报错悄悄返回错误数据后者常发生在浮点数求和顺序不一致时。所以在算子开发里我会强调先写一个CPU或框架侧参考实现把输出结果锁死再去碰GPU实现否则后续优化时根本没法定位是谁引入了误差。4.2 LLVM算子自发现编译器在帮你做模式匹配近几年的编译器领域算子自发现成了一个很热的方向尤其在LLVM及其上层生态MLIR、TVM等里。它做的事情简单说就是编译器自动扫描一段代码的计算模式发现其中重复出现的小算子结构然后推导出等价但更高效的整体实现。在LLVM里这通常是靠pattern matching和循环优化实现的。比如它对循环进行分析发现循环体可以被向量化就自动生成SIMD指令把原来四条标量指令合成一条向量指令。又比如它发现一组乘法和加法结构符合FMA乘加融合模式就自动把它替换成一条FMA硬件指令。这个自发现的思路跟上面说的算子融合异曲同工。在大模型时代算子数量动辄上千全靠手工去调kernel是不可能的。LLVM让编译器去看代码然后自动识别算子模式至少在CPU和通用GPU上缓解了大量人工负担。但我也要说句实在话LLVM自发现在CPU上效果很好到了专有AI加速器昇腾、NPU、GPU的Tensor Core上光靠通用编译器还不够。因为AI加速器有非规则的存储层级、DMA搬移、多核同步、特殊指令等硬件特性需要专门的上层IR比如MLIR、AscendC的tiling和调度框架来承载。所以你会看到昇腾场景里有专门的AscendC算子开发框架而不是只靠LLVM那套东西。4.3 AscendC融合算子以matmulprelu为例在AscendC上做算子开发一个我很常见的例子是融合MatMul和PReLU两个算子。原始模型结构往往是先做矩阵乘法再做参数化ReLU激活。如果不融合中间结果要经过一次全量写显存、再读回来耗时一下就上去了。融合思路是把PReLU的计算逻辑直接并入MatMul的kernel里其中一个block算完一部分矩阵乘结果后在片上寄存器里马上接着做PReLU激活再把结果写回全局内存。这样中间结果根本不出SM或者不出AI Core的片上缓冲省掉的访存就是实打实的性能收益。代码层面的简化示意示意性伪代码非完整AscendC实现// 伪代码示意融合 matmul prelu for (int i block_row_start; i block_row_end; i) { for (int j block_col_start; j block_col_end; j) { acc 0.f; for (int k 0; k K; k) { acc a[i * K k] * b[k * N j]; } // prelu: y x0 ? x : slope*x, 这里融合进去了 out[i * N j] acc 0 ? acc : slope * acc; } }真正写AscendC时还需要处理tiling数据切分、多核分配、double buffer乒乓缓冲这些细节但核心思想就是这个把多个算子的计算逻辑编排进一个kernel减少数据搬运次数。5. 算子的结构性挑战当模型里同时跑上万个算子时跳到更高一层来看。一个现代大模型比如一个标准的Transformerforward过程要执行多少次算子调用我随便拿个小模型来统计就有几百个算子节点大模型那就是几千上万个。这里的问题是纯结构性的跟单个算子没有关系。5.1 首个挑战启动开销的复利效应如果每个算子单独启动一次每个启动都有固定开销几千个算子叠加起来单看startup就会积少成多。更麻烦的是GPU本身有流水线机制理想情况下上一个kernel刚结束下一个kernel立刻顶上。但由于执行依赖关系很多算子必须等前一个算子的输出完全落盘后才能开始。这种串行依赖一旦形成GPU的空闲率就会飙升。我在做部署优化时见过最极端的案例一个BERT小模型在未做图优化时kernel利用率只有30%上下把相邻的若干小算子融合成5个大算子之后利用率直接跳到70%。同一张图跑的算子数量大幅减少端到端延迟降了接近10倍。5.2 第二个挑战显存带宽的一次性消费除了启动开销还有显存带宽问题。每个算子一张输入一张输出数据总在全局内存里兜圈子。一个只有两层的模型还好说一个上百层的模型中间要产生的中间张量会非常可观。算子融合的主要目的说穿了就是把数据在片上多算一会少回一次全局内存。这里有个判断公式我可以分享如果算子的算术强度计算量/数据搬移量低于硬件平台平衡点就属于访存密集型算子先考虑融合或让它少搬数据如果算术强度高于平衡点才需要考虑提升计算本身的吞吐。这句话救了我好几次。因为很多人一看算子慢就急着调线程块大小、调循环展开结果在访存瓶颈下怎么也调不动。5.3 第三个挑战硬件单元的结构性浪费AI加速器通常配备多种计算单元向量单元、矩阵单元、标量单元各有所长。一个算子如果只用了某种单元其他单元就闲置。为了避免这种浪费现在算子开发的趋势是融合异构算子让一个kernel能同时调用向量指令和矩阵指令把矩阵乘的核心计算交给专用单元把激活、归一化这些逐元素操作交给向量单元两者并行跑。这是比单纯减少若干个kernel更高阶的优化思路。5.4 实测中常用到的优化策略清单下面这张表是我在算子性能调优过程中用得最多的几个抓手列出来供参考。它们是长期实践沉淀下来的通用方法不是某个平台专利。优化方向手段适用场景融合多个算子合成一个kernel相邻且存在数据依赖的算子链向量化用SIMD指令批量处理数据逐元素操作、访存带宽受限算子Tiling切分把大张量切成适配片上存储的块大矩阵乘、大卷积算子双缓冲计算当前块时预取下一块数据长计算链、访存延迟敏感的算子共享内存复用把反复使用的数据块留在片上卷积、矩阵乘、窗口类算子降低精度用FP16/BF16/INT8替代FP32对精度不敏感的推理场景隐藏启动开销用CUDA Graph/昇腾Task复用捕获箭头算子数量大、单算子耗时短的模型6. 想入算子开发一条从数学到AscendC的实践路径很多人看到算子开发动辄提到CUDA、LLVM、AscendC第一反应是门槛太高。但我的经验是算子开发更像一门先能用、再玩精的手艺入门路径比想象中平滑关键是按对的顺序学。6.1 第一步把深度学习算子的数学脸认齐先把核心算子的数学定义过一遍矩阵乘法、卷积、池化、归一化、各种激活函数、softmax、注意力机制。这个阶段不需要学多深重点是能回答两个问题这个算子输入输出是什么它每一步做了什么能用一段Python/PyTorch代码手把手把逻辑模拟出来就算过关。这个阶段我建议把Sobel、Canny、拉普拉斯这些传统图像算子的实现也过一遍它们和卷积算子的关系会帮你建立很强的直观感觉。6.2 第二步学会看profiling数据如果你还不了解性能瓶颈长什么样去优化算子就像蒙眼开车。建议先在真实模型上跑一遍nsys或ncuGPU或昇腾的msprof把拿到的时间线pipeline对着算子列表看几遍逐渐学会判断某个kernel慢是慢在访存还是慢在计算。看明白了再动手优化会顺利得多。6.3 第三步手写属于自己的第一个kernel不一定是GPUCPU SIMD也可以用。从向量加法开始再到逐元素实现expscale再到一个小矩阵乘法。每走一步都做两件事和参考实现对比正确性用profile工具测量性能有没有变化。这一步的核心是建立实测-反馈-修改的循环没有这个循环看再多优化文章都只是纸上谈兵。6.4 第四步接触AI编译器框架和昇腾AscendC当你对kernel和硬件的关系有了手感再去接触LLVM、MLIR、TVM以及AscendC这些框架就会舒服很多。AscendC的开发流程虽然有自己的tiling和同步语法但思路跟CUDA非常相似都是先理解数据布局和计算切分再写计算逻辑再做double buffer之类优化。想系统走这条路线的可以直接关注昇腾推出的算子开发能力认证里面覆盖的学习内容和实际工程高度相关考完基本具备了独立开发一个融合算子的能力。结尾的建议最后说一点个人体会。我见过太多人在学算子时要么一头扎进数学公式出不来要么直接跳到调优工具里乱试。如果你也只能记住一件事我建议是先搞清楚你想优化的算子在硬件上把数据搬了几次、算了多少次、等了几次再去谈怎么改代码。算子的本质就是数据在硬件上的流动和变换搞明白数据怎么走算子是什么这个问题的答案就会自然浮现。你现在如果手边有能跑GPU的环境不妨直接打开profiling随便找几个算子看看它们的时间和访存数据比我在这里讲一万字都管用。