1.RoPE 旋转位置编码  核心算子

// ============================================================================
//  PairwiseRotary 算子 Kernel 实现（Ascend C）
// ----------------------------------------------------------------------------
//  算子功能：把一行 64 个数劈成前后半（各 32 个），做二维旋转（RoPE 的 rotate-half）
//      y[0:32]  = xLo * cos - xHi * sin
//      y[32:64] = xLo * sin + xHi * cos
//
//  输入：x   [B, N, S, 64]  fp16
//        cos [B, 1, 1, 32]  fp16   （每个 batch 一组，长度是 64 的一半）
//        sin [B, 1, 1, 32]  fp16
//  输出：y   [B, N, S, 64]  fp16
//
//  并行方式：按「行」把数据平均分给多个 AI Core，
//            每个核处理若干行，核内逐行「搬入 → 计算 → 搬出」。
// ============================================================================

// ---- 头文件保护：防止本文件被 include 两次导致类重复定义 ----
#ifndef _PAIRWISE_ROTARY_H_
#define _PAIRWISE_ROTARY_H_

// ---- 三个必要的头文件 ----
#include "kernel_operator.h"                  // Ascend C 所有 API（DataCopy/Cast/Mul/Add/PipeBarrier...）
#include "kernel_tiling/kernel_tiling.h"      // tiling 相关宏（GET_TILING_DATA 等）
#include "pairwise_rotary_tiling_data.h"      // 本算子自己的 tiling 结构体（host↔kernel 的数据契约）

// ---- 所有 Ascend C API 都在 AscendC 命名空间下，展开后可以直接写 DataCopy(...) ----
using namespace AscendC;

// ---- 编译期常量：写死在本文件里，kernel 侧不需要从 host 传 ----
constexpr int64_t HDIM = 64;       // 隐藏维度：一行有 64 个元素（输入输出最后一维）
constexpr int64_t HDIM_HALF = 32;  // 一行的一半：前后各 32 个元素（cos/sin 的长度）

// ---- Kernel 类：每个核上会构造一个对象，所有逻辑都封装在类里 ----
class KernelPairwiseRotary {
public:
    // ========================================================================
    // 构造函数：把外部传进来的 pipe 和 tiling 保存成成员变量
    //   pipe   : 流水线管理器，负责统一分配 UB 空间
    //   tiling : host 侧算好的切分参数（只读）
    // 参数名后面的下划线是成员变量的命名习惯。
    // 冒号后面是「初始化列表」，比在函数体里赋值更高效。
    // ========================================================================
    __aicore__ inline KernelPairwiseRotary(TPipe* pipe, const PairwiseRotaryTilingData* tiling)
        : pipe_(pipe), tiling_(tiling) {}     // 把两个指针存进成员变量

    // ========================================================================
    // Init：初始化。只做三件事：① 绑定 GM 地址 ② 算本核负责哪一段 ③ 分配 UB
    // 参数是 4 个裸地址，由 kernel 入口函数（pairwise_rotary_arch22.cpp）传进来
    // ========================================================================
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR cos, GM_ADDR sin, GM_ADDR y)
    {
        // ---- ① 把裸地址包装成 GlobalTensor ----
        // GM_ADDR 本质是 __gm__ uint8_t*（不知道类型的显存地址），
        // 所以必须强制转换成 __gm__ half*，再交给 SetGlobalBuffer。
        // 包装之后就能用 xGm[off] 这样按下标访问显存了。
        xGm.SetGlobalBuffer((__gm__ half*)x);    // 输入 x: [B,N,S,64]
        yGm.SetGlobalBuffer((__gm__ half*)y);    // 输出 y: [B,N,S,64]
        cosGm.SetGlobalBuffer((__gm__ half*)cos); // cos: [B,1,1,32]（第 0 维是 batch，其余是 1）
        sinGm.SetGlobalBuffer((__gm__ half*)sin); // sin: [B,1,1,32]

        // ---- ② 算本核负责哪一段 ----
        // 总元素数 = B × N × S × 64。这些 shape 值 kernel 拿不到，
        // 全部由 host 侧的 tiling.cpp 算好、通过 tiling 结构体传进来。
        int64_t totalElem = tiling_->batchSize * tiling_->numHead * tiling_->seqLength * HDIM;

        // 每核负责的元素个数（elements per core）。
        // host 侧按 totalElem / coreNum 算，再向上取整到 64 的倍数（保证不切断行）。
        int64_t epc = tiling_->elemPerCore;

        // GetBlockIdx() 返回当前核号：0, 1, 2, ..., coreNum-1
        // 第 k 个核负责从 k*epc 开始的 epc 个元素
        start_ = GetBlockIdx() * epc;            // 当前 block 起始偏移

        // 本核的结束位置（不含）
        int64_t end = start_ + epc;
        // 最后一个核的 end 可能超出总长度，必须截断到 totalElem，否则越界访问
        if (end > totalElem) end = totalElem;
        // 把「元素个数」换算成「行数」：每行 64 个元素
        curRows_ = (end - start_) / HDIM;        // 当前 block 需要处理的行数

        // ---- ③ 分配 UB 空间 ----
        // InitBuffer(队列对象, 深度, 每块字节数)
        //
        // qX：x 的双缓冲队列。深度 2 表示能同时存在 2 块：
        //     一块正在被 Compute 读，另一块正在被 MTE2 搬运 → 实现流水并行。
        //     每块 64 个 half = 128 字节。
        pipe_->InitBuffer(qX, 2, HDIM * sizeof(half));              // x 双缓冲：预加载下一行
        //
        // qOut：输出队列，深度 1。这块 buffer 要「一物两用」：
        //     前 64 个 float（256 字节）放计算结果；
        //     后 64 个 float 的空间被 ReinterpretCast 成 half，当搬运落地区。
        //     所以申请了 HDIM*sizeof(float)*2 = 512 字节。
        pipe_->InitBuffer(qOut, 1, HDIM * sizeof(float) * 2);       // 输出队列（64 float + 64 half）
        //
        // qCos/qSin：同样是「一物两用」，前 32 个 float 不用，
        //     后 32 个 float 的空间当 half 落地区（详见 LoadCS 的注释）。
        pipe_->InitBuffer(qCos, 1, HDIM_HALF * sizeof(float) * 2);  // cos 加载队列（32 float + 32 half）
        pipe_->InitBuffer(qSin, 1, HDIM_HALF * sizeof(float) * 2);  // sin 加载队列（32 float + 32 half）
        //
        // bXf/bCos/bSin：TBuf 是不带同步的裸 buffer，适合放「写一次、读多次」的中间结果。
        //     它们是 half 转成 float 之后的存放处，后续所有计算都在 float 上做。
        pipe_->InitBuffer(bXf, HDIM * sizeof(float));               // x 转 float 后存储（64 个 float）
        pipe_->InitBuffer(bCos, HDIM_HALF * sizeof(float));         // cos 转 float 后存储（32 个 float）
        pipe_->InitBuffer(bSin, HDIM_HALF * sizeof(float));         // sin 转 float 后存储（32 个 float）
    }

    // ========================================================================
    // Process：主流程。逐行处理本核负责的所有行。
    // 每一行的动作：按需加载 cos/sin → 加载 x → 计算 → 写回。
    //
    // 两个关键优化：
    //   1. 双缓冲：把「下一行 x 的搬运」提前到「当前行计算」之前，
    //      让 MTE2（搬运）和 Vector（计算）并行；
    //   2. 数据复用：cos/sin 每个 batch 只有 32 个数，同 batch 内所有行共用，
    //      所以只在跨 batch 边界时才重新加载。
    // ========================================================================
    __aicore__ inline void Process()
    {
        // 边界保护：如果本核没有分到任何行（核数比数据量多），直接退出。
        // 不写这一句，后面的循环和下标计算会越界。
        if (curRows_ <= 0) return;

        // 取两个 shape 值备用：nh = 每 batch 的 head 数，sl = 序列长度。
        // 因为内存是 [B][N][S][D] 一维展开的，
        // 「行号 ÷ (N×S)」就能反推出这一行属于哪个 batch。
        int64_t nh = tiling_->numHead, sl = tiling_->seqLength;

        // lastBi 记录「当前已经加载到 bCos/bSin 里的 cos/sin 属于哪个 batch」。
        // 初值 -1 只是占位，下面马上会被赋成真实值。
        int64_t lastBi = -1;

        // off0：本核第一行在 GM 中的元素偏移；
        // bi0 ：本核第一行所属的 batch 号 = (元素偏移 ÷ 64) ÷ (N×S)
        int64_t off0 = start_, bi0 = (off0 / HDIM) / (nh * sl);
        // 把 lastBi 更新为真实的 batch 号
        lastBi = bi0;

        // ---- 首行预加载 ----
        // 先把第一行需要的 cos/sin 和 x 搬进来，这样循环体里就可以直接算。
        // 注意这里只 EnQue 不出队，出队在 Compute 里做（生产者-消费者分离）。
        LoadCS(bi0);    // 加载第 bi0 个 batch 的 cos/sin 到 bCos/bSin
        LoadX(off0);    // 加载第 0 行 x 到 qX 队列

        // ---- 逐行主循环 ----
        for (int64_t row = 0; row < curRows_; row++) {
            // 当前行在 GM 中的元素偏移 = 本核起点 + 行号 × 每行元素数
            int64_t off = start_ + row * HDIM;

            // ---- 双缓冲的核心：先搬下一行，再算当前行 ----
            // 因为 qX 深度是 2，可以同时存在 2 块：
            //   - 当前行还在 qX 里等着被 Compute 消费；
            //   - 下一行已经在往另一块 buffer 里搬了。
            // 这样 MTE2（搬运单元）和 Vector（计算单元）就能同时干活。
            // row + 1 < curRows_ 是边界判断：最后一行没有「下一行」，不能预加载。
            if (row + 1 < curRows_) {
                LoadX(off + HDIM);   // 下一行的偏移 = 当前行偏移 + 64
            }

            // ---- 计算当前行 ----
            // Compute 内部会 DeQue qX（如果搬运还没完成会自动阻塞等待）
            Compute();
            // ---- 把当前行的结果写回 GM ----
            Store(off);

            // ---- 预判下一行是否需要重新加载 cos/sin ----
            // cos/sin 的形状是 [B,1,1,32]，每个 batch 只有 32 个值，
            // 但同一个 batch 内有 N×S 行都要用它们。
            // 所以只在「下一行属于新 batch」时才重新加载，避免重复读 GM。
            if (row + 1 < curRows_) {
                int64_t nextOff = off + HDIM;                    // 下一行的元素偏移
                int64_t nextBi = (nextOff / HDIM) / (nh * sl);   // 下一行属于哪个 batch
                if (nextBi != lastBi) {                          // 换了 batch 才重新加载
                    LoadCS(nextBi);                              // 加载新 batch 的 cos/sin
                    lastBi = nextBi;                             // 更新记录
                }
            }
        }
    }

