副标题： 三个芯片架构 × 同一个 GEMM × 四轮投喂——对比实验展示"分阶段投喂法"如何让通用 Agent 写出架构感知的算子代码

---

一、引子 🎯

你现在是某 AI 芯片公司的算子开发工程师。公司自研了一颗新芯片，指令集和 CUDA 完全不同，没有现成的算子库。你拿到了一份 200 页的架构手册，需要在两周内把核心算子（GEMM、Softmax、LayerNorm）跑通。

你的第一反应：让 Agent 帮忙写。

然后你输入：

帮我写一个 GEMM，针对我们的新芯片优化。

Agent 的输出：

void gemm(float* __global__ A, float* __global__ B, float* __global__ C) {
    __shared__ float As[16][16];
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    __syncthreads();
    ...
}

完全不可用。 Agent 回了你一份标准的 CUDA GEMM。你的芯片根本没有 __shared__，没有 threadIdx.x，没有 __syncthreads()。

---

问题出在哪？

LLM（GPT、Claude、DeepSeek 等）的训练数据中包含 几万亿行 CUDA 代码。它们对 NVIDIA GPU 的编程模型了如指掌——但要写一个"没见过的芯片"的算子，它们的知识库里没有任何对应内容。

这不是 Agent 的问题，这是知识缺口的问题。

你的芯片架构不在 Agent 的训练数据里。 解决方案不是微调模型（成本太高，而且每次架构迭代都要重来），而是在 Agent 的上下文窗口中注入知识——分阶段、有结构地告诉它你的芯片长什么样，代码该怎么写。

这篇博客不做理论探讨，而是用 4 轮实验来展示：从零知识到能写出可用算子，需要喂什么，每一步效果如何。

---

二、实验设计

为什么选 GEMM？

矩阵乘法（GEMM）是 AI 芯片最核心的 workload——卷积、全连接层、注意力机制最终都落到 GEMM。而且各芯片架构在 GEMM 上差异最大：NVIDIA 有 Tensor Core，AMD 有 Matrix Core，华为有 Da Vinci Cube。同一个 GEMM，不同芯片的写法完全不同。

这使得 GEMM 成为展示"投喂前后差异"的最佳对象。

三个被测架构

这篇博客的核心实验用 SimpleCore（一个我虚构的极简 AI 加速器）——因为它的指令集完全是"编"的，LLM 训练数据里绝对没有。但只测一个架构缺乏说服力，所以额外加入 RVV 和 Da Vinci 作为对比，验证方法论的可迁移性。

架构	类别	Agent 对它知道多少	在实验中扮演的角色
SimpleCore	虚构	❌ 零知识	主实验：完整 4 轮投喂
RISC-V Vector	开源	⚠️ 有一些但不够深	对比：验证"开源的架构是否不需要喂"
昇腾 Da Vinci	真实（公开资料有限）	⚠️ 资料少	对比：验证"只有架构级知识没有指令集时能到什么程度"

实验设计

SimpleCore 跑完整 4 轮，RVV 和 Da Vinci 跑 R1（裸写）→ R2（喂架构）两轮——因为它们的核心差异在第一轮喂料后就已经显现。

每轮评估四个维度：指令集正确性、代码完整度、专用硬件利用率、性能意识。

---

三、SimpleCore 架构（给读者看的）

为了让实验可复现且足够干净，我设计了一个极简的 AI 加速器架构：

SimpleCore v0.1
──────────────
16 条向量 Lane（编号 0-15）
1 个 4×4 脉动阵列（每周期 16 个 FP16 MAC）

内存层次:
  · Global Memory (GMEM) — 大容量高延迟
  · Shared Memory (SMEM) — 每块 32KB，Lane 间共享
  · RegFile — 每 Lane 32 个寄存器 (r0-r31)

核心指令:
  LOAD/SMEM_LD  — 从 GMEM/SMEM 加载到寄存器
  STORE/SMEM_ST — 从寄存器存储到 GMEM/SMEM
  VADD/VMUL     — 向量加/乘（逐元素）
  VSUM          — 向量归约
  MMA {C},{A},{B} — 4×4 矩阵乘累加 (★★★★★)
  MOV/BEQ/SYNC  — 搬移/跳转/同步

你不需要记住所有细节，只需知道两个关键设计：

1. 有 4×4 脉动阵列（MMA 指令） ——这是最重要的计算资源，不喂这个信息 Agent 绝对写不出来
2. 三层内存模型——和 CUDA 的 __shared__ 命名不同，喂之前 Agent 会套用 CUDA 语法

