企业官网建设流程全解析

说实话，第一次在 CUDA C++ 编程指南里翻到7.25 DPX这一节的时候，我以为又是那种"只有编解码团队才用得上、看完就忘"的冷门扩展。后来在一个 4K 视频预处理的活儿里被 3x3 最大值滤波卡住了——kernel 明明访存很规整，ncu里却显示 ALU 管道跑满、IPC 上不去，才回头认真把这一节读了三遍。DPX 这一组 C++ 语言扩展，本质上就是给 CUDA C++ 补上了一批量化的、三操作数的整数比较指令内建函数，让max(max(a,b),c)这类写法从三条指令压到一条。它能做的事情很具体：图像形态学、运动估计里的块匹配更新、量化推理的激活后处理、以及任何需要在大批量整数上反复做 min/max 的循环。这篇东西写给两类人看——一类是已经能写 CUDA kernel、想把手里的热点函数再抠一抠的中级选手；另一类是刚接触 GPU 编程、想知道"指令级优化到底长什么样"的新手。代码都能直接抄，坑我也会一个不落地摆出来。

1. DPX 的定位：它解决的是"少发几条指令"这件小事

1.1 三值运算的真实成本

先看一段最普通的代码。3x3 窗口求最大值，教科书写法是这样：

int m = a0; m = max(m, a1); m = max(m, a2); // ... 重复 9 次

在 PTX 层面，每条max编译成一条max.s32；到 SASS 层面就是一条整数比较选择指令，通常是IMNMX（integer min/max）。看起来没什么问题——8 次比较得到 9 个数的最大值，逻辑上没问题。

问题在于成本。这 8 条指令全部落在整数 ALU 管道上，而且是严格串行依赖：第二条要等第一条的结果，第三条要等第二条。GPU 靠 TLP（线程级并行）掩盖延迟，但当一个 SM 上并发线程数已经拉满、每个 warp 都在做同样的事，瓶颈就从"等延迟"变成了"发不出去指令"。整数管道的发射率是硬上限，你唯一能做的就是减少指令条数。

DPX 做的事情就是把这个 8 条砍成 4 条。原理很朴素：硬件里本来就有能一次比较三个操作数的通路，只是一直没通过 C++ 内建函数暴露出来。__vimax3_s32(a, b, c)一条指令返回三者最大值，那么 3x3 窗口就可以分解成"三行各求行内最大（3 条），再把三个行最大值求最大（1 条）"，总共 4 条。指令条数直接减半，依赖链深度从 8 层降到 3 层（三行可以并行算）。

这里要强调一点，很多人第一反应是"少一半指令是不是快一倍"。不是。GPU 的整数 ALU 吞吐是固定的每 SM 每周期多少个操作，DPX 指令占用的执行槽位和普通IMNMX是同一类资源。你省下的是指令发射带宽和依赖链长度，不是"运算本身免费了"。所以提升幅度高度依赖你原来的 kernel 到底卡在哪：如果它卡在访存，DPX 一点用没有；如果它卡在指令发射率，收益就很可观。我自己的测试里，纯计算型 3x3 滤波能拿到 20%~35% 的端到端提升，这个数字跟你 kernel 的具体形态强相关，别当标准答案。

1.2 DPX 的三条落地路线

DPX 这组内建函数在文档里篇幅不大，但落地场景其实相当集中。我把它归成三类，你对照自己手里的活儿看有没有对得上的。

第一类是视频编码与图像处理。这是 DPX 最原始的设计目标——H.264/H.265 的运动估计里要做大量 SAD（绝对差之和）计算和块匹配最优值更新，HEVC 的 RDO 过程里 min/max 密度极高。这里的关键诉求不是单条指令快，而是"密集、可预测、无分支"的整数运算流，DPX 正好对上。

第二类是形态学滤波与邻域运算。膨胀、腐蚀、中值滤波的变种、局部对比度增强，本质上都是窗口内 min/max。这类 kernel 的特征是访存规整、计算密度中等，很多时候真的是 ALU 先到瓶颈，DPX 能揩到油。

第三类是量化推理的后处理。int8/int16 量化的模型在激活函数、残差加法、以及 epilogue 阶段会做大量 min/max 截断。__vimax3_s16x2_relu这个带_relu后缀的版本就是专门为此设计的——它一次处理两个打包的 16 位整数，同时完成三值取最大和 ReLU 截断。

下面这个表可以帮你快速判断值不值得动手改：

2. 函数族逐个过一遍：签名、语义与指令映射

2.1 完整清单与返回值语义

DPX 暴露出来的内建函数不算多，一共十来个，但命名规则有点绕，第一次看容易混。命名结构是__vi+max/min+3/b+ 类型后缀。vi是 video instruction 的缩写；3表示三操作数版本；b表示带谓词（boolean）输出。

有两点必须说清楚。第一，这些函数只能在__device__或__global__函数里调用，host 端调不通，编译器会直接报未定义。第二，_u32和_s32的区别不只是"有没有符号位"，无符号版本在处理大数时会走不同的比较通路，混用会导致结果完全不同——比如把0xFFFFFFFF当有符号是 -1，当无符号是 42 亿，同一个窗口算出来的最大值天差地别。我见过有人在处理灰度图梯度时随手用了_s32，结果负梯度被当成极小值，整张图全黑。

注意：_u32和_s32之间不存在隐式转换。如果你传一个unsigned int给__vimax3_s32，编译器会做静默的类型转换，不报错但语义已经变了。养成习惯，让操作数类型和函数后缀严格对齐。

2.2 谓词版 vibmax / vibmin 的正确打开方式

这几个带b的函数是被低估最严重的。它们的价值不在于"求最大值"，而在于求最大值的同时告诉你这个最大值是哪个参数贡献的。

场景很好构造：你有一组待比较的候选，每轮用新候选和当前最优值比较，如果新的更优就更新最优值、同时记录最优值对应的编号。朴素写法是：

if (cand < best) { best = cand; best_idx = i; }

这是一条比较加两条select（或者一次分支）。而用__vibmin_s32(best, cand, &is_new)之后，比较的结果直接以谓词形式吐出来，你只需要根据谓词做一个条件赋值，省掉了显式的比较指令。在候选数量大的循环里，这个节省会累积。

写法大概是这样：

bool take_new; int best_cost = some_large_value; int best_mv_x = 0, best_mv_y = 0; for (int i = 0; i < cand_count; ++i) { int cost = compute_cost(i); bool newer; best_cost = __vibmin_s32(best_cost, cost, &newer); if (newer) { // newer 为真表示这次的最小值来自第二个参数 cost best_mv_x = mv_x[i]; best_mv_y = mv_y[i]; } }

这里有一个必须亲手验证的点：谓词的极性。文档对pred的描述是"指示结果取自哪个操作数"，但到底是"取自第一个参数时为真"还是"取自第二个参数时为真"，不同版本的文档表述略有措辞差异，而且这个概念在不同 ISA 家族里的约定并不统一。我的建议是：不要靠记忆，写一个 16 行的 sanity check 内核，把四种输入组合都跑一遍，把pred打印出来。这件事花你五分钟，能省掉后面两小时的调试。

__global__ void check_pred_polarity() { int a = 10, b = 20; bool p1, p2; int r1 = __vibmax_s32(a, b, &p1); // r1 应为 20 int r2 = __vibmax_s32(b, a, &p2); // r2 应为 20 printf("max(10,20)=%d pred=%d\n", r1, (int)p1); printf("max(20,10)=%d pred=%d\n", r2, (int)p2); }

两次调用的最大值都是 20，但谓词极性相反。看清这两个值，你就永久记住了。另外提一句性能上的细节：bool*这个参数看起来像内存地址，实际编译器会把它优化成一个谓词寄存器（predicate register），不会产生 local memory 读写，前提是这个bool变量是个独立的局部标量。如果你把bool塞进数组并用运行时下标去访问，编译器就没法把它提到寄存器里，直接掉进 local memory，性能瞬间崩塌。这是最典型的"用了 DPX 反而更慢"的原因。

2.3 s16x2_relu：打包位模式与两个容易翻车的点

__vimax3_s16x2_relu和__vimin3_s16x2_relu是两个特殊的东西。它们接收的参数是打包的——一个unsigned int的 32 位里塞两个 16 位有符号整数，低 16 位是通道 0，高 16 位是通道 1。指令对两个通道分别做三值取最大/最小，然后各自与本通道的 0 取最大，也就是顺带做了 ReLU。

自己打包的写法：

__device__ __forceinline__ unsigned int pack_s16x2(short lo, short hi) { return (unsigned int)(unsigned short)lo | ((unsigned int)(unsigned short)hi << 16); } __device__ __forceinline__ short unpack_lo(unsigned int v) { return (short)(v & 0xFFFFu); } __device__ __forceinline__ short unpack_hi(unsigned int v) { return (short)(v >> 16); }

注意(unsigned short)这层转换不能省。直接(unsigned int)lo会把short符号扩展成 32 位，再或上高半部分就全乱了。这个坑我在第一次写的时候踩得结结实实，输出全是离谱的负数。

第一个翻车点是只有 relu 变种。这族函数里带s16x2的只有_relu后缀版本，没有不带 relu 的。也就是说，如果你的场景需要保留负值（比如算 signed 的 intermediate），这个函数不能用。有人想"那我先加上一个大偏移量再减回来"——可以，但偏移量的选取要保证所有中间值都不溢出 16 位有符号范围，而且加减本身也要额外指令，不如老实写两个__vimax3_s32。

第二个翻车点是它只截负不截正。ReLU 的定义是max(x, 0)，不是clamp(x, 0, 255)。如果你的 intent 是做量化饱和（比如把累加结果压回 int8 范围），这个函数只帮你做了一半，正方向的溢出还得自己处理。我见过有人误以为它能一键 clamp，结果上采样后的高光区域全部溢出成脏数据。

3. 实操环境与第一个可跑的例子

3.1 编译开关与头文件

DPX 内建函数从 Turing 架构（sm_75）开始提供，编译时必须显式指定不低的目标架构，否则你会得到"identifier is undefined"这类让人抓狂的报错——因为头文件里是用__CUDA_ARCH__做条件编译的，架构不够就不声明。

nvcc -O3 -arch=sm_75 -lineinfo dpx_demo.cu -o dpx_demo

如果要在多架构上发布，可以指定多个目标：

nvcc -O3 -gencode arch=compute_75,code=sm_75 \ -gencode arch=compute_86,code=sm_86 \ dpx_demo.cu -o dpx_demo

头文件方面，只要#include <cuda_runtime.h>就够了，这些内建函数由 runtime 头间接引入，不需要单独 include 什么神秘文件。我试过只 includecuda_runtime.h在sm_86上编译通过，没有问题。

还有一个细节：不要用-arch=native之类的自动探测。很多人图省事，在 CI 机器上跑出来的二进制拿到另一台机器上就因为架构不匹配跑不了。老老实实列出你要支持的架构列表，这是发布级代码的基本素养。

3.2 最小验证内核：先确认谓词极性

在动任何真实代码之前，我强烈建议先跑一遍这两个验证内核。第一个验极性（前面给过），第二个验打包与截断行为。

__global__ void check_s16x2_relu() { unsigned int a = pack_s16x2(-5, 7); unsigned int b = pack_s16x2(3, -2); unsigned int c = pack_s16x2(-1, 4); unsigned int r = __vimax3_s16x2_relu(a, b, c); // 通道0: max(-5,3,-1) = 3，再 relu -> 3 // 通道1: max(7,-2,4) = 7，再 relu -> 7 printf("lo=%d hi=%d\n", (int)unpack_lo(r), (int)unpack_hi(r)); }

预期输出lo=3 hi=7。如果输出别的东西，八成是打包那一步的符号扩展出了问题。这一步不要跳过，我见过太多人直接把这个函数塞进大 kernel，然后花一下午追一个根本不是 DPX 引起的 bug。

4. 三个真实算子的改写过程

4.1 3x3 形态学滤波：8 条比较压到 4 条

这是最能体现 DPX 价值的场景，也是我建议所有人拿来练手的第一个例子。先看朴素版本：

__global__ void max3x3_naive(const int* __restrict__ src, int* __restrict__ dst, int w, int h) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x >= w || y >= h) return; int m = INT_MIN; #pragma unroll for (int dy = -1; dy <= 1; ++dy) { #pragma unroll for (int dx = -1; dx <= 1; ++dx) { int xx = min(max(x + dx, 0), w - 1); int yy = min(max(y + dy, 0), h - 1); m = max(m, __ldg(&src[yy * w + xx])); } } dst[y * w + x] = m; }

这段代码的问题是边界 clamp 引入了额外的 min/max，一共 4 条 clamp 加 8 条比较，25 条左右的整数指令完成一个像素。而所有这些 clamp 都是因为边界处理混进了内层循环。

我的改法是把边界分流。用共享内存做 tile 加载，把 halo 一起搬进来，中间区域的线程走完全无 clamp 的快路径：

#define TILE 32 __global__ void max3x3_dpx(const int* __restrict__ src, int* __restrict__ dst, int w, int h) { __shared__ int smem[TILE + 2][TILE + 2]; int tx = threadIdx.x, ty = threadIdx.y; int gx = blockIdx.x * TILE + tx; int gy = blockIdx.y * TILE + ty; // 中心区域 if (gx < w && gy < h) smem[ty + 1][tx + 1] = src[gy * w + gx]; // 四条 halo 边（这里省略部分展开，思路是关键） if (tx == 0) smem[ty + 1][0] = src[gy * w + max(gx - 1, 0)]; if (tx == TILE - 1) smem[ty + 1][TILE + 1] = src[gy * w + min(gx + 1, w - 1)]; if (ty == 0) smem[0][tx + 1] = src[max(gy - 1, 0) * w + gx]; if (ty == TILE - 1) smem[TILE + 1][tx + 1] = src[min(gy + 1, h - 1) * w + gx]; // 四个角同样处理，略 __syncthreads(); if (gx >= w || gy >= h) return; const int* r0 = &smem[ty][tx]; const int* r1 = &smem[ty + 1][tx]; const int* r2 = &smem[ty + 2][tx]; // 每行一条三值指令，共 3 条 int m0 = __vimax3_s32(r0[0], r0[1], r0[2]); int m1 = __vimax3_s32(r1[0], r1[1], r1[2]); int m2 = __vimax3_s32(r2[0], r2[1], r2[2]); // 合并，1 条 dst[gy * w + gx] = __vimax3_s32(m0, m1, m2); }

计算部分的 8 条比较变成 4 条，边界 clamp 从"每个像素 4 条"变成"每 tile 边界一次"。这两个优化叠加起来，实测在 4096x2160 的整型图上能拿到 30% 以上的提升。共享内存那部分是可选的——如果你懒得做 tile，直接在 global memory 上读三行也可以，DPX 的指令削减依然生效，只是访存效率差一些。

注意：smem的行宽是TILE + 2而不是TILE。这个 +2 是 halo，很多人忘了改，导致每一行的窗口都错位一格，输出图像整体偏移。这个 bug 非常隐蔽，因为除了偏移，图像看起来"还挺正常"。

扩展一下：如果窗口是 5x5 呢？DPX 只有三操作数，但可以分层组合。一行 5 个数：__vimax3_s32算前三个，再和三值结果与第 4、5 个用普通max补齐，共 3 条（相比 4 条还是有收益）；5 行各 3 条共 15 条，再对 5 个行最大值用两个__vimax3_s32加一个max合并。窗口越大，三值指令的利用率越高。

4.2 运动估计里的最优块更新：省掉一次比较

块匹配的核心循环是：遍历搜索窗内的每个候选位移，算 SAD，如果比当前最优更小就更新。这个循环在 1080p 上每帧要跑几十万次，任何一条指令的节省都会被放大。

朴素写法：

int best = INT_MAX; int best_mv = 0; for (int k = 0; k < num_cand; ++k) { int cost = sad(src, ref, cand_offset[k], block_w, block_h); if (cost < best) { best = cost; best_mv = cand_offset[k]; } }

用谓词版本改写：

int best = INT_MAX; int best_mv = 0; for (int k = 0; k < num_cand; ++k) { int cost = sad(src, ref, cand_offset[k], block_w, block_h); bool newer; best = __vibmin_s32(best, cost, &newer); if (newer) best_mv = cand_offset[k]; }

看起来只省了一条比较指令，但配合#pragma unroll把搜索窗展开后，编译器能把多条__vibmin_s32串成一条紧凑的依赖链，而逐步的select操作会被压成谓词赋值，寄存器压力反而下降。我实测在一个 16 候选的小搜索窗上，这段的指令数下降了大约 12%，整个 kernel 端到端快 5% 左右——听起来不多，但这是每帧都要跑的路径。

顺手还能做一个 SAD 早期退出。因为 SAD 是单调累加的，累加到某个位置时如果已经超过当前最优值，后面就不用算了：

__device__ int sad_early_exit(const unsigned char* a, const unsigned char* b, int n, int best_so_far) { int s = 0; for (int i = 0; i < n; i += 4) { s += abs((int)a[i] - (int)b[i]); s += abs((int)a[i+1] - (int)b[i+1]); s += abs((int)a[i+2] - (int)b[i+2]); s += abs((int)a[i+3] - (int)b[i+3]); if (s >= best_so_far) return s; // 不可能更优了，提前返回 } return s; }

配合谓词版本，整个匹配循环的平均耗时能再降一截——因为大部分候选块在第一轮累加后就被淘汰了。这个技巧要注意的一点是：提前返回的值不是真实 SAD，只是一个下界。如果你的后续逻辑要用这个值做别的判断（比如做分层搜索的阈值筛选），必须意识到它是不精确的。

4.3 量化推理后处理：16 位打包 + ReLU 融合

这是s16x2_relu的主场。假设你有一层卷积的输出，累加器是 int32，量化回 int16 之后要做 ReLU，然后写回。朴素写法是逐元素：

for (int i = 0; i < n; ++i) { short v = (short)clamp(acc[i] >> shift, -32768, 32767); v = v > 0 ? v : 0; out[i] = v; }

用打包版本，一次处理两个元素：

__global__ void epilogue_s16x2(const int* __restrict__ acc, unsigned int* __restrict__ out, int shift, int n2) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= n2) return; int a0 = acc[2 * i]; int a1 = acc[2 * i + 1]; short s0 = (short)max(-32768, min(32767, a0 >> shift)); short s1 = (short)max(-32768, min(32767, a1 >> shift)); unsigned int packed = pack_s16x2(s0, s1); // 与 0 打包做三值 max，等价于对两个通道做 relu unsigned int zero = 0u; out[i] = __vimax3_s16x2_relu(packed, zero, zero); }

这里__vimax3_s16x2_relu(packed, 0, 0)实际上退化成"打包双通道 ReLU"，三操作数能力没用满，但胜在一条指令搞定两个通道。真正能吃掉三操作数优势的场景是残差连接：主路径、旁路、以及一个可选的偏置三路取最大/最小。

有一点必须清楚：_relu版本的截断发生在 16 位有符号范围内，也就是负数截到 0，正数原样保留。如果你需要的是"截断到 [0, 32767]"之外的范围，这个函数帮不上，得自己补。另外，左移和右移的shift如果是个运行时变量，编译器没法做常量折叠，性能会打折，能用模板参数固定就固定。

5. 性能验证：怎么确认它真落地了

5.1 SASS 反汇编与指令计数

改写完之后，第一件事是确认编译器真的把内建函数翻译成了单条指令，而不是"看着像优化其实拆成一堆"。方法是对比编译产物的 SASS。

nvcc -O3 -arch=sm_86 -cubin maxfilter.cu -o maxfilter.cubin cuobjdump -sass maxfilter.cubin | grep -E "IMNMX|VIMNMX|VIMAX|VIMIN"

这里的心态要摆正：不同架构的 SASS 助记符不完全一样，sm_75、sm_86、sm_90上同一条语义的指令可能在反汇编里长得略有差别。不要死记名字，看操作数个数就行——三操作数的 min/max 指令就是 DPX 落地了，两操作数的就是没落地。

如果发现内建函数被拆成了一串IMNMX，常见原因是：

• 编译的目标架构低于sm_75；
• 参数类型和函数后缀不匹配，触发了隐式转换路径；
• 你把内建函数包在了一层__host__ __device__的通用函数里，host 编译路径把整个函数排除了。

顺便看下寄存器用量：

nvcc -O3 -arch=sm_86 -Xptxas -v maxfilter.cu -o maxfilter

输出里的Used N registers是关键。用了 DPX 之后寄存器数小幅上升是正常的（三操作数指令需要同时驻留三个源操作数），但如果暴涨到 64 以上导致 occupancy 掉档，那就得不偿失，得手动拆分循环或者减少 unroll 深度。

5.2 计时方法与瓶颈判断

计时用 CUDA event，不要用 CPU 时钟包着 kernel 调用——那样测出来的主要是启动开销。标准范式：

cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); // 预热，让缓存和频率稳定 for (int i = 0; i < 20; ++i) kernel<<<grid, block>>>(...); cudaDeviceSynchronize(); cudaEventRecord(start); for (int i = 0; i < 200; ++i) kernel<<<grid, block>>>(...); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms = 0; cudaEventElapsedTime(&ms, start, stop); printf("avg %.3f ms\n", ms / 200.0f);

热身这一步别省。GPU 的频率调节有滞后，冷启动的第一轮往往慢 15% 以上，不预热直接测出来的数字基本没有参考价值。

判断瓶颈在哪，看两个指标就够了。sm__inst_executed的下降幅度告诉你省了多少指令；sm__throughput.avg.pct_of_peak_sustained_elapsed和gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed的对比告诉你卡在计算还是访存。如果改造后指令数掉了一半但耗时没动，说明你原来是访存 bound，这时候该去折腾__ldg、共享内存和访存合并，而不是继续扣指令。

5.3 DPX 帮不上忙的三种情况

说句实在的，DPX 不是银弹，遇到下面三种情况它会让你白忙一场。

第一种是访存主导的 kernel。比如一个简单的图像阈值化，每个像素读一次写一次，计算只有一条比较。这种情况下你把那条比较换成什么都不会快，因为时间全花在 DRAM 上。判断方法很简单：算一下算术强度（每字节访存对应多少次运算），低于 1 的基本没戏。

第二种是控制流复杂、分支发散严重的 kernel。DPX 的优势建立在"无分支的整数流"上，如果每个线程走不同路径，warp 内部的指令复用率本来就低，省下的几条指令会被 divergent 的开销完全淹没。

第三种是算法层面还有数量级优化空间的情况。典型例子是 3x3 最大值滤波，如果你的图像很大且允许近似，用可分离的两次一维滤波（先横向后纵向）本身就是 O(n) 而不是 O(9n)，比抠指令划算得多。先把算法选对，再谈指令级优化，顺序反了就是自欺欺人。

6. 踩坑记录与速查表

6.1 编译期与运行期的典型症状

编译期最常见的三个报错，我按出现频率排一下。第一个是identifier "__vimax3_s32" is undefined，九成是-arch设低了，检查你的编译命令里有没有sm_75及以上。第二个是no instance of overloaded function matches the argument list，这是参数类型不匹配，把unsigned和signed对齐就好。第三个是calling a __host__ function from a __global__ function is not allowed，说明你把内建函数包在了一个只标了__host__的函数里，加个__device__就行。

运行期的问题更麻烦，因为不报错、只是结果不对。我整理了一下自己踩过的和帮别人查过的：

输出整体偏移一格。共享内存 tile 的步长忘了加 halo，smem[TILE][TILE]应该是smem[TILE+2][TILE+2]。这个 bug 最阴险，因为图像看着"只是稍微歪了点"，很容易被当成正常现象放过。

负数结果变成极大的正数。类型混用，int传给了_u32版本。检查一下操作数的声明。

打包结果高半部分乱码。打包时没做(unsigned short)强转，short的符号扩展污染了高位。

性能没提升甚至下降。两个方向查：一是bool*参数指向了数组元素导致 local memory 溢出；二是寄存器暴涨导致 occupancy 掉档。用-Xptxas -v看寄存器数，用 SASS 确认指令真的合并了。

ReLU 之后还是负数。检查你用的是不是_relu版本；另外确认打包的两个通道确实都是你想要的通道，高低位别搞反了。

6.2 DPX 使用问题速查表

最后说一个我个人在实际操作中的体会。DPX 这组内建函数最容易让人上头的地方，是它"看起来很高级"，于是有人一上来就想把整个 kernel 全换成 DPX 版本，结果改完发现哪儿都没快。正确的顺序永远是反过来的：先用 profiler 确认你的 kernel 是计算密集且卡在指令发射上，再去替换热点循环里那三五处 min/max，然后立刻测。我自己的经验是，一个 kernel 里真正值得用 DPX 的地方通常不超过三处，找到它们比写它们更花时间。至于那些"顺手也改一下"的部分，十有八九是负收益——多出来的寄存器压力和降低的可读性，换不来那 1% 的速度。

还有一个小技巧分享给需要做定点推理的朋友：如果你手上有大量max(x, 0)形式的后处理，先看看能不能把整个 epilogue 重排成"两个通道并行"的形式，再上s16x2_relu。单通道调用它只能吃一半的吞吐，而重排数据布局的成本往往比你想的低，收益却是实打实翻倍的。