企业官网建设流程全解析

1. Softmax算子为什么值得单独写一篇优化文章

做AI算子开发的朋友应该都有体会，Softmax这个算子看着人畜无害，实际上特别“刁钻”。它的数学形式极其简单，但优化空间和踩坑概率在常用算子中绝对排得上号。尤其到了DCU这种国产加速卡上，想把Softmax跑出接近理论带宽的性能，需要操心的事情远比想象中多。

先说清楚Softmax在干什么。给定一个输入向量，Softmax把每个元素映射成一个概率值，公式长这样：

[ \text{Softmax}(x_i) = \frac{e^{x_i - \max(x)}}{\sum_{j} e^{x_j - \max(x)}} ]

减去最大值那一步不是可有可无，是为了数值稳定性。如果不减，当输入里有大数时，(e^{x_i})直接溢出成inf，后面全完蛋。这在Transformer的注意力层里尤其致命，因为QK^T的点积结果动辄几十上百，不稳定的Softmax会让训练直接发散。

但Softmax真正的麻烦不在数学，而在访存特性。这个算子对每个元素只做几次浮点运算，计算密度极低，属于典型的访存密集型任务。也就是说，性能瓶颈几乎完全取决于你能多快把数据从显存搬进寄存器、把结果写回去。在DCU上做优化，本质上是在跟内存带宽较劲。

这篇文章我从一个实际优化案例出发，完整走一遍从朴素实现到深度调优的过程。案例背景是某推理模型中的注意力模块，Softmax算子的输入shape为[batch, heads, seq_len, seq_len]，其中batch=4，heads=32，seq_len=512，数据类型为FP16。优化目标是把单次调用延迟从2.34ms压到1ms以内。

这个场景非常有代表性。它既有行主序二维矩阵的常规Softmax特征，又有大批量小矩阵的并行特性，还涉及FP16的精度处理，几乎把DCU上算子优化能遇到的典型问题都覆盖了。

适合读这篇文章的人有三类：一是正在做算子移植，需要把PyTorch或CUDA代码迁移到DCU上的工程师；二是做推理优化，被Attention里Softmax耗时困扰的算法工程师；三是对并行编程有兴趣，想看看国产加速卡生态实际长什么样的技术爱好者。

先说结论，最终优化后的Softmax算子单次调用耗时从2.34ms降到了0.62ms，加速比3.77倍，通过了一系列精度对比测试。这个结果不是靠某个单一技巧拿到的，而是一整套策略组合的效果。下面我把每一步的思路、实现细节和踩坑经历都拆开讲。

2. DCU硬件结构与并行编程模型的核心认知

2.1 DCU的硬件架构与计算单元布局

要优化DCU上的算子，首先得搞清楚硬件长什么样。DCU（Deep Computing Unit）是基于GPGPU架构设计的加速卡，核心计算单元叫计算引擎（Compute Engine，CE），每个CE内部包含多个SIMD单元，每个SIMD单元又由若干个线程插槽（Thread Slot）组成。

以我使用的DCU型号为例，单卡有60个CE，每个CE包含4个SIMD单元，每个SIMD单元可以同时驻留一定数量的wavefront（DCU的调度单位，等价于CUDA里的warp，通常为64个线程）。这意味着一个CE上同时活跃的硬件线程数非常可观，能够用大量并行线程来掩盖访存延迟。

DCU的存储层次和主流GPGPU类似，从快到慢依次是：寄存器文件、L1缓存、共享内存（Local Data Share，LDS）、L2缓存、全局显存。每个SIMD单元有独立的L1缓存和LDS，所有CE共享L2缓存和全局显存。

有个关键区别需要特别强调：DCU没有像CPU那样依赖大缓存来加速随机访问，而是靠海量线程的并行切换来隐藏访存延迟。这决定了优化思路的核心——你要做的不是减少线程数，而是尽可能多地创建可以并行执行的线程，让它们在等待内存数据时切换执行其他计算。

