CANN学习资源开源仓的算子开发二Tiling和CMake

自定义结构体AddCustomTilingData。在启动核函数之前，主机侧（Host）需要明确告知设备侧（Device）如何对总数据进行划分，以实现多核并行与单核内流水并行。其通过一个自定义的结构体 AddCustomTilingData 来传递，作为核函数的参数，它将主机侧决定的全局数据视图（totalLength）和单核并行度（tileNum）传递给每个核上执行的核函数实例。

• totalLength 与系统启动的核数（GetBlockNum()，例如8）共同决定了第一级多核并行的粒度 blockLength（即totalLength / GetBlockNum()）。

• tileNum 决定了第二级单核内流水并行的粒度，指导核函数内部如何为TPipe和TQue分配内存，以及循环计算多少次。

这里面存在两级并行：

• 第一级：多核并行，数据在不同AI Core间切分，并行计算
• 第二级：单核内流水并行，数据在单个核内进一步切块，支持流水线并行执行

我们使用CMake作为构建系统，通过特定的配置来调用Ascend C编译工具链。下面提供了一个最小的完整编译框架，实际项目中可能还需要添加头文件路径、链接库、以及其他针对不同构建类型（Debug/Release）的优化选项。

cmake_minimum_required(VERSION 3.16)
# CMake中用于查找和配置Ascend C编译工具链的命令
find_package(ASC REQUIRED)
# 指定项目支持的语言包括ASC和CXX，ASC表示支持使用毕昇编译器对Ascend C编程语言进行编译
project(kernel_samples LANGUAGES ASC CXX)

add_executable(demo
    add_custom.asc
)
# 通过编译选项设置NPU架构
target_compile_options(demo PRIVATE
    $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>
)

• 工具链依赖： find_package(ASC)自动配置了将 .asc 文件编译为NPU可执行代码所需的所有规则。

• 双语言支持：通过 project(… LANGUAGES ASC CXX) 声明，CMake能分别用正确的编译器处理Host侧的C++代码和Device侧的Ascend C核函数代码。

• 架构指定：–npu-arch 必须与运行环境的NPU型号匹配，以确保生成的二进制指令能够正确执行。

在运行cmake之前，设置了 ASC_DIR 环境变量，于是find_package(ASC REQUIRED) 会在此环境变量即$ASCEND_HOME_PATH/aarch64-linux/tikcpp/ascendc_kernel_cmake/中查找 ASCConfig.cmake（就是它了 CANN8.5版本才有） 或 asc-config.cmake 配置文件。ASCConfig.cmake为 CMake 添加了对 ASC 语言的支持，它通过 ASC_CMake 目录中的文件实现了 ASC 语言的完整 CMake 集成

ASC_CMake/
├── CMakeDetermineASCCompiler.cmake    # 检测 ASC 编译器
├── CMakeTestASCCompiler.cmake        # 测试编译器是否可用
├── CMakeASCCompiler.cmake.in         # 编译器配置模板
├── CMakeASCInformation.cmake         # 编译器信息设置
└── FindASC.cmake                      # 查找 ASC 工具链

这些文件是 CMake 支持一门新编程语言的标准结构，类似于 CMake 内置的 CMakeCCompiler.cmake、CMakeCXXCompiler.cmake。

至此我们有了两个核心文件：

1. add_custom.asc

   核心源码文件，采用.asc扩展名以支持Ascend C语法。它实际上是一个混合了设备侧（Device）核函数与主机侧（Host）调用代码的源文件，具体包含：

   • 核函数实现：__global__ __aicore__ void add_custom(…) 函数及其内部调用的算子类 KernelAdd，这部分的代码将在AI Core上执行。

   • Host侧应用程序：封装调用逻辑的 kernel_add 函数、验证函数 VerifyResult 以及程序入口 main 函数，这部分的代码在CPU上执行。

2. CMakeLists.txt

   项目的构建配置文件，它定义了如何将 add_custom.asc 中的混合代码正确地分别编译（为设备代码生成NPU指令，为主机代码生成CPU指令）并链接成一个完整的可执行程序（如 demo）。

add_custom.asc 和 CMakeLists.txt 共同构成了一个完整、可交付的算子工程。编译运行如下：

atomgit@676982cc8b2dd94c4aac433a:~/cann-learning-hub/tutorials/ascendc_operator_development/02_AscendC_basic/src/build (master)$ cmake ..
-- System processer: aarch64
-- CMAKE_ASC_COMPILER: /usr/local/Ascend/cann-8.5.0/bin/bisheng
-- ASCEND_CANN_PACKAGE_LINUX_PATH: /usr/local/Ascend/cann-8.5.0/aarch64-linux
-- CMAKE_ASC_LLD_LINKER: /usr/local/Ascend/cann-8.5.0/aarch64-linux/ccec_compiler/bin/ld.lld
-- The CXX compiler identification is GNU 11.4.0
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Check for working CXX compiler: /usr/bin/c++ - skipped
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Configuring done (0.4s)
-- Generating done (0.0s)
-- Build files have been written to: /opt/atomgit/cann-learning-hub/tutorials/ascendc_operator_development/02_AscendC_basic/src/build

atomgit@676982cc8b2dd94c4aac433a:~/cann-learning-hub/tutorials/ascendc_operator_development/02_AscendC_basic/src/build (master)$ make
[ 50%] Building ASC object CMakeFiles/demo.dir/add_custom.asc.o
[100%] Linking ASC executable demo
[100%] Built target demo

atomgit@676982cc8b2dd94c4aac433a:~/cann-learning-hub/tutorials/ascendc_operator_development/02_AscendC_basic/src/build (master)$ ls
CMakeCache.txt  CMakeFiles  Makefile  autogen  cmake_install.cmake  demo

atomgit@676982cc8b2dd94c4aac433a:~/cann-learning-hub/tutorials/ascendc_operator_development/02_AscendC_basic/src/build (master)$ ./demo
Output: 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 ...
Golden: 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 3.5 ...
[Success] Case accuracy is verification passed.