下面是一个简单的Ascend C的"Hello World"样例，展示了一个Ascend C核函数（设备侧实现的入口函数）的基本写法，及其如何被调用的流程

HelloWorldSample样例参考链接

• 包含核函数的Kernel实现文件hello_world.cpp代码如下：核函数hello_world的核心逻辑为打印"Hello World"字符串。hello_world_do封装了核函数的调用程序，通过<<<>>>内核调用符对核函数进行调用。

#include "kernel_operator.h"
using namespace AscendC;
extern "C" __global__ __aicore__ void hello_world() {
    printf("Hello World!!!\n");
}

void hello_world_do(uint32_t blockDim, void* stream) {
    hello_world<<<blockDim, nullptr, stream>>>();
}

• 调用核函数的应用程序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);
    // 运行管理资源申请
    aclrtContext context;
    int32_t deviceId = 0;
    aclrtSetDevice(deviceId);
    aclrtCreateContext(&context,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);
    aclrtDestroyContext(context);
    aclrtResetDevice(deviceId);
    aclFinalize();
    return 0;
}

• 代码工程目录结构如下

|-- CMakeLists.txt                                        // CMake编译配置文件
|-- hello_world.cpp                                       // Kernel实现
|-- main.cpp                                              // 调用核函数的主程序
|-- run.sh                                                // 一键式编译运行脚本          

• 编译和运行：

bash run.sh -v Ascend310P1

本样例共调度八个核，分别打印了每个核的核号和"Hello World"信息。

root@k8s-node1:/tmp/HelloWorldSample# bash run.sh -v Ascend310P1
-- The C compiler identification is GNU 9.4.0
-- The CXX compiler identification is GNU 9.4.0
-- Check for working C compiler: /usr/bin/cc
-- Check for working C compiler: /usr/bin/cc -- works
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Detecting C compile features
-- Detecting C compile features - done
-- Check for working CXX compiler: /usr/bin/c++
-- Check for working CXX compiler: /usr/bin/c++ -- works
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Configuring done
-- Generating done
-- Build files have been written to: /tmp/HelloWorldSample/build
Scanning dependencies of target kernels_preprocess
[  3%] Creating directories for 'kernels_preprocess'
[  7%] No download step for 'kernels_preprocess'
[ 14%] No update step for 'kernels_preprocess'
[ 14%] No patch step for 'kernels_preprocess'
[ 17%] Performing configure step for 'kernels_preprocess'
-- The C compiler identification is GNU 9.4.0
-- The CXX compiler identification is GNU 9.4.0
-- Check for working C compiler: /usr/bin/cc
-- Check for working C compiler: /usr/bin/cc -- works
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Detecting C compile features
-- Detecting C compile features - done
-- Check for working CXX compiler: /usr/bin/c++
-- Check for working CXX compiler: /usr/bin/c++ -- works
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Configuring done
-- Generating done
-- Build files have been written to: /tmp/HelloWorldSample/build/kernels_preprocess-prefix/src/kernels_preprocess-build
[ 21%] Performing build step for 'kernels_preprocess'
Scanning dependencies of target preprocess_obj
[100%] Building CXX object CMakeFiles/preprocess_obj.dir/tmp/HelloWorldSample/hello_world.cpp.o
[100%] Built target preprocess_obj
Scanning dependencies of target _host_cpp
[100%] Built target _host_cpp
[ 25%] No install step for 'kernels_preprocess'
[ 28%] Completed 'kernels_preprocess'
[ 28%] Built target kernels_preprocess
Scanning dependencies of target kernels_device
Scanning dependencies of target kernels_host
[ 32%] Creating directories for 'kernels_device'
[ 35%] Creating directories for 'kernels_host'
[ 39%] No download step for 'kernels_device'
[ 42%] No download step for 'kernels_host'
[ 46%] No patch step for 'kernels_device'
[ 50%] No update step for 'kernels_device'
[ 57%] No patch step for 'kernels_host'
[ 57%] No update step for 'kernels_host'
[ 60%] Performing configure step for 'kernels_device'
[ 64%] Performing configure step for 'kernels_host'
-- The C compiler identification is GNU 9.4.0
-- The C compiler identification is GNU 9.4.0
-- The CXX compiler identification is GNU 9.4.0
-- The CXX compiler identification is GNU 9.4.0
-- Check for working C compiler: /usr/bin/cc
-- Check for working C compiler: /usr/bin/cc
-- Check for working C compiler: /usr/bin/cc -- works
-- Detecting C compiler ABI info
-- Check for working C compiler: /usr/bin/cc -- works
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Detecting C compile features
-- Detecting C compiler ABI info - done
-- Detecting C compile features - done
-- Check for working CXX compiler: /usr/bin/c++
-- Detecting C compile features
-- Detecting C compile features - done
-- Check for working CXX compiler: /usr/bin/c++
-- Check for working CXX compiler: /usr/bin/c++ -- works
-- Detecting CXX compiler ABI info
-- Check for working CXX compiler: /usr/bin/c++ -- works
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Configuring done
-- Generating done
CMake Warning:
  Manually-specified variables were not used by the project:

    ASCEND_PYTHON_EXECUTABLE