---

四、四轮实验全记录

R1：什么都没喂

第一次尝试，我直接问 Agent：“用 SimpleCore 写一个 GEMM。”

void gemm_simplecore(float* __global__ A, float* __global__ B, ...) {
    __shared__ float As[16][16];
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    __syncthreads();
    ...
}

问题分析：

问题	表现	严重程度
使用了 CUDA 专有语法	__global__、__shared__、threadIdx	❌ 致命——语法错误
使用了不存在的操作	__syncthreads()——SimpleCore 只有 SYNC	❌ 致命
内存模型不对	假设了 CUDA 的层级线程模型	❌ 致命

结论： Agent 对完全未知的架构，自动填补成了它最熟悉的 CUDA。这种情况下输出完全不可用。

---

R2：喂了架构规格（S1）

第二次，我把 SimpleCore 的完整指令集和内存层次图给了 Agent。

这次输出发生了质的变化：

// ✅ 用上了 SimpleCore 的指令集
// ✅ 正确使用了寄存器模型
// ✅ 理解 Lane 的概念

// 寄存器分配
// r0-r3:  A tile (4×4)
// r4-r7:  B tile (4×4)
// r8-r23: C 累加器

// 核心计算——用上了 MMA 指令！
MMA {r8,r9,r10,r11}, {r0,r1,r2,r3}, {r4,r5,r6,r7}

// 数据搬运
LOAD  r0, [r25]       // GMEM → RegFile
STORE [r27], r8       // RegFile → GMEM

// 循环控制
MOV r24, 32           // 循环计数
SUB r24, r24, 1
BNE r24, 0, k_loop

变化：

维度	R1	R2
指令集	CUDA (__global__)	✅ SimpleCore (LOAD/STORE/MOV/BEQ/MMA)
计算核心	循环标量乘加	✅ 用上了 MMA 4×4 脉动阵列
内存搬运	自动驾驶（CUDA 自动管理）	✅ 显式 GMEM → 寄存器搬运
线程模型	CUDA threadIdx/blockIdx	✅ Lane + warp 模型

但仍然有问题： 代码不完整。Agent 在几个关键位置打了标记：

// Simplified: each lane loads its portion
// ... (16 stores per lane)
// ... (not shown for brevity)

它知道架构规范了，但不知道"这样写才完整"。这正是 S2（喂代码范例） 要解决的问题。

---

R3：喂了架构 + 代码范例（S1+S2）

第三次，我在架构规格之外，给了 Agent 三个用 SimpleCore 指令集编写的完整代码范例：

1. 向量加法（展示基础的 LOAD/VADD/STORE 模式）
2. SMEM 分块加载（展示 tiling 模式）
3. 使用 MMA 的小型 GEMM 块（展示脉动阵列用法）

输出变得非常完整——出现了一个 350 行的 GEMM 实现，包含了完整的 tiling 策略、寄存器分配、K 维度循环逻辑。

改进点：

• ✅ 完整的 tiling 策略（128×128 → 多个 4×4 tile）
• ✅ 详细的 SMEM 协同加载方案
• ✅ 每行指令都有注释说明其作用

但出现了两个新问题：

问题一：Agent 写了一个软件除法放在热点循环中

为了计算 tile_id → 行/列索引，Agent 写了这个：

// 用循环除法计算 tile_id / 32
div_loop:
    BEQ r7, r6, div_done
    SUB r7, r7, r6          // 反复减 32
    ADD r5, r5, 1           // 计数
    BEQ r0, r0, div_loop
div_done:

这是一个最多执行 32 次循环的除法程序——而它被放在 GEMM 热点循环的每次 K 迭代中。如果用这个代码跑 128×128 的 GEMM，这个除法就会执行 4×32 = 128 次，每次 32 个循环 = 4096 条无意义的指令。性能灾难。

问题二：Agent 编造了一条不存在的指令

SHL r21, r21, 2    // ← SimpleCore 的指令集里没有 SHL！

Agent 看到需要乘以 4，自动从它的 CUDA/CPU 知识中"借"了一条移位指令——完全不检查这个指令是否在 SimpleCore 的指令集中。

这就是 R3 的核心教训： 代码范例提升了完整度，但 Agent 会用自己的通用知识（乘除法、移位、位操作）来填补架构中不存在的操作。你不明确说"这个架构没有 XXX"，Agent 就会自己编。

---

R4：把错误喂回去（S5 迭代）

