tiled_partition【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况Ascend 950PR/Ascend 950DT支持Atlas A3 训练系列产品/Atlas A3 推理系列产品不支持Atlas A2 训练系列产品/Atlas A2 推理系列产品不支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品AI Core不支持Atlas 推理系列产品Vector Core不支持Atlas 训练系列产品不支持功能说明tiled_partitionAPI用于将一个线程组划分为多个更小、固定大小的子组以便线程在以更精细的粒度上进行协作。提供模板和非模板两个版本的接口分别用于编译时确定划分大小以及运行时确定划分大小的场景。函数原型template unsigned int Size, typename ParentT thread_block_tileSize, ParentT tiled_partition(const ParentT g)thread_group tiled_partition(const thread_group parent, unsigned int tilesz)thread_group tiled_partition(const thread_block parent, unsigned int tilesz)coalesced_group tiled_partition(const coalesced_group parent, unsigned int tilesz)参数说明表1模板版本参数说明参数名输入/输出描述g输入被划分的父组类型只能是thread_block或thread_block_tile。Size输入模板参数指定划分出的thread_block_tile组大小。表2非模板版本参数说明参数名输入/输出描述parent输入被划分的父组类型只能是thread_block或coalesced_group。tilesz输入指定划分出的子组大小。返回值说明返回划分出的子组对象。约束说明Size必须是$2^n$并且必须小于等于32warpSize当前可选值范围1、2、4、8、16、32。对于模板版本的接口父组中的线程数必须能被Size整除。并且如果父组是thread_block_tile则Size必须小于父组大小。调用示例SIMT编程场景using namespace cooperative_groups; __global__ void simt_kernel(...) { ... thread_block block this_thread_block(); auto tile4 tiled_partition4(block); ... }SIMD与SIMT混合编程场景using namespace cooperative_groups; __simt_vf__ inline void simt_kernel(...) { ... thread_block block this_thread_block(); auto tile4 tiled_partition4(block); ... }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
