[AI][昇腾950]SIMT 编程

DaVinci 950 SIMT 编程介绍

1. 概述

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

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 关键参数

2.3 线程索引

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

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 是常量零，S1 是 16’h0。S60-S63 作为循环计数器自动清零。

5. 内存访问

5.1 内存层次

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 概念映射

7.2 执行流程对比

// Host 端kernel<<<grid,block,shared_mem>>>(args);// Device 端__global__voidkernel(float*in,float*out){inttid=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 算术指令可用，但每个线程操作独立的数据元素：

8.2 数据搬移

8.3 控制流

9. 性能优化指南

9.1 Warp 调度优化

9.2 共享内存优化

9.3 DCache 优化

9.4 计算强度优化

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

10. 约束清单

11. SIMD vs SIMT 选择指南