NPU的编程模型和数据流驱动使用方法

最新推荐文章于 2025-05-14 09:44:42 发布

盖世灬英雄z

最新推荐文章于 2025-05-14 09:44:42 发布

阅读量866

点赞数 29

CC 4.0 BY-SA版权

分类专栏：毛虫学习笔记文章标签：人工智能

本文链接：https://blog.csdn.net/qq_30407405/article/details/147306596

毛虫学习笔记专栏收录该内容

4 篇文章

订阅专栏

编程模型

NPU的编程模型从顶层到底层分为Dataflow programming model、High-Level Programming model、Low-level programming model、CT Kernel programming，如下图所示：

其中Dataflow programming model为8系列架构主要面向的编程模型，基于该编程模型，NPU(Tile)将在空间维度划分为多个Tile区域，每个Tile区域负责应用程序中的一段任务，再通过Tile区域的数据流通信来将各段任务串通起来，然后各个Tile区域在运行时将以数据流水的方式并行执行。

High-level programming model为各个Tile区域计算的编程，所以它包含在Dataflow programming model框架内。该编程模型软件可见全局存储（Gmem）和 NPU 的 L1，编程方式类似 SPMD（Single Programming Multiple Data），即一个编程通过多数据并行方式映射到各个 Tile 中并行执行，相当于 Triton 编程模型。不过 High-level programming model 将以 intrinsic 配合张量加减乘除运算符的方式对 L1 的数据块做 Gmem 访问、NPU 通信和 NPU 计算，同时辅以 CPU Scalar 的编程来计算数据的索引和分支跳转控制。实际上，当所有 Tile 看成一个大统一的区域时，数据流编程模型将完全退化到与 CUDA、Triton 完全类似的编程方式。外部用户编程时，一般建议使用到 High-level programming model 即可，因为它具有较强的易用性和后向兼容性。

Low-Level programming model 为 High-level programming model 中的 intrinsic 接口的编程。该编程模型软件可见 NPU 的 L1 和 LO，以直接调用 NPU 指令的方式控制 LSU 搬数、NE 计算、CT Kernel 接口和 KCore Scalar。Low-level 编程时，需要考虑到 LO 存储空间有的有限。

CT Kernel Programming 为 Low-level programming model 中的 CT Kernel 接口做编程。该编程模型软件可见 NPU 的 LOA 和 CT 中的寄存器。若是汇编级编程，软件可见 CT 寄存器，甚至 Scalar Core 的寄存器；若是比汇编更高级的编程，比如 C 编程，则软件不可见 CT 寄存器，由编译器（此时为 LLVM 或者 GCC 编译器）来分配 CT 寄存器资源。CT Kernel 编程时，使用向量 Load、向量 Store、向量计算操作的微指令进行编程。

Low-Level programming model 和 CT Kernel Programming 一般建议内部软件开发和外部高级开发者使用。

Triton编程模型

Triton（可能被误写为“Tricon”）是一种面向GPU的高性能张量计算编程模型和编译器，专为机器学习、高性能计算（HPC）等领域设计，旨在简化GPU编程的复杂性，同时提供接近手写CUDA内核的性能。以下是其核心特点：

1. 核心设计理念

类Python的语法：
Triton使用Python-like语法编写计算内核（如@triton.jit装饰器），用户无需深入掌握CUDA即可实现高效的GPU并行代码。
基于块的编程（Tile-based）：
类似SPMD（Single Program Multiple Data）模型，程序逻辑被映射到GPU的多个线程块（Tiles）上并行执行，每个块处理数据的一个子集（如矩阵分块计算）。
自动内存管理：
编译器自动优化全局内存（Gmem）与共享内存（L1 Cache/SMEM）之间的数据搬运，减少手动调优负担。

2. 关键特性

张量级操作：
直接支持张量的加减乘除等运算符，结合intrinsic函数（如tl.dot()）高效调用GPU硬件指令。
动态并行与索引计算：
允许在GPU内核中动态计算数据索引和分支控制（类似CPU Scalar编程），灵活性高。
跨硬件兼容性：
通过编译器抽象硬件细节，兼容不同GPU架构（如NVIDIA/AMD），类似文中提到的“后向兼容性”。

3. 与文中模型的关联

与High-level programming model的对比：
文中提到High-level模型“相当于Triton编程模型”，均采用SPMD思想，通过高层次抽象（如张量运算符）隐藏底层并行细节，但Triton更专注于GPU，而文中模型可能针对NPU（神经网络处理器）设计。
与CUDA的对比：
Triton比CUDA更易用（无需显式管理线程/内存），但性能接近手写CUDA；文中提到当Tile统一时，数据流模型会退化为类似CUDA/Triton的模式。

4. 典型应用场景

机器学习算子优化：
如矩阵乘法（GEMM）、卷积、注意力机制等，适合需要高效张量计算的场景。
替代部分CUDA开发：
当需要快速实现高性能内核且不想深入CUDA时，Triton是理想选择。

5. 与其他模型的区别

