[AI][昇腾950]SIMT 编程

mlidongfeng

21人浏览 · 2026-06-29 08:33:24

mlidongfeng · 2026-06-29 08:33:24 发布

DaVinci 950 SIMT 编程介绍

1. 概述

DaVinci 950 VecCore（AIV）支持两种向量执行模式：

模式	全称	编程范式
SIMD	Single Instruction Multiple Data	一条指令操作整条向量寄存器（VL宽）
SIMT	Single Instruction Multiple Threads	CUDA风格的细粒度线程并行

SIMT 模式通过 FORK 指令从 SIMD 模式切换进入，提供与 CUDA 高度相似的编程模型，适合不规则并行、分支密集的算法。

2. 硬件线程模型

2.1 线程层次

┌─────────────────────────────────────────────────────────┐
│                    VecCore (AIV)                        │
│                                                         │
│  ┌─ Thread Block ─────────────────────────────────────┐ │
│  │  2048 threads maximum per core                     │ │
│  │                                                    │ │
│  │  ┌─ Warp 0 ──┐  ┌─ Warp 1 ──┐       ┌─ Warp 63┐  │ │
│  │  │ 32 threads │  │ 32 threads │  ...  │ 32 thrd │  │ │
│  │  └────────────┘  └────────────┘       └─────────┘  │ │
│  │                                                    │ │
│  │  Shared Memory: UB 分区                             │ │
│  │  Register File: 128KB                               │ │
│  └────────────────────────────────────────────────────┘ │
│                                                         │
│  4 Warp Schedulers · DCache (32-128KB)                  │
└─────────────────────────────────────────────────────────┘

2.2 关键参数

参数	值
最大线程数/Core	2048
Warp 大小	32 threads
最大 Warp 数	64
Warp 调度器	4 个
寄存器文件	128 KB
DCache	32-128 KB（可编程）
Divergence Stack	126 entries/warp
共享内存	UB 分区

2.3 线程索引

每个 SIMT 线程拥有唯一的三维索引 (thread_x, thread_y, thread_z)，类似 CUDA 的 threadIdx：

维度	来源	说明
thread_x	`[11:0]`	X 维度线程数
thread_y	`[27:16]`	Y 维度线程数
thread_z	`[11:0]`	Z 维度线程数

Warp 内线程按 (z, y, x) 的行优先顺序排列，每个线程获得 1-based 的 thread_id。

3. 进入 SIMT 模式

 void simt_code(void* gmAddr, __ubuf__ int* ubAddr, int scalarValue){
    ....
 }
 VF_CALL<simt_code>( Dim3(1024,1,1), gmAddr=0x1234, ubAddr=0x12, scalarValue=10);

4. SIMT 寄存器模型

4.1 寄存器分配

每个线程拥有独立的寄存器视图：

┌─────────────────────────────────────────┐
│  128 KB Register File                   │
│                                         │
│  Thread 0:  R0  R1  R2 ... R(N-1)      │
│  Thread 1:  R0  R1  R2 ... R(N-1)      │
│  ...                                    │
│  Thread T:  R0  R1  R2 ... R(N-1)      │
│                                         │
│  N = reg_per_thread (from Sm[23:16])    │
│  T = thread_x × thread_y × thread_z    │
│  Constraint: N × T ≤ Total regs        │
└─────────────────────────────────────────┘

4.2 S 寄存器

范围	属性	说明
S0-S63	Read-Only	所有线程共享，由 Parameter Buffer 传入
S64-S95	Read-Write	辅助标量寄存器，SIMT 模式下硬件管理

S0 是常量零，S1 是 16’h0。S60-S63 作为循环计数器自动清零。

5. 内存访问

5.1 内存层次

区域	指令	带宽	用途
Global Memory (GM/OUT)	`VLD` / `VST`	经 DCache	外部 DDR/HBM
Shared Memory (UB)	`VLDS` / `VSTS`	128B/cycle	线程块内共享
Register (VREG)	寄存器操作	最快	线程私有

5.2 共享内存 (Shared Memory)

共享内存由 UB（Unified Buffer）分区提供：

所有同一 Thread Block 的线程共享同一块 UB 区域
通过 VLDS / VSTS 指令访问
32B 对齐，Bank 交织设计
需 MEMBAR 栅栏保证写入可见性

5.3 DCache

大小可编程：32KB - 128KB
Cache Line：128B
2-way 组相联
写回策略
跨 Warp 共享

6. 同步机制

6.1 Thread Block Barrier

┌─────────┐    ┌─────────┐    ┌─────────┐
│ Warp 0  │    │ Warp 1  │    │ Warp N  │
│ arriving│    │ arriving│    │ arriving│
└────┬────┘    └────┬────┘    └────┬────┘
     │              │              │
     └──────────────┼──────────────┘
                    ▼
         ┌─────────────────────┐
         │  Thread Block       │
         │  Barrier            │
         │  (all warps arrive) │
         └─────────┬───────────┘
                   ▼
         All warps resume

所有 Warp 到达 barrier 后才能继续执行
保证 barrier 之前的内存操作对所有线程可见
类似 CUDA 的 __syncthreads()

6.2 MEMBAR (Memory Fence)

MEMBAR.{scope}

保证 fence 之前的内存访问在 fence 之后的内存访问之前完成
Scope：thread/block/core 级别
用于确保共享内存写入对其他 Warp 可见

6.3 Divergence Stack

每 Warp 126 个 entry
处理同一 Warp 内线程分支
类似 CUDA 的 SIMT 执行模型：分支时 mask 掉不活跃线程， divergence stack 记录返回点

Warp (32 threads):
  if (thread_id < 16) {
    // Path A: threads 0-15 active, 16-31 masked
    ...
  } else {
    // Path B: threads 16-31 active, 0-15 masked
    ...
  }
  // Reconvergence: all 32 threads active again