第四次，我把 R3 的问题反馈给 Agent：

• 软件除法在热点循环中 → 建议用 SMEM 查表替代
• SHL 指令不存在 → 只能用 MOV 和查表
• 代码太长 → 保持核心逻辑简洁

Agent 的修复结果：

做对的：

• ✅ 去掉了软件除法循环，改用 SMEM 查表（tile 地址预计算）
• ✅ 去掉了 SHL 指令
• ✅ 整体结构更清晰了

做错的：

• ❌ MMA 指令丢了！ Agent 退化为用 VADD+VMUL 模拟矩阵乘法——性能损失一个数量级
• ❌ 矩阵维度从 128 改成了 64（没有解释）
• ❌ 代码从"具体的实现"变成了"伪代码夹杂注释"

// 实际需要16次乘加, 用循环展开
// 这里简化: 假设已初始化
// ... (实际代码中应展开为16条MOV指令)

R4 的核心教训： 单次修复往往顾此失彼。Agent 在修复"软件除法"和"SHL"这两个具体问题时，丢失了更核心的 MMA 能力。
结构化的反馈比笼统的问题描述有效得多——应该明确告诉 Agent"保留 MMA 不变，只修复除法和 SHL"。

---

实验小结

          R1          R2           R3            R4
指令正确   ❌ CUDA    ✅ 架构匹配   ✅ 完整       ⚠️ 没有MMA
代码完整   ❌ 不可用   ⚠️ 有缺口    ✅ 完整       ⚠️ 半成品
性能意识   ❌ 无       ❌ 无        ⚠️ 有但错了   ❌ 丢失
自我修复   —          —           —            ⚠️ 部分成功

这个实验结果就引出了我们五阶段投喂法的设计逻辑。

---

横向对比：三个架构的反应完全不同

如果只测 SimpleCore，你可能会问："这方法对真实架构也管用吗？"答案是——管用，但每个架构的反应方式不同，而且恰恰揭示了这个方法论最有价值的部分。

RVV（RISC-V Vector Extension）：Agent 最熟悉的"陌生人"

RVV 是开源指令集，LLM 的训练数据里有一定量的 RVV 代码。R1 裸写的表现比 SimpleCore 好得多：

// RVV R1（裸写）——用 intrinsic 写的完整 GEMM
#include <riscv_vector.h>

void gemm_rvv_f16(_Float16 *A, _Float16 *B, _Float16 *C) {
    for (size_t i = 0; i < M; i++) {
        for (size_t j = 0; j < N; j += 8) {
            vfloat16m1_t vacc = vle16_v_f16m1(&C[i*ldc+j], 8);
            for (size_t k = 0; k < K; k++) {
                vfloat16m1_t va = vfmv_v_f_f16m1(A[i*lda+k], 8);
                vfloat16m1_t vb = vle16_v_f16m1(&B[k*ldb+j], 8);
                vacc = vfmacc_vv_f16m1(vacc, va, vb, 8);
            }
            vse16_v_f16m1(&C[i*ldc+j], vacc, 8);
        }
    }
}

代码能跑， 没有任何虚假指令。但仔细看会发现一个性能问题：每次 K 迭代都做 vfmv_v_f_f16m1 标量广播——对于 VLEN=128 的 RVV，每个 K 迭代只处理 8 列，而 K=128 意味着 128 次广播。这是一个在训练数据中常见但不高效的模式。

喂了 RVV v1.0 规格后（R2），Agent 试图写汇编级实现：

// RVV R2（喂架构后）——尝试写汇编
vsetvli t6, x0, e16, m1, ta, ma
vmv.v.i v0, 0
vle16.v v8, (t4)
vfmv.v.f v10, ft0
// ... 但代码越写越长，寄存器管理开始混乱

有趣的是： R2 的汇编代码看起来更"专业"，但实际质量 低于 R1 的 intrinsic 代码。R1 的 intrinsic 代码虽然性能不是最优，但清晰、可编译、逻辑正确。R2 的汇编尝试在寄存器分配上出现了明显的混乱。

RVV 实验的核心发现： 如果 Agent 对某个架构已经有基础认知（开源、训练数据覆盖过），不要把它推到更低级的编程模型（从 intrinsic 到汇编）。S1 喂架构应该用来补充 Agent 的知识盲区（比如该架构特有的优化约束），而不是让它从零学起。