特性	Triton	CUDA	文中High-level模型
编程复杂度	低（Python-like）	高（需管理线程/内存）	中（类似SPMD+intrinsic）
硬件目标	GPU	GPU	NPU
内存管理	自动优化	手动控制	可能部分自动（依赖intrinsic）
适用场景	通用张量计算	通用GPU计算	NPU专用计算

总结

Triton是一种面向GPU的高生产力编程模型，通过抽象并行计算和内存管理细节，平衡了易用性与性能。文中提到的NPU编程模型（High-level）与其理念相似，但针对不同硬件（NPU）设计，可能融合了更多数据流和专用加速器特性。若需进一步探讨NPU模型的细节（如intrinsic或CT Kernel），可继续深入分析。

Intrinsic

在上下文提到的 NPU（神经网络处理器）编程模型 中，intrinsic 指的是由硬件或底层系统提供的、可直接调用的高效原子操作或指令接口，用于实现对特定硬件功能（如内存访问、计算单元、通信等）的底层控制。以下是具体解析：

1. Intrinsic 的核心作用

硬件功能的抽象接口：
Intrinsic 是对 NPU 硬件指令（如LSU数据搬运、NE张量计算、核间通信等）的封装，开发者无需手写汇编即可调用这些底层操作。
平衡效率与易用性：
相比直接写汇编，intrinsic 通过高级语言（如C/Python）暴露硬件能力，既保持性能，又简化编程。

2. 文中与 Intrinsic 相关的场景

High-level 模型中的使用：
- 通过 intrinsic 配合张量运算符（如加减乘除）访问全局内存（Gmem）或 NPU 的 L1 缓存。
- 例如：调用 load_intrinsic 从 Gmem 加载数据块到 L1，再通过 mul_intrinsic 触发乘计算。
- 意义：将数据流控制（如搬移、计算、通信）抽象为高级接口，类似 CUDA 中的 __ldg 或 Triton 的 tl.load。
Low-level 模型中的实现：
- Low-level 模型直接实现这些 intrinsic 的底层逻辑（如生成 NPU 指令），供 High-level 调用。
- 例如：High-level 的 add_intrinsic 在 Low-level 中可能对应一条 NPU 的向量加法指令 VADD。

3. 类比其他领域的 Intrinsic

CUDA：
NVIDIA 提供的 __sinf、__shfl_xor_sync 等函数，直接映射到 GPU 指令。
CPU SIMD：
Intel 的 SSE/AVX 指令集（如 _mm256_add_ps）也是 intrinsic，封装了向量运算。
Triton：
tl.dot()、tl.where() 等函数本质是 GPU 张量操作的 intrinsic。

4. 文中 Intrinsic 的具体功能

根据上下文，NPU 的 intrinsic 可能包括以下类型：

内存操作：
- load/store_intrinsic：在 Gmem 和 L1 之间搬运数据块。
- 需考虑 L1/LO 存储空间限制（如 Low-level 中提到的 LO 容量有限）。
计算操作：
- 张量加减乘除、激活函数等，直接调用 NPU 计算单元（NE）。
通信操作：
- Tile 间数据同步或核间通信的 intrinsic（如 send/receive_intrinsic）。
控制操作：
- 与 CPU Scalar 协作的分支控制（如条件判断的 intrinsic）。

5. 为什么需要 Intrinsic？

性能关键路径优化：
直接调用硬件指令避免高级语言的开销（如循环/函数调用）。
硬件特性暴露：
利用 NPU 的专用计算单元（如矩阵乘法加速器）、定制内存层级（L1/LO/LOA）。
跨层级编程桥梁：
High-level 模型通过 intrinsic 依赖 Low-level 实现，内部开发者可定制优化。

6. 与 “CT Kernel 微指令” 的关系

层级差异：
- Intrinsic 属于 High/Low-level 接口，而 CT Kernel 编程（如向量 Load/Store）是更底层的微指令。
- 例如：High-level 调用 mul_intrinsic → Low-level 生成 NPU 指令 → CT Kernel 将其拆解为微指令（如 VLOAD + VMUL + VSTORE）。

总结

在文中 NPU 编程模型中，intrinsic 是连接高层抽象（如张量运算）与底层硬件指令的关键接口，使开发者既能保持对硬件的控制力，又无需深入汇编细节。这种设计在专用加速器（如NPU/GPU）中极为常见，平衡了性能与开发效率。

Dataflow Programming Model

数据流编程模型（Dataflow programming model），也称空间编程模型，即将应用程序的任务划分到不同的 Tile 区域中。如图 6.2 所示，4 段任务（Layer）被划分到 4 个 Tile 区域中，Tile 区域与 Tile 区域之间以数据流通的方式串通起来。所以，数据流编程模型需要做 Tile 在空间维度的划分。然后每个 Tile 区域内的 Layer 又分为数据流接收（Wait/Recv Func）、计算（Compute Func）、数据流发送（Push/Send Func）三个 Function 进行编程。Wait/Recv Func 需要对 Tile 区域内各 Tile 接收数据的控制进行编程；Compute Func 需要对 Tile 区域内多 Tile 完成计算任务进行编程，这里不仅包括计算，还包括全局存储的访问以及 Tile 区域内的通信；Push/Send Func 需要对 Tile 区域内各 Tile 发送数据的控制进行编程。

