大模型中经常出现depthwith卷积，本质是个分组卷积，这类卷积的特点是groups=cin=cout，这类卷积如果按照group作为外循环的普通方案实现，cube的利用率会很低，因为每个group内的卷积的cin核cout只有1，转换为矩阵乘后，K和N非常小，而且昇腾硬件上每cycle可计算161616大小的矩阵乘，最坏情况每次算力只能发挥1/256。我们提出了group-opt的方案，一次性计算多个groups，每次计算前对weight进行预处理改造，把weight扩充成对角矩阵，只有对角线上值有效，其他点都为0，大幅度提高cube利用率，对比group作为外循环的方案平均性能提升30%

当前昇腾为了满足端侧/云侧/大模型推理/训练已经推出了多种芯片型号，soc的核数也越来越多样化，之前我们矩阵类算子的分核策略亲核于2的幂次方，对于非2的幂次方的核经常出现绑不满核或者分核不均的情况，无法充分发挥算力，借此背景我们推出任意核上负载均衡的策略。把MN的大矩阵按M0N0分成多个小矩阵，每个核按顺序处理一小块，当块数是核数的倍数时可以很好的负载均衡，但如果不是倍数就会存在拖尾。假设一共5块，核数为4，每个核处理完一个后，会剩一个尾块需要一个核处理，这样会造成拖尾，导致负载不均衡。我们会对最后一个尾块再一次进行切分，动态调整基本快大小，尽可能最后一轮每个核上处理的数据量一致。我们还进行了蛇形走位的l2缓存优化，如果一行处理完后我们可以选择从下一行的0号位置开始或者从下一行的尾部开始，我们选择从尾部开始处理，这样可以复用上一行的A矩阵或者B矩阵的l2缓存，提高核间的l2缓存命中率。

搬运bound做对应的优化：
（1）消除重复搬运
（2）消除搬运流水中指令间的空隙
（3）消除bank冲突
（4）cacheLine对齐
（5）增大L2/L1/L0利用率，减少GM搬运
（6）整网中数据预取
（7）量化输入数据，减少搬运数据量，计算换搬运
（8）更快的硬件指令
（9）每次指令搬运的数据量要足够多，实现连续搬运

计算bound做对应的优化：
（1）消除计算流水中指令间的空隙
（2）使用算力更强的计算单元
（3）计算指令配置数据量对齐
（4）增大每次tile块的有效计算量
（5）量化
（6）数学上等价变换

scalar bound做对应的优化：
（1）减少LD和ST
（2）常量传播
（3）头开销优化
（4）一次多载，比如l0多载几份数据进来

无bound
（1）打开乒乓
（2）preload策略
（3）指令合并计算

多核
（1）是否用满核
（2）每个核处理数据量是否均分

mix算子理论时间估计
（1）按vector时间、搬运时间、cube时间最大的来估计理论时间

一、 搬运 Bound (Memory Bound) 优化迁移

GPU 是极其依赖带宽的架构，这一部分的迁移率几乎是 100%。

• （1）消除重复搬运：对应 GPU 中的 Kernel Fusion（算子融合）。减少中间结果写回 Global Memory (GM) 再读出的过程。
• （3）消除 Bank 冲突：在 GPU 中特指 Shared Memory Bank Conflict。GPU 共享内存分为 32 个 bank，如果多个线程访问同一个 bank 的不同地址，会造成串行化。
• （4）CacheLine 对齐：对应 GPU 的 Memory Coalescing（显存合并访问）。确保一个 Warp 内的 32 个线程访问的地址落在同一个 128-byte 的显存块内。
• （5）增大 L2/L1 利用率：GPU 开发者通常通过 Tiling（分块） 策略，将数据留在 Shared Memory（相当于程序员控制的 L1/L0）中重复使用。
• （7）量化输入数据：对应使用 FP16/BF16/INT8。减少比特数直接翻倍带宽利用率，同时 GPU 的 Tensor Core 对低精度有极高的吞吐量。
• （9）连续搬运（Vectorized Load/Store）：在 CUDA 中使用 float4 或 uint4 指令，一次读取 128 bit，能更有效地利用内存总线带宽。

二、 计算 Bound (Compute Bound) 优化迁移

GPU 的计算优化重心在于“喂饱”众多的流处理器（SM）。

• （1）消除计算流水空隙：对应 GPU 的 Latency Hiding（延迟隐藏）。GPU 不靠高频率，而是靠高并行度（足够的 Warp 数量）来掩盖指令延迟。
• （2）更强的计算单元：对应使用 Tensor Cores。在 GPU 上，如果能用 Tensor Core 跑矩阵运算，效率会比普通的 CUDA Core（Vector Unit）高出数倍。
• （3）&（4）数据量对齐与 Tile 块大小：在 GPU 中，Tile 的大小直接影响 Occupancy（活跃线程束比例）。Tile 过大会导致寄存器资源枯竭，过小则无法隐藏搬运延迟，需要寻找平衡点。
• （5）量化：GPU 的低精度计算（如 INT8/FP8）吞吐量远高于 FP32。