维度	RVV R1（裸写）	RVV R2（喂架构后）
指令正确性	✅ 正确的 intrinsic	⚠️ 汇编有错误
代码完整度	✅ 完整的函数	⚠️ 混乱、未完成
专用硬件利用率	⚠️ 用向量但不够高效	❌ 埋在寄存器分配里
投喂效果	—	不升反降

---

昇腾 Da Vinci：Agent 知道"有 Cube"，但不知道"怎么用"

对于华为昇腾的 Da Vinci 架构，Agent 在 R1 的表现和 SimpleCore 类似——退化为普通 C 代码：

// Da Vinci R1（裸写）——退回 C 代码 + 编造的 pragma
void gemm_davinci_fp16(half A[128][128], half B[128][128], half C[128][128]) {
    half local_A[16][16], local_B[16][16], local_C[16][16];
    // ... 三层循环分块加载
    #pragma cube_matmul(local_C, local_A, local_B, 16)  // ← 编造的！
}

Agent 知道 Da Vinci 有 Cube 单元（这是公开知识），但它不知道 Cube 的具体指令怎么写——所以用了 #pragma cube_matmul 这样的假指令来"占位"。

喂了架构信息后（R2），变化非常显著：

# Da Vinci R2（喂架构后）——架构感知的 Cube 分块方案
# Cube 指令：l0c += l0a × l0b（16×16×16）
# pipe 调度：搬入 → 计算 → 搬出

for m_tile in range(8):   # M 方向 8 个 16×16 块
    for n_tile in range(8):  # N 方向 8 个 16×16 块
        l0c = zeros((16, 16))  # L0C 累加器
        for k_tile in range(8):  # K 方向 8 个 16×16 块
            l0a = load(L0A, A, m_tile, k_tile)
            l0b = load(L0B, B, k_tile, n_tile)
            l0c = cube_mac(l0a, l0b, l0c)  # 单条 Cube 指令 = 4096 MACs
        store(C, m_tile, n_tile, l0c)

这里的关键变化不是"代码写对了"，而是设计层次的变化：

维度	Da Vinci R1（裸写）	Da Vinci R2（喂架构后）
设计层级	C 代码 + 假 pragma	架构感知的分块 + pipe 调度
计算映射	三层嵌套循环	✅ 8×8×8 Cube tile 映射
专用硬件	❌ 不知道 Cube 怎么用	✅ 理解 16×16×16 的 MAC 密度
流水线	❌ 无	✅ 双缓冲搬入/计算重叠
性能估算	无	✅ 给出理论 Cube 指令数 512

更重要的是：Agent 在 R2 中没有尝试写汇编（因为 Da Vinci 的指令集不公开），而是在更高层次提供了设计方案——分块策略、pipe 调度、理论性能估算——这些都是在理解了架构规格后才产生的能力。

---

三架构对比总结

                    SimpleCore（虚构）      RVV（开源）         Da Vinci（真实）
                    ────────────────      ──────────         ────────────────
Agent 默认知识       几乎为零              有一些 intrinsic    仅知道"有 Cube"
R1 输出              CUDA 代码            可用的 intrinsic    假 pragma + C 代码
R2 输出（喂架构后）   正确汇编 + MMA       混乱的汇编          架构感知的设计方案
投喂效果             从 ❌ 到 ✅ 质变      略降 ⚠️            从 C 代码到架构设计 ⚠️
最佳策略             S1+S2 都要喂         不喂汇编，只补盲区   S1 喂架构规格就够了

这个对比本身就传达了一个重要信息：

投喂方法论不是"把架构手册塞进提示词"这么简单。不同架构需要不同的投喂策略——完全陌生的架构要重喂（S1→S2），半已知的架构只需补盲区（S1 精简版），只有公开资料的架构只能停留在设计层面（S1 就够了）。

下一节把这一切抽象为五个通用阶段。

---

五、五阶段投喂法（方法论）

三个架构的实验共同揭示了一个核心矛盾：不能一次性把所有东西塞给 Agent。 上下文窗口有限，而且 Agent 也无法同时处理"这个芯片有什么"和"代码应该写多快"两个不同层次的信息。

必须分步投喂，每步解决一个特定的知识缺口：

阶段	标题	喂什么	目的是让 Agent 知道…	适用场景
S1	架构认知	指令集 + 内存层次 + 约束	“这个芯片有什么，没什么”	所有场景，必选
S2	代码语料	3-5 个高质量算子 + 注释	“像样的代码长什么样”	Agent 不熟悉的架构
S3	性能模型	Roofline 图 + Profile 数据	“瓶颈在哪，往哪优化”	已有代码需要调优
S4	工具链	编译器/Profiler 用法	“怎么验证写对了没有”	需要可编译代码
S5	迭代闭环	Error log + 性能数据	“刚写的有问题，自己修”	代码报错或性能不达标