private:
    // ========================================================================
    // LoadCS：加载第 bi 个 batch 的 cos 和 sin，转成 float 后存进 bCos/bSin。
    //
    // 为什么要绕一圈 ReinterpretCast？
    //   DataCopy 要求源和目标的元素类型完全一致：cosGm 是 half，目标也必须是 half。
    //   但队列 qCos 是按 float 申请的（因为后面 Cast 的输出是 float，需要 float 空间）。
    //   解决办法：同一块 buffer 前后分段，后半段当 half 的「落地区」。
    //   即：把 32 个 half 搬到 c[32] 开始的字节位置，Cast 时再从那里读出来。
    // ========================================================================
    __aicore__ inline void LoadCS(int64_t bi)
    {
        // 从 qCos 队列申请一块 UB。申请成 float 类型，所以容量是 64 个 float（256 字节）。
        LocalTensor<float> c = qCos.AllocTensor<float>();

        // DataCopy(目标, 源, 元素个数)
        //   目标 c[HDIM_HALF] → c[32]，即从第 32 个 float 开始的位置（字节偏移 32×4=128）；
        //        .ReinterpretCast<half>() 把这段内存「重新解释」为 half 类型，
        //        128 字节的空间当成 64 个 half 用（实际只用 32 个 = 64 字节）。
        //   源   cosGm[bi * HDIM_HALF] → cos 第 bi 个 batch 的起始位置（偏移 bi×32）；
        //   长度 HDIM_HALF = 32 个 half。
        DataCopy(c[HDIM_HALF].ReinterpretCast<half>(), cosGm[bi * HDIM_HALF], HDIM_HALF);

        // EnQue：告诉硬件「这块 buffer 已交给 MTE2 去搬，别动它」；
        // DeQue：等硬件说「搬完了」，拿回 buffer 使用权。
        // 两句合起来 = 等待 DataCopy 完成（跨 MTE2 → Scalar 的同步）。
        qCos.EnQue(c); c = qCos.DeQue<float>();

        // Cast(目标, 源, 舍入模式, 元素个数)
        //   目标 bCos.Get<float>()：TBuf bCos 的 float 视图（32 个 float）；
        //   源   c[32] 开始的那 32 个 half（刚搬进来的 cos 原始数据）；
        //   RoundMode::CAST_NONE：不做特殊舍入，按 IEEE 规则直接转换。
        //        half → float 是「升精度」，完全无损，所以不需要指定舍入模式。
        Cast(bCos.Get<float>(), c[HDIM_HALF].ReinterpretCast<half>(), RoundMode::CAST_NONE, HDIM_HALF);

        // 归还 buffer。Alloc 和 Free 必须严格配对，
        // 否则队列槽位耗尽，下一次循环 Alloc 会卡死。
        qCos.FreeTensor(c);

        // ---- sin 的处理与 cos 完全同理 ----
        LocalTensor<float> s = qSin.AllocTensor<float>();       // 从 qSin 申请一块（64 个 float 空间）
        DataCopy(s[HDIM_HALF].ReinterpretCast<half>(), sinGm[bi * HDIM_HALF], HDIM_HALF);  // 搬 32 个 half 进后半段
        qSin.EnQue(s); s = qSin.DeQue<float>();                 // 入队再出队 = 等搬运完成
        Cast(bSin.Get<float>(), s[HDIM_HALF].ReinterpretCast<half>(), RoundMode::CAST_NONE, HDIM_HALF);  // half → float
        qSin.FreeTensor(s);                                     // 归还 buffer
    }

    // ========================================================================
    // LoadX：把 GM 上从 off 开始的 64 个 half 搬进 qX 队列。
    // 这是一个「生产者」函数：只 EnQue，不出队。
    // 出队（消费）由 Compute() 完成。
    // ========================================================================
    __aicore__ inline void LoadX(int64_t off)
    {
        // 从 qX 队列申请一块 UB。这次是 half 类型，容量 64 个 half（128 字节）。
        // 因为 x 本来就是 half，不需要 float 落地区，所以直接按 half 申请。
        LocalTensor<half> xl = qX.AllocTensor<half>();  // 从 qX 分配 half 空间

        // 搬运：目标 xl（UB），源 xGm[off]（GM），长度 64 个元素。
        // 因为 qX 深度是 2，这里可以连续调用两次而不会覆盖前一块数据。
        DataCopy(xl, xGm[off], HDIM);                    // 从 GM 搬入 64 个 half

        // 入队。注意：这里没有 DeQue，消费者在 Compute() 里。
        qX.EnQue(xl);                                    // 送入队列供 Compute 使用
    }

    // ========================================================================
    // Compute：PairwiseRotary 的核心计算。
    //   y[0:32]  = xLo * cos - xHi * sin
    //   y[32:64] = xLo * sin + xHi * cos
    // 其中 xLo = x[0:32]，xHi = x[32:64]，cos/sin 各 32 个。
    //
    // 输入：x 在 qX 队列里（half，64 个）；
    //       cos/sin 在 bCos/bSin 里（float，各 32 个）。
    // 输出：结果通过 qOut.EnQue 交给 Store() 搬出。
    //
    // 技巧：用 out 的后半段（out[32:64]）当临时变量，省一块 buffer。
    // ========================================================================
    __aicore__ inline void Compute()
    {
        // 出队：等 LoadX 的搬运完成，拿回当前行的 x。
        // 如果 DataCopy 还没搬完，这里会自动阻塞等待。
        LocalTensor<half> x = qX.DeQue<half>();

        // 拿到三个 TBuf 的类型化视图（只是句柄，不涉及数据拷贝）
        LocalTensor<float> xf = bXf.Get<float>();          // 64 个 float：x 转 float 后的存放处
        LocalTensor<float> cosLocal = bCos.Get<float>();   // 32 个 float：cos
        LocalTensor<float> sinLocal = bSin.Get<float>();   // 32 个 float：sin

        // 申请输出 buffer（从输出队列）。这块有 64 个 float 的空间 + 后面 64 个 half 的落地区。
        LocalTensor<float> out = qOut.AllocTensor<float>();

        // ---- 第一步：half → float ----
        // 把 64 个 half 的 x 转成 64 个 float 存进 xf。
        // 后面所有计算都在 float 上做，避免 half 精度不够。
        Cast(xf, x, RoundMode::CAST_NONE, HDIM);

        // PipeBarrier<PIPE_V>()：Vector 指令屏障。
        // 告诉 Vector 单元「等前面所有 Vector 指令执行完再往下走」。
        // 必须加，因为下一步 Mul 要读 xf，而 Cast 还没写完就会读到脏数据。
        // 判断规则：后一条指令读的 buffer，是不是前一条写的？是 → 加屏障。
        PipeBarrier<PIPE_V>();

        // ================= y[0:32] = xLo * cos - xHi * sin =================

        // Mul(目标, 源1, 源2, 个数)：逐元素相乘
        //   out[0..31] = xf[0..31] × cosLocal[0..31] = xLo × cos
        Mul(out, xf, cosLocal, HDIM_HALF);
        PipeBarrier<PIPE_V>();

        // Muls(目标, 源, 标量, 个数)：乘一个标量（Mul with scalar）
        //   out[32..63] = xf[32..63] × (-1) = -xHi
        //   out[HDIM_HALF] 就是 out[32]，表示从第 32 个元素开始写
        Muls(out[HDIM_HALF], xf[HDIM_HALF], -1.0f, HDIM_HALF);
        PipeBarrier<PIPE_V>();

        //   out[32..63] = out[32..63] × sinLocal = (-xHi) × sin
        // 源和目标可以是同一块内存，API 允许这种原地运算
        Mul(out[HDIM_HALF], out[HDIM_HALF], sinLocal, HDIM_HALF);
        PipeBarrier<PIPE_V>();

        // Add(目标, 源1, 源2, 个数)：逐元素相加
        //   out[0..31] = out[0..31] + out[32..63]
        //              = xLo×cos + (-xHi×sin) = xLo×cos - xHi×sin  ✔
        // 安全性：out[0..31] 和 out[32..63] 是两段不重叠的内存，源和目标不冲突
        Add(out, out, out[HDIM_HALF], HDIM_HALF);
        PipeBarrier<PIPE_V>();

        // ================ y[32:64] = xLo * sin + xHi * cos ================

        //   out[32..63] = xf[0..31] × sinLocal = xLo × sin
        // 注意源写的是 xf（从下标 0 开始），因为 xLo 存在 xf 的前半段
        Mul(out[HDIM_HALF], xf, sinLocal, HDIM_HALF);
        PipeBarrier<PIPE_V>();

        //   xf[32..63] = xf[32..63] × cosLocal = xHi × cos
        // 这里直接改写了 xf 的后半段。因为 xHi 已经用不到了
        // （上面算 -xHi 时已经用完），复用这块内存省空间。
        Mul(xf[HDIM_HALF], xf[HDIM_HALF], cosLocal, HDIM_HALF);
        PipeBarrier<PIPE_V>();

        //   out[32..63] = out[32..63] + xf[32..63]
        //               = xLo×sin + xHi×cos  ✔
        Add(out[HDIM_HALF], out[HDIM_HALF], xf[HDIM_HALF], HDIM_HALF);
        PipeBarrier<PIPE_V>();

        // 归还输入 buffer。此时 x 已经彻底用完（所有 Cast 都做完了）。
        qX.FreeTensor(x);
        // 把结果交给 Store() 消费（出队在 Store 里做）
        qOut.EnQue(out);
    }

    // ========================================================================
    // Store：把 qOut 里的计算结果从 UB 写回 GM 的 y[off .. off+63]。
    //
    // 同样用到「一物两用」的技巧：
    //   out[0..63] 是 64 个 float 的结果；
    //   out[64] 开始的空间被 ReinterpretCast 成 half，当搬运落地区。
    //   这就是 InitBuffer(qOut, 1, HDIM*sizeof(float)*2) 为什么要 ×2。
    // ========================================================================
    __aicore__ inline void Store(int64_t off)
    {
        // 出队：等 Compute 把结果算完（Compute 里 EnQue 了）
        auto out = qOut.DeQue<float>();                       // 获取计算结果

        // Cast(目标, 源, 舍入模式, 元素个数)：float → half
        //   目标 out[HDIM] = out[64]，字节偏移 64×4 = 256，
        //        ReinterpretCast<half>() 把这里开始的 128 字节当 64 个 half 用；
        //   源   out 本身（64 个 float）；
        //   RoundMode::CAST_RINT = Round to nearest INTeger，四舍五入。
        //        这是「降精度」，会丢信息，所以必须显式指定舍入模式（升精度才能用 CAST_NONE）。
        Cast(out[HDIM].ReinterpretCast<half>(), out,          // float → half
             RoundMode::CAST_RINT, HDIM);

        // 搬运：目标 yGm[off]（GM），源 out[64] 开始的 64 个 half（UB），长度 64 个元素
        DataCopy(yGm[off], out[HDIM].ReinterpretCast<half>(), // 写回 GM
                 HDIM);

        // 归还输出 buffer，让下一轮循环可以复用
        qOut.FreeTensor(out);                                  // 释放输出队列
    }

    // ============ 成员变量 ============
    TPipe* pipe_;                                // Ascend C 流水线管理器（外部传入，负责分配 UB）
    const PairwiseRotaryTilingData* tiling_;     // Host 传入的 Tiling 参数（只读）
    GlobalTensor<half> xGm, yGm, cosGm, sinGm;  // 4 个 GM 张量句柄（输入 x/cos/sin + 输出 y）

    // 输入队列。QuePosition::VECIN 表示「数据从 GM 搬入 UB 给 Vector 用」（MTE2 方向）。
    // qX 深度 = 2：双缓冲，能同时存在 2 块，实现「搬下一行」和「算当前行」并行。
    TQue<QuePosition::VECIN, 2> qX;             // x 输入队列，允许连续预加载 2 行
    // qCos/qSin 深度 = 1 就够：它们不是逐行变化的（一个 batch 才换一次），不需要双缓冲。
    TQue<QuePosition::VECIN, 1> qCos, qSin;     // cos/sin 加载队列

    // 输出队列。QuePosition::VECOUT 表示「Vector 算完搬出到 GM」（MTE3 方向）。
    // 深度 = 1：算完立刻搬走，搬出期间没有「下一块」要同时进来，所以不需要双缓冲。
    TQue<QuePosition::VECOUT, 1> qOut;          // 计算结果输出队列

    // TBuf：不带同步的裸 UB 内存，适合放「写一次、读多次」的中间结果。
    // TPosition::VECCALC 表示这是给 Vector 计算用的临时区。
    TBuf<TPosition::VECCALC> bXf;               // x 转 float 后的暂存 (64 floats)
    TBuf<TPosition::VECCALC> bCos;              // cos 转 float 后的暂存 (32 floats)
    TBuf<TPosition::VECCALC> bSin;              // sin 转 float 后的暂存 (32 floats)

    // 两个普通的整数状态变量（存在标量寄存器里，不占 UB）
    int64_t start_ = 0;                         // 当前 block 的起始元素偏移
    int64_t curRows_ = 0;                       // 当前 block 要处理的行数
};

