![[AI][昇腾950]SIMT 编程](images/news1.jpg)
DaVinci 950 SIMT 编程介绍1. 概述DaVinci 950 VecCoreAIV支持两种向量执行模式模式全称编程范式SIMDSingle Instruction Multiple Data一条指令操作整条向量寄存器VL宽SIMTSingle Instruction Multiple ThreadsCUDA风格的细粒度线程并行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 关键参数参数值最大线程数/Core2048Warp 大小32 threads最大 Warp 数64Warp 调度器4 个寄存器文件128 KBDCache32-128 KB可编程Divergence Stack126 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_CALLsimt_code( Dim3(1024,1,1), gmAddr0x1234, ubAddr0x12, scalarValue10);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-S63Read-Only所有线程共享由 Parameter Buffer 传入S64-S95Read-Write辅助标量寄存器SIMT 模式下硬件管理S0 是常量零S1 是 16’h0。S60-S63 作为循环计数器自动清零。5. 内存访问5.1 内存层次区域指令带宽用途Global Memory (GM/OUT)VLD/VST经 DCache外部 DDR/HBMShared Memory (UB)VLDS/VSTS128B/cycle线程块内共享Register (VREG)寄存器操作最快线程私有5.2 共享内存 (Shared Memory)共享内存由 UBUnified Buffer分区提供所有同一 Thread Block 的线程共享同一块 UB 区域通过VLDS/VSTS指令访问32B 对齐Bank 交织设计需MEMBAR栅栏保证写入可见性5.3 DCache大小可编程32KB - 128KBCache Line128B2-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 之后的内存访问之前完成Scopethread/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 again7. 编程模式 — 与 CUDA 对照7.1 概念映射CUDADaVinci 950 SIMTthreadIdx.x/y/zblockDim.x/y/z__global__ void kernel(...)软件封装出的模式__shared__ float s[256]UB 分区通过 VLDS/VSTS 访问__syncthreads()Thread Block Barrier__threadfence()MEMBAR7.2 执行流程对比// Host 端kernelgrid,block,shared_mem(args);// Device 端__global__voidkernel(float*in,float*out){inttidthreadIdx.xblockIdx.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 执行 SEND8. SIMT 指令子集8.1 算术指令SIMT 模式下大部分 SIMD 算术指令可用但每个线程操作独立的数据元素类别指令类型加减VADD, VSUBu8/s8/u16/s16/u32/s32/f16/f32/bf16乘法VMUL, VMULAu16/s16/u32/s32/f16/f32/bf16乘加VFMA, VFMS, VFNMA, VFNMSf16/f32/bf16比较VCMP (EQ/NE/LT/GT/GE/LE)all int f16/f32/bf16最大最小VMAX, VMIN全类型标量-向量VADDS, VMULS, VMAXS, VMINS全类型激活函数VRELU, VLRELU, VPRELUf16/f32数学函数VEXP, VLN, VSQRTf16/f32类型转换VCVTFF, VCVTFI, VCVTII多种格式8.2 数据搬移指令源 → 目的说明VLDGM → VREG全局内存加载经 DCacheVSTVREG → GM全局内存存储VLDSUB → VREG共享内存加载S 寄存器偏移VSTSVREG → UB共享内存存储VLDIGM → VREG立即偏移加载VSTIVREG → GM立即偏移存储VGATHER2GM → VREG间接索引加载VSCATTERVREG → GM间接索引存储8.3 控制流机制说明分支Warp 级 SIMT 分支divergence stack 管理BarrierThread Block 同步屏障MEMBAR内存栅栏保证可见性ENDSIMT 线程结束9. 性能优化指南9.1 Warp 调度优化策略说明Warp 级并行保持 ≥4 个活跃 Warp隐藏延迟减少 Divergence尽量让同一 Warp 的线程走相同路径连续访存相邻线程访问连续地址合并为单次事务寄存器平衡更多线程 vs 更多寄存器/线程的取舍9.2 共享内存优化策略说明Bank 感知避免 Warp 内多线程访问同一 BankPadding在数组维度间插入 padding 避免 Bank 冲突Double BufferPing-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 线程/Core2每线程寄存器数必须是 2 的幂3总寄存器不能超限4Warp 内 SIMT 执行32 线程 lock-step分支导致 mask5Divergence stack 126 entries深度嵌套分支可能溢出6Barrier 必须成对所有 Warp 必须到达同一 barrier7DCache 128B 对齐Cache line 对齐获得最佳性能8无精确异常OOO 执行异常时标记当前周期范围9GM 访问经 DCache使用 MEMBAR 保证一致性11. SIMD vs SIMT 选择指南场景推荐模式原因规则数据并行矩阵乘、卷积SIMDVLOOPv2 VREG 高效批量处理规则归约SIMDVCGMAX → VCMAX 两阶段归约不规则并行稀疏、图SIMT线程独立索引灵活分支条件密集计算SIMT分支 mask 自然处理需要共享内存协作SIMTThread block barrier 原生支持简单元素级操作SIMD单条向量指令覆盖全部数据小计算量 kernelSIMDSIMT 启动开销不值得
[AI][昇腾950]SIMT 编程
DaVinci 950 SIMT 编程介绍1. 概述DaVinci 950 VecCoreAIV支持两种向量执行模式模式全称编程范式SIMDSingle Instruction Multiple Data一条指令操作整条向量寄存器VL宽SIMTSingle Instruction Multiple ThreadsCUDA风格的细粒度线程并行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 关键参数参数值最大线程数/Core2048Warp 大小32 threads最大 Warp 数64Warp 调度器4 个寄存器文件128 KBDCache32-128 KB可编程Divergence Stack126 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_CALLsimt_code( Dim3(1024,1,1), gmAddr0x1234, ubAddr0x12, scalarValue10);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-S63Read-Only所有线程共享由 Parameter Buffer 传入S64-S95Read-Write辅助标量寄存器SIMT 模式下硬件管理S0 是常量零S1 是 16’h0。S60-S63 作为循环计数器自动清零。5. 内存访问5.1 内存层次区域指令带宽用途Global Memory (GM/OUT)VLD/VST经 DCache外部 DDR/HBMShared Memory (UB)VLDS/VSTS128B/cycle线程块内共享Register (VREG)寄存器操作最快线程私有5.2 共享内存 (Shared Memory)共享内存由 UBUnified Buffer分区提供所有同一 Thread Block 的线程共享同一块 UB 区域通过VLDS/VSTS指令访问32B 对齐Bank 交织设计需MEMBAR栅栏保证写入可见性5.3 DCache大小可编程32KB - 128KBCache Line128B2-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 之后的内存访问之前完成Scopethread/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 again7. 编程模式 — 与 CUDA 对照7.1 概念映射CUDADaVinci 950 SIMTthreadIdx.x/y/zblockDim.x/y/z__global__ void kernel(...)软件封装出的模式__shared__ float s[256]UB 分区通过 VLDS/VSTS 访问__syncthreads()Thread Block Barrier__threadfence()MEMBAR7.2 执行流程对比// Host 端kernelgrid,block,shared_mem(args);// Device 端__global__voidkernel(float*in,float*out){inttidthreadIdx.xblockIdx.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 执行 SEND8. SIMT 指令子集8.1 算术指令SIMT 模式下大部分 SIMD 算术指令可用但每个线程操作独立的数据元素类别指令类型加减VADD, VSUBu8/s8/u16/s16/u32/s32/f16/f32/bf16乘法VMUL, VMULAu16/s16/u32/s32/f16/f32/bf16乘加VFMA, VFMS, VFNMA, VFNMSf16/f32/bf16比较VCMP (EQ/NE/LT/GT/GE/LE)all int f16/f32/bf16最大最小VMAX, VMIN全类型标量-向量VADDS, VMULS, VMAXS, VMINS全类型激活函数VRELU, VLRELU, VPRELUf16/f32数学函数VEXP, VLN, VSQRTf16/f32类型转换VCVTFF, VCVTFI, VCVTII多种格式8.2 数据搬移指令源 → 目的说明VLDGM → VREG全局内存加载经 DCacheVSTVREG → GM全局内存存储VLDSUB → VREG共享内存加载S 寄存器偏移VSTSVREG → UB共享内存存储VLDIGM → VREG立即偏移加载VSTIVREG → GM立即偏移存储VGATHER2GM → VREG间接索引加载VSCATTERVREG → GM间接索引存储8.3 控制流机制说明分支Warp 级 SIMT 分支divergence stack 管理BarrierThread Block 同步屏障MEMBAR内存栅栏保证可见性ENDSIMT 线程结束9. 性能优化指南9.1 Warp 调度优化策略说明Warp 级并行保持 ≥4 个活跃 Warp隐藏延迟减少 Divergence尽量让同一 Warp 的线程走相同路径连续访存相邻线程访问连续地址合并为单次事务寄存器平衡更多线程 vs 更多寄存器/线程的取舍9.2 共享内存优化策略说明Bank 感知避免 Warp 内多线程访问同一 BankPadding在数组维度间插入 padding 避免 Bank 冲突Double BufferPing-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 线程/Core2每线程寄存器数必须是 2 的幂3总寄存器不能超限4Warp 内 SIMT 执行32 线程 lock-step分支导致 mask5Divergence stack 126 entries深度嵌套分支可能溢出6Barrier 必须成对所有 Warp 必须到达同一 barrier7DCache 128B 对齐Cache line 对齐获得最佳性能8无精确异常OOO 执行异常时标记当前周期范围9GM 访问经 DCache使用 MEMBAR 保证一致性11. SIMD vs SIMT 选择指南场景推荐模式原因规则数据并行矩阵乘、卷积SIMDVLOOPv2 VREG 高效批量处理规则归约SIMDVCGMAX → VCMAX 两阶段归约不规则并行稀疏、图SIMT线程独立索引灵活分支条件密集计算SIMT分支 mask 自然处理需要共享内存协作SIMTThread block barrier 原生支持简单元素级操作SIMD单条向量指令覆盖全部数据小计算量 kernelSIMDSIMT 启动开销不值得
