已合并
add simd_vf printf/dump example #2558
ChenZhoujie创建于 5月25日
add simd_vf printf/dump example #2558
已合并
共 27 个文件变更+936-248
| @@ -1,107 +1,9 @@ | |||
| 1 | -# Printf接口功能说明 | 1 | +# Printf样例介绍 |
| 2 | - | ||
| 3 | ## 概述 | 2 | ## 概述 |
| 3 | +本样例展示了Ascend C printf接口和simd_vf侧printf接口的基本使用方法。 | ||
A | |||
| 4 | 4 | ||
| 5 | -本样例介绍printf接口使用方法，通过该接口打印核函数相关信息。 | 5 | +## 样例列表 |
 [代码检视建议] 顶层 README 建议补充 simple_printf 与 simd_vf_printf 的差异说明和选择建议。   | |||
| 6 | - | 6 | +| 目录名称 | 功能描述 | 支持的产品 | |
| 7 | -## 支持的产品 | 7 | +| ------------------------------------------------------------ | ---------------------------------------------------- | --- | |
| 8 | - | 8 | +| [simple_printf](./simple_printf) | 使用静态Tensor编程模式实现矩阵乘法，展示AscendC::printf接口的基本使用方法 | Ascend 950PR/Ascend 950DT<br>Atlas A3训练系列产品/Atlas A3推理系列产品<br>Atlas A2训练系列产品/Atlas A2推理系列产品 | |
| 9 | -- Ascend 950PR/Ascend 950DT | 9 | +| [simd_vf_printf](./simd_vf_printf) | 使用vector编程模式，展示simd_vf侧printf接口的基本使用方法 | Ascend 950PR/Ascend 950DT | |
| 10 | -- Atlas A3 训练系列产品/Atlas A3 推理系列产品 | ||
| 11 | -- Atlas A2 训练系列产品/Atlas A2 推理系列产品 | ||
| 12 | - | ||
| 13 | -## 目录结构介绍 | ||
| 14 | - | ||
| 15 | -``` | ||
| 16 | -├── 00_printf | ||
| 17 | -│ ├── scripts | ||
| 18 | -│ │ ├── gen_data.py // 输入数据和真值数据生成脚本 | ||
| 19 | -│ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 | ||
| 20 | -│ ├── CMakeLists.txt // 编译工程文件 | ||
| 21 | -│ ├── data_utils.h // 数据读入写出函数 | ||
| 22 | -│ └── printf.asc // Ascend C样例实现 & 调用样例 | ||
| 23 | -``` | ||
| 24 | - | ||
| 25 | -## 样例描述 | ||
| 26 | - | ||
| 27 | -- 样例功能： | ||
| 28 | - | ||
| 29 | - 使用静态Tensor编程模式实现矩阵乘法，展示printf接口的基本使用方法。 | ||
| 30 | - | ||
| 31 | - 矩阵乘法的计算公式为： | ||
| 32 | - | ||
| 33 | - ``` | ||
| 34 | - C = A * B | ||
| 35 | - ``` | ||
| 36 | - | ||
| 37 | -- 样例规格： | ||
| 38 | - | ||
| 39 | - 样例参数为：M = 256, N = 256, K = 64，shape信息如下表所示： | ||
| 40 | - <table> | ||
| 41 | - <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center">Matmul</td></tr> | ||
| 42 | - </tr> | ||
| 43 | - <tr><td rowspan="3" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr> | ||
| 44 | - <tr><td align="center">a</td><td align="center">[M, K]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 45 | - <tr><td align="center">b</td><td align="center">[K, N]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 46 | - </tr> | ||
| 47 | - </tr> | ||
| 48 | - <tr><td rowspan="1" align="center">样例输出</td><td align="center">c</td><td align="center">[M, N]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 49 | - </tr> | ||
| 50 | - <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">mmad_custom</td></tr> | ||
| 51 | - </table> | ||
| 52 | - | ||
| 53 | -- printf接口支持格式： | ||
| 54 | - | ||
| 55 | - | 格式符 | 说明 | 示例 | | ||
| 56 | - |--------|------|------| | ||
| 57 | - | `%p` | 指针打印 | `AscendC::printf("pointer %p\n", ptr);` | | ||
| 58 | - | `%d` | 整型/bool打印 | `AscendC::printf("value is %d\n", 10);` | | ||
| 59 | - | `%u` | 无符号整型打印 | `AscendC::printf("idx is %u\n", idx);` | | ||
| 60 | - | `%x` | 十六进制打印 | `AscendC::printf("value is %x\n", 255);` | | ||
| 61 | - | `%f` | 浮点型打印(half/float) | `AscendC::printf("half %f\n", val);` | | ||
| 62 | - | `%s` | 字符串打印 | `AscendC::printf("name %s\n", "test");` | | ||
| 63 | - | ||
| 64 | -- 调用实现 | ||
| 65 | - | ||
| 66 | - 使用内核调用符<<<>>>调用核函数。 | ||
| 67 | - | ||
| 68 | -## 编译运行 | ||
| 69 | - | ||
| 70 | -- 配置环境变量 | ||
| 71 | -在本样例根目录下执行如下步骤，编译并执行样例。 | ||
| 72 | - 请根据当前环境上CANN开发套件包的[安装方式](../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。 | ||
| 73 | - - 默认路径，root用户安装CANN软件包 | ||
| 74 | - ```bash | ||
| 75 | - source /usr/local/Ascend/cann/set_env.sh | ||
| 76 | - ``` | ||
| 77 | - | ||
| 78 | - - 默认路径，非root用户安装CANN软件包 | ||
| 79 | - ```bash | ||
| 80 | - source $HOME/Ascend/cann/set_env.sh | ||
| 81 | - ``` | ||
| 82 | - | ||
| 83 | - - 指定路径install_path，安装CANN软件包 | ||
| 84 | - ```bash | ||
| 85 | - source ${install_path}/cann/set_env.sh | ||
| 86 | - ``` | ||
| 87 | - | ||
| 88 | -- 样例执行 | ||
| 89 | - ```bash | ||
| 90 | - mkdir -p build && cd build; # 创建并进入build目录 | ||
| 91 | - cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程 | ||
| 92 | - python3 ../scripts/gen_data.py # 生成测试输入数据 | ||
| 93 | - ./demo # 执行编译生成的可执行程序，执行样例 | ||
| 94 | - python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确 | ||
| 95 | - ``` | ||
| 96 | - | ||
| 97 | -- 编译选项说明 | ||
| 98 | - | ||
| 99 | -| 选项 | 可选值 | 说明 | | ||
| 100 | -|------|--------|------| | ||
| 101 | -| `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品和Atlas A3 训练系列产品/Atlas A3 推理系列产品，dav-3510 对应 Ascend 950PR/Ascend 950DT | | ||
| 102 | - | ||
| 103 | -- 执行结果 | ||
| 104 | - 最终执行结果如下，说明精度对比成功。 | ||
| 105 | - ```bash | ||
| 106 | - test pass! | ||
| 107 | - ``` | ||
| @@ -0,0 +1,12 @@ | |||
| 1 | +# Printf Samples | ||
| 2 | + | ||
| 3 | +## Overview | ||
| 4 | + | ||
| 5 | +This sample shows how to use the Ascend C printf interface and the simd_vf printf interface. | ||
[代码检视建议] 英文 README 建议同步中文文档中的环境准备和运行步骤，避免中英文说明信息不一致。 | |||
| 6 | + | ||
| 7 | +## Sample List | ||
| 8 | + | ||
| 9 | +| Directory | Description | Supported Products | | ||
| 10 | +| --- | --- | --- | | ||
[代码检视建议] Sample List 中建议补充每个样例对应的主要 API 名称，便于用户从目录直接判断覆盖范围。 | |||
| 11 | +| [simple_printf](./simple_printf) | Uses static Tensor programming mode to implement matrix multiplication, showing basic usage of the AscendC::printf interface | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series/Atlas A3 Inference Series<br>Atlas A2 Training Series/Atlas A2 Inference Series | | ||
| 12 | +| [simd_vf_printf](./simd_vf_printf) | Uses vector programming mode, showing basic usage of the simd_vf printf interface | Ascend 950PR/Ascend 950DT | | ||
| @@ -0,0 +1,26 @@ | |||
| 1 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 2 | +# Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | +# CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | +# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 10 | + | ||
| 11 | + | ||
| 12 | +cmake_minimum_required(VERSION 3.16) | ||
| 13 | + | ||
| 14 | +set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510") | ||
| 15 | + | ||
| 16 | +find_package(ASC REQUIRED) | ||
| 17 | + | ||
| 18 | +project(kernel_samples LANGUAGES ASC CXX) | ||
| 19 | + | ||
| 20 | +add_executable(demo | ||
| 21 | + simd_vf_printf.asc | ||
| 22 | +) | ||
| 23 | + | ||
| 24 | +target_compile_options(demo PRIVATE | ||
| 25 | + $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}> | ||
| 26 | +) | ||
| @@ -0,0 +1,88 @@ | |||
| 1 | +# SimdVF Printf接口功能说明 | ||
| 2 | + | ||
| 3 | +## 概述 | ||
| 4 | + | ||
| 5 | +本样例介绍simd_vf侧printf接口使用方法。核函数使用固定输入（不依赖外部数据），通过`printf`接口在vector函数中打印调试信息及计算结果。 | ||
| 6 | + | ||
| 7 | +## 支持的产品 | ||
| 8 | + | ||
| 9 | +- Ascend 950PR/Ascend 950DT | ||
| 10 | + | ||
[代码检视建议] 支持产品章节只列出 Ascend 950PR/950DT，建议说明其他产品不支持的原因或限制，避免用户误用。 | |||
| 11 | +## 目录结构介绍 | ||
| 12 | + | ||
| 13 | +``` | ||
| 14 | +├── simd_vf_printf | ||
| 15 | +│ ├── CMakeLists.txt // 编译工程文件 | ||
| 16 | +│ └── simd_vf_printf.asc // Ascend C样例实现&调用样例 | ||
[代码检视建议] 目录结构处建议补充各脚本的输入输出文件约定，方便测试人员定位生成数据和验证结果。 | |||
| 17 | +``` | ||
| 18 | + | ||
| 19 | +## 样例描述 | ||
| 20 | + | ||
| 21 | +- 样例功能： | ||
| 22 | + | ||
| 23 | + 使用vector编程模式，核函数内使用固定输入展示simd_vf printf接口的基本使用方法。核函数依次调用simd_vf侧printf接口打印不同类型的数据和计算结果。 | ||
| 24 | + | ||
| 25 | + 本样例展示以下printf功能： | ||
| 26 | + 1. `printf`打印int/uint/float/string格式数据 | ||
| 27 | + 2. `printf`打印hex/ptr格式数据 | ||
| 28 | + 3. simd_vf侧`printf`打印纯字符串 | ||
| 29 | + 4. 在`.asc`内使用固定输入写入静态`LocalTensor`，通过AICore侧`AscendC::Add`完成tensor加法，并通过simd_vf侧`printf`打印计算结果 | ||
| 30 | + 5. aicore侧`AscendC::printf`打印核函数开始、blockIdx和结束信息 | ||
| 31 | + | ||
| 32 | + | 格式符 | 说明 | 示例 | | ||
| 33 | + |--------|------|------| | ||
| 34 | + | `%d` | 整型打印 | `printf(fmt, 10);` | | ||
| 35 | + | `%u` | 无符号整型打印 | `printf(fmt, 20U);` | | ||
| 36 | + | `%x` | 十六进制打印 | `printf(fmt, 255);` | | ||
| 37 | + | `%f` | 浮点型打印(float) | `printf(fmt, 3.14f);` | | ||
| 38 | + | `%s` | 字符串打印 | `printf(fmt, "test");` | | ||
[代码检视建议] 运行步骤中建议补充环境变量检查，例如 ASCEND_HOME_PATH/CANN 包路径，减少用户首次运行失败概率。 | |||
| 39 | + | `%p` | 指针打印 | `printf(fmt, ptr);` | | ||
| 40 | + | ||
| 41 | +- 调用实现 | ||
| 42 | + | ||
| 43 | + 在`__global__ __vector__`核函数内使用`asc_vf_call<FuncName>()`调用simd_vf函数。核函数使用固定输入值，不依赖外部数据或host侧数据搬运。Add演示中，AICore侧将固定数据写入静态UB `LocalTensor`，调用`AscendC::Add`生成结果，再将UB地址传给simd_vf函数打印。 | ||
| 44 | + | ||
| 45 | +## 编译运行 | ||
| 46 | + | ||
| 47 | +- 配置环境变量 | ||
| 48 | + 在本样例根目录下执行如下步骤，编译并执行样例。 | ||
| 49 | + 请根据当前环境上CANN开发套件包的[安装方式](https://gitcode.com/cann/asc-devkit/blob/master/docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。 | ||
| 50 | + - 默认路径，root用户安装CANN软件包 | ||
| 51 | + ```bash | ||
| 52 | + source /usr/local/Ascend/cann/set_env.sh | ||
| 53 | + ``` | ||
| 54 | + | ||
| 55 | + - 默认路径，非root用户安装CANN软件包 | ||
| 56 | + ```bash | ||
| 57 | + source $HOME/Ascend/cann/set_env.sh | ||
| 58 | + ``` | ||
[代码检视建议] 建议在预期结果部分给出一段典型 printf 输出，便于区分样例无输出和输出格式异常。 | |||
| 59 | + | ||
| 60 | + - 指定路径install_path，安装CANN软件包 | ||
| 61 | + ```bash | ||
| 62 | + source ${install_path}/cann/set_env.sh | ||
| 63 | + ``` | ||
| 64 | + | ||
| 65 | +- 样例执行 | ||
| 66 | + ```bash | ||
| 67 | + mkdir -p build && cd build; # 创建并进入build目录 | ||
| 68 | + cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程 | ||
| 69 | + ./demo # 执行样例 | ||
| 70 | + ``` | ||
| 71 | + | ||
| 72 | +- 编译选项说明 | ||
| 73 | + | ||
| 74 | + | 选项 | 可选值 | 说明 | | ||
| 75 | + |------|--------|------| | ||
| 76 | + | `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认） | NPU架构：dav-3510对应Ascend 950PR/Ascend 950DT | | ||
| 77 | + | ||
| 78 | +- 执行结果 | ||
| 79 | + 执行日志中可以看到aicore侧和simd_vf侧printf输出。 | ||
| 80 | + ```bash | ||
| 81 | + [AIV Block 0/1] [aicore] simd_vf printf demo start ... | ||
| 82 | + [AIV Block 0/1] [aicore] blockIdx=0 | ||
| 83 | + [simd_vf] int=10, uint=20, float=3.140000, string=hello | ||
| 84 | + [simd_vf] hex=ff, ptr=0x12ff | ||
| 85 | + [simd_vf] This is a simd_vf printf demo string. | ||
| 86 | + [simd_vf] add[0]: 1.250000 + 10.000000 = 11.250000 | ||
| 87 | + [AIV Block 0/1] [aicore] simd_vf printf demo end ... | ||
| 88 | + ``` | ||
| @@ -0,0 +1,101 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/*! | ||
| 12 | + * \file simd_vf_printf.asc | ||
| 13 | + * \brief 使用vector编程模式，展示simd_vf printf接口的基本使用方法 | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | +#include "acl/acl.h" | ||
| 17 | +#include "kernel_operator.h" | ||
| 18 | +#include "utils/debug/asc_printf.h" | ||
| 19 | + | ||
| 20 | +constexpr uint32_t PRINT_COUNT = 1; | ||
| 21 | +constexpr uint32_t ADD_DATA_LEN = 16; | ||
| 22 | + | ||
| 23 | +// simd_vf printf演示: int, uint, float, string格式打印 | ||
| 24 | +__simd_vf__ inline void SimdVfPrintFormat() | ||
| 25 | +{ | ||
| 26 | + __ubuf__ const char* fmt = "[simd_vf] int=%d, uint=%u, float=%f, string=%s\n"; | ||
C 是否构造一个简单的add场景说明一下？稍微带一些业务场景 | |||
| 27 | + printf(fmt, 10, 20U, 3.14f, "hello"); | ||
| 28 | +} | ||
| 29 | + | ||
| 30 | +// simd_vf printf演示: 十六进制和指针格式打印 | ||
| 31 | +__simd_vf__ inline void SimdVfPrintHexPtr() | ||
| 32 | +{ | ||
| 33 | + __ubuf__ const char* fmt = "[simd_vf] hex=%x, ptr=%p\n"; | ||
| 34 | + printf(fmt, 255, (void*)(uintptr_t)0x12ff); | ||
| 35 | +} | ||
| 36 | + | ||
| 37 | +// simd_vf printf演示: 纯字符串打印 | ||
| 38 | +__simd_vf__ inline void SimdVfPrintStr() | ||
[代码检视建议] 建议在 vector 函数中 printf 调用前补充注释说明打印的数据来源和 lane 对应关系。 | |||
| 39 | +{ | ||
| 40 | + __ubuf__ const char* fmt = "[simd_vf] This is a simd_vf printf demo string.\n"; | ||
| 41 | + printf(fmt); | ||
| 42 | +} | ||
| 43 | + | ||
| 44 | +// simd_vf printf演示: 打印AICore侧AscendC::Add计算后的UB tensor结果 | ||
| 45 | +__simd_vf__ inline void SimdVfPrintAdd(__ubuf__ float* x, __ubuf__ float* y, __ubuf__ float* z, uint32_t count) | ||
| 46 | +{ | ||
| 47 | + __ubuf__ const char* fmt = "[simd_vf] add[%u]: %f + %f = %f\n"; | ||
| 48 | + for (uint32_t i = 0; i < count; i++) { | ||
| 49 | + printf(fmt, i, x[i], y[i], z[i]); | ||
| 50 | + } | ||
| 51 | +} | ||
| 52 | + | ||
| 53 | +// AICore侧演示: 固定输入写入LocalTensor，调用AscendC::Add，再交给simd_vf打印 | ||
| 54 | +__aicore__ inline void SimdVfAddTensorDemo() | ||
| 55 | +{ | ||
| 56 | + AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator; | ||
| 57 | + AscendC::LocalTensor<float> xLocal = ubAllocator.Alloc<float, ADD_DATA_LEN>(); | ||
| 58 | + AscendC::LocalTensor<float> yLocal = ubAllocator.Alloc<float, ADD_DATA_LEN>(); | ||
| 59 | + AscendC::LocalTensor<float> zLocal = ubAllocator.Alloc<float, ADD_DATA_LEN>(); | ||
| 60 | + | ||
| 61 | + AscendC::Duplicate<float>(xLocal, 1.25f, ADD_DATA_LEN); | ||
| 62 | + AscendC::Duplicate<float>(yLocal, 10.00f, ADD_DATA_LEN); | ||
| 63 | + | ||
| 64 | + AscendC::Add(zLocal, xLocal, yLocal, ADD_DATA_LEN); | ||
| 65 | + AscendC::PipeBarrier<PIPE_V>(); | ||
| 66 | + | ||
| 67 | + asc_vf_call<SimdVfPrintAdd>((__ubuf__ float*)xLocal.GetPhyAddr(), (__ubuf__ float*)yLocal.GetPhyAddr(), | ||
| 68 | + (__ubuf__ float*)zLocal.GetPhyAddr(), PRINT_COUNT); | ||
| 69 | +} | ||
| 70 | + | ||
| 71 | +extern "C" __global__ __vector__ void simd_vf_printf_kernel() | ||
| 72 | +{ | ||
| 73 | + AscendC::InitSocState(); | ||
| 74 | + | ||
| 75 | + AscendC::printf("[aicore] simd_vf printf demo start ...\n"); | ||
| 76 | + AscendC::printf("[aicore] blockIdx=%u\n", AscendC::GetBlockIdx()); | ||
| 77 | + | ||
| 78 | + asc_vf_call<SimdVfPrintFormat>(); | ||
| 79 | + asc_vf_call<SimdVfPrintHexPtr>(); | ||
| 80 | + asc_vf_call<SimdVfPrintStr>(); | ||
| 81 | + SimdVfAddTensorDemo(); | ||
| 82 | + | ||
[代码检视建议] 建议确认 printf 调用不会在大规模循环中产生过量日志，必要时增加示例中的打印范围限制说明。 | |||
| 83 | + AscendC::printf("[aicore] simd_vf printf demo end ...\n"); | ||
| 84 | +} | ||
| 85 | + | ||
| 86 | +int32_t main(int32_t argc, char* argv[]) | ||
| 87 | +{ | ||
| 88 | + aclInit(nullptr); | ||
| 89 | + int32_t deviceId = 0; | ||
| 90 | + aclrtSetDevice(deviceId); | ||
| 91 | + aclrtStream stream = nullptr; | ||
| 92 | + aclrtCreateStream(&stream); | ||
| 93 | + | ||
| 94 | + simd_vf_printf_kernel<<<1, nullptr, stream>>>(); | ||
| 95 | + | ||
| 96 | + aclrtSynchronizeStream(stream); | ||
| 97 | + aclrtDestroyStream(stream); | ||
| 98 | + aclrtResetDevice(deviceId); | ||
| 99 | + aclFinalize(); | ||
| 100 | + return 0; | ||
| 101 | +} | ||
Rexamples/01_simd_cpp_api/01_utilities/00_printf/CMakeLists.txt→examples/01_simd_cpp_api/01_utilities/00_printf/simple_printf/CMakeLists.txt+0-0
文件重命名但无更改。
| @@ -0,0 +1,107 @@ | |||
| 1 | +# Printf接口功能说明 | ||
| 2 | + | ||
| 3 | +## 概述 | ||
| 4 | + | ||
| 5 | +本样例介绍printf接口使用方法，通过该接口打印核函数相关信息。 | ||
| 6 | + | ||
| 7 | +## 支持的产品 | ||
| 8 | + | ||
| 9 | +- Ascend 950PR/Ascend 950DT | ||
| 10 | +- Atlas A3训练系列产品/Atlas A3推理系列产品 | ||
| 11 | +- Atlas A2训练系列产品/Atlas A2推理系列产品 | ||
| 12 | + | ||
| 13 | +## 目录结构介绍 | ||
| 14 | + | ||
| 15 | +``` | ||
| 16 | +├── 00_printf | ||
| 17 | +│ ├── scripts | ||
| 18 | +│ │ ├── gen_data.py // 输入数据和真值数据生成脚本 | ||
| 19 | +│ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 | ||
| 20 | +│ ├── CMakeLists.txt // 编译工程文件 | ||
| 21 | +│ ├── data_utils.h // 数据读入写出函数 | ||
| 22 | +│ └── printf.asc // Ascend C样例实现 & 调用样例 | ||
| 23 | +``` | ||
| 24 | + | ||
| 25 | +## 样例描述 | ||
| 26 | + | ||
| 27 | +- 样例功能： | ||
| 28 | + | ||
| 29 | + 使用静态Tensor编程模式实现矩阵乘法，展示printf接口的基本使用方法。 | ||
| 30 | + | ||
| 31 | + 矩阵乘法的计算公式为： | ||
| 32 | + | ||
| 33 | + ``` | ||
| 34 | + C = A * B | ||
| 35 | + ``` | ||
| 36 | + | ||
| 37 | +- 样例规格： | ||
| 38 | + | ||
| 39 | + 样例参数为：M = 256, N = 256, K = 64，shape信息如下表所示： | ||
| 40 | + <table> | ||
| 41 | + <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center">Matmul</td></tr> | ||
| 42 | + </tr> | ||
| 43 | + <tr><td rowspan="3" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr> | ||
| 44 | + <tr><td align="center">a</td><td align="center">[M, K]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 45 | + <tr><td align="center">b</td><td align="center">[K, N]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 46 | + </tr> | ||
| 47 | + </tr> | ||
| 48 | + <tr><td rowspan="1" align="center">样例输出</td><td align="center">c</td><td align="center">[M, N]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 49 | + </tr> | ||
| 50 | + <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">mmad_custom</td></tr> | ||
| 51 | + </table> | ||
| 52 | + | ||
| 53 | +- printf接口支持格式： | ||
| 54 | + | ||
| 55 | + | 格式符 | 说明 | 示例 | | ||
| 56 | + |--------|------|------| | ||
| 57 | + | `%p` | 指针打印 | `AscendC::printf("pointer %p\n", ptr);` | | ||
| 58 | + | `%d` | 整型/bool打印 | `AscendC::printf("value is %d\n", 10);` | | ||
| 59 | + | `%u` | 无符号整型打印 | `AscendC::printf("idx is %u\n", idx);` | | ||
| 60 | + | `%x` | 十六进制打印 | `AscendC::printf("value is %x\n", 255);` | | ||
| 61 | + | `%f` | 浮点型打印(half/float) | `AscendC::printf("half %f\n", val);` | | ||
| 62 | + | `%s` | 字符串打印 | `AscendC::printf("name %s\n", "test");` | | ||
| 63 | + | ||
| 64 | +- 调用实现 | ||
| 65 | + | ||
| 66 | + 使用内核调用符<<<>>>调用核函数。 | ||
| 67 | + | ||
| 68 | +## 编译运行 | ||
| 69 | + | ||
| 70 | +- 配置环境变量 | ||
| 71 | +在本样例根目录下执行如下步骤，编译并执行样例。 | ||
| 72 | + 请根据当前环境上CANN开发套件包的[安装方式](https://gitcode.com/cann/asc-devkit/blob/master/docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。 | ||
| 73 | + - 默认路径，root用户安装CANN软件包 | ||
| 74 | + ```bash | ||
| 75 | + source /usr/local/Ascend/cann/set_env.sh | ||
| 76 | + ``` | ||
| 77 | + | ||
| 78 | + - 默认路径，非root用户安装CANN软件包 | ||
| 79 | + ```bash | ||
| 80 | + source $HOME/Ascend/cann/set_env.sh | ||
| 81 | + ``` | ||
| 82 | + | ||
| 83 | + - 指定路径install_path，安装CANN软件包 | ||
| 84 | + ```bash | ||
| 85 | + source ${install_path}/cann/set_env.sh | ||
| 86 | + ``` | ||
| 87 | + | ||
| 88 | +- 样例执行 | ||
| 89 | + ```bash | ||
| 90 | + mkdir -p build && cd build; # 创建并进入build目录 | ||
| 91 | + cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程 | ||
| 92 | + python3 ../scripts/gen_data.py # 生成测试输入数据 | ||
| 93 | + ./demo # 执行编译生成的可执行程序，执行样例 | ||
| 94 | + python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确 | ||
| 95 | + ``` | ||
| 96 | + | ||
| 97 | +- 编译选项说明 | ||
| 98 | + | ||
| 99 | +| 选项 | 可选值 | 说明 | | ||
| 100 | +|------|--------|------| | ||
| 101 | +| `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU架构：dav-2201对应Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品，dav-3510对应Ascend 950PR/Ascend 950DT | | ||
| 102 | + | ||
| 103 | +- 执行结果 | ||
| 104 | + 最终执行结果如下，说明精度对比成功。 | ||
| 105 | + ```bash | ||
| 106 | + test pass! | ||
| 107 | + ``` | ||
Rexamples/01_simd_cpp_api/01_utilities/00_printf/data_utils.h→examples/01_simd_cpp_api/01_utilities/00_printf/simple_printf/data_utils.h+0-0
文件重命名但无更改。
Rexamples/01_simd_cpp_api/01_utilities/00_printf/printf.asc→examples/01_simd_cpp_api/01_utilities/00_printf/simple_printf/printf.asc+0-0
文件重命名但无更改。
Rexamples/01_simd_cpp_api/01_utilities/00_printf/scripts/gen_data.py→examples/01_simd_cpp_api/01_utilities/00_printf/simple_printf/scripts/gen_data.py+0-0
文件重命名但无更改。
Rexamples/01_simd_cpp_api/01_utilities/00_printf/scripts/verify_result.py→examples/01_simd_cpp_api/01_utilities/00_printf/simple_printf/scripts/verify_result.py+0-0
文件重命名但无更改。
| @@ -1,113 +1,9 @@ | |||
| 1 | -# Dump接口功能说明 | 1 | +# Dump样例介绍 |
| 2 | - | ||
| 3 | ## 概述 | 2 | ## 概述 |
| 3 | +本样例展示了Ascend C asc_dump系列接口和simd_vf侧asc_dump接口的基本使用方法。 | ||
| 4 | 4 | ||
| 5 | -本样例介绍asc_dump_xxx系列接口使用方法，通过该接口实现核函数中不同物理位置的张量数据可视化。 | 5 | +## 样例列表 |
[代码检视建议] 顶层 Dump README 建议补充 simple_dump 与 simd_vf_dump 的适用场景差异，便于用户选择样例。 | |||
| 6 | - | 6 | +| 目录名称 | 功能描述 | 支持的产品 | |
| 7 | -## 支持的产品 | 7 | +| ------------------------------------------------------------ | ---------------------------------------------------- | --- | |
| 8 | - | 8 | +| [simple_dump](./simple_dump) | 使用静态Tensor编程模式实现矩阵乘法，展示asc_dump_gm/asc_dump_l1buf/asc_dump_cbuf/asc_dump_ubuf接口的基本使用方法 | Ascend 950PR/Ascend 950DT<br>Atlas A3训练系列产品/Atlas A3推理系列产品<br>Atlas A2训练系列产品/Atlas A2推理系列产品 | |
| 9 | -- Ascend 950PR/Ascend 950DT | 9 | +| [simd_vf_dump](./simd_vf_dump) | 使用vector编程模式，展示simd_vf侧asc_dump_ubuf/asc_dump/asc_dump_reg接口的基本使用方法 | Ascend 950PR/Ascend 950DT | |
| 10 | -- Atlas A3 训练系列产品/Atlas A3 推理系列产品 | ||
| 11 | -- Atlas A2 训练系列产品/Atlas A2 推理系列产品 | ||
| 12 | - | ||
| 13 | -## 目录结构介绍 | ||
| 14 | - | ||
| 15 | -``` | ||
| 16 | -├── 02_dump | ||
| 17 | -│ ├── scripts | ||
| 18 | -│ │ ├── gen_data.py // 输入数据和真值数据生成脚本 | ||
| 19 | -│ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 | ||
| 20 | -│ ├── CMakeLists.txt // 编译工程文件 | ||
| 21 | -│ ├── data_utils.h // 数据读入写出函数 | ||
| 22 | -│ └── dump.asc // Ascend C样例实现 & 调用样例 | ||
| 23 | -``` | ||
| 24 | - | ||
| 25 | -## 样例描述 | ||
| 26 | - | ||
| 27 | -- 样例功能： | ||
| 28 | - | ||
| 29 | - 使用静态Tensor编程模式实现矩阵乘法，展示asc_dump_xxx系列接口的基本使用方法。 | ||
| 30 | - 该系列接口兼容AscendC::DumpTensor接口。但后续开发中，建议优先使用asc_dump系列接口；若需Dump指定偏移位置的数据，由于asc_dump系列暂不支持该能力，可继续使用DumpAccChkPoint接口。 | ||
| 31 | - | ||
| 32 | - 矩阵乘法的计算公式为： | ||
| 33 | - | ||
| 34 | - ``` | ||
| 35 | - C = A * B | ||
| 36 | - ``` | ||
| 37 | - | ||
| 38 | -- 样例规格： | ||
| 39 | - | ||
| 40 | - 样例参数为：M = 256, N = 256, K = 64，shape信息如下表所示： | ||
| 41 | - <table> | ||
| 42 | - <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center">Matmul</td></tr> | ||
| 43 | - </tr> | ||
| 44 | - <tr><td rowspan="3" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr> | ||
| 45 | - <tr><td align="center">a</td><td align="center">[M, K]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 46 | - <tr><td align="center">b</td><td align="center">[K, N]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 47 | - </tr> | ||
| 48 | - </tr> | ||
| 49 | - <tr><td rowspan="1" align="center">样例输出</td><td align="center">c</td><td align="center">[M, N]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 50 | - </tr> | ||
| 51 | - <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">mmad_custom</td></tr> | ||
| 52 | - </table> | ||
| 53 | - | ||
| 54 | -- asc_dump接口用法： | ||
| 55 | - | ||
| 56 | - | 接口 | 内存层级 | 地址类型 | 示例 | | ||
| 57 | - |------|----------|----------|------| | ||
| 58 | - | `asc_dump_gm` | GM (Global Memory) | `__gm__` | `asc_dump_gm<half>((__gm__ half*)addr, id, size);` | | ||
| 59 | - | `asc_dump_l1buf` | L1 Buffer | `__cbuf__` | `asc_dump_l1buf<half>((__cbuf__ half*)addr, id, size);` | | ||
| 60 | - | `asc_dump_cbuf` | L0C Buffer | `__cc__` | `asc_dump_cbuf<float>((__cc__ float*)addr, id, size);` | | ||
| 61 | - | `asc_dump_ubuf` | UB | `__ubuf__` | `asc_dump_ubuf<half>((__ubuf__ half*)addr, id, size);` | | ||
| 62 | - | ||
| 63 | - 参数说明： | ||
| 64 | - - 模板参数：指定数据类型（half/float等） | ||
| 65 | - - id：自定义标识符，用于区分不同dump数据 | ||
| 66 | - - dumpSize：dump数据元素个数 | ||
| 67 | - | ||
| 68 | - 注意：`asc_dump_l1buf`接口仅支持Atlas A2训练系列产品/Atlas A2推理系列产品/Atlas A3训练系列产品/Atlas A3推理系列产品。 | ||
| 69 | - | ||
| 70 | -- 调用实现 | ||
| 71 | - | ||
| 72 | - 使用内核调用符<<<>>>调用核函数。 | ||
| 73 | - | ||
| 74 | -## 编译运行 | ||
| 75 | - | ||
| 76 | -在本样例根目录下执行如下步骤，编译并执行样例。 | ||
| 77 | -- 配置环境变量 | ||
| 78 | - 请根据当前环境上CANN开发套件包的[安装方式](../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。 | ||
| 79 | - - 默认路径，root用户安装CANN软件包 | ||
| 80 | - ```bash | ||
| 81 | - source /usr/local/Ascend/cann/set_env.sh | ||
| 82 | - ``` | ||
| 83 | - | ||
| 84 | - - 默认路径，非root用户安装CANN软件包 | ||
| 85 | - ```bash | ||
| 86 | - source $HOME/Ascend/cann/set_env.sh | ||
| 87 | - ``` | ||
| 88 | - | ||
| 89 | - - 指定路径install_path，安装CANN软件包 | ||
| 90 | - ```bash | ||
| 91 | - source ${install_path}/cann/set_env.sh | ||
| 92 | - ``` | ||
| 93 | - | ||
| 94 | -- 样例执行 | ||
| 95 | - ```bash | ||
| 96 | - mkdir -p build && cd build; # 创建并进入build目录 | ||
| 97 | - cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程 | ||
| 98 | - python3 ../scripts/gen_data.py # 生成测试输入数据 | ||
| 99 | - ./demo # 执行编译生成的可执行程序，执行样例 | ||
| 100 | - python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确 | ||
| 101 | - ``` | ||
| 102 | - | ||
| 103 | -- 编译选项说明 | ||
| 104 | - | ||
| 105 | -| 选项 | 可选值 | 说明 | | ||
| 106 | -|------|--------|------| | ||
| 107 | -| `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品和Atlas A3 训练系列产品/Atlas A3 推理系列产品，dav-3510 对应 Ascend 950PR/Ascend 950DT | | ||
| 108 | - | ||
| 109 | -- 执行结果 | ||
| 110 | - 最终执行结果如下，说明精度对比成功。 | ||
| 111 | - ```bash | ||
| 112 | - test pass! | ||
| 113 | - ``` | ||
| @@ -0,0 +1,12 @@ | |||
| 1 | +# Dump Samples | ||
| 2 | + | ||
| 3 | +## Overview | ||
| 4 | + | ||
| 5 | +This sample shows how to use the Ascend C asc_dump interface series and the simd_vf asc_dump interfaces. | ||
[代码检视建议] 英文 Dump README 建议补充与中文文档一致的运行步骤和输出说明，避免英文用户缺少关键操作信息。 | |||
| 6 | + | ||
| 7 | +## Sample List | ||
| 8 | + | ||
| 9 | +| Directory | Description | Supported Products | | ||
| 10 | +| --- | --- | --- | | ||
| 11 | +| [simple_dump](./simple_dump) | Uses static Tensor programming mode to implement matrix multiplication, showing basic usage of the asc_dump_gm/asc_dump_l1buf/asc_dump_cbuf/asc_dump_ubuf interfaces | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series/Atlas A3 Inference Series<br>Atlas A2 Training Series/Atlas A2 Inference Series | | ||
| 12 | +| [simd_vf_dump](./simd_vf_dump) | Uses vector programming mode, showing basic usage of the simd_vf asc_dump_ubuf/asc_dump/asc_dump_reg interfaces | Ascend 950PR/Ascend 950DT | | ||
| @@ -0,0 +1,26 @@ | |||
| 1 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 2 | +# Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | +# CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | +# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 10 | + | ||
| 11 | + | ||
| 12 | +cmake_minimum_required(VERSION 3.16) | ||
| 13 | + | ||
| 14 | +set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510") | ||
| 15 | + | ||
| 16 | +find_package(ASC REQUIRED) | ||
| 17 | + | ||
| 18 | +project(utils_debug LANGUAGES ASC CXX) | ||
| 19 | + | ||
| 20 | +add_executable(demo | ||
| 21 | + simd_vf_dump.asc | ||
| 22 | +) | ||
| 23 | + | ||
| 24 | +target_compile_options(demo PRIVATE | ||
| 25 | + $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}> | ||
| 26 | +) | ||
| @@ -0,0 +1,92 @@ | |||
| 1 | +# SimdVF Dump接口功能说明 | ||
| 2 | + | ||
| 3 | +## 概述 | ||
| 4 | + | ||
| 5 | +本样例介绍simd_vf侧asc_dump系列接口使用方法，在vector函数中通过`asc_dump_ubuf`、`asc_dump_reg`、`asc_dump`等接口实现UB和寄存器数据的可视化打印。核函数从二进制文件读取输入数据，执行Adds计算后写出输出数据，通过比较输出与真值进行校验。 | ||
| 6 | + | ||
| 7 | +## 支持的产品 | ||
| 8 | + | ||
| 9 | +- Ascend 950PR/Ascend 950DT | ||
| 10 | + | ||
| 11 | +## 目录结构介绍 | ||
| 12 | + | ||
[代码检视建议] 支持产品章节建议明确 simd_vf dump 对 SoC/运行模式的限制，避免在不支持环境中执行失败。 | |||
| 13 | +``` | ||
| 14 | +├── simd_vf_dump | ||
| 15 | +│ ├── scripts | ||
| 16 | +│ │ ├── gen_data.py // 输入数据和真值数据生成脚本 | ||
| 17 | +│ │ └── verify_result.py // 验证输出数据和真值数据是否一致 | ||
| 18 | +│ ├── CMakeLists.txt // 编译工程文件 | ||
| 19 | +│ ├── data_utils.h // 数据读入写出函数 | ||
| 20 | +│ └── simd_vf_dump.asc // Ascend C样例实现&调用样例 | ||
| 21 | +``` | ||
| 22 | + | ||
| 23 | +## 样例描述 | ||
| 24 | + | ||
| 25 | +- 样例功能： | ||
[代码检视建议] 目录结构建议补充 dump 输出目录和验证脚本读取路径，便于排查文件路径错误。 | |||
| 26 | + | ||
| 27 | + 使用vector编程模式，展示simd_vf侧asc_dump系列接口的基本使用方法。核函数通过`ReadFile`读取二进制输入数据，使用`asc_vf_call`调用simd_vf函数进行dump打印，执行Adds计算后通过`WriteFile`写出二进制输出数据，最终通过校验脚本比对输出与真值。 | ||
| 28 | + | ||
| 29 | + 1. `asc_dump_ubuf<T>`-打印UB数据 | ||
| 30 | + 2. `asc_dump_reg<T>`-打印寄存器数据 | ||
| 31 | + 3. `asc_dump<T>`-打印UB/Reg数据（通用接口） | ||
| 32 | + | ||
| 33 | +- simd_vf dump接口用法： | ||
| 34 | + | ||
| 35 | + | 接口 | 地址类型 | 示例 | | ||
| 36 | + |------|----------|------| | ||
| 37 | + | `asc_dump_ubuf<T>` | `__ubuf__` | `asc_dump_ubuf<float>(input, desc, dump_size);` | | ||
| 38 | + | `asc_dump_reg<T>` | Reg | `asc_dump_reg<float>(input, desc, dump_size);` | | ||
| 39 | + | `asc_dump<T>` | `__ubuf__` / Reg | `asc_dump<float>(input, desc, dump_size);` | | ||
| 40 | + | ||
| 41 | + 参数说明： | ||
| 42 | + - 模板参数`T`：指定数据类型（float/half等） | ||
| 43 | + - `input`：UB数据地址或Reg寄存器对象 | ||
| 44 | + - `desc`：打印控制描述符 | ||
| 45 | + - `dump_size`：单次dump的数据元素个数 | ||
| 46 | + | ||
| 47 | + | ||
| 48 | + | ||
| 49 | +- 调用实现 | ||
| 50 | + | ||
[代码检视建议] 运行步骤建议明确是否需要清理历史 dump 文件，避免重复运行时旧结果干扰本次验证。 | |||
| 51 | + 在`__global__ __vector__`核函数内使用`asc_vf_call<FuncName>()`调用simd_vf函数完成dump打印，随后通过`AscendC::Adds`执行计算并将结果写回GM。Host侧通过`ReadFile`/`WriteFile`完成二进制数据的读写。 | ||
| 52 | + | ||
| 53 | +## 编译运行 | ||
| 54 | + | ||
| 55 | +- 配置环境变量 | ||
| 56 | + 在本样例根目录下执行如下步骤，编译并执行样例。 | ||
| 57 | + 请根据当前环境上CANN开发套件包的[安装方式](https://gitcode.com/cann/asc-devkit/blob/master/docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。 | ||
| 58 | + - 默认路径，root用户安装CANN软件包 | ||
| 59 | + ```bash | ||
| 60 | + source /usr/local/Ascend/cann/set_env.sh | ||
| 61 | + ``` | ||
| 62 | + | ||
| 63 | + - 默认路径，非root用户安装CANN软件包 | ||
| 64 | + ```bash | ||
| 65 | + source $HOME/Ascend/cann/set_env.sh | ||
| 66 | + ``` | ||
| 67 | + | ||
| 68 | + - 指定路径install_path，安装CANN软件包 | ||
| 69 | + ```bash | ||
| 70 | + source ${install_path}/cann/set_env.sh | ||
[代码检视建议] 预期结果建议展示 dump 数据片段或校验通过日志，方便测试人员快速判定执行结果。 | |||
| 71 | + ``` | ||
| 72 | + | ||
| 73 | +- 样例执行 | ||
| 74 | + ```bash | ||
| 75 | + mkdir -p build && cd build; # 创建并进入build目录 | ||
| 76 | + cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程 | ||
| 77 | + python3 ../scripts/gen_data.py # 生成测试输入数据 | ||
| 78 | + ./demo # 执行编译生成的可执行程序，执行样例 | ||
| 79 | + python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确 | ||
| 80 | + ``` | ||
| 81 | + | ||
| 82 | +- 编译选项说明 | ||
| 83 | + | ||
| 84 | + | 选项 | 可选值 | 说明 | | ||
| 85 | + |------|--------|------| | ||
| 86 | + | `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认） | NPU架构：dav-3510对应Ascend 950PR/Ascend 950DT | | ||
| 87 | + | ||
| 88 | +- 执行结果 | ||
| 89 | + 最终执行结果如下，说明精度对比成功。 | ||
| 90 | + ```bash | ||
| 91 | + test pass! | ||
| 92 | + ``` | ||
| @@ -0,0 +1,87 @@ | |||
| 1 | +/** | ||
| 2 | +* Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | +* This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | +* CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | +* Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | +* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | +* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | +* See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | +*/ | ||
| 10 | + | ||
| 11 | + | ||
| 12 | +/* ! | ||
| 13 | + * \file data_utils.h | ||
| 14 | + * \brief | ||
| 15 | + */ | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | +bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize) | ||
| 27 | +{ | ||
| 28 | + struct stat sBuf; | ||
| 29 | + int fileStatus = stat(filePath.data(), &sBuf); | ||
| 30 | + if (fileStatus == -1) { | ||
| 31 | + ERROR_LOG("failed to get file"); | ||
| 32 | + return false; | ||
| 33 | + } | ||
| 34 | + if (S_ISREG(sBuf.st_mode) == 0) { | ||
| 35 | + ERROR_LOG("%s is not a file, please enter a file", filePath.c_str()); | ||
| 36 | + return false; | ||
| 37 | + } | ||
| 38 | + | ||
| 39 | + std::ifstream file; | ||
| 40 | + file.open(filePath, std::ios::binary); | ||
| 41 | + if (!file.is_open()) { | ||
| 42 | + ERROR_LOG("Open file failed. path = %s", filePath.c_str()); | ||
| 43 | + return false; | ||
| 44 | + } | ||
| 45 | + | ||
| 46 | + std::filebuf *buf = file.rdbuf(); | ||
| 47 | + size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in); | ||
| 48 | + if (size == 0) { | ||
| 49 | + ERROR_LOG("file size is 0"); | ||
| 50 | + file.close(); | ||
| 51 | + return false; | ||
| 52 | + } | ||
| 53 | + if (size > bufferSize) { | ||
| 54 | + ERROR_LOG("file size is larger than buffer size"); | ||
| 55 | + file.close(); | ||
| 56 | + return false; | ||
| 57 | + } | ||
| 58 | + buf->pubseekpos(0, std::ios::in); | ||
| 59 | + buf->sgetn(static_cast<char *>(buffer), size); | ||
| 60 | + fileSize = size; | ||
| 61 | + file.close(); | ||
| 62 | + return true; | ||
| 63 | +} | ||
| 64 | + | ||
| 65 | +bool WriteFile(const std::string &filePath, const void *buffer, size_t size) | ||
| 66 | +{ | ||
| 67 | + if (buffer == nullptr) { | ||
| 68 | + ERROR_LOG("Write file failed. buffer is nullptr"); | ||
| 69 | + return false; | ||
| 70 | + } | ||
| 71 | + | ||
| 72 | + int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE); | ||
| 73 | + if (fd < 0) { | ||
| 74 | + ERROR_LOG("Open file failed. path = %s", filePath.c_str()); | ||
| 75 | + return false; | ||
| 76 | + } | ||
| 77 | + | ||
| 78 | + size_t writeSize = write(fd, buffer, size); | ||
| 79 | + (void)close(fd); | ||
| 80 | + if (writeSize != size) { | ||
| 81 | + ERROR_LOG("Write file Failed."); | ||
| 82 | + return false; | ||
| 83 | + } | ||
| 84 | + | ||
| 85 | + return true; | ||
| 86 | +} | ||
| 87 | + | ||
| @@ -0,0 +1,30 @@ | |||
| 1 | +#!/usr/bin/python3 | ||
| 2 | +# coding=utf-8 | ||
| 3 | + | ||
| 4 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 5 | +# Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 6 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 7 | +# CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 8 | +# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 9 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 10 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 11 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 12 | +# ---------------------------------------------------------------------------------------------------------- | ||
[代码检视建议] 数据生成脚本建议固定随机种子或说明数据生成规则，保证样例验证结果可复现。 | |||
| 13 | + | ||
| 14 | + | ||
| 15 | +import os | ||
| 16 | +import numpy as np | ||
| 17 | + | ||
C 文件不是必须，可以删除，我感觉verify_result.py也不是必须，建议就用一个.asc文件就可以了 | |||
| 18 | + | ||
| 19 | +def gen_data(): | ||
| 20 | + data_len = 32 | ||
| 21 | + os.makedirs("input", exist_ok=True) | ||
| 22 | + os.makedirs("output", exist_ok=True) | ||
| 23 | + x = np.linspace(0, data_len - 1, data_len, dtype=np.float32) | ||
| 24 | + x.tofile("./input/x_gm.bin") | ||
| 25 | + golden = x + 1.0 | ||
| 26 | + golden.tofile("./output/golden.bin") | ||
| 27 | + | ||
| 28 | + | ||
| 29 | +if __name__ == "__main__": | ||
| 30 | + gen_data() | ||
| @@ -0,0 +1,57 @@ | |||
| 1 | +#!/usr/bin/python3 | ||
| 2 | +# coding=utf-8 | ||
| 3 | + | ||
| 4 | +# ---------------------------------------------------------------------------------------------------------- | ||
| 5 | +# Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 6 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 7 | +# CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 8 | +# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 9 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 10 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 11 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 12 | +# ---------------------------------------------------------------------------------------------------------- | ||
[代码检视建议] 数据生成脚本建议补充输出目录不存在时的创建逻辑或明确报错，避免依赖外部目录准备。 | |||
| 13 | + | ||
| 14 | + | ||
| 15 | +import sys | ||
| 16 | +import numpy as np | ||
| 17 | + | ||
| 18 | + | ||
| 19 | +RELATIVE_TOL = 1e-3 | ||
| 20 | +ABSOLUTE_TOL = 1e-4 | ||
| 21 | +ERROR_TOL = 1e-4 | ||
| 22 | + | ||
| 23 | + | ||
| 24 | +def verify_result(output, golden): | ||
| 25 | + output = np.fromfile(output, dtype=np.float32).reshape(-1) | ||
| 26 | + golden = np.fromfile(golden, dtype=np.float32).reshape(-1) | ||
| 27 | + different_element_results = np.isclose(output, | ||
| 28 | + golden, | ||
| 29 | + rtol=RELATIVE_TOL, | ||
| 30 | + atol=ABSOLUTE_TOL, | ||
| 31 | + equal_nan=True) | ||
| 32 | + different_element_indexes = np.where(different_element_results == False)[0] | ||
| 33 | + for index in range(len(different_element_indexes)): | ||
| 34 | + real_index = different_element_indexes[index] | ||
| 35 | + golden_data = golden[real_index] | ||
| 36 | + output_data = output[real_index] | ||
| 37 | + print( | ||
| 38 | + "data index: %06d, expected: %-.9f, actual: %-.9f, rdiff: %-.6f" % | ||
| 39 | + (real_index, golden_data, output_data, | ||
| 40 | + abs(output_data - golden_data) / golden_data)) | ||
| 41 | + if index == 100: | ||
| 42 | + break | ||
| 43 | + error_ratio = float(different_element_indexes.size) / golden.size | ||
| 44 | + print("error ratio: %.4f, tolerance: %.4f" % (error_ratio, ERROR_TOL)) | ||
| 45 | + return error_ratio <= ERROR_TOL | ||
| 46 | + | ||
| 47 | + | ||
| 48 | +if __name__ == '__main__': | ||
| 49 | + try: | ||
| 50 | + res = verify_result(sys.argv[1], sys.argv[2]) | ||
| 51 | + if not res: | ||
| 52 | + raise ValueError("[ERROR] result error") | ||
| 53 | + else: | ||
| 54 | + print("test pass!") | ||
| 55 | + except Exception as e: | ||
| 56 | + print(e) | ||
| 57 | + sys.exit(1) | ||
| @@ -0,0 +1,131 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/*! | ||
| 12 | + * \file simd_vf_dump.asc | ||
| 13 | + * \brief 使用vector编程模式，展示simd_vf侧asc_dump_ubuf/asc_dump/asc_dump_reg接口的基本使用方法 | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | +#include "acl/acl.h" | ||
| 17 | +#include "kernel_operator.h" | ||
| 18 | +#include "utils/debug/asc_dump.h" | ||
| 19 | +#include "data_utils.h" | ||
| 20 | + | ||
| 21 | +namespace { | ||
| 22 | + constexpr uint32_t DATA_LEN = 32; | ||
| 23 | + constexpr uint32_t REG_COUNT = 32; | ||
| 24 | +} | ||
| 25 | + | ||
| 26 | +// simd_vf dump演示: 使用asc_dump_ubuf打印UB数据 | ||
| 27 | +__simd_vf__ inline void DumpUbufData(__ubuf__ float* x) | ||
| 28 | +{ | ||
| 29 | + asc_dump_ubuf<float>(x, 5, 32); | ||
| 30 | + asc_dump<float>(x, 5, 16); | ||
| 31 | +} | ||
| 32 | + | ||
| 33 | +// simd_vf dump演示: 使用asc_dump_reg打印Reg数据 | ||
[代码检视建议] 建议在 asc_dump 调用附近补充注释说明 dump 的物理位置和期望数据含义。 | |||
| 34 | +__simd_vf__ inline void DumpRegData(__ubuf__ float* x) | ||
| 35 | +{ | ||
| 36 | + AscendC::Reg::RegTensor<float> srcReg; | ||
| 37 | + AscendC::Reg::MaskReg maskReg; | ||
| 38 | + uint32_t count = REG_COUNT; | ||
| 39 | + maskReg = AscendC::Reg::UpdateMask<float>(count); | ||
| 40 | + AscendC::Reg::LoadAlign(srcReg, x); | ||
| 41 | + AscendC::Reg::Duplicate<float>(srcReg, (float)3); | ||
| 42 | + | ||
| 43 | + asc_dump_reg<float>(srcReg, 5, 32); | ||
| 44 | + AscendC::Reg::Duplicate<float>(srcReg, (float)4); | ||
| 45 | + asc_dump<float>(srcReg, 5, 16); | ||
| 46 | +} | ||
| 47 | + | ||
| 48 | +// simd_vf dump演示: 使用asc_dump打印修改后的UB数据 | ||
| 49 | +__simd_vf__ inline void DumpModifiedUbuf(__ubuf__ float* x) | ||
| 50 | +{ | ||
| 51 | + asc_dump<float>(x, 5, 32); | ||
| 52 | +} | ||
| 53 | + | ||
| 54 | +__global__ __vector__ void simd_vf_dump_kernel(__gm__ uint8_t* x, __gm__ uint8_t* z) | ||
| 55 | +{ | ||
| 56 | + uint32_t xLength = DATA_LEN; | ||
| 57 | + | ||
| 58 | + AscendC::GlobalTensor<float> gmInput; | ||
| 59 | + AscendC::GlobalTensor<float> gmOutput; | ||
| 60 | + gmInput.SetGlobalBuffer((__gm__ float*)x, xLength); | ||
| 61 | + gmOutput.SetGlobalBuffer((__gm__ float*)z, xLength); | ||
| 62 | + | ||
| 63 | + AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator; | ||
| 64 | + AscendC::LocalTensor<float> ubInput = ubAllocator.Alloc<float, DATA_LEN>(); | ||
| 65 | + | ||
| 66 | + AscendC::DataCopy(ubInput, gmInput, xLength); | ||
| 67 | + AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); | ||
| 68 | + AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); | ||
| 69 | + | ||
| 70 | + AscendC::printf("[aicore] simd_vf dump demo start ...\n"); | ||
| 71 | + | ||
| 72 | + // 1. 调用asc_dump_ubuf接口：打印UB数据 | ||
| 73 | + AscendC::printf("========== UB data (asc_dump_ubuf + asc_dump) ==========\n"); | ||
| 74 | + asc_vf_call<DumpUbufData>((__ubuf__ float *)ubInput.GetPhyAddr()); | ||
| 75 | + | ||
| 76 | + // 2. 调用asc_dump_reg接口：从UB加载到Reg后打印Reg数据 | ||
| 77 | + AscendC::printf("========== Register data (asc_dump_reg + asc_dump) ==========\n"); | ||
| 78 | + asc_vf_call<DumpRegData>((__ubuf__ float *)ubInput.GetPhyAddr()); | ||
| 79 | + | ||
| 80 | + // 3. 调用asc_dump接口：打印UB数据 | ||
| 81 | + AscendC::printf("========== UB data (asc_dump) ==========\n"); | ||
| 82 | + asc_vf_call<DumpModifiedUbuf>((__ubuf__ float *)ubInput.GetPhyAddr()); | ||
| 83 | + | ||
| 84 | + // 4. 执行Adds计算（x + 1.0），将结果写回GM，用于后续校验 | ||
| 85 | + AscendC::Adds(ubInput, ubInput, 1.0f, xLength); | ||
| 86 | + AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); | ||
| 87 | + AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); | ||
| 88 | + AscendC::DataCopy(gmOutput, ubInput, xLength); | ||
| 89 | + | ||
| 90 | + AscendC::printf("[aicore] simd_vf dump demo end ...\n"); | ||
| 91 | +} | ||
| 92 | + | ||
| 93 | +int32_t main(int32_t argc, char* argv[]) | ||
| 94 | +{ | ||
| 95 | + aclInit(nullptr); | ||
| 96 | + int32_t deviceId = 0; | ||
| 97 | + aclrtSetDevice(deviceId); | ||
[代码检视建议] 建议对 dump 数据量进行适当控制，避免示例在默认配置下生成过大调测文件。 | |||
| 98 | + aclrtStream stream = nullptr; | ||
| 99 | + aclrtCreateStream(&stream); | ||
| 100 | + | ||
| 101 | + constexpr uint32_t xTotalLength = DATA_LEN; | ||
| 102 | + size_t fileSize = xTotalLength * sizeof(float); | ||
| 103 | + | ||
| 104 | + uint8_t* xHost; | ||
| 105 | + uint8_t* xDevice; | ||
| 106 | + aclrtMallocHost((void**)(&xHost), fileSize); | ||
| 107 | + aclrtMalloc((void**)&xDevice, fileSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 108 | + ReadFile("./input/x_gm.bin", fileSize, xHost, fileSize); | ||
| 109 | + aclrtMemcpy(xDevice, fileSize, xHost, fileSize, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 110 | + | ||
| 111 | + uint8_t* zHost; | ||
| 112 | + uint8_t* zDevice; | ||
| 113 | + aclrtMallocHost((void**)(&zHost), fileSize); | ||
| 114 | + aclrtMalloc((void**)&zDevice, fileSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 115 | + | ||
| 116 | + constexpr uint32_t numBlocks = 1; | ||
| 117 | + simd_vf_dump_kernel<<<numBlocks, nullptr, stream>>>(xDevice, zDevice); | ||
| 118 | + aclrtSynchronizeStream(stream); | ||
| 119 | + | ||
| 120 | + aclrtMemcpy(zHost, fileSize, zDevice, fileSize, ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 121 | + WriteFile("./output/output.bin", zHost, fileSize); | ||
| 122 | + | ||
| 123 | + aclrtFree(xDevice); | ||
| 124 | + aclrtFreeHost(xHost); | ||
| 125 | + aclrtFree(zDevice); | ||
| 126 | + aclrtFreeHost(zHost); | ||
| 127 | + aclrtDestroyStream(stream); | ||
| 128 | + aclrtResetDevice(deviceId); | ||
| 129 | + aclFinalize(); | ||
| 130 | + return 0; | ||
| 131 | +} | ||
Rexamples/01_simd_cpp_api/01_utilities/02_dump/CMakeLists.txt→examples/01_simd_cpp_api/01_utilities/02_dump/simple_dump/CMakeLists.txt+0-0
文件重命名但无更改。
| @@ -0,0 +1,113 @@ | |||
| 1 | +# Dump接口功能说明 | ||
| 2 | + | ||
| 3 | +## 概述 | ||
| 4 | + | ||
| 5 | +本样例介绍asc_dump_xxx系列接口使用方法，通过该接口实现核函数中不同物理位置的张量数据可视化。 | ||
| 6 | + | ||
| 7 | +## 支持的产品 | ||
| 8 | + | ||
| 9 | +- Ascend 950PR/Ascend 950DT | ||
| 10 | +- Atlas A3训练系列产品/Atlas A3推理系列产品 | ||
| 11 | +- Atlas A2训练系列产品/Atlas A2推理系列产品 | ||
| 12 | + | ||
| 13 | +## 目录结构介绍 | ||
| 14 | + | ||
| 15 | +``` | ||
| 16 | +├── 02_dump | ||
| 17 | +│ ├── scripts | ||
| 18 | +│ │ ├── gen_data.py // 输入数据和真值数据生成脚本 | ||
| 19 | +│ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 | ||
| 20 | +│ ├── CMakeLists.txt // 编译工程文件 | ||
| 21 | +│ ├── data_utils.h // 数据读入写出函数 | ||
| 22 | +│ └── dump.asc // Ascend C样例实现 & 调用样例 | ||
| 23 | +``` | ||
| 24 | + | ||
| 25 | +## 样例描述 | ||
| 26 | + | ||
| 27 | +- 样例功能： | ||
| 28 | + | ||
| 29 | + 使用静态Tensor编程模式实现矩阵乘法，展示asc_dump_xxx系列接口的基本使用方法。 | ||
| 30 | + 该系列接口兼容AscendC::DumpTensor接口。但后续开发中，建议优先使用asc_dump系列接口；若需Dump指定偏移位置的数据，由于asc_dump系列暂不支持该能力，可继续使用DumpAccChkPoint接口。 | ||
| 31 | + | ||
| 32 | + 矩阵乘法的计算公式为： | ||
| 33 | + | ||
| 34 | + ``` | ||
| 35 | + C = A * B | ||
| 36 | + ``` | ||
| 37 | + | ||
| 38 | +- 样例规格： | ||
| 39 | + | ||
| 40 | + 样例参数为：M = 256, N = 256, K = 64，shape信息如下表所示： | ||
| 41 | + <table> | ||
| 42 | + <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center">Matmul</td></tr> | ||
| 43 | + </tr> | ||
| 44 | + <tr><td rowspan="3" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr> | ||
| 45 | + <tr><td align="center">a</td><td align="center">[M, K]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 46 | + <tr><td align="center">b</td><td align="center">[K, N]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 47 | + </tr> | ||
| 48 | + </tr> | ||
| 49 | + <tr><td rowspan="1" align="center">样例输出</td><td align="center">c</td><td align="center">[M, N]</td><td align="center">half</td><td align="center">ND</td></tr> | ||
| 50 | + </tr> | ||
| 51 | + <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">mmad_custom</td></tr> | ||
| 52 | + </table> | ||
| 53 | + | ||
| 54 | +- asc_dump接口用法： | ||
| 55 | + | ||
| 56 | + | 接口 | 内存层级 | 地址类型 | 示例 | | ||
| 57 | + |------|----------|----------|------| | ||
| 58 | + | `asc_dump_gm` | GM (Global Memory) | `__gm__` | `asc_dump_gm<half>((__gm__ half*)addr, id, size);` | | ||
| 59 | + | `asc_dump_l1buf` | L1 Buffer | `__cbuf__` | `asc_dump_l1buf<half>((__cbuf__ half*)addr, id, size);` | | ||
| 60 | + | `asc_dump_cbuf` | L0C Buffer | `__cc__` | `asc_dump_cbuf<float>((__cc__ float*)addr, id, size);` | | ||
| 61 | + | `asc_dump_ubuf` | UB | `__ubuf__` | `asc_dump_ubuf<half>((__ubuf__ half*)addr, id, size);` | | ||
| 62 | + | ||
| 63 | + 参数说明： | ||
| 64 | + - 模板参数：指定数据类型（half/float等） | ||
| 65 | + - id：自定义标识符，用于区分不同dump数据 | ||
| 66 | + - dumpSize：dump数据元素个数 | ||
| 67 | + | ||
| 68 | + 注意：`asc_dump_l1buf`接口仅支持Atlas A2训练系列产品/Atlas A2推理系列产品/Atlas A3训练系列产品/Atlas A3推理系列产品。 | ||
| 69 | + | ||
| 70 | +- 调用实现 | ||
| 71 | + | ||
| 72 | + 使用内核调用符<<<>>>调用核函数。 | ||
| 73 | + | ||
| 74 | +## 编译运行 | ||
| 75 | + | ||
| 76 | +在本样例根目录下执行如下步骤，编译并执行样例。 | ||
| 77 | +- 配置环境变量 | ||
| 78 | + 请根据当前环境上CANN开发套件包的[安装方式](https://gitcode.com/cann/asc-devkit/blob/master/docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。 | ||
| 79 | + - 默认路径，root用户安装CANN软件包 | ||
| 80 | + ```bash | ||
| 81 | + source /usr/local/Ascend/cann/set_env.sh | ||
| 82 | + ``` | ||
| 83 | + | ||
| 84 | + - 默认路径，非root用户安装CANN软件包 | ||
| 85 | + ```bash | ||
| 86 | + source $HOME/Ascend/cann/set_env.sh | ||
| 87 | + ``` | ||
| 88 | + | ||
| 89 | + - 指定路径install_path，安装CANN软件包 | ||
| 90 | + ```bash | ||
| 91 | + source ${install_path}/cann/set_env.sh | ||
| 92 | + ``` | ||
| 93 | + | ||
| 94 | +- 样例执行 | ||
| 95 | + ```bash | ||
| 96 | + mkdir -p build && cd build; # 创建并进入build目录 | ||
| 97 | + cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程 | ||
| 98 | + python3 ../scripts/gen_data.py # 生成测试输入数据 | ||
| 99 | + ./demo # 执行编译生成的可执行程序，执行样例 | ||
| 100 | + python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确 | ||
| 101 | + ``` | ||
| 102 | + | ||
| 103 | +- 编译选项说明 | ||
| 104 | + | ||
| 105 | +| 选项 | 可选值 | 说明 | | ||
| 106 | +|------|--------|------| | ||
| 107 | +| `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU架构：dav-2201对应Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品，dav-3510对应Ascend 950PR/Ascend 950DT | | ||
| 108 | + | ||
| 109 | +- 执行结果 | ||
| 110 | + 最终执行结果如下，说明精度对比成功。 | ||
| 111 | + ```bash | ||
| 112 | + test pass! | ||
| 113 | + ``` | ||
Rexamples/01_simd_cpp_api/01_utilities/02_dump/data_utils.h→examples/01_simd_cpp_api/01_utilities/02_dump/simple_dump/data_utils.h+0-0
文件重命名但无更改。
Rexamples/01_simd_cpp_api/01_utilities/02_dump/dump.asc→examples/01_simd_cpp_api/01_utilities/02_dump/simple_dump/dump.asc+0-0
文件重命名但无更改。
Rexamples/01_simd_cpp_api/01_utilities/02_dump/scripts/gen_data.py→examples/01_simd_cpp_api/01_utilities/02_dump/simple_dump/scripts/gen_data.py+0-0
文件重命名但无更改。
Rexamples/01_simd_cpp_api/01_utilities/02_dump/scripts/verify_result.py→examples/01_simd_cpp_api/01_utilities/02_dump/simple_dump/scripts/verify_result.py+0-0
文件重命名但无更改。
| @@ -26,35 +26,29 @@ | |||
| 26 | 26 | ||
| 27 | ### 调用实现 | 27 | ### 调用实现 |
| 28 | 28 | ||
| 29 | -调用aclrtc接口系列在运行时编译并执行核函数，完整链路分为两大阶段： | 29 | +调用aclrtc接口系列在运行时编译并执行核函数，完整链路如下： |
| 30 | - | ||
| 31 | -#### 编译 | ||
| 32 | - | ||
| 33 | -通过aclrtc接口将核函数源码字符串编译为deviceELF二进制。 | ||
| 34 | 30 | ||
| 31 | +#### 编译阶段 | ||
| 35 | 1. `aclrtcCreateProg` — 创建编译程序对象，传入核函数源码字符串 | 32 | 1. `aclrtcCreateProg` — 创建编译程序对象，传入核函数源码字符串 |
| 36 | -2. `aclrtcCompileProg` — 执行运行时编译，通过options传入`--npu-arch`指定NPU架构 | 33 | +2. `aclrtcAddNameExpr` — 注册需要导出的核函数名（含模板参数，如`Kernel::add_custom<float>`） |
| 37 | -3. `aclrtcGetCompileLogSize`/`aclrtcGetCompileLog` — 编译失败时获取编译错误日志，用于问题定位 | 34 | +3. `aclrtcCompileProg` — 执行运行时编译，通过options传入`--npu-arch`指定NPU架构 |
| 38 | -4. `aclrtcGetBinDataSize`/`aclrtcGetBinData` — 编译成功后获取编译产物的二进制大小及数据（deviceELF） | 35 | +4. `aclrtcGetBinDataSize`/`aclrtcGetBinData` — 获取编译产物的二进制大小及数据（deviceELF） |
| 36 | +5. `aclrtcGetLoweredName` — 获取核函数编译后的mangledname，用于后续查找 | ||
| 39 | 37 | ||
| 40 | -#### 加载与执行 | 38 | +#### 加载阶段 |
| 39 | +1. `aclrtBinaryLoadFromData` — 将编译产物的二进制加载到设备（通过`ACL_RT_BINARY_MAGIC_ELF_AICORE`标记为AICore可执行） | ||
| 40 | +2. `aclrtBinaryGetFunction` — 从加载的二进制中获取核函数句柄（`funcHandle`） | ||
| 41 | 41 | ||
| 42 | -通过aclrt运行时接口加载deviceELF并启动核函数计算，执行完毕后销毁编译程序对象。 | 42 | +#### 参数配置阶段 |
| 43 | +1. `aclrtKernelArgsInit` — 初始化核函数参数句柄 | ||
| 44 | +2. `aclrtKernelArgsAppend` — 逐个追加参数（核函数为`__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z`，对应传入三个Device内存指针） | ||
| 45 | +3. `aclrtKernelArgsFinalize` — 完成参数配置 | ||
| 43 | 46 | ||
| 44 | -- **二进制加载** | 47 | +#### 执行阶段 |
| 45 | - 1. `aclrtBinaryLoadFromData` — 将编译产物的二进制加载到设备（通过`ACL_RT_BINARY_MAGIC_ELF_AICORE`标记为AICore可执行） | 48 | +1. `aclrtLaunchKernelWithConfig` — 启动核函数，指定block数量、stream等执行配置 |
| 46 | - 2. `aclrtBinaryGetFunction` — 从加载的二进制中获取核函数句柄（`funcHandle`） | ||
| 47 | 49 | ||
| 48 | -- **参数配置** | 50 | +#### 资源清理 |
| 49 | - 1. `aclrtKernelArgsInit` — 初始化核函数参数句柄 | 51 | +1. `aclrtcDestroyProg` — 销毁编译程序对象 |
| 50 | - 2. `aclrtKernelArgsAppend` — 逐个追加参数（核函数为`__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z`，对应传入三个Device内存指针） | ||
| 51 | - 3. `aclrtKernelArgsFinalize` — 完成参数配置 | ||
| 52 | - | ||
| 53 | -- **核函数启动** | ||
| 54 | - 1. `aclrtLaunchKernelWithConfig` — 启动核函数，指定block数量、stream等执行配置 | ||
| 55 | - | ||
| 56 | -- **资源清理** | ||
| 57 | - 1. `aclrtcDestroyProg` — 销毁编译程序对象 | ||
| 58 | 52 | ||
| 59 | 数据生成与精度校验均在Host侧C++内完成，不依赖外部脚本。 | 53 | 数据生成与精度校验均在Host侧C++内完成，不依赖外部脚本。 |
| 60 | 54 | ||
| @@ -42,21 +42,35 @@ | |||
| 42 | 42 | ||
| 43 | ### 样例实现 | 43 | ### 样例实现 |
| 44 | 44 | ||
| 45 | -#### Kernel实现 | 45 | +### Kernel实现 |
| 46 | 46 | ||
| 47 | 核函数源码以字符串形式嵌入Host代码中，通过LocalMemAllocator接口完成内存管理，通过SetFlag/WaitFlag接口完成事件同步，实现数据搬运、计算和结果搬出的完整流程。 | 47 | 核函数源码以字符串形式嵌入Host代码中，通过LocalMemAllocator接口完成内存管理，通过SetFlag/WaitFlag接口完成事件同步，实现数据搬运、计算和结果搬出的完整流程。 |
| 48 | 48 | ||
| 49 | -#### 调用实现 | 49 | +### 调用实现 |
| 50 | 50 | ||
| 51 | -调用aclrtc接口系列在运行时编译并执行核函数，完整链路分为两大阶段： | 51 | +调用aclrtc接口系列在运行时编译并执行核函数，完整链路如下： |
| 52 | - | ||
| 53 | -##### 编译 | ||
| 54 | - | ||
| 55 | -通过aclrtc接口将核函数源码字符串编译为deviceELF二进制。 | ||
| 56 | 52 | ||
| 53 | +#### 编译阶段 | ||
| 57 | 1. `aclrtcCreateProg` — 创建编译程序对象，传入核函数源码字符串 | 54 | 1. `aclrtcCreateProg` — 创建编译程序对象，传入核函数源码字符串 |
| 58 | -2. `aclrtcCompileProg` — 执行运行时编译，通过options传入`--npu-arch`指定NPU架构 | 55 | +2. `aclrtcAddNameExpr` — 注册需要导出的核函数名（含模板参数，如`Kernel::add_custom<float>`） |
| 59 | -3. `aclrtcGetBinDataSize`/`aclrtcGetBinData` — 获取编译产物的二进制大小及数据（deviceELF） | 56 | +3. `aclrtcCompileProg` — 执行运行时编译，通过options传入`--npu-arch`指定NPU架构 |
| 57 | +4. `aclrtcGetBinDataSize`/`aclrtcGetBinData` — 获取编译产物的二进制大小及数据（deviceELF） | ||
| 58 | +5. `aclrtcGetLoweredName` — 获取核函数编译后的mangledname，用于后续查找 | ||
| 59 | + | ||
| 60 | +#### 加载阶段 | ||
| 61 | +1. `aclrtBinaryLoadFromData` — 将编译产物的二进制加载到设备（通过`ACL_RT_BINARY_MAGIC_ELF_AICORE`标记为AICore可执行） | ||
| 62 | +2. `aclrtBinaryGetFunction` — 从加载的二进制中获取核函数句柄（`funcHandle`） | ||
| 63 | + | ||
| 64 | +#### 参数配置阶段 | ||
| 65 | +1. `aclrtKernelArgsInit` — 初始化核函数参数句柄 | ||
| 66 | +2. `aclrtKernelArgsAppend` — 逐个追加参数（核函数为`GM_ADDR x, GM_ADDR y, GM_ADDR z`，对应传入三个Device内存指针） | ||
| 67 | +3. `aclrtKernelArgsFinalize` — 完成参数配置 | ||
| 68 | + | ||
| 69 | +#### 执行阶段 | ||
| 70 | +1. `aclrtLaunchKernelWithConfig` — 启动核函数，指定block数量、stream等执行配置 | ||
| 71 | + | ||
| 72 | +#### 资源清理 | ||
| 73 | +1. `aclrtcDestroyProg` — 销毁编译程序对象 | ||
| 60 | 74 | ||
| 61 | 对于模板核函数（如本样例），还需： | 75 | 对于模板核函数（如本样例），还需： |
| 62 | 76 | ||
目录readme需要加英文版本