#endif   // _PAIRWISE_ROTARY_H_   头文件保护结束

2.tanh(x) = (exp(x) - exp(-x)) / (exp(x) + exp(-x))   核心算子

#include "kernel_operator.h"
#include "tanh_custom_tiling.h"

// 每个队列的缓冲块数：2 表示双缓冲。
// 双缓冲让"搬运下一块数据(DMA)"与"计算当前块数据(Vector)"重叠执行，提升流水效率。
constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue

// Tanh 算子 kernel 实现。
// 整体流程：Init 切核/切块并申请 UB -> Process 循环(CopyIn->Compute->CopyOut)。
class KernelTanh {
public:
    __aicore__ inline KernelTanh() {}

    // 初始化：完成按核数据切分、GlobalTensor 绑定、UB(TBuf/TQue) 空间分配。
    // x/y        : 输入、输出在 GM(全局内存) 上的裸地址
    // totalLength: 总元素个数（来自 tiling）
    // tileNum    : 每核分块份数（来自 tiling）
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t totalLength, uint32_t tileNum)
    {
        // 每个 AI Core 平均分到的元素个数（本算子按元素均分，无尾部处理）。
        // GetBlockNum() 返回 host 侧 SetBlockDim 设置的核数（此处为 8）。
        this->blockLength = totalLength / AscendC::GetBlockNum();

        this->tileNum = tileNum;

        // 单次搬运/计算的元素个数。
        // 每核数据被切成 tileNum*BUFFER_NUM 块，因此需再除以 BUFFER_NUM。
        // 与 Process 中 loopCount = tileNum*BUFFER_NUM 严格对应：
        //   loopCount * tileLength == blockLength，正好覆盖本核全部数据。
        this->tileLength = this->blockLength / tileNum / BUFFER_NUM;

        // 将裸地址包装为 GlobalTensor，并按核号做连续切片：
        // 第 GetBlockIdx() 个核处理 [blockLength*idx, blockLength*idx + blockLength)。
        // 各核区间互不重叠，实现数据并行。
        xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
        yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);

        // 分配 UB 空间，大小单位为字节。
        // 输入/输出用 TQue：数据需跨 GM<->UB 搬运，依赖队列的 EnQue/DeQue 做同步。
        // DTYPE_X/DTYPE_Y 由 host 注册的 DataType 决定，此处为 half(fp16)。
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y));

        // 计算用的临时张量使用 TBuf<VECCALC>：只在 UB 内部参与计算、不参与搬运，
        // 因此无需队列管理。统一用 float 存放，保证中间计算精度。
        pipe.InitBuffer(tmpBuf0, this->tileLength * sizeof(float));
        pipe.InitBuffer(tmpBuf1, this->tileLength * sizeof(float));
        pipe.InitBuffer(tmpBuf2, this->tileLength * sizeof(float));
    }

    // 主处理流程：循环处理本核的所有分块。
    __aicore__ inline void Process()
    {
        // 循环次数 = 分块份数 * 双缓冲倍数，保证覆盖本核全部数据。
        int32_t loopCount = this->tileNum * BUFFER_NUM;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);   // 搬入第 i 块
            Compute(i);  // 计算第 i 块
            CopyOut(i);  // 搬出第 i 块
        }
    }