-- Build files have been written to: /tmp/HelloWorldSample/build/kernels_host-prefix/src/kernels_host-build
-- Detecting CXX compiler ABI info - done
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Configuring done
-- Generating done
-- Build files have been written to: /tmp/HelloWorldSample/build/kernels_device-prefix/src/kernels_device-build
[ 67%] Performing build step for 'kernels_host'
[ 71%] Performing build step for 'kernels_device'
Scanning dependencies of target host_bisheng_obj
Scanning dependencies of target device_obj
[100%] Building CXX object CMakeFiles/host_bisheng_obj.dir/tmp/HelloWorldSample/hello_world.cpp.o
[100%] Building CXX object CMakeFiles/device_obj.dir/tmp/HelloWorldSample/build/auto_gen/kernels/auto_gen_hello_world.cpp.o
[100%] Built target host_bisheng_obj
[ 75%] Performing install step for 'kernels_host'
[100%] Built target host_bisheng_obj
Install the project...
-- Install configuration: "Debug"
-- Installing: /tmp/HelloWorldSample/build/kernels_host_dir/./objects-Debug/host_bisheng_obj/tmp/HelloWorldSample/hello_world.cpp.o
[ 78%] Completed 'kernels_host'
[ 78%] Built target kernels_host
In file included from /tmp/HelloWorldSample/build/auto_gen/kernels/auto_gen_hello_world.cpp:7:
In file included from /tmp/HelloWorldSample/hello_world.cpp:10:
In file included from /usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/kernel_operator.h:27:
In file included from /usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/interface/kernel_operator_intf.h:25:
In file included from /usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/interface/kernel_operator_dump_tensor_intf.h:28:
/usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/impl/dav_m200/kernel_operator_dump_tensor_impl.h:116:14: warning: 'SetAddr' is deprecated: NOTICE: SetAddr has been deprecated and will be removed in the next version. Please do not use it! [-Wdeprecated-declarations]
    tmpLocal.SetAddr(tbuf_tmpLocal);
             ^
/usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/impl/dav_m200/kernel_operator_dump_tensor_impl.h:159:5: note: in instantiation of function template specialization 'AscendC::InitTmpTensor<unsigned char>' requested here
    InitTmpTensor(tmpLocal, (uint8_t)QuePosition::VECIN);
    ^
/usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/impl/kernel_tensor_base.h:117:7: note: 'SetAddr' has been explicitly marked deprecated here
    [[deprecated("NOTICE: SetAddr has been deprecated and will be removed in the next version. "
      ^
In file included from /tmp/HelloWorldSample/build/auto_gen/kernels/auto_gen_hello_world.cpp:7:
In file included from /tmp/HelloWorldSample/hello_world.cpp:10:
In file included from /usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/kernel_operator.h:27:
In file included from /usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/interface/kernel_operator_intf.h:25:
In file included from /usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/interface/kernel_operator_dump_tensor_intf.h:28:
In file included from /usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/impl/dav_m200/kernel_operator_dump_tensor_impl.h:29:
/usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/impl/kernel_pop_stack_buffer.h:57:14: warning: 'SetAddr' is deprecated: NOTICE: SetAddr has been deprecated and will be removed in the next version. Please do not use it! [-Wdeprecated-declarations]
    popLocal.SetAddr(addr);
             ^
/usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/lib/filter/../../impl/filter/dropout/dropout_membase_impl.h:43:16: note: in instantiation of function template specialization 'AscendC::PopStackBuffer<short, AscendC::TPosition::VECCALC>' requested here
    bool ans = PopStackBuffer<int16_t, TPosition::LCM>(stackBuffer);
               ^
/usr/local/Ascend/ascend-toolkit/latest/tools/tikcpp/tikcfw/impl/kernel_tensor_base.h:117:7: note: 'SetAddr' has been explicitly marked deprecated here
    [[deprecated("NOTICE: SetAddr has been deprecated and will be removed in the next version. "
      ^
2 warnings generated.
[100%] Built target device_obj
Scanning dependencies of target merge_device_obj
/usr/local/Ascend/ascend-toolkit/latest/tools/ccec_compiler/bin/ld.lld  -m aicorelinux  -n -Ttext=0 /tmp/HelloWorldSample/build/kernels_device-prefix/src/kernels_device-build/CMakeFiles/device_obj.dir/tmp/HelloWorldSample/build/auto_gen/kernels/auto_gen_hello_world.cpp.o -static -o /tmp/HelloWorldSample/build/kernels_merge_obj_dir/device.o
ld.lld: warning: -n (--nmagic) and -N (--omagic) will be ignored
[100%] Built target merge_device_obj
[ 82%] No install step for 'kernels_device'
[ 85%] Completed 'kernels_device'
[ 85%] Built target kernels_device
Scanning dependencies of target kernels_merge_obj
[ 85%] Built target kernels_merge_obj
Scanning dependencies of target kernels_host_stub_obj
[ 89%] Building CXX object CMakeFiles/kernels_host_stub_obj.dir/auto_gen/kernels/host_stub.cpp.o
[ 89%] Built target kernels_host_stub_obj
Scanning dependencies of target kernels
[ 92%] Linking CXX static library lib/libkernels.a
/usr/local/Ascend/ascend-toolkit/latest/bin/ascendc_pack_kernel /tmp/HelloWorldSample/build/CMakeFiles/kernels_host_stub_obj.dir/auto_gen/kernels/host_stub.cpp.o /tmp/HelloWorldSample/build/kernels_merge_obj_dir/device.o 0 /tmp/HelloWorldSample/build/CMakeFiles/kernels_host_stub_obj.dir/auto_gen/kernels/host_stub.cpp.o
recompile: /usr/bin/ar qc lib/libkernels.a  CMakeFiles/kernels_host_stub_obj.dir/auto_gen/kernels/host_stub.cpp.o /tmp/HelloWorldSample/build/kernels_host_dir/objects-Debug/host_bisheng_obj/tmp/HelloWorldSample/hello_world.cpp.o  elf_tool.c.o ascendc_runtime.cpp.o
recompile: /usr/bin/ranlib lib/libkernels.a
[ 92%] Built target kernels
Scanning dependencies of target main
[ 96%] Building CXX object CMakeFiles/main.dir/main.cpp.o
[100%] Linking CXX executable main
[100%] Built target main
-- Install configuration: "Debug"
-- Installing: /tmp/HelloWorldSample/out/lib/libkernels.a
-- Installing: /tmp/HelloWorldSample/out/include
-- Installing: /tmp/HelloWorldSample/out/include/kernels
-- Installing: /tmp/HelloWorldSample/out/include/kernels/aclrtlaunch_triple_chevrons_func.h
-- Installing: /tmp/HelloWorldSample/out/include/kernels/aclrtlaunch_hello_world.h
opType=hello_world, DumpHead: AIC-0, block_id=0, block dim=8, block_remain_len=1048496, block_initial_space=1048576, rsv=0, magic=5aa5bccd
Hello World!!!
opType=hello_world, DumpHead: AIC-1, block_id=1, block dim=8, block_remain_len=1048496, block_initial_space=1048576, rsv=0, magic=5aa5bccd
Hello World!!!
opType=hello_world, DumpHead: AIC-2, block_id=2, block dim=8, block_remain_len=1048496, block_initial_space=1048576, rsv=0, magic=5aa5bccd
Hello World!!!
opType=hello_world, DumpHead: AIC-3, block_id=3, block dim=8, block_remain_len=1048496, block_initial_space=1048576, rsv=0, magic=5aa5bccd
Hello World!!!
opType=hello_world, DumpHead: AIC-4, block_id=4, block dim=8, block_remain_len=1048496, block_initial_space=1048576, rsv=0, magic=5aa5bccd
Hello World!!!
opType=hello_world, DumpHead: AIC-5, block_id=5, block dim=8, block_remain_len=1048496, block_initial_space=1048576, rsv=0, magic=5aa5bccd
Hello World!!!
opType=hello_world, DumpHead: AIC-6, block_id=6, block dim=8, block_remain_len=1048496, block_initial_space=1048576, rsv=0, magic=5aa5bccd
Hello World!!!
opType=hello_world, DumpHead: AIC-7, block_id=7, block dim=8, block_remain_len=1048496, block_initial_space=1048576, rsv=0, magic=5aa5bccd
Hello World!!!