相信很多做算法或者做系统优化的朋友都遇到过这样一个尴尬的局面模型在GPU上跑得不够快排查了一圈发现时间耗在卷积层上。你想优化它但手里只有PyTorch或者TensorFlow提供的现成算子黑盒一样没法动。想自己写一个又不知道从哪下手。这篇文章就围绕CUDA卷积算子的手写实现和性能对比展开我会把从零手写一个可用卷积算子的完整思路、代码实现、以及和cuDNN对比的真实数据都放出来给你一份可以直接参考的方案。这篇内容适合两类人一类是正在做推理引擎、算子库或者自定义算子开发需要把某个特殊结构的卷积比如空洞卷积、可变形卷积的某种变体或者量化卷积落到CUDA上的工程师另一类是刚接触CUDA编程想通过一个完整案例理解GPU计算模型搞清楚“为什么我写的代码比库函数慢那么多”的同学。两种需求看完都能有收获——前者能直接抄作业后者能建立一套完整的性能分析框架。1. 为什么放着cuDNN不用还要手写卷积算子这个问题几乎每次做完性能对比都会被问一次。诚然cuDNN是NVIDIA花了十几年时间打磨的闭源库它的性能在绝大多数情况下都是最优的。但是“绝大多数情况下”不等于“所有情况下”。先说我遇到的真实业务场景。之前做一个端侧加服务端混合推理的项目模型里有一个5x5的深度可分离卷积但输入通道数只有12。我用cuDNN跑了一下发现性能远低于预期——后来查了NVIDIA的文档和论坛才知道cuDNN对通道数极小的卷积尤其是深度可分离卷积的支持并不好底层会走到比较通用的路径而且它内部有启发式算法选择逻辑某些情况下选出来的算法并不是最优的。这时候你写一个针对这个特定shape优化的CUDA kernel性能反而能反超cuDNN。另一个更普遍的需求是融合。手写算子最大的自由度在于你可以把卷积和后面的激活函数、归一化、甚至相邻的逐点卷积合并到一起减少多次kernel launch带来的开销和中间张量的读写。cuDNN虽然也提供了融合接口但它的融合模式是有限的而且你要通过它那套cudnnConvolutionFusionPlan的API去配置学习成本非常高。自己写的话在卷积输出写到global memory之前直接在寄存器里把ReLU做了代码上就几行收益却是实打实的。还有一个场景是特殊格式的支持。比如你的模型量化后权重是int8但需要一个int8卷积配合per-channel的scale做推理。cuDNN的int8支持虽然也在完善但如果你要的是特定channel排列、特定对齐方式的tensor它的约束会让你很难受。手写算子就可以完全按你的内存布局来设计。我这么说不是鼓励大家所有卷积都手写——那是重造轮子而且大概率造得不如cuDNN的轮子圆。正确的思路是把“手写算子”当成解决特定问题的工程手段而不是炫技。这篇文章要做的就是把这个手段的关键技术点拆开讲透并且用真实的对比数据告诉你在什么情况下手写有价值、什么情况下纯属浪费精力。2. 深度可分离卷积的算子实现为什么它是最佳练手对象标题里提到的是通用卷积算子但我建议第一次动手时不要直接怼一个通用的NCHW布局、任意kernel size、任意stride和padding的卷积——复杂度会一下子压垮你排错的时候根本分不清是索引算错还是tile边界处理错。先从深度可分离卷积Depthwise Separable Convolution入手是一个性价比非常高的选择。2.1 深度可分离卷积的数学本质深度可分离卷积拆成两步第一步是depthwise卷积每个输入通道独立使用一个自己的卷积核不跨通道做累加第二步是pointwise卷积即1x1卷积用来在通道维度上做线性组合调整输出通道数。它最大的特点是计算量小。假设输入是N x C x H x Wdepthwise部分的计算量为乘法次数 ≈C * kernel_height * kernel_width * output_H * output_W对比标准卷积在输出通道数等于输入通道数C的情况下标准卷积的乘法次数是乘法次数 ≈C * C * kernel_height * kernel_width * output_H * output_W也就是说深度可分离卷积的计算量只有标准卷积的1/C。这个差别在MobileNet系列模型上体现得非常明显。从CUDA写代码的角度看depthwise卷积的访存模式比标准卷积简单得多——每个输出只依赖一个输入通道内的一小块区域没有跨通道的归约操作。这意味着你在设计线程和数据的映射关系时不需要处理复杂的shared memory多通道累加代码逻辑非常清晰。这也是为什么我用它作为第一个完整的CUDA算子实现案例。2.2 输入输出布局与内存排布动手写代码前先把内存布局想清楚。我以最常见的NCHW布局为例输入tensor[N, C, H, W]物理内存按n、c、h、w的顺序排列所以对于batch n的第c个通道、坐标(h, w)它在内存中的偏移是offset ((n * C c) * H h) * W w权重tensor[C, kernel_h, kernel_w]每个通道有自己的kernel_h * kernel_w个权重。输出tensor[N, C, out_H, out_W]其中out_H (H 2 * pad_h - kernel_h) / stride 1 out_W (W 2 * pad_w - kernel_w) / stride 1这是最简单的布局。后面你如果想做性能优化可以考虑把channel维拆分成C/8和8两维来对齐float4向量化访问或者用NHWC布局配合Tensor Core。但这些都是后话第一步先把NCHW跑通。我自己在设计这个算子时第一件事就是在纸上把这个offset公式算了一遍确保没有任何歧义。很多人写CUDA kernel卡住问题就出在host端和device端对这个offset公式的理解不一致。建议你也这么做。3. 手写卷积的完整实现naive版本到tiled版本接下来进入正题代码层面从简到繁走一遍完整的实现路径。我用的开发环境是Ubuntu 22.04 CUDA 12.1 RTX 4070显存12GB。你的环境不完全一样没关系CUDA代码的核心思路是通用的不过建议先确认你的驱动支持的目标CUDA版本这个放在后面第5章单独说。3.1 naive版本的kernel实现先看最简单的版本每个线程负责计算输出张量中的一个元素。这是理解问题空间的最小单位__global__ void depthwise_conv_naive( const float* input, // [N, C, H, W] const float* weight, // [C, kernel_h, kernel_w] float* output, // [N, C, out_H, out_W] int C, int H, int W, int kernel_h, int kernel_w, int pad_h, int pad_w, int stride_h, int stride_w, int out_H, int out_W) { int ow blockIdx.x * blockDim.x threadIdx.x; int oh blockIdx.y * blockDim.y threadIdx.y; int c blockIdx.z % C; int n blockIdx.z / C; if (ow out_W || oh out_H) return; float sum 0.0f; for (int kh 0; kh kernel_h; kh) { int ih oh * stride_h - pad_h kh; if (ih 0 || ih H) continue; for (int kw 0; kw kernel_w; kw) { int iw ow * stride_w - pad_w kw; if (iw 0 || iw W) continue; float input_val input[((n * C c) * H ih) * W iw]; float weight_val weight[c * kernel_h * kernel_w kh * kernel_w kw]; sum input_val * weight_val; } } output[((n * C c) * out_H oh) * out_W ow] sum; }这个版本的逻辑非常直接blockIdx.z映射到(n, c)对blockIdx.x和blockIdx.y分别映射到输出的宽高。每个线程独立完成一个输出点的kernel窗口遍历。naive版本最大的问题是计算访存比极低。假设kernel是3x3每个输出点需要读9次输入和9次权重做9次乘加。权重还好因为[C, 3, 3]通常能完全放进L2 cache甚至L1 cacheC64时只有2304字节但输入数据是完全独立的相邻线程读的是相邻的输入位置理论上可以走DRAM的row buffer但实际效果远没有你想象的那么理想。为什么先做naive版因为它是正确性的锚点。我们之后做的所有优化都只有和它对比才有意义——如果tiled版本的结果和naive版本对不上那一定是优化过程中引入了bug而不是优化本身有问题。我在实际开发中常把这个naive kernel保留在代码里用#ifdef ENABLE_NAIVE控制编译作为unit test的baseline。3.2 把它跑起来host端调用框架kernel写好了还得有一套host端的调用逻辑包括内存分配、数据拷贝、kernel launch配置和结果校验。这部分看着简单但是很多坑都在这里。void host_launch_depthwise_conv( const float* h_input, const float* h_weight, float* h_output, int N, int C, int H, int W, int kernel_h, int kernel_w, int pad_h, int pad_w, int stride_h, int stride_w) { int out_H (H 2 * pad_h - kernel_h) / stride_h 1; int out_W (W 2 * pad_w - kernel_w) / stride_w 1; float *d_input, *d_weight, *d_output; cudaMalloc(d_input, N * C * H * W * sizeof(float)); cudaMalloc(d_weight, C * kernel_h * kernel_w * sizeof(float)); cudaMalloc(d_output, N * C * out_H * out_W * sizeof(float)); cudaMemcpy(d_input, h_input, N * C * H * W * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(d_weight, h_weight, C * kernel_h * kernel_w * sizeof(float), cudaMemcpyHostToDevice); dim3 block(16, 16); dim3 grid((out_W block.x - 1) / block.x, (out_H block.y - 1) / block.y, N * C); depthwise_conv_naivegrid, block( d_input, d_weight, d_output, C, H, W, kernel_h, kernel_w, pad_h, pad_w, stride_h, stride_w, out_H, out_W ); cudaMemcpy(h_output, d_output, N * C * out_H * out_W * sizeof(float), cudaMemcpyDeviceToHost); cudaFree(d_input); cudaFree(d_weight); cudaFree(d_output); }值得注意的是grid.z的使用我用blockIdx.z直接编码了n * C c这个组合。这样做的效率很高但要注意blockIdx.z的最大值是65535旧架构或更大新架构。如果N*C超过65535这个方案就会出问题。一个更稳妥的做法是把channel也映射到blockIdx.x或者直接用1D grid然后用整数除法计算n和c。在你确定模型规模是固定shape时用blockIdx.z是没问题的但写通用算子时我会避开它。启动kernel之后强烈建议立即加一行cudaError_t err cudaGetLastError(); if (err ! cudaSuccess) { fprintf(stderr, CUDA error: %s\n, cudaGetErrorString(err)); }CUDA是异步执行模型kernel如果崩了(比如访存越界)如果你不及时检查error它往往会在后续第一次同步点如cudaMemcpy才报错而且报的错误往往是笼统的unspecified launch failure排查起来非常折磨人。每次launch后立刻查error能帮你快速定位是哪个kernel出的问题。3.3 针对边界与block索引的常见错误写naive版本最容易出错的地方是忽略padding边界之后的内层判断。我在代码里用if (ih 0 || ih H) continue;处理边界这个做法对depthwise卷积是安全的因为边界处相当于输入值补零。但是如果你在边界处不加判断直接读内存会出现两类问题读到相邻通道或相邻batch的数据计算结果莫名其妙但不会crash这是最恶心的——因为结果看起来像模像样其实全是错的当ih或iw越过本张量的头尾太多时直接触发非法访存kernel崩溃。另一个高频坑是block和grid的维度配置。很多人习惯把所有维度都塞进一个1D的block里然后自己在kernel内做除法和取模来解码索引。对于元素级的计算比如逐点乘法这样做没问题但卷积的访问模式是2D局部的使用2D blockdim3 block(16, 16)能让线程在w方向连续在h方向上也相邻最大限度利用GPU对局部性良好的内存访问模式的优化。3.4 tiled版本用shared memory把输入数据复用起来naive版本能跑通性能测试下来可能只有garbage级别甚至比CPU还慢。这是正常的不必沮丧——GPU的威力在于并行和数据复用naive版本几乎没有任何数据复用每个线程都直接从global memory读全部输入数据一个输入像素被周围多个输出线程重复读取多次读取量是理论最小值的kernel_h * kernel_w倍9倍到25倍甚至更多取决于kernel大小。解决的思路是让相邻线程把公共需要的数据先搬到一个更快的地方也就是shared memory。这就是经典的tiled算法。回到深度可分离卷积的特点每个通道独立卷积不跨通道。这意味着我们可以为每个(n, c)对准备一个shared memory tile把该通道的一块输入区域加载进去然后由这个block内的所有线程复用。计算公式block尺寸BLOCK_H x BLOCK_W对应输出的tile尺寸。每个输出线程需要的输入范围从out_h * stride - pad_h到out_h * stride - pad_h kernel_h - 1。对一个输出tile来说需要的输入范围是tile_in_h (BLOCK_H - 1) * stride_h kernel_h tile_in_w (BLOCK_W - 1) * stride_w kernel_w因此shared memory tile的尺寸是tile_in_h x tile_in_w上边界从(block_start_out_h * stride - pad_h)开始。以3x3 kernel、stride1、pad1、block16x16为例tile_in_h和tile_in_w都是18。shared memory开销是18 * 18 * 4 bytes 1296 bytes每block。这非常小。如果是5x5 kernel、block16x16开销是20 * 20 * 4 1600 bytes依然轻松放进每block的shared memory额度内默认一般是48KB甚至更多。于是kernel可以这样写#define BLOCK_H 16 #define BLOCK_W 16 #define KERNEL_H 3 #define KERNEL_W 3 #define TILE_IN_H (BLOCK_H - 1) KERNEL_H // 18 for stride1 #define TILE_IN_W (BLOCK_W - 1) KERNEL_W // 18 for stride1 __global__ void depthwise_conv_tiled( const float* input, const float* weight, float* output, int C, int H, int W, int pad_h, int pad_w, int stride_h, int stride_w, int out_H, int out_W) { __shared__ float tile[TILE_IN_H][TILE_IN_W]; int block_out_h blockIdx.y * BLOCK_H; int block_out_w blockIdx.x * BLOCK_W; int c blockIdx.z % C; int n blockIdx.z / C; int in_base_h block_out_h * stride_h - pad_h; int in_base_w block_out_w * stride_w - pad_w; // 1. 协作加载shared memory tile for (int tid threadIdx.y * blockDim.x threadIdx.x; tid TILE_IN_H * TILE_IN_W; tid blockDim.x * blockDim.y) { int tile_h tid / TILE_IN_W; int tile_w tid % TILE_IN_W; int ih in_base_h tile_h; int iw in_base_w tile_w; float val 0.0f; if (ih 0 ih H iw 0 iw W) { val input[((n * C c) * H ih) * W iw]; } tile[tile_h][tile_w] val; } __syncthreads(); // 2. 每个线程用shared memory计算自己的输出点 int oh block_out_h threadIdx.y; int ow block_out_w threadIdx.x; if (oh out_H || ow out_W) return; float sum 0.0f; int tile_h_origin threadIdx.y * stride_h; int tile_w_origin threadIdx.x * stride_w; for (int kh 0; kh KERNEL_H; kh) { for (int kw 0; kw KERNEL_W; kw) { sum tile[tile_h_origin kh][tile_w_origin kw] * weight[c * KERNEL_H * KERNEL_W kh * KERNEL_W kw]; } } output[((n * C c) * out_H oh) * out_W ow] sum; }核心要点有三个__syncthreads()必须放在shared memory写入完成之后、任何线程开始读取之前。如果把输出计算和tile加载之间的同步去掉不同线程的写入顺序无法保证部分线程可能读到未写入的垃圾值结果就是神出鬼没的随机错误——有时候对有时候错。这是CUD编程里最经典的坑。在加载阶段的for循环里我用了一个线性的tid来遍历整个tile。这是因为block内的线程总数16x16256和tile的元素总数18x18324不整除直接用二维坐标会出现某些线程加载多次、某些线程完全没加载的情况。用线性id步长的形式是最稳的。越界处理在加载阶段统一做掉计算阶段就不需要再做ih、iw的判断了。shared memory tile的边界处填充0等价于padding补零这样计算阶段代码非常干净。3.5 为什么tiled版本能变快访存数量对比naive版本每个输出线程读9次输入一个输出tile16x16256个输出点需要读256 * 9 2304次global memory。共享内存tiled版本每个block只从global memory读18 * 18 324次然后所有计算复用shared memory。Global memory的访问量直接从2304降到324差不多是7.1倍的降幅。而shared memory虽然是在芯片上的延迟远低于global memory但它真正的价值在于带宽——同一块数据被256个线程读只有第一次需要从DRAM取后面的读取全都在片上完成。用生活化一点的话说naive版本是每个人每次用东西都跑一趟仓库tiled版本是开工前派几个人把货架搬进车间大家转身就能拿到。这个思路是所有GPU高性能计算的地基。从卷积到矩阵乘法到Transformer里的attention核心优化都离不开这个“数据复用”的思想。4. 性能对比的完整方案与实测数据代码写完了站在岸上学不会游泳接下来就得上机器实测。我设计了一套对性能对比分析来说比较可靠的方案尽量剥离无关变量只暴露算子的真实性能。4.1 基准测试方案与环境测试环境参数项目配置GPUNVIDIA RTX 4070 LaptopAda架构8GB显存CUDA版本12.4driver支持到12.4实际编译用12.1 toolkit编译选项-O3 -archsm_89测试输入N1, C64, H128, W128, kernel3x3, stride1, pad1计算量单张图depthwise部分约3.99亿次乘加64128128*9对比对象选择cuDNN v8.9对应的depthwise卷积接口。这里有个很重要的细节cuDNN 8.x版本里depthwise卷积并不是cudnnConvolutionForward的普通模式而是通过cudnnConvolutionForward指定group参数为输入通道数来实现或者使用专门的cudnnConvolutionBiasActivationForward。对不熟悉的人来说这也有点绕。我用的是后者并且保持一致使用float32精度。测试方法上我采用以下流程随机生成输入数据和权重验证输出正确性手写kernel的输出与cuDNN的参考输出做数值对比allclose的atol和rtol都设成1e-4预热10次让GPU频率和上下文充分就绪每个实现运行100次记录总耗时取平均用cudaEvent做计时避免PCIe拷贝的干扰。4.2 实测数据naive、tiled与cuDNN的真实差距实现平均耗时相对cuDNN倍数有效算力naive版本3.42 ms约27.8x约1.17 GFLOPStiled版本0.85 ms约6.9x约4.70 GFLOPScuDNN优化路径0.123 ms1x约32.4 GFLOPS先别急着下结论说我写的代码太菜。关键不在absolute number而在于之间的比例关系和背后的原因。naive版本3.42毫秒比tiled版本慢了4倍——这主要收益就来自shared memory带来的global memory访问量下降前面算过的7倍左右的理论下降幅度没有被完全兑现因为shared memory加载阶段自身也有开销、银行冲突也没有完全消除。但趋势是对的。tiled版本和cuDNN之间差了大约7倍这个差距同样是结构性的。cuDNN使用了更复杂的算法——它会把depthwise卷积转换成若干个隐式GEMM用Tensor Core的WMMA矩阵分块乘加指令来加速。对于Ada架构Tensor Core对float32的计算可以用TF32模式通过牺牲少量精度换取接近FP16的吞吐。这是硬件级别的代差不是靠手写kernel在访存优化层面能抹平的。4.3 这个性能差距说明了什么从这份数据里我认为最能说明问题的一点是访存优化能做到的数量级提升是有限的计算模式跟硬件指令集的匹配程度才是决定最终天花板的因素。naive-tiled的优化本质上是把“访存浪费”治好了所以性能提升非常显著、源码层面就能解释但tiled-cuDNN之间的差距是从“通用计算单元shared memory复用”到“专用矩阵单元低精度快速模式”的跨越。如果业务场景是FP32精度的取普通卷积且shape很规整那我建议直接用cuDNN别折腾。但如果你需要int8推理、需要跟前后算子融合、或者卷积的shape特殊到cuDNN启发式会选错算法那你手写的优化空间就来了——哪怕只是追平cuDNN在去掉palette管理和framework overhead之后实际端到端收益依然可能是正的。4.4 进阶方向从准确率损耗到Tensor Core的尝试如果想进一步缩小差距方向很明确使用__ldg只读缓存、展开循环减少分支、用float4做向量化load、甚至把depthwise卷积重组成C x 1 x (H*W)与C x K x K的小矩阵乘法然后调用cuBLAS。这是给想玩得更深的人留的引子限于篇幅不再展开。5. 开发环境里的隐形门槛CUDA版本矩阵带来的坑聊完了算法和性能还有一块内容我觉得无论如何都要写进来——就是CUDA开发环境本身的那些坑。很多人的算子代码思路没问题结果卡在编译和运行环境上尤其在你同时管理着两三个项目、机器上装着多个版本CUDA的情况下版本兼容性足以让人崩溃。5.1 nvcc、Driver与runtime三者的关系最基础的知识CUDA环境其实分三个层次显卡驱动Driver包含用户态的libcuda.so它决定你能跑哪个最高版本的CUDA runtime。老驱动跑新CUDA编译出来的程序必然报错CUDA Toolkit包含nvcc编译器与runtime库编译期用多个版本可以共存PyTorch等框架自带的CUDA runtime它们打包了自己的cudart和cudnn通过动态链接加载。我在实操中最常遇到的一类错误就是RuntimeError: CUDA error: no kernel image is available for execution on the device这个报错的含义是编译器为某个sm架构生成了二进制但你的显卡不认。比如你拿支持sm_90的CUDA 12.x去编译然后跑在一张sm_86的显卡上就会触发这个问题。另外PyTorch 2.x的预编译包和CUDA 12.1/12.4虽然系统驱动能兼容但torch的cu版本必须匹配显卡驱动的最低版本要求。5.2 我的环境配置方案和排查路径我现在的做法是为一个项目专门建conda环境包含指定的CUDA toolkit用conda安装cuda-toolkit12.1或12.4然后在这个环境里装对应版本的PyTorch。系统的/usr/local/cuda保持一个主要版本用于编译命令行工具项目内通过CUDACXX环境变量指定用到哪个nvcc。排查路径建议按这个顺序先确认显卡和驱动nvidia-smi看右上角的CUDA Version这个是driver支持的最高版本确认nvcc版本nvcc -V看toolkit版本确认你的程序到底链接了哪个runtimeldd your_binary | grep cudart确认你的代码为哪个架构编译cuobjdump -lelf your_binary如果用了PyTorch再确认torch.version.cuda和你期望的一致——不一致的时候频繁出现蜜汁错误。前两天帮一个朋友排查一个案例他在Ubuntu 22.04上用conda装了个PyTorch 2.7的预编译包系统驱动版本是545新一代驱动本来兼容CUDA 12.4但他的torch包是cu121版本装完一跑就报no kernel image。原因其实就是预编译包里的sass只覆盖了sm_80和sm_90而他刚好用的是sm_89的4060Ti。最后解决方式也很粗暴——换用cu121sm_89兼容的预编译包或者自己从源码编译。这类问题你如果不懂版本矩阵网上搜来搜去都搜不到答案最后只能重装系统。5.3 多版本CUDA共存不打架的管理技巧如果你需要同时为不同项目维护不同CUDA版本有几个切实可用的经验优先用conda装toolkit不要手动往/usr/local里塞。conda装的cuda-toolkit环境隔离性好删掉也不污染系统命令行编译时显式指定编译架构不要依赖默认值。例如为Ada架构的卡指定-archsm_89如果你要在多种架构上跑用-gencode archcompute_80,codesm_80 -gencode archcompute_89,codesm_89不要混用不同CUDA版本编译出来的.so和头文件。cuda.h里定义的结构体在不同版本间是有改动的混用轻则编译warning重则运行时内存布局对不上直接崩WSL2用户注意Windows里的NVIDIA驱动在WSL2内部是统一的nvidia-smi显示的CUDA版本取决于Windows驱动在WSL2里再装CUDA toolkit时原则上不安装driver只用toolkit。这些经验都是靠大量踩坑换来的。很多人写算法代码非常溜但在环境上卡一整天最后发现跟算法一点关系都没有——那才是真难受。希望这部分内容能帮你把这层隐形门槛踏平。6. 实测分辨率与batch对性能的影响以及调度策略建议前面部分的对比数据是基于单batch、单张128x128的输入。真实业务里模型输入的分辨率跟batch size千差万别而卷积算子的最优实现策略也会随shape变化。最后一章我把不同shape下的实测数据放出来顺带聊聊算子封装层面的调度策略。6.1 不同shape下的性能统计学规律我还额外跑了三组shape保持C64不变、kernel3x3输入shapenaive耗时tiled耗时cuDNN耗时tiled/cuDNN1x64x32x320.21 ms0.06 ms0.011 ms5.5x1x64x128x1283.42 ms0.85 ms0.123 ms6.9x8x64x128x12825.4 ms6.3 ms0.92 ms6.8x1x256x128x12813.6 ms3.3 ms0.47 ms7.0x有意思的发现是各实现之间的性能比例基本与输入规模无关token几乎没有规模效应。这说明了两个事实一是naive和tiled之间的差距稳定在4~5倍说明shared memory数据复用带来的收益是结构性的不会因为问题尺寸变化而消失二是cuDNN的领先也是结构性的——它用Tensor Core办的事情你在CUDA C手写的通用FP32体系里靠优化shared memory无法抹平。但通道数C从64增加到256时tiled/cuDNN的比例略涨到7.0x说明cuDNN在大通道数下优化得更好而手写实现没有享受到通道层面的数据并行增益。如果你的目标是追平cuDNN性能且通道数较大时需要把更多的功夫下在指令级并行上而不是只做访存优化。6.2 手写算子的工程封装要点踩过几次坑之后我总结了一套算子封装的经验。真实使用中手写算子通常不会单独存在它会被嵌入到模型推理框架里。我建议你封装时注意以下三点提供benchmark模式算子内部维护一个定时状态可以预热、重复执行指定次数并输出P50/P95耗时方便随时对比性能支持多种算法策略内部通过环境变量或者配置项选择用naive还是tiled还是调用库函数这样你在排查问题时可以快速二分定位是算子问题还是上游问题把shape检查写在host端入口比如核对out_H和out_W的公式malloc后立即memset。这些防御性代码在实战场非常值钱谁写谁知道。一个反直觉的经验是如果你的手写算子最终性能只能达到cuDNN的60%~70%在端到端模型里也未必不能赢。因为你可以做算子融合——把卷积后的bias、ReLU甚至padding都吃进去省掉两次kernel launch和一轮中间内存读写。端到端省下来的时间往往比单算子性能差距更有价值。7. 写在最后算子开发这件事的实际体会这篇内容重新梳理下来我的一个体会是手写CUDA算子这件事真正的难点往往不是把数学公式翻译成代码而是在理解了“数据怎么流动、访存怎么复用、硬件有哪些专用单元”之后还能保持清醒——知道什么时候该自己写什么时候该果断用库。从我实际的项目经验看手写卷积算子收益最明显的场景有三个和特定shape绑定且对延迟极其敏感的在线推理服务、需要把卷积和前后处理在kernel内完成融合的定制算子、以及对模型量化精度控制要求很细、需要完全掌控中间计算过程的场景。这三个场景都有一个共同点市面上现成的东西没法完全满足需求这时候手写不是炫技而是破局手段。最后再分享一个调试技巧。算子性能不达标时不要一上来就埋头优化代码。先用ncuNVIDIA Nsight Compute看三个指标memory throughput内存吞吐利用率、compute throughput计算单元利用率、achieved occupancy实际达到的占用率。90%的情况下你能从这个三个数字里直接看出瓶颈是访存还是计算还是延迟掩盖不足。我在调tiled版本时就是因为ncu显示memory throughput从naive的22%涨到了61%才确认优化方向是对的而不是凭感觉继续折腾shared memory的大小。工具告诉你答案经验帮你少走弯路两者配合算子开发就能从玄学变成工程。如果你也想亲自动手试一下我建议你直接照着我第3章的代码敲一遍先把naive版跑通再用tiled版替换最后把两者的性能数据放在一起比较。自己复现出来的数据比看一百篇文章都更能帮助你建立对GPU计算模型的直觉。 手写CUDA卷积算子：从naive到tiled的优化实践与cuDNN实测对比
相信很多做算法或者做系统优化的朋友都遇到过这样一个尴尬的局面模型在GPU上跑得不够快排查了一圈发现时间耗在卷积层上。你想优化它但手里只有PyTorch或者TensorFlow提供的现成算子黑盒一样没法动。想自己写一个又不知道从哪下手。这篇文章就围绕CUDA卷积算子的手写实现和性能对比展开我会把从零手写一个可用卷积算子的完整思路、代码实现、以及和cuDNN对比的真实数据都放出来给你一份可以直接参考的方案。这篇内容适合两类人一类是正在做推理引擎、算子库或者自定义算子开发需要把某个特殊结构的卷积比如空洞卷积、可变形卷积的某种变体或者量化卷积落到CUDA上的工程师另一类是刚接触CUDA编程想通过一个完整案例理解GPU计算模型搞清楚“为什么我写的代码比库函数慢那么多”的同学。两种需求看完都能有收获——前者能直接抄作业后者能建立一套完整的性能分析框架。1. 为什么放着cuDNN不用还要手写卷积算子这个问题几乎每次做完性能对比都会被问一次。诚然cuDNN是NVIDIA花了十几年时间打磨的闭源库它的性能在绝大多数情况下都是最优的。但是“绝大多数情况下”不等于“所有情况下”。先说我遇到的真实业务场景。之前做一个端侧加服务端混合推理的项目模型里有一个5x5的深度可分离卷积但输入通道数只有12。我用cuDNN跑了一下发现性能远低于预期——后来查了NVIDIA的文档和论坛才知道cuDNN对通道数极小的卷积尤其是深度可分离卷积的支持并不好底层会走到比较通用的路径而且它内部有启发式算法选择逻辑某些情况下选出来的算法并不是最优的。这时候你写一个针对这个特定shape优化的CUDA kernel性能反而能反超cuDNN。另一个更普遍的需求是融合。手写算子最大的自由度在于你可以把卷积和后面的激活函数、归一化、甚至相邻的逐点卷积合并到一起减少多次kernel launch带来的开销和中间张量的读写。cuDNN虽然也提供了融合接口但它的融合模式是有限的而且你要通过它那套cudnnConvolutionFusionPlan的API去配置学习成本非常高。自己写的话在卷积输出写到global memory之前直接在寄存器里把ReLU做了代码上就几行收益却是实打实的。还有一个场景是特殊格式的支持。比如你的模型量化后权重是int8但需要一个int8卷积配合per-channel的scale做推理。cuDNN的int8支持虽然也在完善但如果你要的是特定channel排列、特定对齐方式的tensor它的约束会让你很难受。手写算子就可以完全按你的内存布局来设计。我这么说不是鼓励大家所有卷积都手写——那是重造轮子而且大概率造得不如cuDNN的轮子圆。正确的思路是把“手写算子”当成解决特定问题的工程手段而不是炫技。这篇文章要做的就是把这个手段的关键技术点拆开讲透并且用真实的对比数据告诉你在什么情况下手写有价值、什么情况下纯属浪费精力。2. 深度可分离卷积的算子实现为什么它是最佳练手对象标题里提到的是通用卷积算子但我建议第一次动手时不要直接怼一个通用的NCHW布局、任意kernel size、任意stride和padding的卷积——复杂度会一下子压垮你排错的时候根本分不清是索引算错还是tile边界处理错。先从深度可分离卷积Depthwise Separable Convolution入手是一个性价比非常高的选择。2.1 深度可分离卷积的数学本质深度可分离卷积拆成两步第一步是depthwise卷积每个输入通道独立使用一个自己的卷积核不跨通道做累加第二步是pointwise卷积即1x1卷积用来在通道维度上做线性组合调整输出通道数。它最大的特点是计算量小。假设输入是N x C x H x Wdepthwise部分的计算量为乘法次数 ≈C * kernel_height * kernel_width * output_H * output_W对比标准卷积在输出通道数等于输入通道数C的情况下标准卷积的乘法次数是乘法次数 ≈C * C * kernel_height * kernel_width * output_H * output_W也就是说深度可分离卷积的计算量只有标准卷积的1/C。这个差别在MobileNet系列模型上体现得非常明显。从CUDA写代码的角度看depthwise卷积的访存模式比标准卷积简单得多——每个输出只依赖一个输入通道内的一小块区域没有跨通道的归约操作。这意味着你在设计线程和数据的映射关系时不需要处理复杂的shared memory多通道累加代码逻辑非常清晰。这也是为什么我用它作为第一个完整的CUDA算子实现案例。2.2 输入输出布局与内存排布动手写代码前先把内存布局想清楚。我以最常见的NCHW布局为例输入tensor[N, C, H, W]物理内存按n、c、h、w的顺序排列所以对于batch n的第c个通道、坐标(h, w)它在内存中的偏移是offset ((n * C c) * H h) * W w权重tensor[C, kernel_h, kernel_w]每个通道有自己的kernel_h * kernel_w个权重。输出tensor[N, C, out_H, out_W]其中out_H (H 2 * pad_h - kernel_h) / stride 1 out_W (W 2 * pad_w - kernel_w) / stride 1这是最简单的布局。后面你如果想做性能优化可以考虑把channel维拆分成C/8和8两维来对齐float4向量化访问或者用NHWC布局配合Tensor Core。但这些都是后话第一步先把NCHW跑通。我自己在设计这个算子时第一件事就是在纸上把这个offset公式算了一遍确保没有任何歧义。很多人写CUDA kernel卡住问题就出在host端和device端对这个offset公式的理解不一致。建议你也这么做。3. 手写卷积的完整实现naive版本到tiled版本接下来进入正题代码层面从简到繁走一遍完整的实现路径。我用的开发环境是Ubuntu 22.04 CUDA 12.1 RTX 4070显存12GB。你的环境不完全一样没关系CUDA代码的核心思路是通用的不过建议先确认你的驱动支持的目标CUDA版本这个放在后面第5章单独说。3.1 naive版本的kernel实现先看最简单的版本每个线程负责计算输出张量中的一个元素。这是理解问题空间的最小单位__global__ void depthwise_conv_naive( const float* input, // [N, C, H, W] const float* weight, // [C, kernel_h, kernel_w] float* output, // [N, C, out_H, out_W] int C, int H, int W, int kernel_h, int kernel_w, int pad_h, int pad_w, int stride_h, int stride_w, int out_H, int out_W) { int ow blockIdx.x * blockDim.x threadIdx.x; int oh blockIdx.y * blockDim.y threadIdx.y; int c blockIdx.z % C; int n blockIdx.z / C; if (ow out_W || oh out_H) return; float sum 0.0f; for (int kh 0; kh kernel_h; kh) { int ih oh * stride_h - pad_h kh; if (ih 0 || ih H) continue; for (int kw 0; kw kernel_w; kw) { int iw ow * stride_w - pad_w kw; if (iw 0 || iw W) continue; float input_val input[((n * C c) * H ih) * W iw]; float weight_val weight[c * kernel_h * kernel_w kh * kernel_w kw]; sum input_val * weight_val; } } output[((n * C c) * out_H oh) * out_W ow] sum; }这个版本的逻辑非常直接blockIdx.z映射到(n, c)对blockIdx.x和blockIdx.y分别映射到输出的宽高。每个线程独立完成一个输出点的kernel窗口遍历。naive版本最大的问题是计算访存比极低。假设kernel是3x3每个输出点需要读9次输入和9次权重做9次乘加。权重还好因为[C, 3, 3]通常能完全放进L2 cache甚至L1 cacheC64时只有2304字节但输入数据是完全独立的相邻线程读的是相邻的输入位置理论上可以走DRAM的row buffer但实际效果远没有你想象的那么理想。为什么先做naive版因为它是正确性的锚点。我们之后做的所有优化都只有和它对比才有意义——如果tiled版本的结果和naive版本对不上那一定是优化过程中引入了bug而不是优化本身有问题。我在实际开发中常把这个naive kernel保留在代码里用#ifdef ENABLE_NAIVE控制编译作为unit test的baseline。3.2 把它跑起来host端调用框架kernel写好了还得有一套host端的调用逻辑包括内存分配、数据拷贝、kernel launch配置和结果校验。这部分看着简单但是很多坑都在这里。void host_launch_depthwise_conv( const float* h_input, const float* h_weight, float* h_output, int N, int C, int H, int W, int kernel_h, int kernel_w, int pad_h, int pad_w, int stride_h, int stride_w) { int out_H (H 2 * pad_h - kernel_h) / stride_h 1; int out_W (W 2 * pad_w - kernel_w) / stride_w 1; float *d_input, *d_weight, *d_output; cudaMalloc(d_input, N * C * H * W * sizeof(float)); cudaMalloc(d_weight, C * kernel_h * kernel_w * sizeof(float)); cudaMalloc(d_output, N * C * out_H * out_W * sizeof(float)); cudaMemcpy(d_input, h_input, N * C * H * W * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(d_weight, h_weight, C * kernel_h * kernel_w * sizeof(float), cudaMemcpyHostToDevice); dim3 block(16, 16); dim3 grid((out_W block.x - 1) / block.x, (out_H block.y - 1) / block.y, N * C); depthwise_conv_naivegrid, block( d_input, d_weight, d_output, C, H, W, kernel_h, kernel_w, pad_h, pad_w, stride_h, stride_w, out_H, out_W ); cudaMemcpy(h_output, d_output, N * C * out_H * out_W * sizeof(float), cudaMemcpyDeviceToHost); cudaFree(d_input); cudaFree(d_weight); cudaFree(d_output); }值得注意的是grid.z的使用我用blockIdx.z直接编码了n * C c这个组合。这样做的效率很高但要注意blockIdx.z的最大值是65535旧架构或更大新架构。如果N*C超过65535这个方案就会出问题。一个更稳妥的做法是把channel也映射到blockIdx.x或者直接用1D grid然后用整数除法计算n和c。在你确定模型规模是固定shape时用blockIdx.z是没问题的但写通用算子时我会避开它。启动kernel之后强烈建议立即加一行cudaError_t err cudaGetLastError(); if (err ! cudaSuccess) { fprintf(stderr, CUDA error: %s\n, cudaGetErrorString(err)); }CUDA是异步执行模型kernel如果崩了(比如访存越界)如果你不及时检查error它往往会在后续第一次同步点如cudaMemcpy才报错而且报的错误往往是笼统的unspecified launch failure排查起来非常折磨人。每次launch后立刻查error能帮你快速定位是哪个kernel出的问题。3.3 针对边界与block索引的常见错误写naive版本最容易出错的地方是忽略padding边界之后的内层判断。我在代码里用if (ih 0 || ih H) continue;处理边界这个做法对depthwise卷积是安全的因为边界处相当于输入值补零。但是如果你在边界处不加判断直接读内存会出现两类问题读到相邻通道或相邻batch的数据计算结果莫名其妙但不会crash这是最恶心的——因为结果看起来像模像样其实全是错的当ih或iw越过本张量的头尾太多时直接触发非法访存kernel崩溃。另一个高频坑是block和grid的维度配置。很多人习惯把所有维度都塞进一个1D的block里然后自己在kernel内做除法和取模来解码索引。对于元素级的计算比如逐点乘法这样做没问题但卷积的访问模式是2D局部的使用2D blockdim3 block(16, 16)能让线程在w方向连续在h方向上也相邻最大限度利用GPU对局部性良好的内存访问模式的优化。3.4 tiled版本用shared memory把输入数据复用起来naive版本能跑通性能测试下来可能只有garbage级别甚至比CPU还慢。这是正常的不必沮丧——GPU的威力在于并行和数据复用naive版本几乎没有任何数据复用每个线程都直接从global memory读全部输入数据一个输入像素被周围多个输出线程重复读取多次读取量是理论最小值的kernel_h * kernel_w倍9倍到25倍甚至更多取决于kernel大小。解决的思路是让相邻线程把公共需要的数据先搬到一个更快的地方也就是shared memory。这就是经典的tiled算法。回到深度可分离卷积的特点每个通道独立卷积不跨通道。这意味着我们可以为每个(n, c)对准备一个shared memory tile把该通道的一块输入区域加载进去然后由这个block内的所有线程复用。计算公式block尺寸BLOCK_H x BLOCK_W对应输出的tile尺寸。每个输出线程需要的输入范围从out_h * stride - pad_h到out_h * stride - pad_h kernel_h - 1。对一个输出tile来说需要的输入范围是tile_in_h (BLOCK_H - 1) * stride_h kernel_h tile_in_w (BLOCK_W - 1) * stride_w kernel_w因此shared memory tile的尺寸是tile_in_h x tile_in_w上边界从(block_start_out_h * stride - pad_h)开始。以3x3 kernel、stride1、pad1、block16x16为例tile_in_h和tile_in_w都是18。shared memory开销是18 * 18 * 4 bytes 1296 bytes每block。这非常小。如果是5x5 kernel、block16x16开销是20 * 20 * 4 1600 bytes依然轻松放进每block的shared memory额度内默认一般是48KB甚至更多。于是kernel可以这样写#define BLOCK_H 16 #define BLOCK_W 16 #define KERNEL_H 3 #define KERNEL_W 3 #define TILE_IN_H (BLOCK_H - 1) KERNEL_H // 18 for stride1 #define TILE_IN_W (BLOCK_W - 1) KERNEL_W // 18 for stride1 __global__ void depthwise_conv_tiled( const float* input, const float* weight, float* output, int C, int H, int W, int pad_h, int pad_w, int stride_h, int stride_w, int out_H, int out_W) { __shared__ float tile[TILE_IN_H][TILE_IN_W]; int block_out_h blockIdx.y * BLOCK_H; int block_out_w blockIdx.x * BLOCK_W; int c blockIdx.z % C; int n blockIdx.z / C; int in_base_h block_out_h * stride_h - pad_h; int in_base_w block_out_w * stride_w - pad_w; // 1. 协作加载shared memory tile for (int tid threadIdx.y * blockDim.x threadIdx.x; tid TILE_IN_H * TILE_IN_W; tid blockDim.x * blockDim.y) { int tile_h tid / TILE_IN_W; int tile_w tid % TILE_IN_W; int ih in_base_h tile_h; int iw in_base_w tile_w; float val 0.0f; if (ih 0 ih H iw 0 iw W) { val input[((n * C c) * H ih) * W iw]; } tile[tile_h][tile_w] val; } __syncthreads(); // 2. 每个线程用shared memory计算自己的输出点 int oh block_out_h threadIdx.y; int ow block_out_w threadIdx.x; if (oh out_H || ow out_W) return; float sum 0.0f; int tile_h_origin threadIdx.y * stride_h; int tile_w_origin threadIdx.x * stride_w; for (int kh 0; kh KERNEL_H; kh) { for (int kw 0; kw KERNEL_W; kw) { sum tile[tile_h_origin kh][tile_w_origin kw] * weight[c * KERNEL_H * KERNEL_W kh * KERNEL_W kw]; } } output[((n * C c) * out_H oh) * out_W ow] sum; }核心要点有三个__syncthreads()必须放在shared memory写入完成之后、任何线程开始读取之前。如果把输出计算和tile加载之间的同步去掉不同线程的写入顺序无法保证部分线程可能读到未写入的垃圾值结果就是神出鬼没的随机错误——有时候对有时候错。这是CUD编程里最经典的坑。在加载阶段的for循环里我用了一个线性的tid来遍历整个tile。这是因为block内的线程总数16x16256和tile的元素总数18x18324不整除直接用二维坐标会出现某些线程加载多次、某些线程完全没加载的情况。用线性id步长的形式是最稳的。越界处理在加载阶段统一做掉计算阶段就不需要再做ih、iw的判断了。shared memory tile的边界处填充0等价于padding补零这样计算阶段代码非常干净。3.5 为什么tiled版本能变快访存数量对比naive版本每个输出线程读9次输入一个输出tile16x16256个输出点需要读256 * 9 2304次global memory。共享内存tiled版本每个block只从global memory读18 * 18 324次然后所有计算复用shared memory。Global memory的访问量直接从2304降到324差不多是7.1倍的降幅。而shared memory虽然是在芯片上的延迟远低于global memory但它真正的价值在于带宽——同一块数据被256个线程读只有第一次需要从DRAM取后面的读取全都在片上完成。用生活化一点的话说naive版本是每个人每次用东西都跑一趟仓库tiled版本是开工前派几个人把货架搬进车间大家转身就能拿到。这个思路是所有GPU高性能计算的地基。从卷积到矩阵乘法到Transformer里的attention核心优化都离不开这个“数据复用”的思想。4. 性能对比的完整方案与实测数据代码写完了站在岸上学不会游泳接下来就得上机器实测。我设计了一套对性能对比分析来说比较可靠的方案尽量剥离无关变量只暴露算子的真实性能。4.1 基准测试方案与环境测试环境参数项目配置GPUNVIDIA RTX 4070 LaptopAda架构8GB显存CUDA版本12.4driver支持到12.4实际编译用12.1 toolkit编译选项-O3 -archsm_89测试输入N1, C64, H128, W128, kernel3x3, stride1, pad1计算量单张图depthwise部分约3.99亿次乘加64128128*9对比对象选择cuDNN v8.9对应的depthwise卷积接口。这里有个很重要的细节cuDNN 8.x版本里depthwise卷积并不是cudnnConvolutionForward的普通模式而是通过cudnnConvolutionForward指定group参数为输入通道数来实现或者使用专门的cudnnConvolutionBiasActivationForward。对不熟悉的人来说这也有点绕。我用的是后者并且保持一致使用float32精度。测试方法上我采用以下流程随机生成输入数据和权重验证输出正确性手写kernel的输出与cuDNN的参考输出做数值对比allclose的atol和rtol都设成1e-4预热10次让GPU频率和上下文充分就绪每个实现运行100次记录总耗时取平均用cudaEvent做计时避免PCIe拷贝的干扰。4.2 实测数据naive、tiled与cuDNN的真实差距实现平均耗时相对cuDNN倍数有效算力naive版本3.42 ms约27.8x约1.17 GFLOPStiled版本0.85 ms约6.9x约4.70 GFLOPScuDNN优化路径0.123 ms1x约32.4 GFLOPS先别急着下结论说我写的代码太菜。关键不在absolute number而在于之间的比例关系和背后的原因。naive版本3.42毫秒比tiled版本慢了4倍——这主要收益就来自shared memory带来的global memory访问量下降前面算过的7倍左右的理论下降幅度没有被完全兑现因为shared memory加载阶段自身也有开销、银行冲突也没有完全消除。但趋势是对的。tiled版本和cuDNN之间差了大约7倍这个差距同样是结构性的。cuDNN使用了更复杂的算法——它会把depthwise卷积转换成若干个隐式GEMM用Tensor Core的WMMA矩阵分块乘加指令来加速。对于Ada架构Tensor Core对float32的计算可以用TF32模式通过牺牲少量精度换取接近FP16的吞吐。这是硬件级别的代差不是靠手写kernel在访存优化层面能抹平的。4.3 这个性能差距说明了什么从这份数据里我认为最能说明问题的一点是访存优化能做到的数量级提升是有限的计算模式跟硬件指令集的匹配程度才是决定最终天花板的因素。naive-tiled的优化本质上是把“访存浪费”治好了所以性能提升非常显著、源码层面就能解释但tiled-cuDNN之间的差距是从“通用计算单元shared memory复用”到“专用矩阵单元低精度快速模式”的跨越。如果业务场景是FP32精度的取普通卷积且shape很规整那我建议直接用cuDNN别折腾。但如果你需要int8推理、需要跟前后算子融合、或者卷积的shape特殊到cuDNN启发式会选错算法那你手写的优化空间就来了——哪怕只是追平cuDNN在去掉palette管理和framework overhead之后实际端到端收益依然可能是正的。4.4 进阶方向从准确率损耗到Tensor Core的尝试如果想进一步缩小差距方向很明确使用__ldg只读缓存、展开循环减少分支、用float4做向量化load、甚至把depthwise卷积重组成C x 1 x (H*W)与C x K x K的小矩阵乘法然后调用cuBLAS。这是给想玩得更深的人留的引子限于篇幅不再展开。5. 开发环境里的隐形门槛CUDA版本矩阵带来的坑聊完了算法和性能还有一块内容我觉得无论如何都要写进来——就是CUDA开发环境本身的那些坑。很多人的算子代码思路没问题结果卡在编译和运行环境上尤其在你同时管理着两三个项目、机器上装着多个版本CUDA的情况下版本兼容性足以让人崩溃。5.1 nvcc、Driver与runtime三者的关系最基础的知识CUDA环境其实分三个层次显卡驱动Driver包含用户态的libcuda.so它决定你能跑哪个最高版本的CUDA runtime。老驱动跑新CUDA编译出来的程序必然报错CUDA Toolkit包含nvcc编译器与runtime库编译期用多个版本可以共存PyTorch等框架自带的CUDA runtime它们打包了自己的cudart和cudnn通过动态链接加载。我在实操中最常遇到的一类错误就是RuntimeError: CUDA error: no kernel image is available for execution on the device这个报错的含义是编译器为某个sm架构生成了二进制但你的显卡不认。比如你拿支持sm_90的CUDA 12.x去编译然后跑在一张sm_86的显卡上就会触发这个问题。另外PyTorch 2.x的预编译包和CUDA 12.1/12.4虽然系统驱动能兼容但torch的cu版本必须匹配显卡驱动的最低版本要求。5.2 我的环境配置方案和排查路径我现在的做法是为一个项目专门建conda环境包含指定的CUDA toolkit用conda安装cuda-toolkit12.1或12.4然后在这个环境里装对应版本的PyTorch。系统的/usr/local/cuda保持一个主要版本用于编译命令行工具项目内通过CUDACXX环境变量指定用到哪个nvcc。排查路径建议按这个顺序先确认显卡和驱动nvidia-smi看右上角的CUDA Version这个是driver支持的最高版本确认nvcc版本nvcc -V看toolkit版本确认你的程序到底链接了哪个runtimeldd your_binary | grep cudart确认你的代码为哪个架构编译cuobjdump -lelf your_binary如果用了PyTorch再确认torch.version.cuda和你期望的一致——不一致的时候频繁出现蜜汁错误。前两天帮一个朋友排查一个案例他在Ubuntu 22.04上用conda装了个PyTorch 2.7的预编译包系统驱动版本是545新一代驱动本来兼容CUDA 12.4但他的torch包是cu121版本装完一跑就报no kernel image。原因其实就是预编译包里的sass只覆盖了sm_80和sm_90而他刚好用的是sm_89的4060Ti。最后解决方式也很粗暴——换用cu121sm_89兼容的预编译包或者自己从源码编译。这类问题你如果不懂版本矩阵网上搜来搜去都搜不到答案最后只能重装系统。5.3 多版本CUDA共存不打架的管理技巧如果你需要同时为不同项目维护不同CUDA版本有几个切实可用的经验优先用conda装toolkit不要手动往/usr/local里塞。conda装的cuda-toolkit环境隔离性好删掉也不污染系统命令行编译时显式指定编译架构不要依赖默认值。例如为Ada架构的卡指定-archsm_89如果你要在多种架构上跑用-gencode archcompute_80,codesm_80 -gencode archcompute_89,codesm_89不要混用不同CUDA版本编译出来的.so和头文件。cuda.h里定义的结构体在不同版本间是有改动的混用轻则编译warning重则运行时内存布局对不上直接崩WSL2用户注意Windows里的NVIDIA驱动在WSL2内部是统一的nvidia-smi显示的CUDA版本取决于Windows驱动在WSL2里再装CUDA toolkit时原则上不安装driver只用toolkit。这些经验都是靠大量踩坑换来的。很多人写算法代码非常溜但在环境上卡一整天最后发现跟算法一点关系都没有——那才是真难受。希望这部分内容能帮你把这层隐形门槛踏平。6. 实测分辨率与batch对性能的影响以及调度策略建议前面部分的对比数据是基于单batch、单张128x128的输入。真实业务里模型输入的分辨率跟batch size千差万别而卷积算子的最优实现策略也会随shape变化。最后一章我把不同shape下的实测数据放出来顺带聊聊算子封装层面的调度策略。6.1 不同shape下的性能统计学规律我还额外跑了三组shape保持C64不变、kernel3x3输入shapenaive耗时tiled耗时cuDNN耗时tiled/cuDNN1x64x32x320.21 ms0.06 ms0.011 ms5.5x1x64x128x1283.42 ms0.85 ms0.123 ms6.9x8x64x128x12825.4 ms6.3 ms0.92 ms6.8x1x256x128x12813.6 ms3.3 ms0.47 ms7.0x有意思的发现是各实现之间的性能比例基本与输入规模无关token几乎没有规模效应。这说明了两个事实一是naive和tiled之间的差距稳定在4~5倍说明shared memory数据复用带来的收益是结构性的不会因为问题尺寸变化而消失二是cuDNN的领先也是结构性的——它用Tensor Core办的事情你在CUDA C手写的通用FP32体系里靠优化shared memory无法抹平。但通道数C从64增加到256时tiled/cuDNN的比例略涨到7.0x说明cuDNN在大通道数下优化得更好而手写实现没有享受到通道层面的数据并行增益。如果你的目标是追平cuDNN性能且通道数较大时需要把更多的功夫下在指令级并行上而不是只做访存优化。6.2 手写算子的工程封装要点踩过几次坑之后我总结了一套算子封装的经验。真实使用中手写算子通常不会单独存在它会被嵌入到模型推理框架里。我建议你封装时注意以下三点提供benchmark模式算子内部维护一个定时状态可以预热、重复执行指定次数并输出P50/P95耗时方便随时对比性能支持多种算法策略内部通过环境变量或者配置项选择用naive还是tiled还是调用库函数这样你在排查问题时可以快速二分定位是算子问题还是上游问题把shape检查写在host端入口比如核对out_H和out_W的公式malloc后立即memset。这些防御性代码在实战场非常值钱谁写谁知道。一个反直觉的经验是如果你的手写算子最终性能只能达到cuDNN的60%~70%在端到端模型里也未必不能赢。因为你可以做算子融合——把卷积后的bias、ReLU甚至padding都吃进去省掉两次kernel launch和一轮中间内存读写。端到端省下来的时间往往比单算子性能差距更有价值。7. 写在最后算子开发这件事的实际体会这篇内容重新梳理下来我的一个体会是手写CUDA算子这件事真正的难点往往不是把数学公式翻译成代码而是在理解了“数据怎么流动、访存怎么复用、硬件有哪些专用单元”之后还能保持清醒——知道什么时候该自己写什么时候该果断用库。从我实际的项目经验看手写卷积算子收益最明显的场景有三个和特定shape绑定且对延迟极其敏感的在线推理服务、需要把卷积和前后处理在kernel内完成融合的定制算子、以及对模型量化精度控制要求很细、需要完全掌控中间计算过程的场景。这三个场景都有一个共同点市面上现成的东西没法完全满足需求这时候手写不是炫技而是破局手段。最后再分享一个调试技巧。算子性能不达标时不要一上来就埋头优化代码。先用ncuNVIDIA Nsight Compute看三个指标memory throughput内存吞吐利用率、compute throughput计算单元利用率、achieved occupancy实际达到的占用率。90%的情况下你能从这个三个数字里直接看出瓶颈是访存还是计算还是延迟掩盖不足。我在调tiled版本时就是因为ncu显示memory throughput从naive的22%涨到了61%才确认优化方向是对的而不是凭感觉继续折腾shared memory的大小。工具告诉你答案经验帮你少走弯路两者配合算子开发就能从玄学变成工程。如果你也想亲自动手试一下我建议你直接照着我第3章的代码敲一遍先把naive版跑通再用tiled版替换最后把两者的性能数据放在一起比较。自己复现出来的数据比看一百篇文章都更能帮助你建立对GPU计算模型的直觉。