private:
    // 搬入：从 GM 拷贝一块数据到 UB 的输入队列。
    __aicore__ inline void CopyIn(int32_t progress)
    {
        // 从输入队列申请一块 UB 空间，再按偏移搬运第 progress 块。
        AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
        AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
        // 入队：通知后续消费方（Compute）该块数据已就绪，并触发必要的同步。
        inQueueX.EnQue(xLocal);
    }

    // 计算：在 UB 上完成 tanh(x) = (e^x - e^-x) / (e^x + e^-x)。
    // 使用 fp32 中间精度，最后再转回 fp16 输出。
    __aicore__ inline void Compute(int32_t progress)
    {
        // 从输入队列取出已就绪的数据块；从输出队列申请结果空间。
        AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
        AscendC::LocalTensor<DTYPE_Y> yLocal = outQueueY.AllocTensor<DTYPE_Y>();

        // 取出三块临时计算 buffer（复用，故 3 块足够，见下方推导）。
        AscendC::LocalTensor<float> tmp0 = tmpBuf0.Get<float>();
        AscendC::LocalTensor<float> tmp1 = tmpBuf1.Get<float>();
        AscendC::LocalTensor<float> tmp2 = tmpBuf2.Get<float>();

        // tmp0 = float(x)：fp16 精度不足，先升到 fp32 再做 Exp/Div。
        AscendC::Cast(tmp0, xLocal, AscendC::RoundMode::CAST_NONE, this->tileLength);
        // tmp1 = e^x
        AscendC::Exp(tmp1, tmp0, this->tileLength);
        // tmp2 = 1.0
        AscendC::Duplicate(tmp2, (float)1.0, this->tileLength);
        // tmp0 = 1 / e^x = e^-x （此时 tmp0 中的 x 已不再需要，可安全覆盖）
        AscendC::Div(tmp0, tmp2, tmp1, this->tileLength);
        // tmp2 = e^x - e^-x （分子）
        AscendC::Sub(tmp2, tmp1, tmp0, this->tileLength);
        // tmp1 = e^x + e^-x （分母）
        AscendC::Add(tmp1, tmp1, tmp0, this->tileLength);
        // tmp0 = 分子 / 分母 = tanh(x)
        AscendC::Div(tmp0, tmp2, tmp1, this->tileLength);
        // 结果转回 fp16，CAST_RINT 表示就近取整，误差小于直接截断。
        AscendC::Cast(yLocal, tmp0, AscendC::RoundMode::CAST_RINT, this->tileLength);

        // 结果入队，交给 CopyOut；同时释放输入 buffer 供下一轮复用。
        outQueueY.EnQue(yLocal);
        inQueueX.FreeTensor(xLocal);
    }

    // 搬出：把计算结果从 UB 的输出去队列拷贝回 GM。
    __aicore__ inline void CopyOut(int32_t progress)
    {
        AscendC::LocalTensor<DTYPE_Y> yLocal = outQueueY.DeQue<DTYPE_Y>();
        AscendC::DataCopy(yGm[progress * this->tileLength], yLocal, this->tileLength);
        // 释放输出 buffer，供双缓冲下一轮复用。
        outQueueY.FreeTensor(yLocal);
    }

private:
    AscendC::TPipe pipe;                                                             // UB 内存管理器
    AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX;                 // 输入队列(VECIN)
    AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueY;               // 输出队列(VECOUT)
    AscendC::TBuf<AscendC::QuePosition::VECCALC> tmpBuf0, tmpBuf1, tmpBuf2;          // 计算临时 buffer
    AscendC::GlobalTensor<DTYPE_X> xGm;                                              // 输入 GM 张量
    AscendC::GlobalTensor<DTYPE_Y> yGm;                                              // 输出 GM 张量
    uint32_t blockLength;                                                            // 每核处理的元素数
    uint32_t tileNum;                                                                // 每核分块份数
    uint32_t tileLength;                                                             // 单次搬运/计算元素数
};

// kernel 入口：由框架在 AI Core 上调用。
// workspace 本算子未使用；tiling 指向 host 下发的切分参数。
extern "C" __global__ __aicore__ void tanh_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    // 注册默认 tiling 结构体类型，并把它从 tiling 内存反序列化出来。
    REGISTER_TILING_DEFAULT(TanhCustomTilingData);
    GET_TILING_DATA(tilingData, tiling);

    // 用 host 计算好的切分参数初始化并执行。
    KernelTanh op;
    op.Init(x, y, tilingData.totalLength, tilingData.tileNum);
    op.Process();
}

3.算子的 Tiling 数据结构定义

 * Tanh 算子的 Tiling 数据结构定义。
 *
 * 该结构体是 host 侧（TilingFunc）与 kernel 侧（tanh_custom 核函数）之间
 * 传递切分参数的唯一契约：
 *   - host 侧计算好参数后写入这块内存，随算子下发一起拷贝到 device；
 *   - kernel 侧通过 GET_TILING_DATA 反序列化后读取使用。
 * 因此两边必须包含同一份定义（host 通过相对路径 include 本文件）。
 *
 * 说明：题目要求直接定义结构体，无需套用 BEGIN_TILING_DATA_DEF 等宏，
 * 与之配套的注册/解析宏为 REGISTER_TILING_DEFAULT + GET_TILING_DATA。
 */

#ifndef TANH_CUSTOM_TILING_H
#define TANH_CUSTOM_TILING_H
#include <cstdint>

struct TanhCustomTilingData {
    // 参与计算的总元素个数（由 host 侧根据输入 shape 求得，如 8*2048=16384）。
    // kernel 只能拿到裸指针，无法自行获知 shape，必须靠该字段确定处理范围。
    uint32_t totalLength;
    // 每个核内的分块份数（不含双缓冲倍数 BUFFER_NUM）。
    // 用于控制 UB 单次搬运/计算的粒度：UB 容量有限，需把每核数据切成小块循环处理。
    // 该值由 host 决定并可随硬件调优，因此通过 tiling 下发而非写死在 kernel 中。
    uint32_t tileNum;
};
#endif