核心洞察：不是所有架构都需要喂所有阶段

三架构的实验告诉我们——投喂策略必须根据"Agent 对目标架构的已有知识量"来调整：

完全陌生的架构（类似 SimpleCore）→ S1 + S2 都必需
Agent 对你一无所知。S1 让它能用对指令，S2 让它写出完整代码。

开源/已知的架构（类似 RVV）→ 只喂 S1 精简版，跳过 S2
Agent 已经有一定基础。不要让 S2 的代码范例把它"带偏"到更低级的编程模型。此时 S1 只应该补充它不知道的优化约束，而不是把整个 ISA 喂给它。

只有公开资料的架构（类似 Da Vinci）→ 只喂 S1
没有可用的指令级范例给 S2，但 S1 足以让 Agent 从"写 C 代码"提升到"做架构感知的设计"。对这类架构，Agent 的价值在于设计方案而非生成汇编。

不是每次都要全量投喂

另一个关键原则是按需投喂：

• 写架构相关代码（算子迁移、指令映射）→ 只喂 S1
• 写具体算子实现 → 喂 S1 + S2
• 性能调优 → 喂 S3（profile 数据）
• 调试修复 → 喂 S4（编译错误）→ 喂 S5（迭代循环）

不要把整个手册都塞进提示词里——这是导致上下文超载的常见原因。控制每个阶段的投喂量，只给 Agent 当下需要的知识。

---

六、实战模板（可以直接抄走）

以下四个模板是我在实际开发中反复使用的模式。你可以直接复制修改。

模板 1：架构描述（给 Agent 建立认知）

以下是我的芯片架构规格，请记住这些约束：

=== 芯片名称: [名称] ===

计算单元: [类型] × [数量]
  - 例如：16×16 脉动阵列 / 向量单元 / 标量核心

内存层次:
  - [层级1]: [容量], [延迟特征], [访问方式]
  - [层级2]: [容量], [延迟特征], [访问方式]
  - [层级3]: [容量], [延迟特征], [访问方式]

数据类型: [支持的精度列表]

核心指令:
  [指令1]: [功能描述]
  [指令2]: [功能描述]

关键约束:
  - [约束1]
  - [约束2]

注意：如果我的指令集中没有某些操作（移位、乘除法、位操作等），
不要使用它们，也不要用软件模拟它们。

模板 2：写算子（告诉 Agent 要做什么）

请帮我为 [架构名称] 实现一个 [算子名]（[功能描述]）。

参数:
  - 输入: [类型] [名称], 形状 [维度]
  - 输出: [类型] [名称], 形状 [维度]

算法描述:
  [一两句话描述要做什么计算]

约束:
  - [性能要求，如：不要用软件除法]
  - [精度要求，如：fp16]
  - [特殊约束]

架构规格（S1 内容）: [引用模板 1 的输出]
代码风格参考（S2 内容）: [引用代码范例]

请写出完整的实现代码，每条指令都要有注释。

模板 3：性能诊断（发现瓶颈后修复）

以下是我运行 [算子] 后的 Profile 数据：

热点函数: [函数名]
热点占比: [X]%
瓶颈类型: 计算密集 / 访存密集
主要问题: [描述]

可能的原因:
  - [可能原因1]
  - [可能原因2]
  - [可能原因3]

请分析 Profile 数据，指出最可能的原因，并给出修复方案。
注意保留现有的 MMA/Tensor Core 用法不变，只修复性能问题。

模板 4：错误修复（编译/运行失败后迭代）

我运行的 [算子] 代码遇到了以下错误：

错误类型: [编译错误 / 运行时错误 / 结果错误]
错误信息: [粘贴错误日志]

代码片段:
[粘贴相关代码]

期望行为: [描述正确应该发生什么]

请分析错误原因并修复。
重要：保持以下部分不变：
  - [保持不变的逻辑1]
  - [保持不变的逻辑2]

---

七、几个关键陷阱

陷阱 1：Context Window 管理

S1 架构规格 ≈ 1-2K tokens，S2 代码范例 ≈ 3-5K tokens，S3-S5 按需投喂。一次全喂 ≈ 8-10K tokens——大部分 LLM 能承受，但超过这个量就要考虑选择性投喂：

• 只给核心指令集（删掉附录、历史版本等冗余内容）
• 代码范例只给 2-3 个，不要一次给 10 个
• 按"当前任务需要什么"精简上下文