Wait/Recv Func 的数据给到 Compute Func，Compute Func 的数据给到 Push/Send Func，它们三者间以管道流水的方式并行执行。Compute Func 是必需的，但是 Wait/Recv Func 和 Push/Send Func 不是必需的。例如图 5.2 中，Layer 0 数据来源不是其它 Tile 区域，则不需要 Wait/Recv Func；Layer 3 数据不需要发送到其它 Tile 区域，所以不需要 Push/Send Func。若当 Layer 的数据来源既不是其它 Tile 区域，也不发送数据给其它 Tile 区域，此时仅需要 Compute Func，如图 6.3 所示。

数据流编程模型下，手工写程序完成Tile区域空间划分比较困难，所以数据流编程模型的程序一般需要由编译器自动化生成。

High Level Programming Model

High-level programming model (高层次编程模型) 用于为上一节提到的Tile区域内的 Wait/Recv Func、Compute Func 和 Push/Send Func 进行编程。在该编程模型下，仅全局存储（Gmem）和NPU内部的 L1缓存 对程序员可见（如图6.4所示）。可使用的语法包括：

Gmem ↔ L1 的数据加载/存储（Load/Store）
L1 ↔ L1 的NPU通信原语
L1数据块计算操作 的Intrinsic接口
CPU Scalar 用于计算索引和分支控制

该模型采用 SPMD（Single Program Multiple Data） 方式编程，即程序员只需编写一个适用于单个Tile（或虚拟化为一个Program ID）的程序，通过ID计算本Tile所需的源数据、完成计算并输出结果数据块。这样，同一程序可部署到多个Tile中并行执行，因此该编程模型与 Triton 类似。

在编程语言设计上，建议对 Gmem 和 L1 进行抽象（例如使用数组、指针、结构体表示，而非直接操作地址）。推荐选用 Triton 语言并在此基础上扩展——Triton是第三方编程语言，目前在深度学习领域逐渐被广泛接受，其语法符合此处高层次编程模型的定义。

特殊情况：
若Tile区域仅需 Compute Func，则编程方式与Triton完全一致。
Triton编程语言将尽可能实现完全的生态兼容性。

SIMT Programming Model

SIMT编程模型（SIMT Programming Model） 同样属于**高层次编程模型（High-level Programming Model）**的范畴。SIMT编程的代表性框架包括：

NVIDIA的CUDA
AMD的HIP
开源的OpenCL

NPU 的 SIMT 编程模型 可参考上述 SIMT 编程语言进行开发。如图6.5所示，在 NPU SIMT 编程模型 中，可见的存储结构包括：

Shared Memory（共享内存）：供一个 Thread Block（线程块） 内所有线程共享，由软件管理。

NPU-SIMT 架构 包含 3 种 Shared Memory：

默认 Shared Memory（LOA Shared Memory）
- 对应 NVIDIA GPU SM 的 Shared Memory
- 对应 AMD GPU 的 LDS（Local Data Share）
- 实际硬件实现为 LOA SPM（Scratchpad Memory）
LOW Shared Memory
- 对应 LOW SPM
L1 Shared Memory
- 对应 L1 SPM

SIMT 编程的访问规则

LOA Shared Memory
- Thread（线程） 可直接以 Scalar（标量）编程方式 对数据进行计算操作。
LOW Shared Memory 和 L1 Shared Memory
- 不允许 Thread 直接计算，必须通过 Warp-level（线程束级） 方式访问：
  - LOW Shared Memory
    - 需通过 Warp-level 操作 NE（计算单元） 进行计算访问，例如 SIMT 的矩阵乘法（matrix multiply）和卷积（conv）Intrinsic 原语。
  - L1 Shared Memory 需通过 Warp-level 的内存拷贝（memcopy） 访问，例如将 L1 数据搬运至计算单元进行处理。这种分层存储访问机制的设计，旨在平衡计算效率与硬件资源利用率：

LOA Shared Memory 支持细粒度线程级操作，适合高并行计算任务；
LOW/L1 Shared Memory 通过 Warp 级批处理访问，可优化带宽密集型操作（如矩阵乘、卷积）。

与前文 High-level 模型的关联

编程层级关系
- High-level 模型（含 SIMT）→ Low-level 模型（指令/Intrinsic）→ CT Kernel（微指令）。
- SIMT 作为 High-level 的实现方式之一，与 Triton 类似，但更贴近硬件线程管理。
存储层级对比

存储类型 SIMT 模型 High-level 模型
全局存储 Gmem Gmem
高速缓存 LOA/LOW/L1 Shared Mem L1 Cache
编程可见性 Thread/Warp 级控制 Tile 级抽象（SPMD）