2.2 HIP编程模型与DTK工具链

DCU的编程模型遵循HIP（Heterogeneous Interface for Portability）规范。HIP的语法和CUDA高度相似，从CUDA代码迁移到HIP通常只需要做少量替换，比如__global__改__global__（HIP里也是__global__），blockIdx.x变hipBlockIdx_x，threadIdx.x变hipThreadIdx_x。但有些细节差异会在后面实操部分讲清楚。

DCU的软件开发工具包是DTK（DCU Toolkit），里面包含HIP编译器（基于LLVM）、HIPify工具（用于自动把CUDA代码转换成HIP代码）、性能分析工具等。在使用过程中，我用的是DTK 24.04版本，编译命令大致如下：

hipcc -O3 -std=c++17 -DNDEBUG -o softmax_bench softmax_bench.cpp

注意：-O3在高性能算子编译中基本是标配，但如果你在代码里用了restrict关键字或者__builtin_assume，一定要检查生成汇编是否真正优化到位，因为DCU编译器在某些场景下对复杂指针别名的处理不如预期，保守的代码写法反而会拖累性能。

2.3 Wavefront、线程块与调度机制

DCU的调度机制和NVIDIA GPU最直观的差异就是wavefront大小为64线程，而CUDA的warp是32线程。这个差异影响深远。

在Softmax算子优化中，我们经常需要做线程间的数据归约（比如求最大值、求和）。在CUDA里你习惯用__shfl_down_sync在32个线程内做shuffle归约，到了DCU/ROCm平台，对应的是__shfl_down，同时因为wavefront有64个线程，归约的步数多了一级。

另一个需要适应的是CE的调度粒度。DCU的硬件调度单元以wavefront为单位，一个wavefront里的线程执行相同的指令，如果出现分支分歧，会出现串行执行不同路径的情况，性能损失明显。所以在设计内核时，尽量保证同一wavefront内的线程走相同分支，或者干脆避免分支。

实际测算下来，同样一份Softmax内核，我最初从CUDA直接搬过来时，性能只有预期的60%左右。排查之后发现主要问题就在wavefront大小的差异上，线程块维度和归约逻辑没有针对64线程重新设计，导致大量线程空转。

3. Softmax算子的基础实现与首轮性能摸底

3.1 朴素实现：一个线程处理一行

在动手优化之前，先把基础版本写出来，作为后续优化的基准线。最简单直观的思路是：让一个线程处理输入矩阵的一行。假设输入是[M, N]的二维矩阵，那就开M个线程，每个线程独立处理一行数据。

朴素实现的伪代码逻辑如下：

