一、CPU 是博士，NPU 是一万个小学生

第一次看昇腾算子课，老师说"我们从 Hello World 开始"。我以为就是 printf。结果看到代码才发现——昇腾的 Hello World 根本不打印字符串。它干的是：在多个 AI Core 上并行把一段数据加 1，结果搬回内存，再让 CPU 打印。一个 hello world，把"数据从哪来、谁来算、算完放哪"整条链路走了一遍。这篇就用它把昇腾芯片讲清楚。

CPU 的设计哲学是"什么都能干"：8~64 个超强核心，分支预测、乱序执行、大缓存，擅长串行和复杂逻辑。NPU 相反：几十个 AI Core，每个内部还有大量并行计算单元，擅长矩阵乘、向量运算这类"简单但海量"的任务。

要让 1000 万个数各自加 1：

• CPU：8 个博士生各分一批，飞快地一个个算
• NPU：几十个 Core 同时上，每个再用向量指令一次算 128 个

深度学习场景下，后者快几个数量级。这就是理解昇腾的第一把钥匙：它不是通用计算芯片，是为"把大量简单计算并行铺开"设计的。

二、达芬奇架构：AI Core 里住着三个角色

┌──────────── AI Core ────────────┐
│  Scalar 单元 → 调度、算地址        │
│        ↓                        │
│  Cube 单元   → 矩阵乘             │
│        ↓                        │
│  Vector 单元 → 向量运算           │
│  UB ← 片上高速缓存,数据从 GM 搬    │
└─────────────────────────────────┘

单元	干什么	类比
Scalar	调度、地址计算、循环控制	工头，指挥不干重活
Cube	矩阵乘加，大模型算力主力	重型机械，一次算一大块
Vector	逐元素运算、激活、归一化	流水线工人

为什么要分三个？因为它们对应深度学习里三类计算：Cube 对应卷积和矩阵乘（占大模型 90% 算力），Vector 对应 ReLU、LayerNorm 这类（数量多计算量小），Scalar 负责把两者串起来。这就是"异构计算"：不是一块芯片干所有事，而是各司其职，用流水线串起来。

三、三个关键字，看懂昇腾代码

__global__ __aicore__ void hello_world(
    __gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)

• __global__ — 这段代码在 AI Core 上跑，CPU 根本不会执行它。它会被编译成 NPU 指令，下载到设备上运行。
• __aicore__ — 运行在 AI Core 内部。
• __gm__ — 指针指向全局内存（片外大内存）。必须显式标注，因为 AI Core 内部还有 UB，编译器靠这个标记决定生成什么访存指令。

四、SPMD：一份代码，几十个核同时跑

uint32_t blockIdx = GetBlockIdx();   // 当前是第几个 AI Core

这是昇腾编程模型的精髓——SPMD（单程序多数据）：

     你写的一份代码
    ┌──────┼──────┐
   Core0   Core1   Core2 ...
  算数据0  算数据1  算数据2
    └──────┴──────┘
        结果汇总回 GM

你只写一份代码，它在几十个 AI Core 上同时执行，每个核拿到的 blockIdx 不同，处理的数据也不同。为什么要这样？一张 4096×4096 矩阵乘，单核算到天荒地老。SPMD 让几十个核同时开工，把大任务切小块并行。理解了 SPMD，就理解了昇腾并行编程的一半。

五、为什么不能直接写 x[i] += 1？

这是新手问得最多的问题。答案很简单：因为那个 x 根本不在 AI Core 手边。

CPU 里内存是平铺的，x[100] 直接访问，缓存自动兜底。但 AI Core 内部只有几百 KB 的 UB，你传进来的指针指向几 GB 外的 GM。所以必须：

1. GM → UB 搬数据
2. 在 UB 上算
3. UB → GM 搬回去

这三步就是 Ascend C 永远避不开的骨架。所有"麻烦"都在解决一个物理问题：如何高效搬运数据，并让搬运和计算重叠。

六、总结

Hello World 把一个完整闭环浓缩在最短的代码里：

概念	体现
异构计算	CPU 调度，AI Core 计算
SPMD 并行	GetBlockIdx()
多级存储	GM ↔ UB
设备侧代码	global aicore
内存空间	gm

它之所以是起点，不是因为它简单，而是因为它把昇腾的设计哲学讲完了。

下一个问题：数据"一次搬多少"该怎么定？这就是 Tiling。想看实操，看我另一篇：《昇腾 Ascend C 入门：手撕 Add 算子，从核函数到 NPU 跑通》。

参考：昇腾社区 CANN 文档 · Ascend C 算子开发入门课程