AscendC从入门到精通系列(一)初步感知AscendC

举报
zjun 发表于 2024/11/08 19:49:28 2024/11/08
【摘要】 Ascend C是CANN针对算子开发场景推出的编程语言,原生支持C和C++标准规范,兼具开发效率和运行性能。基于Ascend C编写的算子程序,通过编译器编译和运行时调度,运行在昇腾AI处理器上。使用Ascend C,开发者可以基于昇腾AI硬件,高效的实现自定义的创新算法。

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 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  // 注意修改为当前环境的路径
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库:

operator/HelloWorldSample/run.sh · Ascend/samples - 码云 - 开源中国 (gitee.com)
​gitee.com/ascend/samples/tree/master/operator/Hello

【版权声明】本文为华为云社区用户原创内容,转载时必须标注文章的来源(华为云社区)、文章链接、文章作者等基本信息, 否则作者和本社区有权追究责任。如果您发现本社区中有涉嫌抄袭的内容,欢迎发送邮件进行举报,并提供相关证据,一经查实,本社区将立刻删除涉嫌侵权内容,举报邮箱: cloudbbs@huaweicloud.com
  • 点赞
  • 收藏
  • 关注作者

评论(0

0/1000
抱歉,系统识别当前为高风险访问,暂不支持该操作

全部回复

上滑加载中

设置昵称

在此一键设置昵称,即可参与社区互动!

*长度不超过10个汉字或20个英文字符,设置后3个月内不可修改。

*长度不超过10个汉字或20个英文字符,设置后3个月内不可修改。