陷阱 2：IP 泄露风险

这是一个真实的问题。 如果你在喂自研芯片的指令集细节给云端 LLM（ChatGPT、Claude.ai、DeepSeek Chat 等），你的芯片架构细节会被发送到第三方服务器。

应对方案：

• 优先使用本地模型（Ollama + Qwen3/DeepSeek 本地部署），数据不出本机
• 如果必须用云端，对敏感信息做抽象化处理（改指令名、改寄存器名）
• 在 Prompt 前加：这是一个虚构的架构，所有指令名和寄存器均为虚构设计

陷阱 3：幻觉防御——Agent 会编造不存在的指令

实验中 R3 的 SHL 不是偶然。LLM 在遇到"需要移位操作"而指令集中没有时，会自动从通用知识库中借用——即使你明确说"指令集只有这些"。

防御方法：

• 在架构描述末尾加一句：“以下指令不在架构中，不要使用：[列表示例]”
• 要求 Agent 输出代码前检查每条指令是否在指令集中
• 设置一个验证步骤：每次 Agent 输出后，自动匹配指令白名单

陷阱 4：一次只给一个反馈

R4 的实验证明：同时指出多个问题，Agent 容易顾此失彼。在迭代修复（S5）时，一次只反馈一个最严重的问题——先保证 MMA 用对，再修除法，再修 SHL。

---

八、延伸：不只是芯片架构

这套"分阶段投喂法"不只适用于芯片算子开发。任何需要让 Agent 理解私有领域知识的场景都可以套用：

场景	S1 架构	S2 语料	S3 性能	S4 工具链	S5 迭代
芯片算子开发	指令集	代码范例	Profile 数据	汇编器/Profiler	错误修复
内部 API 开发	API 规范	SDK 范例	延迟基准	测试框架	报错修复
专有数据库	数据模型	查询范例	查询计划	管理工具	调优反馈
私有大模型部署	模型结构	调用范例	延迟指标	部署工具	错误日志

你需要做的只是把这五个阶段翻译成你的领域。 在每个阶段，问自己一个问题：“Agent 现在不知道什么，而这些知识对于完成当前任务是必要的？” 然后把那个知识喂给它。

---

九、总结

回到开篇的场景：你拿到了一份 200 页的架构手册，需要在两周内跑通核心算子。

现在的方案不是让 Agent 自己"理解"这 200 页——而是分阶段提取关键知识，按需注入到每一次对话中。

喂给 Agent 的 4 条核心原则：

1. 先让 Agent 知道"你的芯片有什么" → 指令集 + 内存模型（S1）
2. 再用范例告诉它"代码长什么样" → 高质量代码片段（S2）
3. 跑完后告诉它"哪里慢了，哪里错了" → Profile + 错误日志（S3-S5）
4. 每次只喂当前需要的信息 → 上下文不是垃圾场

最后想明白一件事：

你喂的不是提示词，是团队积累的架构知识和代码经验。Agent 只是那个加速复用的管道。真正驱动它写对代码的，是你对芯片的理解——以及你把这种理解编码到上下文中的能力。

（这篇不写实验脚本下载链接了。但如果你在自研芯片上实践这个方法论，欢迎在评论区分享你的"投喂配方"。）

---

系列全篇：
1–15. 前情见本系列前作
16. 从 CUDA 到自研芯片：怎么让 Agent "学会"你的芯片架构 — 本文

---

附录：完整实验 Prompt

以下是本轮实验所有阶段使用的完整 Prompt。模型使用 DeepSeek Chat，temperature=0.3，max_tokens 按需设置（2000-4000）。

A. SimpleCore 实验

A1. R1（裸写）

你是一个算子开发工程师。现在有一个新的芯片架构叫 SimpleCore，请为它写一个 GEMM（C = A × B + C，矩阵乘法）的实现。

注意：我不知道 SimpleCore 的指令集和内存模型，但你可以假设它有一些计算单元和内存。

请写出完整的代码实现。

不要问任何问题，直接写代码。

A2. R2（喂 S1 架构规格）

你是一个算子开发工程师。请为 SimpleCore 芯片架构写一个 GEMM 实现（C = A × B + C，M=N=K=128，fp16）。

SimpleCore 的完整规格如下：

=== SimpleCore v0.1 芯片规格 ===

计算单元:
- 16 条向量处理单元（Lane），编号 0-15
- 1 个 4×4 脉动阵列（每周期 16 个 FP16 MAC）