__global__ void softmax_naive(const float* input, float* output, int M, int N) { int row = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; if (row >= M) return; // 1. 找最大值 float max_val = -FLT_MAX; for (int i = 0; i < N; i++) { max_val = fmaxf(max_val, input[row * N + i]); } // 2. 计算指数和 float sum = 0.0f; for (int i = 0; i < N; i++) { sum += expf(input[row * N + i] - max_val); } // 3. 归一化写回 for (int i = 0; i < N; i++) { output[row * N + i] = expf(input[row * N + i] - max_val) / sum; } }

这个实现完全正确，但性能惨不忍睹。原因有三：

第一，三次遍历输入数据。每次都要从全局显存读一遍数据，如果N比较大，访存流量翻了三倍。第二，没有做任何向量化访存，每次只能读一个float。DCU的全局访存带宽是按128字节为单位的cacheline对齐的，标量访问浪费了大量带宽。第三，大量冗余计算，指数函数被调用了两次。

在我测试的shape下，这个版本的耗时是2.34ms。注意，这个数字本身就包含了并行执行的收益，因为M=4×32×512=65536个线程块请求已经足够多，吞吐量主要被访存次数和指令效率限制。

3.2 为什么朴素实现会这么慢

拿2.34ms来算一下有效带宽。输入数据量是4×32×512×512×2字节（FP16）= 64MB，加上输出64MB，总共128MB。2.34ms对应的带宽大约是55GB/s。

看DCU的规格，理论显存带宽通常在几百GB/s到1TB/s以上。55GB/s连理论值的零头都不到。这中间的差距去哪了？

访存模式是罪魁祸首。朴素实现里，每个线程访问的行是连续的128KB数据（512×2字节），但线程间访问的行却是完全离散的。同一时刻，一个wavefront里的64个线程正在访问64个不同行的首地址，这64个地址分布在整个显存地址空间中，每次访存都触发一次完整的cacheline加载，无法合并。实际上有效的带宽利用率非常低，大量总线周期花在了等待不连续地址的数据返回上。

访存合并（Memory Coalescing）是整个优化的基石。正确的做法是让同一wavefront内的线程访问连续地址，也就是把矩阵按列方向拆分给不同线程。后面所有优化方案都建立在这个认知之上。

3.3 性能分析的基线工具使用

在动手改代码之前，先用工具记录一份基线数据，方便后续对照。DCU的生态里能用到的分析工具主要是dcu_prof和hip_analyze，前者做硬件计数器采样的时间线分析，后者做静态代码检查。

常用的操作是：

dcu_prof -t softmax_bench ./softmax_bench

从prof输出里可以读到内核执行时间、占用率、全局访存吞吐、L2命中率等关键指标。第一次跑朴素版本时，L2命中率只有21%，这个数字直接说明访存模式有严重问题。

实操心得：每次改动内核后，先记录L2命中率和全局访存吞吐两个指标，如果它们没有明显变化，说明优化方向不对，不必纠结延迟数字。这两个指标能帮你快速判断瓶颈在访存还是计算。

4. 核心优化策略：从访存模式到并行划分的全面调整

4.1 方案一：行级并行加向量化访存

朴素实现的问题在于线程块内线程处理不同行，导致访存不合并。第一版优化从访存模式下手——让一个wavefront协同处理一行数据，同时每个线程连续读取多个相邻元素，形成向量化访存。

具体做法是：每行数据由64个线程协作处理，每个线程负责连续N/64个元素。这样同一wavefront的线程在某一步访问的地址是连续的，硬件可以把这些访问合并成较少的几次cacheline传输。

用伪代码表示：

__global__ void softmax_v1(const half* input, half* output, int M, int N) { int row = hipBlockIdx_x; int tid = hipThreadIdx_x; // 0..63 int stride = N / 64; const half* row_ptr = input + row * N; half* out_ptr = output + row * N; float local_max = -FLT_MAX; #pragma unroll 4 for (int i = 0; i < stride; i++) { int idx = tid * stride + i; float val = __half2float(row_ptr[idx]); local_max = fmaxf(local_max, val); } // wavefront内归约求全局最大值 for (int offset = 32; offset > 0; offset >>= 1) { local_max = fmaxf(local_max, __shfl_down(local_max, offset)); } // ... 指数求和、归一化类似处理 }

这里__shfl_down是wavefront内部线程间数据交换的关键指令，它允许线程直接读取同一wavefront中其他线程的寄存器值，不需要经过共享内存或全局内存，延迟极低。在DCU上，它的实现效率和CUDA的shuffle指令相当，是归约类操作的利器。

这一版优化后，耗时从2.34ms降到1.48ms。提升明显，但还没达到目标，因为每个线程访问的元素间隔是stride而不是1，向量化程度不够。如果数据宽度允许，应该用float2或float4类型做显式向量化访问。

4.2 方案二：向量化访存与循环展开

DCU的编译器对显式向量类型的支持非常关键。把数据当作float4（16字节）一次性读取，可以显著减少访存指令数量，同时提高cacheline利用率。

修改核心循环：

const float4* input_v4 = reinterpret_cast<const float4*>(row_ptr); float4 vals[4]; float local_max = -FLT_MAX; #pragma unroll 8 for (int i = 0; i < stride / 4; i++) { vals[i % 4] = input_v4[tid * (stride / 4) + i]; local_max = fmaxf(local_max, fmaxf(fmaxf(vals[i % 4].x, vals[i % 4].y), fmaxf(vals[i % 4].z, vals[i % 4].w))); }

这里有个前置条件：N必须能被4整除，N/64也必须能被4整除，否则需要处理边界。好在seq_len=512这个场景完全满足条件。

循环展开的作用是让编译器生成更多独立的访存指令，这些指令可以在等待内存返回时并行执行，提高内存级并行（MLP，Memory Level Parallelism）。从实际效果看，展开因子8比展开因子4性能更好，但继续加大收益就不明显了，猜测是寄存器压力过大导致spill到local memory。

这版优化跑到了0.98ms，终于突破1ms大关，但距离最优还有空间。这时瓶颈开始从访存模式转向计算效率和归约开销。

4.3 方案三：两遍遍历合并成一遍

在基础实现中，需要三次遍历数据：找最大值、算指数和、算结果。但实际上第一遍找最大值和第二遍算指数和是可以合并的，现代GPU上常见做法是分块处理，先对每个块做局部统计，再跨块归约。

这里介绍一种常用技巧：online softmax。它允许你在不知道全局最大值的情况下，边读数据边更新统计量。对于流式数据或无法多次访存的场景非常有用。公式如下：

维护当前最大值m和累加和sum。每读到一个新元素x，计算：

[ m' = \max(m, x) ] [ sum' = sum \cdot e^{m - m'} + e^{x - m'} ]

这样一来，一趟遍历就能同时完成找最大值和求和。虽然不是所有场景都需要这招（多数情况下两遍遍历就够），但理解了它，对设计更复杂的高性能版本会有帮助。

实际优化中，我用的是两遍遍历合并到同一内核的策略，做法是：每个线程读取自己负责的数据块，先求出局部最大值，然后立即在同一段数据上计算部分指数和。因为局部最大值可能不是全局最大值，最后归一化时需要调整，但这个调整可以在最终归约阶段一次完成。

4.4 方案四：LDG缓存策略与__ldg替代

DCU的全局内存读取存在L2缓存，L2命中与否对性能影响极大。对于Softmax这种数据会被多次读取的操作，应该尽量让数据留在L2里。

在HIP中，可以用__ldg内置函数标记只读数据访问路径，编译器会生成非临时加载指令，优先命中缓存。实测在部分shape下，使用__ldg能额外带来10%~15%的性能提升。

float val = __ldg(&input[row * N + idx]);

注意，__ldg只对指针指向的只读数据有效。如果你之后对同一内存地址做了写操作，编译器可能优化掉__ldg，或者更糟，缓存一致性反而拖慢速度。这个场景下，input和output指针完全分离，放心用。

4.5 方案五：FP16数据类型的精度处理

我的案例输入是FP16，而计算过程中用FP32做中间累加是必须的。FP16的有效精度只有10位尾数，直接用它累加几百个指数值，误差会被放大到不可接受。

但这里有个隐蔽的问题：在读取FP16数据并转成FP32时，转换指令本身有开销。DCU的__half2float是一条独立指令，大量的转换会拖慢内核。

优化技巧是使用half2向量类型。一个half2寄存器可以同时装载两个FP16数，一条指令转换成两个FP32。如果输入对齐到4字节，甚至可以用half4。

half2 val2 = *reinterpret_cast<half2*>(row_ptr + idx); float2 valf2 = __half22float2(val2);

这样访存指令数和转换指令数同时减半，整体指令吞吐量大幅提升。

这个优化比较tricky的地方在于对齐。DCU要求half2指针必须4字节对齐，而输入张量是64字节对齐的，只要偏移量是2的倍数就没问题。在实际代码里我加了静态断言，确保编译期就能发现对齐问题。

4.6 方案六：大矩阵场景的分块调度设计

上面的优化针对的是seq_len=512、行数很多的情况。但Softmax还有一个典型的性能陷阱：当单行数据量特别大（比如seq_len=4096或8192），或者矩阵特别小而行的数量特别多时，策略需要完全不同。

针对行数据很大的场景，不能再用“一个wavefront处理一行”的方案，因为每行数据太大，wavefront的寄存器装不下全部数据，也没法一次性归约。正确的做法是把每行拆分成多个列块，先分别计算局部统计量，再用第二个内核或第二次pass做全局归约。

这种两阶段处理有个工程细节要注意：第一次pass结果需要暂存在中间缓冲区，而中间缓冲区需要按块对齐分配，避免bank conflict。我遇到过的最典型问题就是中间缓冲区大小没有按LDS的bank数对齐，导致同时访问时发生大量冲突，性能比朴素实现还差。

针对行数特别多的场景，更应该关注任务调度粒度。索引从一维线性化：把整个矩阵看成一个大一维数组，按固定大小的tile切分给不同线程块。然后每个线程块从tile里提取它需要处理的行的片段。这样做的好处是线程块之间的负载更均衡，不容易出现某些CE空闲的情况。

5. 实战记录：三版迭代完整性能对比与关键参数选择

5.1 线程块尺寸与wavefront对齐的策略选择

线程块尺寸的设计直接决定了并行度上限。DCU的一个CE最多可以驻留一定数量的wavefront，超过之后多余线程只能排队。对于Softmax这种访存密集型算子，并行度要尽量高，但也不能无脑加大。

在我最终方案里，线程块大小设为256，即4个wavefront（256/64）。这样设置的原因有两点：第一，256个线程刚好可以完整覆盖一行512个FP16数据（每个线程处理2个half，刚好组成一个half2向量），不需要额外处理边界；第二，256个线程的寄存器占用不会超过CE的资源限制，允许足够多的线程块同时驻留。

关于每个线程处理多少数据，有个经验公式：理想情况下，每个线程处理4~8个FP32元素（或者8~16个FP16元素），既不会让访存指令太少导致延迟掩盖不足，也不会让寄存器溢出。我的案例中，每行512个FP16，每个线程处理8个元素，即4个half2，分配算式是512/(64×4)=2个half2每线程。这个组合实测效果最佳。

如果行数是奇数乘数（例如N=768），处理方式就会复杂一些，需要最后一个wavefront部分线程处理额外元素，或者调整每个线程的数据量。大多数情况下，选择每个线程处理相同数量的元素，然后对越界部分做mask处理，性能损失可以控制在几个百分点内。

5.2 共享内存与寄存器使用的平衡

在归约阶段，最初版本我用共享内存做跨线程数据交换。每个线程把自己的局部最大值写到共享内存，然后做分阶段同步归约。

但很快发现，共享内存归约有明显的性能代价：每次归约都需要__syncthreads()，这个同步指令会让整个wavefront停下来等待最慢的线程，频繁使用会浪费大量周期。

改用warp shuffle归约后，效果立竿见影。shuffle指令直接在寄存器之间交换数据，不走共享内存，也不需要同步。因为一个wavefront的64个线程在DCU上是同时调度的，天然保证了一致性，shuffle的效率远高于共享内存。

在最终代码里，我彻底移除了共享内存的使用，全部归约都用shuffle完成。这不仅减少了同步开销，还降低了寄存器压力，因为共享内存和寄存器在某些架构上是共享资源的。

5.3 三版方案的性能测量与放大对比

整个调优过程我保留了三个关键版本的测量数据，整理成表格方便对照。

有效带宽的计算方法是：总数据量（输入+输出，FP16各64MB）除以耗时。v2版本的207GB/s依然没到理论峰值，但考虑到FP16转FP32、指数计算、shuffle交换这些额外开销，已经到了非常合理的水平。

这个对比也暴露了一个现象：v1到v2的进步不是靠单项优化，而是simultaneously调整了访存宽度、归约机制、缓存策略和指令混合，每个方向一点点累积，最后叠加出大提升。单靠某一招想要3.77倍加速，基本不可能。

5.4 精度验证与数据比对

性能再漂亮，精度不对就是废代码。Softmax的精度验证我用了三步：

第一步，最大绝对误差检查。在同一输入下，用DCU内核输出与PyTorch CPU的float64参考实现对比，计算每个元素的最大绝对误差。我的实现最终最大误差在2e-4以内，对FP16输出来说完全可接受。

第二步，分布一致性检查。Softmax的输出本质是一个概率分布，检查每行输出是否满足求和接近1。实测最大偏差小于1e-3，符合预期。

第三步，端到端模型验证。把优化后的Softmax接回原推理模型，跑一批真实数据，观察最终输出的top-1准确率和原始实现是否一致。这一步最重要，因为算子层面的微小误差在某些场景下会累积放大。

特别提醒：如果你优化的是训练过程中的Softmax，一定要检查反向传播的表现，因为梯度计算对精度更敏感。我这次只做推理优化，所以反传不在范围内，如果你需要支持训练，建议在反向Softmax的优化上单独投入时间，很多推理场景的优化技巧不能直接套用。

6. 常见问题与性能排查技巧实录

6.1 为什么我的shuffle归约结果不正确

在DCU上使用__shfl_down时最容易踩的坑是忘记处理线程数不等于wavefront大小的情况。如果线程块大小是128，而你只用前64个线程做归约并shuffle到其他线程，就会漏掉数据。

我调试时发现一个更隐蔽的问题：__shfl_down的offset如果小于32，行为正常；但如果offset在32到63之间，某些DCU型号的驱动行为不一致，结果丢数据。排查了很久，最后改成两次循环，先做offset=32的归约，再做offset=16、8、4、2、1的归约，问题消失。

还有个常见错误是忘记把参与shuffle的变量赋初值。如果某线程的局部最大值初始化为0而不是-FLT_MAX，恰好这一行所有输入都是负数，最终归约结果就会错误地变成0。这种bug不会崩溃，只会让精度测试不过，非常隐蔽。

6.2 L2命中率上不去，问题出在哪

用dcu_prof看到L2命中率低，第一反应自然是缓存复用不够。但对于Softmax这种流式访问主导的算子，L2命中率天生不会太高，因为每行数据基本只会被读一次（如果用两遍遍历，会被读两次，但第二次可能已被挤出L2）。

如果L2命中率特别低，另一个排查方向是线程块调度顺序。DCU的线程块调度顺序可以控制L2的空间局部性。尝试修改hipLaunchKernelGGL的grid维度排列方式，把相邻线程块映射到共享L2区域的线程块，能提升命中率。

我实测过把grid从[M]改成[M/8, 8]（二维grid），其中第二维是L2切片索引，结果L2命中率从21%提升到了34%。虽然数字看着不大，但整体耗时降低了约8%。

6.3 编译器自动向量化失败，如何检查

DCU的HIP编译器对自动向量化的能力有限，尤其在循环内有fmaxf这类数学函数时，常常不会自动生成向量访存指令。检查方法很简单：在编译命令里加-S选项生成汇编文件，然后搜索v_load或global_load_dwordx这类指令。

如果发现循环体里都是global_load_dword（4字节标量加载），说明编译器没有自动向量化成功。解决办法是像前面那样显式使用float2或half2类型，把向量化写死在代码逻辑里。编译器对显式向量类型的支持很好，反汇编能看到global_load_dwordx2（8字节）或global_load_dwordx4（16字节）指令。

6.4 性能抖动：为什么内核耗时忽高忽低

内核耗时不稳定，可能是其他任务抢占显存带宽，也可能是时钟频率波动。排查方法是连续运行多次benchmark，观察分布。

更有意思的一个原因是：DCU在同时跑多个上下文时，L2缓存会被共享，如果同一个GPU上有其他kernel在跑，Softmax的L2命中率会骤降。这是硬件层面的资源竞争，代码层面没法完全规避，但可以在任务调度时避免同时运行多个大访存kernel，或使用单独的GPU实例。

6.5 边界条件处理的最优解

当N不能被线程数整除时，不要直接开根号强行整除，那会让代码逻辑臃肿且难以维护。更优雅的方案是让每个线程处理固定数量的元素，最后留一个线程处理尾部数据。

或者使用grid-stride loop，让每个线程循环处理多个元素，循环条件里判断边界。这个模式对不规则shape的适配性最好，性能损失也很小。

6.6 从CUDA代码迁移到DCU的隐藏问题

如果是从CUDA代码直接迁移，hipify工具能完成90%的替换工作，但剩下10%会导致性能剧烈下降。最常见的问题：

第一，__syncthreads()的语义在DCU上不如CUDA严格，某些编译器优化可能导致同步被移除。如果代码里有复杂的共享内存读写依赖，最好检查一下反汇编里是否真的存在barrier指令。

第二，cudaMalloc换成hipMalloc后，分配的大内存默认属性可能与CUDA不同，访问延迟更高。建议尝试hipMallocManaged或调整对齐属性。

第三，__restrict__关键字在DCU编译器上的优化力度不如NVIDIA的nvcc。如果迁移后性能达不到预期，尝试手动把指针加载到局部变量，消除每次访问的指针解引用开销。

7. 关于DCU算子优化生态的一些补充经验

除了Softmax本身，这次优化过程中接触到的DCU工具链和生态也值得说几句。

DTK的hipcc编译器总体质量不错，但对某些优化模式的支持还不成熟。比如我尝试过用内联PTX（DCU上对应的是内联GCN汇编）手写FMA和指数指令的组合，确实能压掉几条指令，但代码可维护性急剧下降。除非为了追求极限性能，否则不建议在工程代码里大量使用。

DCU的性能分析工具这几年进步明显，dcu_prof的硬件计数器覆盖已经比较全。但相比成熟工具还是少了些便利性，比如不能直接在时间线上查某条指令的详细信息。解决方法是自己在代码里插桩，用clock64()记录关键阶段耗时。这个方法土，但有效。

社区方面，虽然DCU生态还比不上CUDA，但近两年的文档和示例代码质量提升很大。遇到问题时，先查/opt/dtk目录下的示例代码，通常能找到对应的内核模板，再结合官方性能优化指南里的架构特性说明，大多数问题都能定位。

对于团队来说，如果要做算子迁移，建议建立一份自己的性能基线库，把每个算子在不同shape下的朴素实现耗时、优化后耗时、理论带宽和实测带宽都记录下来。有了这个基线库，后续新算子的优化进度评估会快很多，也能避免把已经解决过的坑再踩一遍。

8. 最终版本核心代码参考实现

把最终可运行的优化内核主体贴出来，方便参考。这个版本融合了前面所有可行优化，代码故意保留了shuffle归约和half2显式向量化，方便对照理解。

#include <hip/hip_runtime.h> #include <hip/hip_fp16.h> #include <cmath> #include <cfloat> // 假设 M 能被 gridDim.x * blockDim.x 整除（此处 M = 65536） // 每行 N = 512 个 half，由 blockDim.x = 256 个线程处理 // 每个线程处理 1 个 half2（即 2 个 half），一个 block 覆盖 2 行 __global__ void softmax_opt_kernel(const half* __restrict__ input, half* __restrict__ output, int M, int N) { const int tid = hipThreadIdx_x; // 0..255 const int wid = tid >> 6; // 所在 wavefront 编号 0..3 const int lane = tid & 63; // wavefront 内线程编号 0..63 // 每个 block 处理 4 行，间隔 blockDim.y? 此处简化为线性处理 int row = hipBlockIdx_x * 4 + wid; // 一个 wavefront 处理一行 if (row >= M) return; const int half_per_thread = N / 256; // 2 const half2* row_ptr = reinterpret_cast<const half2*>(input + row * N); half2* out_ptr = reinterpret_cast<half2*>(output + row * N); float local_max = -FLT_MAX; // 每线程负责一个 half2 half2 val = __ldg(&row_ptr[lane]); float2 valf = __half22float2(val); local_max = fmaxf(fmaxf(valf.x, valf.y), local_max); // wavefront 内归约最大值 (64 -> 1) #pragma unroll for (int offset = 32; offset > 0; offset >>= 1) { local_max = fmaxf(local_max, __shfl_down(local_max, offset)); } float row_max = __shfl(local_max, 0); // 广播到整个 wavefront // 计算指数与和 float e0 = __expf(valf.x - row_max); float e1 = __expf(valf.y - row_max); float local_sum = e0 + e1; #pragma unroll for (int offset = 32; offset > 0; offset >>= 1) { local_sum += __shfl_down(local_sum, offset); } float row_sum = __shfl(local_sum, 0); // 归一化写回 half2 res; res.x = __float2half(e0 / row_sum); res.y = __float2half(e1 / row_sum); out_ptr[lane] = res; }

这段代码有几个使用前提：行数必须是block数×4的整数倍（这个案例满足），N必须等于256的倍数（512满足）。如果shape不满足，需要加上边界判断，逻辑会复杂一些。

关于__expf和expf的选择，我在优化中特意用了快速版本。__expf的精度比expf低一些，但速度更快。对Softmax的最终输出影响在1e-5量级，完全可接受。在训练等需要高精度的场景，建议换回标准expf。

另一个细节是__shfl和__shfl_down的配合使用。先通过__shfl_down把最大值归约到lane 0，再用__shfl广播给全wavefront，这个模式比每个线程都保存一份最终值更省指令。

9. 后续扩展方向与实践建议

Softmax的优化经验可以直接延伸到其他访存密集型算子。像LayerNorm、RMSNorm、通道均值方差计算，它们的结构都是“读数据-归约-归一化-写回”，和Softmax几乎同构。把这次用的shuffle归约、half2向量化、L2缓存友好的线程块调度这几个技巧迁移过去，通常能快速获得类似幅度的提升。

在更复杂的注意力机制中，Softmax经常和QK^T矩阵乘法、矩阵掩码操作融合。如果你的融合目标是减少kernel launch次数，可以在Softmax内核里加一个参数，判断是否需要应用attention mask。在DCU上，kernel launch的固定开销比NVIDIA高，减少launch次数对端到端性能的影响更明显。这也是我后续在做的工作方向。

从业余时间接触DCU到现在，我最大的体会是：硬件不同，但底层逻辑共通。访存合并、归约效率、指令混合、缓存利用，这些在任何GPU架构上都是核心命题。只不过在DCU上，你需要更主动地管理这些优化，因为工具链和编译器还没成熟到自动帮你搞定一切。

对于刚开始接触DCU优化的朋友，建议从小算子入手，别一上来就挑战大模型端到端调优。Softmax其实是个很好的起点：数学简单，但访存、归约、向量化、缓存全涉及一遍。把这个流程走通，再看其他算子会轻松很多。

这次优化到0.62ms之后，我并没有继续往下压。再往下走就要开始碰汇编手写指数运算和指令级调度了，收益可能还有20%左右，但代码可读性和可维护性会急剧恶化。在工程实践中，一个能维护、能debug、能迁移的内核，比一个快10%但谁也看不懂的内核有价值得多。如果你走的是产品化路线，记住这个判断标准。