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

简介: Ascend C是CANN推出的一种编程语言,专为算子开发设计,支持C/C++标准,旨在提高开发效率与运行性能。通过Ascend C,开发者能在昇腾AI处理器上高效实现自定义算法。本文档介绍了如何使用Ascend C编写和运行一个简单的“Hello World”程序,包括核函数的编写、主程序调用及CMake配置,展示了Ascend C的基本使用流程。

1 什么是AscendC

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

image.png

算子开发学习地图:
image.png

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

相关文章
|
机器学习/深度学习 存储 并行计算
一篇就够:高性能推理引擎理论与实践 (TensorRT)
本文分享了关于 NVIDIA 推出的高性能的深度学习推理引擎 TensorRT 的背后理论知识和实践操作指南。
14821 9
一篇就够:高性能推理引擎理论与实践 (TensorRT)
|
JavaScript 搜索推荐 程序员
Vuepress + gitee五分钟免费搭建个人博客(保姆级教程)
前言 作为一个程序员,没有折腾过个人博客是不算完整的。技术文章的输出是我们程序员能力的一种体现,也是一种非常好的个人总结。 市面上有很多搭建个人博客的工具或框架,包括hexo、wordpress等等。不可否认,市面上有些博客系统做得很好,博客主题也很丰富,但是往往存在一个问题:比较重。 作为一个Vue程序员,我就比较推荐使用vuepress搭建个人博客,因为它毕竟是Vue出品的,大家熟知的vue官方就是利用vuepress搭建的。 本篇文章就从零开始教大家搭建一个免费的博客,零基础小白也可以学习哦!
3911 0
Vuepress + gitee五分钟免费搭建个人博客(保姆级教程)
|
测试技术 开发者 异构计算
AscendC从入门到精通系列(二)基于Kernel直调开发AscendC算子
本文介绍了AscendC算子的开发流程,包括核函数开发、算子类定义及其实现、核函数的CPU和NPU侧运行验证。通过具体示例`add_custom.cpp`,详细展示了如何使用Ascend C完成算子核函数的定义、初始化、数据搬运和计算过程,并提供了完整的CPU和NPU侧调用程序代码,帮助开发者理解和实践AscendC算子的开发。
|
存储 人工智能 JSON
AscendC从入门到精通系列(三)基于自定义算子工程开发AscendC算子
本文介绍了基于Ascend C的自定义算子开发流程,涵盖从工程创建、代码编写、编译部署到运行验证的全过程。以动态shape的AddCustom算子为例,详细描述了如何利用CANN提供的工具msOpGen生成开发工程,实现算子核函数与host侧代码,以及如何编译、部署和测试自定义算子。
|
11月前
|
人工智能 物联网 测试技术
FireRedASR:精准识别普通话、方言和歌曲歌词!小红书开源工业级自动语音识别模型
小红书开源的工业级自动语音识别模型,支持普通话、中文方言和英语,采用 Encoder-Adapter-LLM 和 AED 架构,实现 SOTA 性能。
3464 17
FireRedASR:精准识别普通话、方言和歌曲歌词!小红书开源工业级自动语音识别模型
|
人工智能 自动驾驶 机器人
吴泳铭:AI最大的想象力不在手机屏幕,而是改变物理世界
过去22个月,AI发展速度超过任何历史时期,但我们依然还处于AGI变革的早期。生成式AI最大的想象力,绝不是在手机屏幕上做一两个新的超级app,而是接管数字世界,改变物理世界。
24701 73
吴泳铭:AI最大的想象力不在手机屏幕,而是改变物理世界
|
PyTorch API 算法框架/工具
AscendC从入门到精通系列(四)使用Pybind调用AscendC算子
本文介绍了如何通过Pybind11在PyTorch框架中调用自定义的Ascend C算子。首先,通过编写算子的C++实现和pybind11封装,将算子功能暴露给Python。接着,构建Python调用脚本,利用torch接口生成数据并调用封装好的算子模块。最后,通过CMake配置文件编译整个项目,实现从算子开发到测试的完整流程。
|
人工智能 并行计算 编译器
【AI系统】SIMD & SIMT 与 CUDA 关系
本文深入解析了AI芯片中SIMD和SIMT的计算本质,基于NVIDIA CUDA实现的对比,探讨了不同并行编程模型,包括串行(SISD)、数据并行(SIMD)和多线程(MIMD/SPMD)。文章详细介绍了各模型的特点及应用场景,特别强调了英伟达GPU中的SIMT机制如何通过SPMD编程模型实现高效并行计算,以及SIMD、SIMT、SPMD之间的关系和区别。
731 13
|
11月前
|
存储 人工智能 搜索推荐
Memobase:开源AI长期记忆系统,让AI真正记住每个用户的秘密武器
Memobase 是一个开源的长期记忆系统,专为生成式 AI 应用设计,通过用户画像和时间感知记忆功能,帮助 AI 记住、理解并适应用户需求。
2046 0