1 什么是AscendC
Ascend C是CANN針對算子開發(fā)場景推出的編程語言,原生支持C和C++標準規(guī)范,兼具開發(fā)效率和運行性能?;贏scend C編寫的算子程序,通過編譯器編譯和運行時調度,運行在昇騰AI處理器上。使用Ascend C,開發(fā)者可以基于昇騰AI硬件,高效的實現(xiàn)自定義的創(chuàng)新算法。

image.png
算子開發(fā)學習地圖:

image.png
2 從helloworld出發(fā)感受AscendC
2.1 使用AscendC寫核函數(shù)
包含核函數(shù)的Kernel實現(xiàn)文件hello_world.cpp代碼如下:核函數(shù)hello_world的核心邏輯為打印"Hello World"字符串。hello_world_do封裝了核函數(shù)的調用程序,通過<<<>>>內核調用符對核函數(shù)進行調用。
#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調用核函數(shù)
便攜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);
// 設置參與運算的核數(shù)為8
constexpr uint32_t blockDim = 8;
// 用內核調用符<<<>>>調用核函數(shù),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 requirement
cmake_minimum_required(VERSION 3.16.0)
# project information
project(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 library
ascendc_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 // 注意修改為當前環(huán)境的路徑
rm -rf build
mkdir -p build
cmake -B build \
-DSOC_VERSION=${SOC_VERSION} \
-DASCEND_CANN_PACKAGE_PATH=/usr/local/Ascend/ascend-toolkit/latest
cmake --build build -j
cmake --install build
./build/main
注意:編譯有報錯,確認下CANN版本和Sample的版本是不是匹配的。
比如CANN是8.0RC2,那么Sample庫的版本最好也切到8.0RC2。
如果CANN是8.0RC3這種,Sample中沒有8.0RC2,那就直接用master。
詳細可以Ascend官方gitee庫: