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 * dimz；dimx，dimy，dimz为表示线程的dim3结构体。

• N的取值范围为1到2048。

  最大线程数决定了每个线程可分配的寄存器数量，具体对应关系请见下表，寄存器用于存储线程中的局部变量，若局部变量的个数超出寄存器个数，容易出现栈溢出等问题。建议最大线程数与启动核函数的dim3线程数保持一致。

  表 1__launch_bounds__的Thread数量与每个Thread可用寄存器数

  Thread的个数(个)	每个Thread可用寄存器个数(个)
  1025~2048	16
  513~1024	32
  257~512	64
  1~256	127

<<<>>>调用

在SIMT编程下使用<<<>>>调用__global__限定符修饰的函数时必须指定执行配置，形如：

<<<grid_dim, block_dim, dynamic_mem_size, stream>>>

其中：

• grid_dim：int或dim3类型，用于指定网格（grid）的维度与规模，grid_dim.x * grid_dim.y * grid_dim.z等于启动的线程块总数。
• block_dim：int或dim3类型，用于指定每个线程块（block）的维度与规模，block_dim.x * block_dim.y * block_dim.z等于每个线程块包含的线程数，需要小于等于__launch_bounds__配置。
• dynamic_mem_size：size_t类型，用于指定每个线程块动态分配的共享内存大小，单位为字节。
• stream：aclrtStream类型指针，指定关联的流，用于维护异步操作的执行顺序。

以下示例展示了内核函数的声明与调用方式。

// 声明 __global__ void add_custom(float* x, float* y, float* z, uint64_t total_length); // 调用 add_custom<<<block_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