从 CUDA 对 AI 芯片思考
从技术的角度重新看英伟达生态,有很多值得借鉴的方面。本文将主要从流水编排、SIMT 前端、分支预测和交互方式等方面进行分析,同时对比 DSA 架构,思考可以从英伟达 CUDA 中借鉴的要点。
英伟达生态的思考点
从软件和硬件架构的角度出发,CUDA 和 SIMT 之间存在一定的关系,而目前 AI 芯片采用的 DSA 架构在编程模型和硬件执行模型上还处于较为早期的状态,英伟达强大的生态同样离不开 CUDA 在编程方面的易用性。
面对新的 AI 芯片,在流水隐藏方面,实现架构层面的隐藏流水编排机制,提出一个形式上与 SPMD 没有关系的编程模式,而且易用性堪比 CUDA 的软件是可能的。但是反过来在核心问题上没有得到解决,提出形式上与 CUDA 类似的编程模型也仍然会有易用性的问题,开发者很难获得一个足够好的初始性能。
在软硬件架构方面,对于 DSA 架构而言,一方面需要建立一套开放的软硬件架构,联合其他 DSA 架构一起对抗 CUDA 生态;另一方面需要明确面向不同层级开发者的易用性和软件开发形态。
SIMT 与 CUDA 的关系
英伟达为了维护 CUDA 生态对 SIMT 硬件架构做出了调整和取舍,因此 CUDA 会在一定程度上对英伟达硬件架构产生约束,例如保留 SM、Warp、Thread 等线程分层概念。CUDA 架构在近几年没有做出重大的改变,主要是维护编程体系软件对外的抽象和易用性。
DSA 之所以在硬件架构的指令和设计上比较激进,并非软件体系做得好,而是在刚开始并没有太多地考虑编程体系的问题,自然没有为了实现软硬件协同带来的架构约束。CUDA 的成功之处在于通过 SIMT 架构掩盖了流水编排、并行指令隐藏以及 CUDA 的易用性。
DSA 硬件架构执行方式
DSA 硬件架构一般是指单核单线程,线程内指令可以通过多核共享 Cache 协作。编程模型上缺乏统一的标准,因此需要专门搭建编译器和编程体系,硬件主要以 AI 加速芯片(TPU、NPU 等)为主。
关于 DSA 的硬件执行方式,DSA 硬件目前的裸接口一般是每个核一个线程,每个线程内串行调 DSA 指令集,指令在硬件上通常会分发到不同的指令执行流水线上,正确性部分靠软件同步实现,部分靠硬件保证。
CUDA 客户能力区分
按照使用 CUDA 的难易程度,可以将 CUDA 的使用用户分为三类,分别是初阶、中阶和高阶用户。
-
初阶用户:掌握 CUDA 并行编程能力,了解 NV SIMT 硬件基础架构,可以拿到并行指令、流水掩盖、并行计算三部分性能。
-
中阶用户:进一步运用 CUDA 提供的切块 Tiling、流水 Pipeline 能力,进一步获取更高的性能收益。
-
高阶用户:深入了解 SIMT 微架构细节,解决线程 bank 冲突、精细化流水掩盖、精细化指令使用、极致的切块 Tiling 策略,从而实现极致性能。
CUDA 在开发方面具有很好的易用性,以下是使用 CPU 编写的矩阵加法运算:
void add_matrix(float* a, float* b, float* c, int N) {
int index;
for (int i = 0; i < N; ++i) {
index = i + j * N;
c[index] = a[index] + b[index];
}
}
int main()
{
add_matrix(a, b, c, N);
}
以下是使用 GPU 编写的矩阵加法运算,与 CPU 编程相比,因为使用的是并行计算,所以没有 for 循环:
__global__ void add_matrix(float* a, float* b, f