CANN/asc-devkit核函数配置

核函数配置【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit核函数定义核函数是SIMT编程的Device侧入口函数负责协调整个算子的执行流程。函数定义语法为__global__ void kernel_name(uint32_t* param1, float* param2, ...);关键修饰符说明如下__global__必需修饰符作用为标识核函数表明可在Host侧通过...调用。核函数定义有以下几个约束返回值类型必须是void入参支持基础数据类型和基础数据类型的指针类型如int32_t、float、int32_t*、float*等。指针参数必须是指向Global Memory上的内存地址。__launch_bounds__(N)在多线程并发执行时每个线程使用较少的寄存器可以让更多的线程和线程块驻留在AI处理器上从而提升性能。因此编译器会采用启发式算法将寄存器溢出register spilling和指令数量控制在最低水平同时尽量减少寄存器的使用量。应用程序可以通过在__global__函数定义中使用__launch_bounds__()限定符来限制启动边界launch bounds提供附加信息辅助编译器优化这一过程这属于可选配置。__launch_bounds__()函数标记宏在核函数上可选配置用于在编译期指定核函数启动的最大线程数。若未配置__launch_bounds__最大线程数默认为1024。参数N需要满足N dimx * dimy * dimzdimxdimydimz为表示线程的dim3结构体。N的取值范围为1到2048。最大线程数决定了每个线程可分配的寄存器数量具体对应关系请见下表寄存器用于存储线程中的局部变量若局部变量的个数超出寄存器个数容易出现栈溢出等问题。建议最大线程数与启动核函数的dim3线程数保持一致。表 1__launch_bounds__的Thread数量与每个Thread可用寄存器数Thread的个数(个)每个Thread可用寄存器个数(个)1025~204816513~102432257~512641~256127调用在SIMT编程下使用调用__global__限定符修饰的函数时必须指定执行配置形如grid_dim, block_dim, dynamic_mem_size, stream其中grid_dimint或dim3类型用于指定网格grid的维度与规模grid_dim.x * grid_dim.y * grid_dim.z等于启动的线程块总数。block_dimint或dim3类型用于指定每个线程块block的维度与规模block_dim.x * block_dim.y * block_dim.z等于每个线程块包含的线程数需要小于等于__launch_bounds__配置。dynamic_mem_sizesize_t类型用于指定每个线程块动态分配的共享内存大小单位为字节。streamaclrtStream类型指针指定关联的流用于维护异步操作的执行顺序。以下示例展示了内核函数的声明与调用方式。// 声明 __global__ void add_custom(float* x, float* y, float* z, uint64_t total_length); // 调用 add_customblock_num, thread_num_per_block, dyn_ubuf_size, stream(x, y, z, 1024);【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考