1 什么是 AscendC
Ascend C 是 CANN 针对算子开发场景推出的编程语言，原生支持 C 和 C++标准规范，兼具开发效率和运行性能。基于 Ascend C 编写的算子程序，通过编译器编译和运行时调度，运行在昇腾 AI 处理器上。使用 Ascend C，开发者可以基于昇腾 AI 硬件，高效的实现自定义的创新算法。
算子开发学习地图：
2 从 helloworld 出发感受 AscendC
2.1 使用 AscendC 写核函数
包含核函数的 Kernel 实现文件 hello_world.cpp 代码如下：核函数 hello_world 的核心逻辑为打印"Hello World"字符串。hello_world_do 封装了核函数的调用程序，通过<<<>>>内核调用符对核函数进行调用。
#include "kernel_operator.h"extern "C" __global__ __aicore__ void hello_world(){ AscendC::printf("Hello World!!!\n");}
void hello_world_do(uint32_t blockDim, void* stream){ hello_world<<<blockDim, nullptr, stream>>>();}
复制代码
2.2 通过 main.cpp 调用核函数
便携 main.cpp 进行调用。
#include "acl/acl.h"extern void hello_world_do(uint32_t coreDim, void* stream);
int32_t main(int argc, char const *argv[]){ // AscendCL初始化 aclInit(nullptr); // 运行管理资源申请 int32_t deviceId = 0; aclrtSetDevice(deviceId); aclrtStream stream = nullptr; aclrtCreateStream(&stream);
// 设置参与运算的核数为8 constexpr uint32_t blockDim = 8; // 用内核调用符<<<>>>调用核函数，hello_world_do中封装了<<<>>>调用 hello_world_do(blockDim, stream); aclrtSynchronizeStream(stream); // 资源释放和AscendCL去初始化 aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0;}
复制代码
2.3 添加 CMakeLists 文件
注意修改：SOC_VERSION，一般是: Ascend310P3,Ascend910B3 等，通过 npu-sim info 命令查询。
# Copyright (c) Huawei Technologies Co., Ltd. 2020. All rights reserved.
# CMake lowest version requirementcmake_minimum_required(VERSION 3.16.0)
# project informationproject(Ascend_C)set(SOC_VERSION "Ascend310P3" CACHE STRING "system on chip type")set(ASCEND_CANN_PACKAGE_PATH "/usr/local/Ascend/ascend-toolkit/latest" CACHE PATH "ASCEND CANN package installation directory")set(RUN_MODE "npu" CACHE STRING "run mode: npu")set(CMAKE_BUILD_TYPE "Debug" CACHE STRING "Build type Release/Debug (default Debug)" FORCE)set(CMAKE_INSTALL_PREFIX "${CMAKE_CURRENT_LIST_DIR}/out" CACHE STRING "path for install()" FORCE)
if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake)elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake)elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/ascendc_devkit/tikcpp/samples/cmake) set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/ascendc_devkit/tikcpp/samples/cmake)else() message(FATAL_ERROR "ascendc_kernel_cmake does not exist, please check whether the cann package is installed.")endif()
include(${ASCENDC_CMAKE_DIR}/ascendc.cmake)
# ascendc_library use to add kernel file to generate ascendc libraryascendc_library(kernels STATIC hello_world.cpp)
add_executable(main main.cpp)
target_link_libraries(main PRIVATE kernels)
复制代码
2.4 编译运行
注意修改：SOC_VERSION，一般是: Ascend310P3, Ascend910B3 等，通过 npu-sim info 命令查询。
source /usr/local/Ascend/ascend-toolkit/latest/bin/setenv.bash // 注意修改为当前环境的路径rm -rf buildmkdir -p buildcmake -B build \ -DSOC_VERSION=${SOC_VERSION} \ -DASCEND_CANN_PACKAGE_PATH=/usr/local/Ascend/ascend-toolkit/latestcmake --build build -jcmake --install build
./build/main
复制代码
注意：编译有报错，确认下 CANN 版本和 Sample 的版本是不是匹配的。
比如 CANN 是 8.0RC2，那么 Sample 库的版本最好也切到 8.0RC2。
如果 CANN 是 8.0RC3 这种，Sample 中没有 8.0RC2，那就直接用 master。
详细可以 Ascend 官方 gitee 库：
operator/HelloWorldSample/run.sh · Ascend/samples - 码云 - 开源中国 (gitee.com)gitee.com/ascend/samples/tree/master/operator/Hello
评论