三、 Scalar Bound 优化迁移

GPU 虽然是典型的 SIMT（单指令多线程）架构，但每个 SM 都有标量单元。

• （1）减少 LD 和 ST：在 GPU 中体现为 Register Spilling（寄存器溢出） 优化。避免编译器将局部变量放入慢速的 Local Memory（实则在显存中）。
• （2）常量传播：使用 __constant__ 内存。GPU 有专门的 Constant Cache，对于 Warp 内所有线程访问同一地址的情况有加速。
• （3）头开销优化：对应 Kernel Launch Overhead。通过 CUDA Graph 或把小算子合并来减少 CPU 侧下发指令的开销。

四、 无 Bound (General/Pipeline) 优化迁移

• （1）打开乒乓（Double Buffering）：在 CUDA 中通过 shared memory 配合 cp.async（异步拷贝指令）实现。在计算当前 Tile 时，异步加载下一个 Tile。
• （2）Preload 策略：对应 __prefetch_global_l1 指令，或者在计算循环开始前先做一次异步加载。
• （3）指令合并计算：如 FMA (Fused Multiply-Add)。一个周期完成乘加运算，GPU 编译器通常会自动处理。

五、 多核 (Multi-SM) 优化迁移

GPU 的“多核”体现为数百个 SM。

• （1）是否用满核：对应 Grid Size。如果生成的线程块（Thread Blocks）太少，会导致 GPU 大量 SM 闲置（特别是现在的 H100/H20).
• （2）均分数据（Load Balance）：避免 Tail Effect（长尾效应）。如果某个线程块处理的数据远多于其他块，整个 Kernel 的结束时间取决于最慢的那个块。在 GPU 动态任务分配中常用 Atomic Add 实现任务池（Task Queue）。

总结：GPU 独有的额外优化点

虽然 NPU 的经验迁移性很强，但 GPU 有几个特有的性能陷阱需要额外注意：

1. Warp Divergence（分支分叉）：if-else 逻辑如果导致同一 Warp 内线程走不同分支，会造成严重的计算性能下降（NPU 某些架构对标量逻辑处理较好，GPU 则非常忌讳）。
2. Register Pressure（寄存器压力）：GPU 的寄存器是有限资源，每个线程用的寄存器越多，能并行跑的线程就越少（Occupancy 下降）。
3. Unified Memory / PCIe 瓶颈：如果数据在 CPU 和 GPU 之间频繁同步，搬运 Bound 会变成整个系统的致命瓶颈。

你的总结已经触及了高性能计算的本质，只要把 NPU 的 L0/L1/L2 缓存概念 映射到 GPU 的 Registers/Shared Memory/L2 Cache，逻辑是完全相通的。

如何从头实现一个高性能算子？
1、审视数学公式，进行等价化简，确定对硬件最优的数据排布格式，分析使用哪些指令可以完成功能，选用最高效的指令。根据理论计算指令吞吐、多级缓存带宽与延迟去进行理论建模，根据理论建模评估不同场景下的 bound 类型，简单算子使用roofline模型评估，融合算子按耗时最长的流水进行评估。
2、根据理论建模去推导最优tiling切块。
3、配合tiling实现算子kernel。
4、完成算子开发后使用性能分析工具查看能否达到理论bound类型，如果不能先明确为什么达不到，是否是代码bug？是否某种场景有局限性无法达到理论性能？根据实际情况刷新bound类型，再根据bound类型去优化tiling和kernel。

for batch:
    for N:
        for s1_loop: // q
            lse_max = -inf;
            lse_sum = 0;
            acc_out = 0;
            for s2_loop:
                qk_tile = mm(q, k, transposeB=true); // cube
                // l0c_to_ub
                new_lse_max = max(lse_max, max(qk_tile)); // ub
                p_qk_tile = exp(qk_tile - new_lse_max); // ub
                new_lse_sum = lse_sum * exp(lse_max - new_lse_max) + sum(p_qk_tile); // ub

                new_acc_out = mm(p_qk_tile, v); // cube
                // l0c_to_ub
                acc_out = acc_out * exp(lse_max - new_lse_max) + new_acc_out; // ub

                lse_max = new_lse_max; // ub
                lse_sum = new_lse_sum; // ub
            acc_out /= lse_sum; // ub
            output[offset] = acc_out;// ub_to_gm

for s1_loop:
    lse_max = -inf;
    lse_sum = 0;
    lse_out = 0;
    for d_loop:
        new_lse_max = max(lse_max, max(d_tile));
        new_d_tile = exp(d_tile - new_lse_max);
        new_lse_sum = lse_sum * exp(lse_max - new_lse_max) + sum(new_d_tile);
        lse_out = lse_out * exp(lse_max - new_lse_max);
        lse_out[offset] = new_d_tile;
        lse_max = new_lse_max;
        lse_sum = new_lse_sum;
    lse_out /= lse_sum;
    output[offset] = lse_out;

mte2头开销减少
1、在L1空间够的情况下，weight切K分块载入，但不换出，保持复用。
2、增加重复载入量，强行切K，cubebound场景可做。