4.算子注册

// 引入 kernel 侧的 tiling 结构体定义：host 与 kernel 必须共用同一份定义，
// 这里 host 通过相对路径直接 include op_kernel 目录下的头文件。
#include "../op_kernel/tanh_custom_tiling.h"
#include "register/op_def_registry.h"


// Tiling 实现：在算子下发前运行，负责决定"数据怎么切分"，
// 并把结果写入 tiling 内存，供 device 侧 kernel 读取。
namespace optiling {
const uint32_t BLOCK_DIM = 8;  // 启动的 AI Core 数量（核数），决定并行度
const uint32_t TILE_NUM = 8;   // 每个核内的分块份数（不含双缓冲倍数）

static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    // 框架已预分配一块 raw tiling buffer，这里直接按我们的结构体解释，
    // 避免额外的一次内存拷贝，得到一个可写指针。
    TanhCustomTilingData* tiling = reinterpret_cast<TanhCustomTilingData*>(context->GetRawTilingData()->GetData());

    // 取输入张量的元素总数（返回 int64_t），转成结构体字段类型 uint32_t。
    // 例：shape 8x2048 -> totalLength = 16384。
    uint32_t totalLength = static_cast<uint32_t>(context->GetInputShape(0)->GetOriginShape().GetShapeSize());

    // 告诉运行时本算子启动 BLOCK_DIM 个核。
    // kernel 侧 GetBlockNum() 将返回该值、GetBlockIdx() 返回 0..BLOCK_DIM-1。
    context->SetBlockDim(BLOCK_DIM);

    // 回填切分参数（先写数据，再设置数据大小，顺序不能颠倒）。
    tiling->totalLength = totalLength;
    tiling->tileNum = TILE_NUM;

    // 设置实际有效的 tiling 数据长度，决定往 device 拷贝多少字节。
    context->GetRawTilingData()->SetDataSize(sizeof(TanhCustomTilingData));

    // tiling 接口约定至少要暴露一个 workspace 槽位。
    // Tanh 是纯 element-wise 运算，无需额外显存，故设为 0。
    size_t* currentWorkspace = context->GetWorkspaceSizes(1);
    currentWorkspace[0] = 0;

    return ge::GRAPH_SUCCESS;
}
}


// Shape / 数据类型推导：用于图模式下的形状与类型推断。
namespace ge {
static ge::graphStatus InferShape(gert::InferShapeContext* context)
{
    // Tanh 为逐元素运算，输出形状与输入完全一致，直接拷贝即可。
    const gert::Shape* x1_shape = context->GetInputShape(0);
    gert::Shape* y_shape = context->GetOutputShape(0);
    *y_shape = *x1_shape;
    return GRAPH_SUCCESS;
}
static ge::graphStatus InferDataType(gert::InferDataTypeContext *context)
{
    // 输出数据类型与输入保持一致（此处均为 float16）。
    const auto inputDataType = context->GetInputDataType(0);
    context->SetOutputDataType(0, inputDataType);
    return ge::GRAPH_SUCCESS;
}
}


// 算子原型注册：定义输入/输出的名字、数量、数据类型、格式，
// 并绑定 shape/type 推导函数与 tiling 函数。
namespace ops {
class TanhCustom : public OpDef {
public:
    explicit TanhCustom(const char* name) : OpDef(name)
    {
        // 输入 x：必选，float16，ND 格式（含未知 shape 场景）。
        // 框架据此在编译 kernel 时把 DTYPE_X 替换为 half。
        this->Input("x")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND});
        // 输出 y：必选，float16，ND 格式；对应 kernel 中的 DTYPE_Y。
        this->Output("y")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND});

        // 绑定 shape 与数据类型推导函数。
        this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType);

        // 绑定 Tiling 实现，并声明该算子支持的硬件平台。
        this->AICore()
            .SetTiling(optiling::TilingFunc);
        this->AICore().AddConfig("ascend910b");

    }
};

// 向框架注册该算子。
OP_ADD(TanhCustom);
}

5.逐元素型模板

// ============================================================================
//  模板 1：逐元素型 Vector 算子骨架
//  适用：Add / Sub / Mul / Div / Relu / Abs / Exp / Ln / Sqrt / Gelu / Sigmoid / Cast
//  特征：输出 shape == 输入 shape，且 y[i] 只依赖 x[i]
// ============================================================================

// ---- 头文件保护：防止重复 include ----
#ifndef _ELEMWISE_H_
#define _ELEMWISE_H_

// ---- 三个固定 include ----
#include "kernel_operator.h"                // Ascend C 全部 API
#include "kernel_tiling/kernel_tiling.h"    // tiling 相关宏
#include "elemwise_tiling_data.h"           // host 定义的 tiling 结构体（名字按题目改）

// ---- 展开 AscendC 命名空间，之后可以直接写 DataCopy(...) ----
using namespace AscendC;

// ---- 每次搬进 UB 处理的元素个数 ----
// 两个约束：
//   ① 对齐：TILE × sizeof(T) 必须是 32 字节整数倍
//            half(2B) → TILE 取 16 的倍数；float(4B) → TILE 取 8 的倍数
//   ② 容量：所有 InitBuffer 字节数之和 < UB（910B 约 192KB）
// 拿不准就取 2048
constexpr int64_t TILE = 2048;

class KernelElemwise {
public:
    // ---- 构造函数：保存 pipe 和 tiling ----
    // 每个 __aicore__ 成员函数都必须带 inline（AI Core 上函数调用开销大）
    __aicore__ inline KernelElemwise(TPipe* pipe, const ElemwiseTilingData* tiling)
        : pipe_(pipe), tiling_(tiling) {}    // 初始化列表，比函数体赋值高效

    // ---- Init：① 绑 GM ② 算本核范围 ③ 分配 UB ----
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y)
    {
        // ① 绑 GM。GM_ADDR 是 __gm__ uint8_t*，必须强转成具体类型才能下标访问
        xGm.SetGlobalBuffer((__gm__ half*)x);   // 输入
        yGm.SetGlobalBuffer((__gm__ half*)y);   // 输出

        // ② 算本核负责哪一段
        int64_t total   = tiling_->totalLen;     // 总元素数（host 算好传进来）
        int64_t perCore = tiling_->lenPerCore;   // 每核元素数（host 算好传进来）

        start_ = GetBlockIdx() * perCore;        // 核号 × 每核元素数 = 本核起点
        len_   = total - start_;                 // 从起点到末尾还剩多少
        if (len_ > perCore) len_ = perCore;      // 不能超过配额（只有尾核不满）
        if (len_ < 0) len_ = 0;                  // 核数多于数据量时归零

        // ③ 分配 UB。InitBuffer(队列对象, 深度, 每块字节数)
        pipe_->InitBuffer(qIn,  2, TILE * sizeof(half));  // 深度 2 = 双缓冲
        pipe_->InitBuffer(qOut, 1, TILE * sizeof(half));  // 深度 1 够用
    }

    // ---- Process：主流程 ----
    __aicore__ inline void Process()
    {
        if (len_ <= 0) return;                   // 边界：本核没活干，直接退出

        // ---- 核内分块循环 ----
        // 为什么要分块？因为本核的 len_ 可能远大于 UB 能装下的量
        for (int64_t off = 0; off < len_; off += TILE) {
            // 尾块截断：最后一块可能不足 TILE 个
            // 例：len_=5000, TILE=2048 → 第 3 块只有 5000-4096=904 个
            int64_t cur = (len_ - off > TILE) ? TILE : len_ - off;

            // 三步流水。注意传「全局偏移」= start_ + off，不是循环变量 off
            CopyIn(start_ + off, cur);           // GM → UB
            Compute(cur);                        // UB 内计算
            CopyOut(start_ + off, cur);          // UB → GM
        }
    }

