编程模型【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit适用场景在SIMD与SIMT混合编程中SIMT能够简化复杂算子与不规则控制流的开发而SIMD基于向量寄存器与指令实现高效的数据并行处理即单指令处理多数据提升每周期的吞吐量。SIMD与SIMT混合编程支持开发者根据算子特征进行精细化映射规则的逐元素elementwise操作通过SIMD获得高带宽和高算力利用率不规则或包含分支的计算通过SIMT来缓解发散和控制复杂度。在系统层面这有利于提高硬件利用率和能效同时也更便于进行算子融合和数据复用等优化。同一个算子中既包含SIMD擅长的连续规整计算也包含SIMT擅长的离散访问等任务从而在同一算子中同时利用SIMD和SIMT的优势。线程架构SIMT编程模型采用分层的抽象线程组织结构示意图如下图所示从顶层到底层依次为Grid线程块网格、Thread Block线程块、Warp线程束和Thread线程。这种层次化的设计使得开发者可以方便地将并行计算问题映射到硬件资源上。图 1线程结构示意图Thread线程线程是整个架构中的最小单元每个线程独立完成计算任务拥有独立的寄存器和栈空间。Thread Block线程块Thread Block是Grid的组成单元由若干线程组成最大2048个线程。使用内置变量blockDim来表示一个Thread Block启用的线程数量。在SIMD与SIMT混合编程场景中Thread Block有以下特点使用asc_vf_call调用SIMT VF时需显式指定Thread Block的维度信息同一Thread Block内的线程可以访问共享内存UB实现数据交互Thread Block内的线程可通过同步机制实现协作在SIMT函数定义时可使用__launch_bounds__()配置最大线程数用于编译期底层对寄存器资源进行预分配。Grid线程块网格Grid是SIMT线程层次结构的最顶层由多个Thread Block组成。使用内置变量gridDim来表示Grid中启用的线程块数量。在SIMD与SIMT混合编程场景Grid有以下特征Grid的维度配置由用户设置启动的AIV核数决定核函数执行期间不可更改Grid中的所有线程块具有相同的尺寸和维度配置同一Grid中的线程块相互独立按任意顺序执行限制说明如下Grid中线程块总数不能超过65535。由于gridDim实际就是用户配置的启动AIV核数当设置的核数超过物理核数时头尾开销会增大强烈建议启动的AIV核数小于物理核数。线程索引每个线程都有唯一的标识开发者可以通过内置变量获取线程的信息从而确定每个线程负责处理的数据相关内置变量如下表所示内置变量说明数据类型约束gridDimGrid的维度大小。dim3gridDim.x 65535gridDim.y、gridDim.z必须为1blockDimThread Block的维度大小。dim3blockDim.x * blockDim.y *blockDim.z 2048blockIdx当前Thread Block在Grid中的索引。dim3无threadIdx当前线程在Thread Block内的索引。dim3无各Thread Block可通过线程块索引blockIdx进行标识各线程可通过线程块内线程索引threadIdx进行标识。对于一维Grid和Block索引计算公式为int idx blockIdx.x * blockDim.x threadIdx.x;对于二维Grid和Block索引计算公式为int x_index blockIdx.x * blockDim.x threadIdx.x; int y_index blockIdx.y * blockDim.y threadIdx.y;Warp执行机制Warp是SIMT架构中基本的调度和执行单位。每个Warp包含32个线程这些线程从相同的程序地址开始执行拥有各自的指令地址计数器和寄存器状态并且可以选择分支独立执行。在一个Thread Block内所有线程按线性顺序被硬件自动划分为每32个线程一组的Warp同一Warp内的所有线程执行同一条指令。Warp内的线程虽然执行同一段代码但可通过条件分支进入不同的执行路径这种情况称为分支发散Warp Divergence。当Warp中的32个线程均执行相同的代码分支时硬件利用率最高一旦发生分支发散硬件会串行执行每个分支路径只有进入当前分支的线程即活跃线程会被执行其余线程则被屏蔽从而导致Warp执行效率下降。此外同一Warp内的线程相互独立彼此不能存在依赖关系。Thread Block的线程数建议设置为32的整数倍。若线程数未满足该要求最后一个Warp将包含不足32个线程导致该Warp内存在空闲线程通道进而降低执行效率。UB划分UB即Unified Buffer是同一线程块内所有线程均可访问的内存空间位于每个AIV内部。UB内存空间总大小为256KB参考图2按功能划分为四个主要区域从低地址向高地址依次为静态内存、动态内存、 预留空间 、Data Cache。图 2UB内存分配图内存空间说明静态内存从内存的起始地址分配一段指定大小的内存空间其大小在编译时确定不可动态修改。// 静态内存通过数组分配例如 __ubuf__ char static_buf[1024];动态内存位于静态内存之后通过...中参数dyn_ub_size指定的动态内存大小空间可通过以下方式申请使用使用动态数组分配。// 动态内存通过动态数组分配例如 extern __ubuf__ char dynamic_buf[];通过LocalMemAllocator的Alloc接口申请。由于申请动态内存时均从静态内存结束位置之后开始分配如果同时使用可能会导致地址空间重叠从而引发未定义行为因此只能选择其中一种方法进行申请。预留空间系统预留空间大小固定为8KB。Data CacheSIMT专有的Data Cache空间UB内扣除静态内存、动态内存以及预留空间以后剩余内存大小为Data CacheData Cache最小为32KB剩余空间超过128KB时Data Cache大小固定为128KB具体计算公式为DataCache UB总大小256KB – 静态内存 – 动态内存 – 预留空间(8KB若DataCache小于32KB会出现校验报错。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANN/asc-devkit：混合编程模型
发布时间：2026/5/21 23:08:34
编程模型【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit适用场景在SIMD与SIMT混合编程中SIMT能够简化复杂算子与不规则控制流的开发而SIMD基于向量寄存器与指令实现高效的数据并行处理即单指令处理多数据提升每周期的吞吐量。SIMD与SIMT混合编程支持开发者根据算子特征进行精细化映射规则的逐元素elementwise操作通过SIMD获得高带宽和高算力利用率不规则或包含分支的计算通过SIMT来缓解发散和控制复杂度。在系统层面这有利于提高硬件利用率和能效同时也更便于进行算子融合和数据复用等优化。同一个算子中既包含SIMD擅长的连续规整计算也包含SIMT擅长的离散访问等任务从而在同一算子中同时利用SIMD和SIMT的优势。线程架构SIMT编程模型采用分层的抽象线程组织结构示意图如下图所示从顶层到底层依次为Grid线程块网格、Thread Block线程块、Warp线程束和Thread线程。这种层次化的设计使得开发者可以方便地将并行计算问题映射到硬件资源上。图 1线程结构示意图Thread线程线程是整个架构中的最小单元每个线程独立完成计算任务拥有独立的寄存器和栈空间。Thread Block线程块Thread Block是Grid的组成单元由若干线程组成最大2048个线程。使用内置变量blockDim来表示一个Thread Block启用的线程数量。在SIMD与SIMT混合编程场景中Thread Block有以下特点使用asc_vf_call调用SIMT VF时需显式指定Thread Block的维度信息同一Thread Block内的线程可以访问共享内存UB实现数据交互Thread Block内的线程可通过同步机制实现协作在SIMT函数定义时可使用__launch_bounds__()配置最大线程数用于编译期底层对寄存器资源进行预分配。Grid线程块网格Grid是SIMT线程层次结构的最顶层由多个Thread Block组成。使用内置变量gridDim来表示Grid中启用的线程块数量。在SIMD与SIMT混合编程场景Grid有以下特征Grid的维度配置由用户设置启动的AIV核数决定核函数执行期间不可更改Grid中的所有线程块具有相同的尺寸和维度配置同一Grid中的线程块相互独立按任意顺序执行限制说明如下Grid中线程块总数不能超过65535。由于gridDim实际就是用户配置的启动AIV核数当设置的核数超过物理核数时头尾开销会增大强烈建议启动的AIV核数小于物理核数。线程索引每个线程都有唯一的标识开发者可以通过内置变量获取线程的信息从而确定每个线程负责处理的数据相关内置变量如下表所示内置变量说明数据类型约束gridDimGrid的维度大小。dim3gridDim.x 65535gridDim.y、gridDim.z必须为1blockDimThread Block的维度大小。dim3blockDim.x * blockDim.y *blockDim.z 2048blockIdx当前Thread Block在Grid中的索引。dim3无threadIdx当前线程在Thread Block内的索引。dim3无各Thread Block可通过线程块索引blockIdx进行标识各线程可通过线程块内线程索引threadIdx进行标识。对于一维Grid和Block索引计算公式为int idx blockIdx.x * blockDim.x threadIdx.x;对于二维Grid和Block索引计算公式为int x_index blockIdx.x * blockDim.x threadIdx.x; int y_index blockIdx.y * blockDim.y threadIdx.y;Warp执行机制Warp是SIMT架构中基本的调度和执行单位。每个Warp包含32个线程这些线程从相同的程序地址开始执行拥有各自的指令地址计数器和寄存器状态并且可以选择分支独立执行。在一个Thread Block内所有线程按线性顺序被硬件自动划分为每32个线程一组的Warp同一Warp内的所有线程执行同一条指令。Warp内的线程虽然执行同一段代码但可通过条件分支进入不同的执行路径这种情况称为分支发散Warp Divergence。当Warp中的32个线程均执行相同的代码分支时硬件利用率最高一旦发生分支发散硬件会串行执行每个分支路径只有进入当前分支的线程即活跃线程会被执行其余线程则被屏蔽从而导致Warp执行效率下降。此外同一Warp内的线程相互独立彼此不能存在依赖关系。Thread Block的线程数建议设置为32的整数倍。若线程数未满足该要求最后一个Warp将包含不足32个线程导致该Warp内存在空闲线程通道进而降低执行效率。UB划分UB即Unified Buffer是同一线程块内所有线程均可访问的内存空间位于每个AIV内部。UB内存空间总大小为256KB参考图2按功能划分为四个主要区域从低地址向高地址依次为静态内存、动态内存、 预留空间 、Data Cache。图 2UB内存分配图内存空间说明静态内存从内存的起始地址分配一段指定大小的内存空间其大小在编译时确定不可动态修改。// 静态内存通过数组分配例如 __ubuf__ char static_buf[1024];动态内存位于静态内存之后通过...中参数dyn_ub_size指定的动态内存大小空间可通过以下方式申请使用使用动态数组分配。// 动态内存通过动态数组分配例如 extern __ubuf__ char dynamic_buf[];通过LocalMemAllocator的Alloc接口申请。由于申请动态内存时均从静态内存结束位置之后开始分配如果同时使用可能会导致地址空间重叠从而引发未定义行为因此只能选择其中一种方法进行申请。预留空间系统预留空间大小固定为8KB。Data CacheSIMT专有的Data Cache空间UB内扣除静态内存、动态内存以及预留空间以后剩余内存大小为Data CacheData Cache最小为32KB剩余空间超过128KB时Data Cache大小固定为128KB具体计算公式为DataCache UB总大小256KB – 静态内存 – 动态内存 – 预留空间(8KB若DataCache小于32KB会出现校验报错。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
