核函数【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit正如在概述中所介绍的能够在AI处理器NPU上并行执行、并由主机端Host发起调用的函数被称为核函数Kernel。核函数被设计为由大量线程Threads同时并行运行。每个线程都拥有独立的寄存器和栈它们协同工作以完成海量的数据计算任务。核函数的定义Ascend C支持开发者自定义核函数来扩展C。定义核函数时通过引入特定的函数限定符编译器能够识别该函数需要在设备端Device编译并允许其从主机端被发起调用。核函数的定义示例如下// 核函数定义示例 __global__ void vec_add(float* x, float* y, float* z) { // 具体的计算逻辑 }定义核函数时需要遵循以下规则使用函数类型限定符__global__标识它是一个核函数。核函数必须具有void返回类型。核函数参数列表需遵循函数参数列表限制。核函数的调用算子程序中的函数可以分为三类host侧执行函数、核函数device侧执行、device侧执行函数除核函数之外。下图以Kernel直调算子开发方式为例描述三者的调用关系host侧执行函数可以调用其它host执行函数也就是通用C/C编程中的函数调用也可以通过...调用核函数。核函数可以调用除核函数之外的其它device侧执行函数。device侧执行函数除核函数之外使用类型限定符__aicore__标识可以调用同类的其它device侧执行函数。图1核函数、host侧执行函数、device侧执行函数调用关系Host侧通过核函数调用符...的语法形式调用核函数如下所示kernel_nameblocks_per_grid, threads_per_block, dyn_ubuf_size, stream(argument list)...内的参数为核函数的执行配置由4个参数决定详细用法请参考核函数配置blocks_per_gridint或dim3类型用于指定网格grid的维度与规模blocks_per_grid.x * blocks_per_grid.y * blocks_per_grid.z等于启动的线程块总数不得大于65535。threads_per_blockint或dim3类型用于指定每个线程块block的维度与规模threads_per_block.x * threads_per_block.y * threads_per_block.z等于每个线程块包含的线程数需要小于等于__launch_bounds__配置。dyn_ubuf_sizesize_t类型该参数指定除静态分配的内存外本次调用为每个线程块动态分配的共享内存字节数单位为bytes默认设置为0。具体用法请参考共享内存中的“动态申请”方式streamaclrtStream类型指针指定关联的流用于维护异步操作的执行顺序。Grid与Thread索引内置变量在核函数内Ascend C提供了内置变量来获取执行配置的参数以及线程或线程块的索引。常用的内置索引变量说明如下threadIdx获取当前线程在其所属线程块内的索引。threadIdx.xthreadIdx.ythreadIdx.z分别表示当前线程在3个维度的索引threadIdx.x的范围为[0, blockDim.x)threadIdx.y的范围为[0, blockDim.y)threadIdx.z的范围为[0, blockDim.z)。blockDim获取线程块中配置的三维层次结构即核函数启动时配置的dim3结构体实例值。blockDim.xblockDim.yblockDim.z分别表示线程块中三个维度的线程数。在核函数启动的执行配置中指定对应核函数配置中的threads_per_block参数。blockIdx: 获取当前线程块在其所属网格中的索引表示当前线程所在的线程块在整个网格中的位置坐标。gridDim表示整个计算任务在各个维度上分别由多少个线程块构成。在核函数启动的执行配置中指定对应核函数配置中的blocks_per_grid参数。各个维度上线程块关系需满足gridDim.x * gridDim.y * gridDim.z 65535。开发者可通过组合上述变量计算出当前线程在整个庞大计算网格空间中的全局唯一数据索引Global Index。这使得成千上万个执行相同代码的线程能够精准地定位并处理属于自己的特定数据分片。边界检查在实际一维并行处理开发中待处理的数据总长度往往无法被单一线程块的尺寸完美整除。因此通常需要分配多余的物理线程并在核函数内部引入边界检查机制以防止内存越界。一维数据索引的标准计算公式为int global_index blockIdx.x * blockDim.x threadIdx.x;以下代码片段展示了如何结合索引机制与边界检查实现动态长度的矢量相加// 1. Kernel implementation __global__ void vec_add(float* x, float* y, float* z, int vector_length) { // Calculate the global element index corresponding to the current thread int global_index blockIdx.x * blockDim.x threadIdx.x; // Boundary check: Constrain the computation range to prevent excess threads from accessing out-of-bounds memory if (global_index vector_length) { // Only threads with indices within the valid range will execute the specific addition operation z[global_index] x[global_index] y[global_index]; } } // 2. Host-side invocation example int main() { // Assume memory allocation and initialization on both the Host and Device sides are complete int vector_length 1000; // Empirical configuration: Specify that each block contains 256 threads int threads 256; // Calculate parameters: Calculate the required number of blocks based on the ceiling principle // Algorithm: (length threads - 1) / threads int blocks (vector_length threads - 1) / threads; // The result is 4 // Execute invocation: Launch 4 blocks, with 256 threads per block (a total of 1024 physical threads dispatched) vec_addblocks, threads, 0, stream(dev_x, dev_y, dev_z, vector_length); // ... Subsequent stream synchronization and resource release logic ... return 0; }在本示例中目标处理元素总数为1000个。按单块256线程的策略划分Host侧实际向硬件下发了4个线程块总计启动了$4 \times 256 1024$个物理线程。当程序运行至最后一个线程块即blockIdx.x为3时该块内末尾的24个线程从索引1000到1023推导出的global_index将落在[1000, 1023]区间。此时通过if (global_index vector_length)这道边界拦截这部分超出有效数据范围的“多余”线程将直接结束任务并安全退出。该机制在保证计算结果完备性的同时有效规避了内存越界访问的风险。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
Ascend C核函数编程指南
核函数【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit正如在概述中所介绍的能够在AI处理器NPU上并行执行、并由主机端Host发起调用的函数被称为核函数Kernel。核函数被设计为由大量线程Threads同时并行运行。每个线程都拥有独立的寄存器和栈它们协同工作以完成海量的数据计算任务。核函数的定义Ascend C支持开发者自定义核函数来扩展C。定义核函数时通过引入特定的函数限定符编译器能够识别该函数需要在设备端Device编译并允许其从主机端被发起调用。核函数的定义示例如下// 核函数定义示例 __global__ void vec_add(float* x, float* y, float* z) { // 具体的计算逻辑 }定义核函数时需要遵循以下规则使用函数类型限定符__global__标识它是一个核函数。核函数必须具有void返回类型。核函数参数列表需遵循函数参数列表限制。核函数的调用算子程序中的函数可以分为三类host侧执行函数、核函数device侧执行、device侧执行函数除核函数之外。下图以Kernel直调算子开发方式为例描述三者的调用关系host侧执行函数可以调用其它host执行函数也就是通用C/C编程中的函数调用也可以通过...调用核函数。核函数可以调用除核函数之外的其它device侧执行函数。device侧执行函数除核函数之外使用类型限定符__aicore__标识可以调用同类的其它device侧执行函数。图1核函数、host侧执行函数、device侧执行函数调用关系Host侧通过核函数调用符...的语法形式调用核函数如下所示kernel_nameblocks_per_grid, threads_per_block, dyn_ubuf_size, stream(argument list)...内的参数为核函数的执行配置由4个参数决定详细用法请参考核函数配置blocks_per_gridint或dim3类型用于指定网格grid的维度与规模blocks_per_grid.x * blocks_per_grid.y * blocks_per_grid.z等于启动的线程块总数不得大于65535。threads_per_blockint或dim3类型用于指定每个线程块block的维度与规模threads_per_block.x * threads_per_block.y * threads_per_block.z等于每个线程块包含的线程数需要小于等于__launch_bounds__配置。dyn_ubuf_sizesize_t类型该参数指定除静态分配的内存外本次调用为每个线程块动态分配的共享内存字节数单位为bytes默认设置为0。具体用法请参考共享内存中的“动态申请”方式streamaclrtStream类型指针指定关联的流用于维护异步操作的执行顺序。Grid与Thread索引内置变量在核函数内Ascend C提供了内置变量来获取执行配置的参数以及线程或线程块的索引。常用的内置索引变量说明如下threadIdx获取当前线程在其所属线程块内的索引。threadIdx.xthreadIdx.ythreadIdx.z分别表示当前线程在3个维度的索引threadIdx.x的范围为[0, blockDim.x)threadIdx.y的范围为[0, blockDim.y)threadIdx.z的范围为[0, blockDim.z)。blockDim获取线程块中配置的三维层次结构即核函数启动时配置的dim3结构体实例值。blockDim.xblockDim.yblockDim.z分别表示线程块中三个维度的线程数。在核函数启动的执行配置中指定对应核函数配置中的threads_per_block参数。blockIdx: 获取当前线程块在其所属网格中的索引表示当前线程所在的线程块在整个网格中的位置坐标。gridDim表示整个计算任务在各个维度上分别由多少个线程块构成。在核函数启动的执行配置中指定对应核函数配置中的blocks_per_grid参数。各个维度上线程块关系需满足gridDim.x * gridDim.y * gridDim.z 65535。开发者可通过组合上述变量计算出当前线程在整个庞大计算网格空间中的全局唯一数据索引Global Index。这使得成千上万个执行相同代码的线程能够精准地定位并处理属于自己的特定数据分片。边界检查在实际一维并行处理开发中待处理的数据总长度往往无法被单一线程块的尺寸完美整除。因此通常需要分配多余的物理线程并在核函数内部引入边界检查机制以防止内存越界。一维数据索引的标准计算公式为int global_index blockIdx.x * blockDim.x threadIdx.x;以下代码片段展示了如何结合索引机制与边界检查实现动态长度的矢量相加// 1. Kernel implementation __global__ void vec_add(float* x, float* y, float* z, int vector_length) { // Calculate the global element index corresponding to the current thread int global_index blockIdx.x * blockDim.x threadIdx.x; // Boundary check: Constrain the computation range to prevent excess threads from accessing out-of-bounds memory if (global_index vector_length) { // Only threads with indices within the valid range will execute the specific addition operation z[global_index] x[global_index] y[global_index]; } } // 2. Host-side invocation example int main() { // Assume memory allocation and initialization on both the Host and Device sides are complete int vector_length 1000; // Empirical configuration: Specify that each block contains 256 threads int threads 256; // Calculate parameters: Calculate the required number of blocks based on the ceiling principle // Algorithm: (length threads - 1) / threads int blocks (vector_length threads - 1) / threads; // The result is 4 // Execute invocation: Launch 4 blocks, with 256 threads per block (a total of 1024 physical threads dispatched) vec_addblocks, threads, 0, stream(dev_x, dev_y, dev_z, vector_length); // ... Subsequent stream synchronization and resource release logic ... return 0; }在本示例中目标处理元素总数为1000个。按单块256线程的策略划分Host侧实际向硬件下发了4个线程块总计启动了$4 \times 256 1024$个物理线程。当程序运行至最后一个线程块即blockIdx.x为3时该块内末尾的24个线程从索引1000到1023推导出的global_index将落在[1000, 1023]区间。此时通过if (global_index vector_length)这道边界拦截这部分超出有效数据范围的“多余”线程将直接结束任务并安全退出。该机制在保证计算结果完备性的同时有效规避了内存越界访问的风险。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