private:
    // ---- CopyIn：搬入 ----
    __aicore__ inline void CopyIn(int64_t off, int64_t n)
    {
        LocalTensor<half> x = qIn.AllocTensor<half>();  // 从队列申请一块 UB
        DataCopy(x, xGm[off], n);                       // 搬 n 个元素（内部自动换算字节）
        qIn.EnQue(x);                                   // 入队：交给 MTE2 去搬
    }

    // ---- Compute：★ 唯一需要按题目改的地方 ★ ----
    __aicore__ inline void Compute(int64_t n)
    {
        LocalTensor<half> x = qIn.DeQue<half>();        // 出队：等搬运完成（会阻塞）
        LocalTensor<half> y = qOut.AllocTensor<half>(); // 申请输出 buffer

        // ================= 在这里写题目公式 =================
        //  y = relu(x)   →  Maxs(y, x, (half)0, n);
        //  y = x * 2     →  Muls(y, x, (half)2, n);
        //  y = x + b     →  Add(y, x, bLocal, n);
        //  y = |x|       →  Abs(y, x, n);
        //  y = e^x       →  Exp(y, x, n);
        //
        //  需要精度时先升 float：
        //    Cast(xf, x, RoundMode::CAST_NONE, n);   // half → float
        //    ... 在 float 上算 ...
        //    Cast(y, yf, RoundMode::CAST_RINT, n);   // float → half
        //
        //  多条有依赖的指令之间必须加 PipeBarrier<PIPE_V>();
        //  判断规则：后一条读的 buffer，是不是前一条写的？
        // ====================================================

        qOut.EnQue(y);                                  // 交给 Store 搬出
        qIn.FreeTensor(x);                              // 归还输入（Alloc/Free 必须配对）
    }

    // ---- CopyOut：搬出 ----
    __aicore__ inline void CopyOut(int64_t off, int64_t n)
    {
        LocalTensor<half> y = qOut.DeQue<half>();       // 出队：等算完
        DataCopy(yGm[off], y, n);                       // 写回 GM
        qOut.FreeTensor(y);                             // 归还 buffer
    }

    // ---- 成员变量 ----
    TPipe* pipe_;                          // 流水线管理器（外部传入）
    const ElemwiseTilingData* tiling_;     // tiling 参数（外部传入，只读）
    GlobalTensor<half> xGm, yGm;           // GM 张量句柄
    TQue<QuePosition::VECIN,  2> qIn;      // 输入队列（GM→UB），深度 2
    TQue<QuePosition::VECOUT, 1> qOut;     // 输出队列（UB→GM），深度 1
    int64_t start_ = 0;                    // 本核起始元素偏移
    int64_t len_   = 0;                    // 本核元素个数
};

#endif

6.规约型模板

// ============================================================================
//  模板 2：规约型 Vector 算子骨架
//  适用：ReduceSum / ReduceMax / ReduceMin / ReduceMean / ArgMax
//  特征：x [rows, cols] → y [rows]，输出少一维
//  与模板 1 的三个区别：
//    ① 切分单位是「行」不是「元素」
//    ② 循环粒度是「一次一行」
//    ③ 搬出长度是 1，必须用 DataCopyPad（对齐！）
// ============================================================================

#ifndef _REDUCE_H_
#define _REDUCE_H_

#include "kernel_operator.h"
#include "kernel_tiling/kernel_tiling.h"
#include "reduce_tiling_data.h"

using namespace AscendC;

class KernelReduce {
public:
    __aicore__ inline KernelReduce(TPipe* pipe, const ReduceTilingData* tiling)
        : pipe_(pipe), tiling_(tiling) {}

    // ---- Init：注意切分粒度是「行」 ----
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y)
    {
        xGm.SetGlobalBuffer((__gm__ half*)x);   // 输入 [rows, cols]
        yGm.SetGlobalBuffer((__gm__ half*)y);   // 输出 [rows]

        int64_t rows       = tiling_->rows;        // 总行数
        int64_t cols       = tiling_->cols;        // 每行列数 = 规约长度
        int64_t rowPerCore = tiling_->rowPerCore;  // 每核处理多少行（host 算好）

        // 和模板 1 同样的套路，只是单位从「元素」换成「行」
        rowStart_ = GetBlockIdx() * rowPerCore;    // 本核起始行号
        int64_t rowEnd = rowStart_ + rowPerCore;   // 本核结束行号（不含）
        if (rowEnd > rows) rowEnd = rows;          // 尾核截断
        rowNum_ = rowEnd - rowStart_;              // 本核要处理的行数
        if (rowNum_ < 0) rowNum_ = 0;              // 边界归零

        // 输入队列：每块装一整行（cols 个 half），深度 2 做双缓冲
        pipe_->InitBuffer(qIn,  2, cols * sizeof(half));

        // 输出队列：申请 64 字节而不是 2 字节。
        // 原因见 CopyOut —— 需要一块 32 字节对齐的区域放 half 结果。
        pipe_->InitBuffer(qOut, 1, 64);

        // bXf：half 的 x 转 float 后存放（cols 个 float）。
        // 为什么必须转 float？规约是「累加几千个数」，half 累加精度崩得厉害。
        pipe_->InitBuffer(bXf, cols * sizeof(float));

        // bWork：ReduceSum 需要一个临时工作区放中间结果，长度和输入一样
        pipe_->InitBuffer(bWork, cols * sizeof(float));
    }

    // ---- Process：一行一行处理 ----
    __aicore__ inline void Process()
    {
        if (rowNum_ <= 0) return;               // 边界

        int64_t cols = tiling_->cols;

        // ---- 按行循环，每次处理一整行 ----
        // 注意：这里不是 off += TILE，而是 r++
        for (int64_t r = 0; r < rowNum_; r++) {
            // 本行第一个元素在 GM 中的偏移 = 行号 × 每行列数
            int64_t off = (rowStart_ + r) * cols;

            CopyIn(off, cols);        // 搬入一整行 cols 个元素
            Compute(cols);            // 规约成 1 个数
            CopyOut(rowStart_ + r);   // 写回 1 个数
        }
    }

private:
    // ---- CopyIn：搬入一整行 ----
    __aicore__ inline void CopyIn(int64_t off, int64_t n)
    {
        LocalTensor<half> x = qIn.AllocTensor<half>();  // 申请 UB
        DataCopy(x, xGm[off], n);                       // n = cols，搬一整行
        qIn.EnQue(x);                                   // 入队
    }

    // ---- Compute：把一整行规约成 1 个数 ----
    __aicore__ inline void Compute(int64_t cols)
    {
        LocalTensor<half>  x   = qIn.DeQue<half>();     // 取出本行的 cols 个 half
        LocalTensor<float> out = qOut.AllocTensor<float>();  // 申请输出（只要 1 个数）

        // ---- 第一步：half → float ----
        LocalTensor<float> xf = bXf.Get<float>();
        Cast(xf, x, RoundMode::CAST_NONE, cols);
        // 下一句 ReduceSum 要读 xf，必须等 Cast 写完
        PipeBarrier<PIPE_V>();

        // ---- 第二步：规约 ----
        // ReduceSum(目标, 源, 工作区, 个数)
        // 把 xf 里的 cols 个数全部加起来，结果写 1 个 float 到 out[0]
        ReduceSum(out, xf, bWork.Get<float>(), cols);
        // 下面要 EnQue 交给 MTE3，先等 Vector 算完
        PipeBarrier<PIPE_V>();

        qOut.EnQue(out);
        qIn.FreeTensor(x);      // 输入用完了（必须在 Cast 之后才 Free）
    }

    // ---- CopyOut：把 1 个数写回 GM ----
    __aicore__ inline void CopyOut(int64_t row)
    {
        LocalTensor<float> out = qOut.DeQue<float>();   // 取出结果

        // ★ 关键点：float → half，目标地址必须 32 字节对齐
        //   out 是 float 数组，每个 float 占 4 字节
        //   out[8] 的字节偏移 = 8 × 4 = 32，正好对齐 ✔
        //   ReinterpretCast<half>() 把这块内存重新解释成 half 类型
        LocalTensor<half> outH = out[8].ReinterpretCast<half>();

        // 把 out[0] 那 1 个 float 转成 1 个 half，写进 out[8] 开始的位置
        // 源（byte 0）和目标（byte 32）不重叠，安全
        // CAST_RINT = 四舍五入（降精度必须指定舍入模式）
        Cast(outH, out, RoundMode::CAST_RINT, 1);

        // ★ 为什么不能用 DataCopy？
        //   DataCopy 要求长度是 32 字节整数倍，
        //   这里只写 1 个 half = 2 字节，不满足 → 必须用 DataCopyPad
        // DataCopyExtParams{blockCount, blockLen(字节), srcStride, dstStride, rsv}
        DataCopyExtParams params{
            1,                        // 1 个连续块
            (uint32_t)sizeof(half),   // 块长度 = 2 字节
            0, 0, 0                   // stride 全 0（连续搬运）
        };
        DataCopyPad(yGm[row], outH, params);

        qOut.FreeTensor(out);   // 归还
    }

    // ---- 成员变量 ----
    TPipe* pipe_;
    const ReduceTilingData* tiling_;
    GlobalTensor<half> xGm, yGm;
    TQue<QuePosition::VECIN,  2> qIn;     // 输入队列，深度 2
    TQue<QuePosition::VECOUT, 1> qOut;    // 输出队列，深度 1
    TBuf<TPosition::VECCALC> bXf;         // x 转 float 后的存放处
    TBuf<TPosition::VECCALC> bWork;       // ReduceSum 的临时工作区
    int64_t rowStart_ = 0;                // 本核起始行号
    int64_t rowNum_   = 0;                // 本核要处理的行数
};

