CUDA与AscendC对比（二）

为什么CUDA当时要设计的现在看上去比较复杂？

2006年CUDA刚诞生时，通用GPU编程还是个全新的领域。NVIDIA没有前人的经验可以参考，他们做的事情本质上是：把GPU这种原本只画图的专用硬件，硬生生改造成了通用并行处理器。当时的设计决策：

• 直接暴露硬件特性（Warp、Shared Memory、寄存器）

• 让程序员手动管理一切（线程调度、内存层级、同步）

• 不做过多抽象（性能至上，灵活性第一）

类比：CUDA就像C语言。C语言设计出来时，有人觉得它"复杂"（要手动管理内存、指针很危险），但它成为了系统编程的基石。后来的高级语言（Java、Python）在C之上做了抽象，但代价是性能和可控性。

如果隐藏这些细节会发生什么？

• 自动处理线程束发散→编译器难以优化→性能会不可预测

• 自动管理Shared Memory→你无法控制数据重用→访存延迟暴增

• 隐藏内存层次→你无法做手工调优→计算强度高的算法跑不起来

NVIDIA其实在不断简化CUDA，只是大部分简化发生在库和工具层面，而不是核心语言层面：

趋势：

• 2006年：只能写底层CUDA C

• 2010年：有了Thrust（类似STL的并行算法库）

• 2017年：CUDA 9引入Cooperative Groups（简化线程协作）

• 2020年：CUDA C++增强（Lambda、constexpr支持）

• 2023年：CUDA 12引入更简洁的核函数启动语法（实验性）

Ascend C站在了CUDA的肩膀上，它看到了：

• CUDA暴露太多的线程级控制，但对大多数任务并不必要

• 开发者真正需要的是数据级并行，而不是线程级裸奔

• 可以建立更高的抽象（流水线范式、队列模型）

但这不代表Ascend C更强大：

• CUDA能做的事情（比如手动调Warp shuffle、PTX汇编），Ascend C做不到

• CUDA在奇异硬件行为（如事务内存、异步拷贝）上的控制力，Ascend C无法比拟

如果你觉得CUDA太复杂，一个折中方案是使用 Numba 或 PyTorch的CUDA扩展，它们提供了类似Ascend C的抽象层次，但底层仍然是CUDA。

昇腾NPU里面是按AiCore来并行的，比如910B4有40个vector Core可以并行计算,那么就性能和GPU相比，是不是会差？因为cuda都是可以几百上千个线程同时跑？

用“40个核心 vs 上千个线程”来对比，是混淆了GPU的“核心”(Core)和昇腾AI Core的计算粒度。

打个比方：你是在用工厂的数量和工人的数量做对比，而实际要比较的是产出速度。

• GPU的“CUDA核心”更像是“流水线上的工位”：一个SM（相当于车间）里有64或128个CUDA核心（工位），但它们必须32个一组整齐划一行动（这叫线程束Warp），单独一个工位没法干自己的活。你说“上千个线程”，是指整个芯片里所有工位上的人（线程）加起来。

• 昇腾AI Core更像是“全自动化工位”：一个AI Core内部，不是只有一个工人。它有一个Cube单元（做矩阵乘法，一拍算一个16x16x16的立方体），一个Vector单元（做向量运算，一拍处理256个数据），它们可以像流水线一样同时干活。

所以，40个AI Core是40个高度自动化的车间在并行，每个车间里还有多个流水线在同时运转。

直接用昇腾910和英伟达A100对比，结果更直观。910的AI Core是64个，不是40个，它的对手是A100：

昇腾一个Vector单元（做向量运算，一拍处理256Byte个数据），对应GPU，需要256个线程同时来，而每个线程只处理一个数据（Byte）？

GPU的"单指令多线程"（SIMT）模式

// GPU需要256个线程并行执行
__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;  // idx从0到255
    c[idx] = a[idx] + b[idx];  // 每个线程处理1个元素
}
​
// 启动256个线程
vec_add<<<1, 256>>>(a, b, c, 256);

• 256个线程在SM上被组织成8个Warp（32×8=256）

• 硬件调度器在多个线程束之间切换来隐藏延迟

• 控制开销：需要管理256个PC寄存器、256个栈、复杂的调度逻辑

昇腾的"单指令多数据"（SIMD）模式

// Ascend C在一个Vector Core内完成同样的事情 下面代码和注释应该有错，vector的1个cycle处理的带宽是256Bytes，而不是256个数。
__kernel__ void vec_add(half* a, half* b, half* c) {
    // 假设当前核负责第block_idx个数据块
    LocalTensor<half> aLocal = ...;  // 一次性搬入256个half
    LocalTensor<half> bLocal = ...;
    LocalTensor<half> cLocal = ...;
    
    // 单条Vector指令：一拍完成256个half的加法
    Add(cLocal, aLocal, bLocal, 256);
    
    // 搬出结果
    DataCopy(c + offset, cLocal, ...);
}

• 一条指令 替代了256个线程的循环

• 编译器不需要管理256个执行上下文

• 硬件调度负担极低

把GPU和昇腾比作两种工厂：

GPU：有1000个普通工人，每个工人一次搬一块砖。优点是灵活，可以临时叫某些工人去干别的活。缺点是管理这1000个工人的开销很大，而且如果工人干的活不一样，效率会下降。

昇腾：有40个智能机械臂，每个机械臂一次能抓起256块砖同时搬。优点是搬砖效率极高（一拍完成256块砖）。缺点是所有机械臂都在做相同的工作，无法让机械臂A去焊接同时机械臂B去喷漆。

昇腾的Vector单元确实是GPU"多线程"的硬件加速版——把256个线程的调度和执行，压缩成了一条Vector指令。这就是为什么昇腾能用更少的"核心"，在某些场景下达到与GPU媲美的性能。

既然Vector这么高效，为什么GPU不也做成纯Vector架构？答案：因为GPU需要兼顾图形渲染和通用计算，SIMT模式在处理不规则控制流时更灵活。而昇腾专门为AI计算设计，可以牺牲灵活性来换取极致的吞吐量。