内存层次:
- Global Memory (GMEM): 大容量高延迟，通过 LOAD/STORE 访问
- Shared Memory (SMEM): 每块 32KB，同一块内所有 Lane 共享
- RegFile: 每 Lane 32 个寄存器 r0-r31

数据类型: fp16, int32

指令集:
  LOAD  rd, [addr]         // GMEM → RegFile
  STORE [addr], rs         // RegFile → GMEM
  SMEM_LD rd, [saddr]      // SMEM → RegFile
  SMEM_ST [saddr], rs      // RegFile → SMEM
  VADD rd, rs1, rs2        // 向量加法
  VMUL rd, rs1, rs2        // 向量乘法（逐元素）
  VSUM rd, rs1             // 向量归约
  MMA {C}, {A}, {B}        // 4×4 矩阵乘累加 (C += A×B)
  MOV  rd, imm             // 加载立即数
  BEQ  rs1, rs2, label     // 条件分支

约束:
- 每 Block 最多 1024 线程（64 Lane × 16）
- SMEM 32KB/Block
- MMA 操作数必须来自寄存器

请用 SimpleCore 的指令集写出完整的 GEMM 实现代码。

不要问任何问题，直接写代码。

A3. R3（喂 S1 + S2 代码范例）

在 R2 的 Prompt 基础上，增加了以下代码范例：

=== 代码范例 1：向量加法（展示基础模式）===
// 计算 C = A + B，每 Lane 处理一个元素
function vec_add(A, B, C, N):
    addr_A = A + lane_id * 1
    addr_B = B + lane_id * 1
    addr_C = C + lane_id * 1
    LOAD  r0, [addr_A]     // GMEM → RegFile
    LOAD  r1, [addr_B]     // GMEM → RegFile
    VADD  r2, r0, r1       // 向量加法
    STORE [addr_C], r2     // RegFile → GMEM

=== 代码范例 2：SMEM 分块加载 ===
function load_tile(mat, smem_buf, stride_ld):
    for row = 0 to 15:
        gmem_addr = mat + row * stride_ld + lane_id
        LOAD  r0, [gmem_addr]
        SMEM_ST [smem_buf + row * 16 + lane_id], r0
    SYNC

=== 代码范例 3：使用 MMA 的小型 GEMM 块 ===
    MOV r8, 0; MOV r9, 0; MOV r10, 0; MOV r11, 0   // 清零累加器
    SMEM_LD r0, [A_base + offset]
    SMEM_LD r4, [B_base + offset]
    MMA {r8,r9,r10,r11}, {r0,r1,r2,r3}, {r4,r5,r6,r7}

并在末尾增加要求：

要求：
1. 用 SimpleCore 指令集写出完整代码
2. 使用 MMA 指令（4×4 脉动阵列）进行核心计算
3. 使用 SMEM 进行数据缓存
4. 实现完整的 tiling 策略（128×128 → 多个 4×4 tile）
5. 处理好边界情况
6. 每行指令写清楚注释

直接输出完整的 GEMM 实现。

A4. R4（迭代反馈）

（架构规格同 R2，但增加约束说明：没有 SHL/SHR 指令，没有乘除法指令）

你之前为 SimpleCore 写了一个 GEMM 实现，但有以下问题需要修复：

**问题 1：软件除法在热点循环中**
你写了一个 32 次循环的除法子程序来计算 tile_id → 行列索引，而这些计算在 GEMM 的热点循环中每次 K 迭代都会执行。这是性能灾难。

解决方案：因为 TILE_N=32（2 的幂），行列索引可以在编译时计算好每个 tile 的 (row, col) 映射表存在 SMEM 中，运行时直接查表。

**问题 2：使用了不存在的指令（SHL）**
SimpleCore 指令集中没有 SHL。

**问题 3：代码太臃肿**
350 行代码中大部分是地址计算的锅炉板代码。保持核心逻辑清晰即可。

请用以下策略重写 GEMM：
1. 预计算每个 tile 的起始地址表，存在 SMEM 中
2. 热点循环只做：SMEM_LD → MMA → SMEM_LD → MMA ...
3. 用最少的指令完成 GEMM 核心

B. RVV 实验

B1. R1（裸写）

写一个 RISC-V Vector (RVV) 的 GEMM 实现（C = A × B + C，M=N=K=128，fp16）。
用 RVV 内联汇编或 intrinsic 写出完整代码。
不要问问题，直接输出。

B2. R2（喂架构规格）