#endif

7.广播型（只讲与逐元素型的区别）

// ---- 改动 1：Init 里多分配一块 TBuf 存广播源 ----
//   如果广播源是 [D]（每行一样）或 [B,1,1,D]（每 batch 一样），
//   把它常驻在 TBuf 里，循环里反复用，避免重复读 GM
pipe_->InitBuffer(bBias, BDIM * sizeof(float));

// ---- 改动 2：Process 里循环外先加载一次 ----
LoadBiasOnce();                       // ← 只加载一次
for (int64_t off = 0; off < len_; off += TILE) {
    // 如果是「每 batch 一样」（本题的 cos/sin），这里还要判断是否跨 batch：
    //   int64_t bi = ((start_ + off) / ROWLEN) / (nh * sl);
    //   if (bi != lastBi) { LoadBiasOnce(bi); lastBi = bi; }
    CopyIn(start_ + off, cur);
    Compute(cur);
    CopyOut(start_ + off, cur);
}

// ---- 改动 3：Compute 里直接用常驻的广播源 ----
//   Mul(y, x, bBias.Get<float>(), n);   // bBias 不用每次搬

8.两遍扫描型（以softmax为例）

// ============================================================================
//  模板 4：两遍扫描型 Vector 算子骨架（以 Softmax 为例）
//  适用：Softmax / LayerNorm / RMSNorm / Normalize
//  特征：y[i] 依赖整行的统计量（max / sum / mean / var），必须先扫完一整行
//  结构 = 模板 2（按行循环）+ 额外的第二遍计算
// ============================================================================

// ---- Init 与模板 2 相同，额外多一块 buffer 放统计量 ----
// pipe_->InitBuffer(bRowMax, 32);    // 放整行 max（1 个 float + 对齐余量）
// pipe_->InitBuffer(bRowSum, 32);    // 放整行 sum
// pipe_->InitBuffer(bTmp,    32);    // 放广播后的统计量

// ---- Process 与模板 2 相同：for (r = 0; r < rowNum_; r++) { ... } ----

// ---- Compute：两遍扫描 ----
__aicore__ inline void Compute(int64_t cols)
{
    LocalTensor<half>  x   = qIn.DeQue<half>();
    LocalTensor<float> out = qOut.AllocTensor<float>();

    // ===== 第 0 步：half → float =====
    LocalTensor<float> xf = bXf.Get<float>();
    Cast(xf, x, RoundMode::CAST_NONE, cols);
    PipeBarrier<PIPE_V>();

    // ===== 第一遍：求整行的 max =====
    // 为什么要减 max？因为 exp(x) 在 x 稍大时就会溢出（half 上限 65504）。
    // 减去 max 后最大值为 exp(0)=1，数值稳定。
    LocalTensor<float> rowMax = bRowMax.Get<float>();
    ReduceMax(rowMax, xf, bWork.Get<float>(), cols);   // cols 个数 → 1 个 max
    PipeBarrier<PIPE_V>();

    // ===== 第二遍开始：减 max =====
    // Sub(xf, xf, rowMax, cols) 里 rowMax 只有 1 个数，
    // 但 API 会自动把它广播到整个向量（这是 Sub 的广播语义）
    Sub(xf, xf, rowMax, cols);
    PipeBarrier<PIPE_V>();

    // ===== 求指数 =====
    Exp(xf, xf, cols);
    PipeBarrier<PIPE_V>();

    // ===== 求整行的 sum =====
    LocalTensor<float> rowSum = bRowSum.Get<float>();
    ReduceSum(rowSum, xf, bWork.Get<float>(), cols);   // cols 个数 → 1 个 sum
    PipeBarrier<PIPE_V>();

    // ===== 归一化：除以 sum =====
    // 把 1 个 float 广播成 cols 个相同的值，才能做逐元素除法
    LocalTensor<float> sumVec = bTmp.Get<float>();
    Duplicate(sumVec, rowSum.GetValue(0), cols);   // 标量 → 向量
    PipeBarrier<PIPE_V>();

    Div(out, xf, sumVec, cols);                     // y = exp(x-max) / sum
    PipeBarrier<PIPE_V>();

    qOut.EnQue(out);
    qIn.FreeTensor(x);
}

9.算子原型定义（ op_host/xxx_def.cp）

// ============================================================================
//  算子原型定义：声明「算子长什么样」
//  最常考：改输入个数 / 改数据类型 / 改芯片型号
// ============================================================================

#include "register/op_def_registry.h"   // 提供 OpDef 基类和 OP_ADD 宏

namespace ops {                          // 算子定义必须放在 ops 命名空间

// 类名 = 算子名，必须继承 OpDef
class PairwiseRotary : public OpDef {
public:
    // 参数 name 是算子在框架里的注册名
    explicit PairwiseRotary(const char* name) : OpDef(name)
    {
        // ---- 声明输入 1：x ----
        this->Input("x")                          // ★ 输入名，顺序必须和 kernel 入口参数一致
            .ParamType(REQUIRED)                  // REQUIRED 必选 / OPTIONAL 可选 / DYNAMIC 动态个数
            .DataType({ge::DT_FLOAT16})           // ★ 支持的类型列表（可写多种，框架按序匹配）
            .Format({ge::FORMAT_ND})              // 排布：ND 普通 / NCHW / NHWC ...
            .UnknownShapeFormat({ge::FORMAT_ND})  // 动态 shape 时用的排布
            .AutoContiguous();                    // 自动把非连续输入转成连续，省得自己处理 stride

        // ---- 声明输入 2：cos（结构同上） ----
        this->Input("cos")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND})
            .AutoContiguous();

        // ---- 声明输入 3：sin ----
        this->Input("sin")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND})
            .AutoContiguous();

        // ---- 声明输出：y ----
        // 输出不需要 AutoContiguous（框架分配的显存一定连续）
        this->Output("y")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND});

        // ---- AI Core 配置 ----
        OpAICoreConfig aicoreConfig;
        aicoreConfig.DynamicCompileStaticFlag(true)   // 动态编译静态化（编译期能定就定死）
            .DynamicFormatFlag(false)                 // 不支持动态 format
            .DynamicRankSupportFlag(true)             // 支持动态 rank（维度个数不定）
            .DynamicShapeSupportFlag(true)            // 支持动态 shape（维度大小不定）
            .NeedCheckSupportFlag(false)              // 不做支持度校验（省编译时间）
            .PrecisionReduceFlag(true);               // 允许精度降低（可用快速数学指令）

        // ★ 把配置挂到芯片上。要换芯片就改这个字符串（ascend910b / ascend910_93）
        this->AICore().AddConfig("ascend910b", aicoreConfig);
    }
};

// ---- 注册算子，宏展开生成框架需要的注册代码 ----
OP_ADD(PairwiseRotary);
} // namespace ops

10.op_host/xxx_infershape.cpp —— 形状推导（作用：图编译期告诉框架"输出该分配多大显存"。规约型必须改这个文件。）

// ============================================================================
//  形状推导：图编译期决定输出张量的形状
//  逐元素型：y.shape = x.shape（就一句）
//  规约型：要去掉/改小被规约的那一维
// ============================================================================

