CANN/asc-devkit SIMT线程架构指南

线程架构【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkitSIMT编程模型采用分层的抽象线程组织结构示意图如下图所示从顶层到底层依次为Grid线程块网格、Thread Block线程块和Thread线程。这种层次化的设计使得开发者可以方便地将并行计算问题映射到硬件资源上。线程层次结构Thread线程线程是整个架构中的最小单元每个线程独立完成计算任务拥有独立的寄存器和栈空间。Thread Block线程块Thread Block由若干线程组成最大2048个线程使用内置变量blockDim表示一个线程块在各维度上的线程数量。Thread Block有以下特征同一Thread Block内的线程可以访问共享内存实现数据交互Thread Block内的线程可通过同步机制实现协作Grid线程块网格Grid是SIMT线程层次结构的最顶层由多个Thread Block组成。使用内置变量gridDim来表示Grid在各维度上的线程块数量并有以下特征Grid的维度由用户的执行配置决定核函数执行期间不可更改Grid中的所有线程块具有相同的尺寸和维度配置同一Grid中的线程块相互独立按任意顺序执行基于SIMT编程模型的程序在AIV核上执行多个结构相同的线程块执行的总线程数等于gridDim.x*gridDim.y*gridDim.z*blockDim.x*blockDim.y*blockDim.z。当前Grid中线程块总数不能超过65535。线程索引在SIMT线程架构下每个线程都有唯一的标识开发者可以通过内置变量获取线程的身份信息从而确定每个线程负责处理的数据相关内置变量如下表所示内置变量说明数据类型约束gridDimGrid的维度大小。dim3gridDim.x * gridDim.y * gridDim.z 65535blockDimThread Block的维度大小。dim3blockDim.x * blockDim.y * blockDim.z 2048blockIdx当前Thread Block在Grid中的索引。dim3无threadIdx当前线程在Thread Block内的索引。dim3无各Thread Block可通过线程块索引blockIdx进行标识各线程可通过线程块内线程索引threadIdx进行标识。线程索引的计算示例如下图所示对于一维Grid和Block索引计算公式为int global_id blockIdx.x * blockDim.x threadIdx.x;对于二维Grid和Block索引计算公式为int global_id_x blockIdx.x * blockDim.x threadIdx.x; int global_id_y blockIdx.y * blockDim.y threadIdx.y;Warp执行机制实际上每个线程块内的线程并无法同时并发执行底层是通过Warp对线程块内的线程进行划分在一个Thread Block内所有线程按线性顺序被硬件自动划分为每32个线程一组的Warp同一Warp内的所有线程执行同一条指令。Warp是SIMT架构中基本的调度和执行单位。Warp内的线程从相同的程序地址开始执行拥有各自的指令地址计数器和寄存器状态并且可以选择分支独立执行。基于Warp的划分机制Thread Block的blockDim建议设置为32的整数倍否则最后一个Warp将包含不足32个线程导致该Warp内存在空闲线程通道进而降低执行效率。Warp内的线程虽然执行同一段代码但在分支处可能进入不同的执行路径这种情况称为分支发散Warp Divergence。控制流指令if、switch、do、for、while都可能会引起分支发散。当Warp中的32个线程均执行相同的代码分支时硬件利用率最高一旦发生分支发散只有进入当前分支的线程即活跃线程会被执行其余线程则被掩码屏蔽等待从而导致Warp执行效率下降。例如若Warp中的线程遇到条件分支导致threadIdx.x小于8的线程与大于等于8的线程进入不同执行路径。由于同一Warp内线程必须同步执行同一条指令硬件会串行执行各分支先执行threadIdx.x小于8所在分支大于等于8的线程被掩码屏蔽等待再执行大于等于8所在分支小于8的线程被掩码屏蔽等待分支发散结束后所有线程重新汇合继续执行如下图所示。此外同一Warp内的线程相互独立彼此不能存在依赖关系。配置最大线程数上文提到每个线程块内的线程拥有独立的寄存器和栈空间而每个AIV核总的寄存器是有限的。若算子计算任务越复杂需要使用到的寄存器越多支持启动的线程数越少。当前支持用户在核函数定义时使用__launch_bounds__(N)配置最大线程数N限制算子能启动的线程数个数上限。编译时底层编译器会根据该值对寄存器资源进行分配用户配置的N越大每个线程可用寄存器越少具体关系如下表所示配置的最大线程数每个Thread可用寄存器个数1025~204816513~102432257~512641~256127需要注意的是__launch_bounds__(N)是可选配置默认值为1024。N是编译期确定的值因此必须是静态变量或者常量。算子执行期实际启动的线程数不能超过最大线程数否则会出现运行报错。用户配置的最大线程数应与算子复杂度相匹配若N过大会导致每个线程分配到的寄存器数量过少线程执行时编译器可能会选择将部分数据存放到栈空间称为寄存器溢出register spilling影响计算性能甚至导致精度异常影响算子功能而若N过小会限制实际启动的线程数造成资源浪费。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考