核函数配置【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit核函数定义核函数是SIMD与SIMT混合编程的Device侧入口函数负责协调整个算子的执行流程包括VF的调度和调用。vector计算单元的混合编程场景下函数定义语法为__global__ __vector__ void kernel_name(__gm__ type* param1, __gm__ type* param2, ...);关键修饰符说明如下__global__必需修饰符作用为标识核函数表明可在Host侧通过...调用。__vector__必需修饰符作用为标识函数是在Device侧AIV核上执行。核函数定义有以下几个约束返回值类型必须是void入参支持指针类型需使用__gm__修饰和Ascend C内置数据类型指针参数必须是指向Global Memory上的内存地址使用__gm__修饰。__launch_bounds__(N)在多线程并发执行时每个线程使用较少的寄存器可以让更多的线程和线程块驻留在AI处理器上从而提升性能。因此编译器会采用启发式算法将寄存器溢出register spilling和指令数量控制在最低水平同时尽量减少寄存器的使用量。应用程序可以通过在__global__函数定义中使用__launch_bounds__()限定符来限制启动边界launch bounds提供附加信息辅助编译器优化这一过程这属于可选配置。__launch_bounds__(N)函数标记宏在SIMT VF入口函数上可选配置用于在编译期指定SIMT VF启动的最大线程数。若未配置__launch_bounds__最大线程数默认为1024。参数N需要满足N dimx * dimy * dimzdimxdimydimz为表示线程的dim3结构体。N的取值范围为1到2048。最大线程数决定了每个线程可分配的寄存器数量具体对应关系请见下表寄存器用于存储线程中的局部变量若局部变量的个数超出寄存器个数容易出现栈溢出等问题。建议最大线程数与启动VF任务的dim3线程数保持一致。表 1__launch_bounds__的Thread数量与每个Thread可用寄存器数Thread的个数(个)每个Thread可用寄存器个数(个)1025~204816513~102432257~512641~256127调用核函数的调用是通过...内核调用符在Host侧调用语法如下kernel_nameblock_num, dyn_ub_size, stream(args...);内核调用符内的配置参数说明如下参数类型说明约束block_numuint32_t设置核函数启用的核数取值范围[1, 65535]dyn_ub_sizeuint32_t指定动态内存大小单位为字节不超过最大可配置值256KB - 8KB - 32KB - 静态内存streamaclrtStream用于维护异步操作执行顺序无asc_vf_call调用在SIMD与SIMT混合编程场景需使用asc_vf_call启动SIMT VFVector Function子任务通过参数配置启动指定数目的线程执行指定的SIMT VF函数。其函数原型如下template auto funcPtr, typename... Args __aicore__ inline void asc_vf_call(dim3 threadNums, Args ...args)其中模板参数为指定的SIMT VF函数名以及SIMT VF函数参数具体描述如下表 2模板参数说明参数名描述funcPtr用于指定SIMT入口核函数。Args定义可变参数用于传递实参到SIMT入口核函数。表 3参数说明参数名输入/输出描述threadNums输入dim3结构定义为{dimxdimydimz}用于指定SIMT线程块内线程数量。线程总数为dimx * dimy * dimz该值的大小必须小于等于2048建议为32的倍数。args输入可变参数用于传递实参到SIMT入口核函数。以下示例展示了SIMD与SIMT混合编程场景下如何使用asc_vf_call调用__simt_vf__函数。#include simt_api/common_functions.h __simt_vf__ inline void add_simt( __gm__ float* dst, __gm__ float* src0, __gm__ float* src1) { // simt 代码 int idx threadIdx.x blockIdx.x * blockDim.x; dst[idx] src0[idx] src1[idx]; } __global__ __vector__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { // asc_vf_call启动SIMT VF子任务函数名为add_simt配置blockDim为dim3{1024, 1, 1} asc_vf_calladd_simt(dim3{1024, 1, 1}, z, x, y); }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考