#include "register/op_impl_registry.h"   // IMPL_OP_INFERSHAPE 宏
#include "op_common/log/log.h"           // OP_CHECK_NULL_WITH_CONTEXT / OP_LOGE

using namespace ge;
namespace ops {

// 函数名习惯：InferShape4<算子名>
graphStatus InferShape4PairwiseRotary(gert::InferShapeContext* context)
{
    // 取第 0 个输入的 shape（x）
    const gert::Shape* xInputShape = context->GetInputShape(0);
    OP_CHECK_NULL_WITH_CONTEXT(context, xInputShape);   // 空指针检查

    // 取第 0 个输出的 shape，拿到的是可写的形状对象
    gert::Shape* yShape = context->GetOutputShape(0);
    OP_CHECK_NULL_WITH_CONTEXT(context, yShape);

    // ★ 核心：逐元素型就这一句
    *yShape = *xInputShape;

    return GRAPH_SUCCESS;
}

// 注册到算子上
IMPL_OP_INFERSHAPE(PairwiseRotary).InferShape(InferShape4PairwiseRotary);
} // namespace ops

11.op_host/xxx_tiling.cpp —— Tiling 切分（host 侧最常考）（作用：把大问题切成 coreNum 份，填 tiling 结构体，告诉框架用几个核。）

// ============================================================================
//  Tiling（切分）—— host 侧最重要的一段
//  文件位置：op_host/pairwise_rotary_tiling.cpp
// ============================================================================

#include "register/op_def_registry.h"
#include "op_common/log/log.h"
#include "op_common/op_host/util/math_util.h"         // CeilDiv 等
#include "op_common/op_host/util/platform_util.h"     // PlatformAscendC
#include "../op_kernel/pairwise_rotary_tiling_data.h" // tiling 结构体（host/kernel 共用）

namespace optiling {

using Ops::Base::CeilDiv;   // 向上取整除法：CeilDiv(10,3) = 4

// ---- 常量 ----
constexpr int64_t X_INDEX = 0;                  // 输入 x 的下标
constexpr int64_t COS_INDEX = 1;                // 输入 cos 的下标
constexpr int64_t DEFAULT_WORKSPACE_SIZE = 64;  // workspace 预留 64 字节（本算子用不到）
constexpr int64_t HIDDEN_DIM = 64;              // x 最后一维固定 64
constexpr int64_t COS_DIM = 32;                 // cos/sin 最后一维固定 32
constexpr size_t  WORKSPACE_NUM = 1;            // workspace 块数

// ============================================================================
//  工具函数 1：拿平台信息（核数、UB 大小）
// ============================================================================
static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)
{
    auto platformInfoPtr = context->GetPlatformInfo();   // 拿平台描述
    OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);

    // 包装成 AscendC 平台对象，提供便捷查询接口
    auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);

    *coreNum = ascendcPlatform.GetCoreNumAiv();          // ★ 查询 AI Vector 核数量
    OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);

    // ★ 查询 UB 字节数（用来算 tile 大小能不能装下）
    ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize);
    OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);

    return ge::GRAPH_SUCCESS;
}

// ============================================================================
//  工具函数 2：设置 workspace 大小
// ============================================================================
static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
{
    // workspace 是一块框架分配的临时显存，多核之间交换数据时用。
    // 本算子是纯逐元素/逐行处理，核间不需要通信，所以给最小值即可。
    size_t* currentWorkspace = context->GetWorkspaceSizes(WORKSPACE_NUM);
    OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
    currentWorkspace[0] = DEFAULT_WORKSPACE_SIZE;
    return ge::GRAPH_SUCCESS;
}

// ============================================================================
//  主函数：Tiling 逻辑
// ============================================================================
static ge::graphStatus PairwiseRotaryTilingFunc(gert::TilingContext* context)
{
    // ---- 第一步：拿平台信息 ----
    uint64_t ubSize;
    int64_t  coreNum;
    OP_CHECK_IF(GetPlatformInfo(context, &ubSize, &coreNum) != ge::GRAPH_SUCCESS,
                OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);

    // ---- 第二步：拿输入 shape ----
    auto xShape = context->GetInputShape(X_INDEX);
    OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
    auto cosShape = context->GetInputShape(COS_INDEX);
    OP_CHECK_NULL_WITH_CONTEXT(context, cosShape);

    // GetStorageShape()：拿「实际存储」的 shape（动态 shape 时比 GetOriginShape 更准）
    auto xGs   = xShape->GetStorageShape();
    auto cosGs = cosShape->GetStorageShape();

    // ---- 第三步：校验 shape，不合法立刻报错返回 ----
    int64_t xDimNum = xGs.GetDimNum();     // 维度个数
    OP_CHECK_IF(xDimNum != 4, OP_LOGE(context, "x must be 4D (B,N,S,D)"), return ge::GRAPH_FAILED);

    int64_t batchSize = xGs.GetDim(0);     // B
    int64_t numHead   = xGs.GetDim(1);     // N
    int64_t seqLength = xGs.GetDim(2);     // S
    int64_t hiddenDim = xGs.GetDim(3);     // D
    OP_CHECK_IF(hiddenDim != HIDDEN_DIM, OP_LOGE(context, "hiddenDim must be 64"), return ge::GRAPH_FAILED);

    int64_t cosDimNum = cosGs.GetDimNum();
    OP_CHECK_IF(cosDimNum != 4, OP_LOGE(context, "cos must be 4D"), return ge::GRAPH_FAILED);
    int64_t cosB = cosGs.GetDim(0);
    int64_t cosD = cosGs.GetDim(3);
    OP_CHECK_IF(cosD != COS_DIM, OP_LOGE(context, "cos last dim must be 32"), return ge::GRAPH_FAILED);
    OP_CHECK_IF(cosB != batchSize, OP_LOGE(context, "cos batch must equal x batch"), return ge::GRAPH_FAILED);

    // ---- 第四步：设置 workspace ----
    OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
                OP_LOGE(context, "GetWorkspaceSize error"), return ge::GRAPH_FAILED);

    // ---- 第五步：拿 tiling 结构体的写入位置 ----
    // GetTilingData<T>() 返回指向「框架预分配的 tiling 内存」的指针
    PairwiseRotaryTilingData* tilingData = context->GetTilingData<PairwiseRotaryTilingData>();
    OP_CHECK_NULL_WITH_CONTEXT(context, tilingData);

    // 先清零，避免残留脏数据。memset_s 是安全版本，返回 EOK 表示成功
    OP_CHECK_IF(memset_s(tilingData, sizeof(PairwiseRotaryTilingData), 0, sizeof(PairwiseRotaryTilingData)) != EOK,
                OP_LOGE(context, "memset_s tiling error"), return ge::GRAPH_FAILED);

    // ---- 第六步：填 shape 参数 ----
    tilingData->batchSize = batchSize;
    tilingData->numHead   = numHead;
    tilingData->seqLength = seqLength;

    // ---- 第七步：★ 核心切分逻辑 ----
    int64_t totalElem = batchSize * numHead * seqLength * HIDDEN_DIM;  // 总元素数

    // 平均每核分多少（向上取整，保证覆盖全部元素）
    int64_t rawElemPerCore = CeilDiv(totalElem, coreNum);

    // ★ 向上取整到 HIDDEN_DIM(64) 的倍数。
    //   为什么必须对齐？因为「行」不能被切开——一行 64 个数必须落在同一个核里，
    //   否则 cos/sin 的复用逻辑和双缓冲的按行处理都会错乱。
    int64_t elemPerCore = ((rawElemPerCore + HIDDEN_DIM - 1) / HIDDEN_DIM) * HIDDEN_DIM;

    // 实际需要的核数 = 总元素数 ÷ 每核元素数（向上取整）
    int64_t usedCoreNum = CeilDiv(totalElem, elemPerCore);

    tilingData->numBlocks   = usedCoreNum;
    tilingData->elemPerCore = elemPerCore;

    // ---- 第八步：★ 告诉框架「用几个核」----
    // 不写这一句，kernel 只在 1 个核上跑，性能分全丢
    context->SetBlockDim(usedCoreNum);

    return ge::GRAPH_SUCCESS;
}

// 注册：把函数绑定到 PairwiseRotary
IMPL_OP_OPTILING(PairwiseRotary).Tiling(PairwiseRotaryTilingFunc);

} // namespace optiling