7. 编程模式 — 与 CUDA 对照

7.1 概念映射

CUDA	DaVinci 950 SIMT
`threadIdx.x/y/z`
`blockDim.x/y/z`
`__global__ void kernel(...)`	软件封装出的模式
`__shared__ float s[256]`	UB 分区，通过 VLDS/VSTS 访问
`__syncthreads()`	Thread Block Barrier
`__threadfence()`	MEMBAR

7.2 执行流程对比

// Host 端
kernel<<<grid, block, shared_mem>>>(args);

// Device 端
__global__ void kernel(float* in, float* out) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    out[tid] = in[tid] * 2.0f;
}

DaVinci 950 SIMT 内部实现:

; SIMD 阶段：准备参数
SMOV S2, 0x00000080    ; thread_x = 128, thread_y = 1
SMOV S4, 0x00080001    ; thread_z = 1, reg_per_thread = 8
; S6 = input buffer address
; S8 = output buffer address
FORK S2, S4, 16, 1

; === SIMT 阶段开始 ===
; 每线程获取自己的 tid
; thread_id 可通过 SPR 获取
; 加载输入
VLD V0, [S6], A0, #normal     ; 从 Global Memory 加载
; 计算
VMUL V1, V0, S_alpha          ; 乘以 2.0（标量广播）
; 存储结果
VST V1, [S8], A0, #norm_b32   ; 写回 Global Memory
; 结束
END

; === SIMD 恢复 (PC + St - 4) ===
; 继续 SIMD 执行
SEND

8. SIMT 指令子集

8.1 算术指令

SIMT 模式下大部分 SIMD 算术指令可用，但每个线程操作独立的数据元素：

类别	指令	类型
加减	VADD, VSUB	u8/s8/u16/s16/u32/s32/f16/f32/bf16
乘法	VMUL, VMULA	u16/s16/u32/s32/f16/f32/bf16
乘加	VFMA, VFMS, VFNMA, VFNMS	f16/f32/bf16
比较	VCMP (EQ/NE/LT/GT/GE/LE)	all int + f16/f32/bf16
最大最小	VMAX, VMIN	全类型
标量-向量	VADDS, VMULS, VMAXS, VMINS	全类型
激活函数	VRELU, VLRELU, VPRELU	f16/f32
数学函数	VEXP, VLN, VSQRT	f16/f32
类型转换	VCVTFF, VCVTFI, VCVTII	多种格式

8.2 数据搬移

指令	源 → 目的	说明
`VLD`	GM → VREG	全局内存加载（经 DCache）
`VST`	VREG → GM	全局内存存储
`VLDS`	UB → VREG	共享内存加载（S 寄存器偏移）
`VSTS`	VREG → UB	共享内存存储
`VLDI`	GM → VREG	立即偏移加载
`VSTI`	VREG → GM	立即偏移存储
`VGATHER2`	GM → VREG	间接索引加载
`VSCATTER`	VREG → GM	间接索引存储

8.3 控制流

机制	说明
分支	Warp 级 SIMT 分支，divergence stack 管理
Barrier	Thread Block 同步屏障
MEMBAR	内存栅栏，保证可见性
END	SIMT 线程结束

9. 性能优化指南

9.1 Warp 调度优化

策略	说明
Warp 级并行	保持 ≥4 个活跃 Warp，隐藏延迟
减少 Divergence	尽量让同一 Warp 的线程走相同路径
连续访存	相邻线程访问连续地址，合并为单次事务
寄存器平衡	更多线程 vs 更多寄存器/线程的取舍

9.2 共享内存优化

策略	说明
Bank 感知	避免 Warp 内多线程访问同一 Bank
Padding	在数组维度间插入 padding 避免 Bank 冲突
Double Buffer	Ping-pong 读写隐藏延迟
Barrier 最小化	减少 barrier 次数，增加每个 phase 的计算量

9.3 DCache 优化

策略	说明
空间局部性	连续地址访问充分利用 128B cache line
时间局部性	重用已缓存数据，减少 GM 访问
预取	在计算当前数据时加载下一批数据

9.4 计算强度优化

Arithmetic Intensity = FLOPs / Bytes_Transferred

推荐: 计算访存比 ≥ 1:1
FP16 优先: 吞吐量是 FP32 的 2 倍
FP8/HiF8: 推理场景可达 4 倍吞吐
FMA 链: 一次融合乘加 = 2 FLOP，减少中间舍入

10. 约束清单

#	约束	说明
1	最大 2048 线程/Core
2	每线程寄存器数必须是 2 的幂
3	总寄存器不能超限
4	Warp 内 SIMT 执行	32 线程 lock-step，分支导致 mask
5	Divergence stack 126 entries	深度嵌套分支可能溢出
6	Barrier 必须成对	所有 Warp 必须到达同一 barrier
7	DCache 128B 对齐	Cache line 对齐获得最佳性能
8	无精确异常	OOO 执行，异常时标记当前周期范围
9	GM 访问经 DCache	使用 MEMBAR 保证一致性

11. SIMD vs SIMT 选择指南

场景	推荐模式	原因
规则数据并行（矩阵乘、卷积）	SIMD	VLOOPv2 + VREG 高效批量处理
规则归约	SIMD	VCGMAX → VCMAX 两阶段归约
不规则并行（稀疏、图）	SIMT	线程独立索引，灵活分支
条件密集计算	SIMT	分支 + mask 自然处理
需要共享内存协作	SIMT	Thread block barrier 原生支持
简单元素级操作	SIMD	单条向量指令覆盖全部数据
小计算量 kernel	SIMD	SIMT 启动开销不值得