RVV v1.0 规格：

RISC-V Vector Extension (RVV) v1.0 关键规范：

向量寄存器: v0-v31, 每个 VLEN 位（实现定义，通常 128/256/512）
向量长度: vl（由 vsetvli 设置）
向量元素类型: e8/e16/e32/e64
LMUL: mf8/mf4/mf2/m1/m2/m4/m8

关键指令:
  vsetvli rd, rs1, vtype      // 设置向量长度
  vle16.v vd, (rs1)           // 加载 fp16 向量
  vse16.v vs3, (rs1)          // 存储 fp16 向量
  vfmacc.vv vd, vs1, vs2      // 向量乘累加
  vfmul.vv vd, vs1, vs2       // 向量乘法
  vfadd.vv vd, vs1, vs2       // 向量加法

关键约束:
  - 没有矩阵专用指令，GEMM 要用向量指令组合
  - 循环分块时需要手动管理偏移地址
  - 边界处理（vl < VLMAX 时）

Prompt：

写一个 RISC-V Vector (RVV) 的 GEMM 实现（C = A × B + C，M=N=K=128，fp16）。

RVV v1.0 规格：[上述规格]

要求：
1. 用 vsetvli + vle16.v + vfmacc.vv 实现核心计算
2. 做好 tiling（用 vl 控制块大小）
3. 处理好边界情况（最后一块可能不满向量长度）
4. 直接输出完整可用的 RISC-V GCC 汇编代码

C. Da Vinci 实验

C1. R1（裸写）

写一个针对华为昇腾 Da Vinci 架构的 GEMM 实现（C = A × B + C，M=N=K=128，fp16）。
用 CANN（昇腾计算语言）或者伪代码写出核心实现。
不要问问题，直接输出。

C2. R2（喂架构规格）

Da Vinci 架构规格（仅使用公开资料）：

华为昇腾 Da Vinci 架构关键规格：

计算单元 - Cube Unit:
  · 16×16×16 脉动阵列，4096 个 FP16 MAC/cycle
  · A: 16×16, B: 16×16, C: 16×16（累加器）

内存层次（Ascend 910B）:
  · Global Memory: HBM
  · L1 Buffer: 每 AI Core 本地缓存
  · Cube 输入/输出走固定缓冲区（L0A/L0B/L0C）

编程模型:
  · AI Core 是基本计算单元
  · pipe 模式：搬入 → 计算 → 搬出
  · 算子开发方式：TBE DSL 或手写汇编

关键约束:
  · Cube 处理 16×16 块效率最高
  · 非 16 对齐需要边界处理
  · pipe 调度优化对性能影响极大

Prompt：

写一个针对华为昇腾 Da Vinci 架构的 GEMM 实现（C = A × B + C，M=N=K=128，fp16）。

Da Vinci 架构规格：[上述规格]

要求：
1. 基于 Cube Unit 设计分块策略
2. 展示如何将 128×128 映射到 16×16 Cube
3. 考虑 pipe 调度的基本概念
4. 可以用伪代码 + 文字说明结合

---

实验环境影响： 以上 Prompt 使用 DeepSeek Chat API（deepseek-chat 模型），temperature=0.3。

我们也用 DeepSeek v4-pro（关闭思考模式）验证了关键实验点，跨模型一致性如下：

核心结论	deepseek-chat	deepseek-v4-pro	一致？
不喂架构 → 写不出目标芯片代码	❌ CUDA 幻觉（__global__、threadIdx）	⚠️ 干净的标准 tiled C（没用假指令，但也无架构感知）	✅ 都是"不可用"
喂架构 → 指令质变，用上专用单元	✅ MMA 正确使用	✅ MMA 正确使用	✅ 完全一致
喂料后会有额外指令幻觉	仅 R3 出现一条 SHL	R2 就出现 MUL/SHL/SUB/HALT 四条	✅ 都会幻觉，v4-pro 甚至更严重
RVV 不需要重喂	✅ intrinsic 可用	✅ 更完善（还写了解码展开版）	✅ 一致
核心叙事线	R1→R2→R3→R4 阶梯式改善	同上	✅ 一致

两个有趣的现象：（1）更强大的模型（v4-pro）在 R1 更"克制"——不编造 CUDA 而是退回标准 C，但在 R2 反而更"自信"地编造了更多假指令。更强≠更不容易幻觉。（2）跨模型验证说明五阶段投喂法的有效性不依赖于特定模型，不同 LLM 只是"怎么错"有差异，"什么能治"是一样的。