Star774
1077
代码
代码
Issues57
Pull Requests211
Actions
项目讨论
Wiki
分析
Star774
1077
已合并
add simt example #897
wulinyu创建于 3月14日
add simt example #897
已合并
wulinyu创建于 3月14日
从wulinyu/asc-devkitnewest合入到cann/asc-devkitmaster
共 28 个文件变更+76-1017
MREADME.md+2-2
@@ -50,7 +50,7 @@ Ascend C提供三类接口，均可实现底层的完备编程能力：
50 50 
51<img src="docs/figures/architecture.png" alt="架构图" width="850px" height="580px">51<img src="docs/figures/architecture.png" alt="架构图" width="850px" height="580px">
52 52 
53-- **语言扩展层C API**：纯C接口，支持数组分配内存、基于指针的计算接口，提供与业界一致的C语言编程体验，并开放芯片完备编程能力。Atlas A2/A3支持SIMD的纯C接口；Ascend 950PR/Ascend 950DT将支持与业界类似的纯SIMT编程能力、SIMD/SIMT混合编程能力；53+- **语言扩展层C API**：纯C接口，支持数组分配内存、基于指针的计算接口，提供与业界一致的C语言编程体验，并开放芯片完备编程能力。Atlas A2/A3支持SIMD的纯C接口；Ascend 950PR/Ascend 950DT将支持与业界类似的SIMT编程能力、SIMD/SIMT混合编程能力；
54- **基础API**：单指令抽象的C++类库API，一般基于Tensor编程；逐步基于Layout完善Tensor编程能力；54- **基础API**：单指令抽象的C++类库API，一般基于Tensor编程；逐步基于Layout完善Tensor编程能力；
55- **高阶API**：基于单核对常见算法进行抽象和封装，提供公共算法的实现；55- **高阶API**：基于单核对常见算法进行抽象和封装，提供公共算法的实现；
56- **算子模板库**：基于模板提供算子的完整实现参考，简化Tiling开发，支持用户自定义扩展；56- **算子模板库**：基于模板提供算子的完整实现参考，简化Tiling开发，支持用户自定义扩展；
@@ -172,7 +172,7 @@ Ascend C提供三类接口，均可实现底层的完备编程能力：
172## 📌后续规划172## 📌后续规划
173 173 
174- 基于Altas A2/A3发布语言扩展层纯C接口，提供基于数组分配内存能力，支持基于指针的计算接口，实现与业界类似的纯C编程体验；174- 基于Altas A2/A3发布语言扩展层纯C接口，提供基于数组分配内存能力，支持基于指针的计算接口，实现与业界类似的纯C编程体验；
175-- Ascend 950PR/Ascend 950DT将支持纯SIMT编程、SIMD与SIMT混合编程，并通过Layout进一步强化Tensor编程能力；175+- Ascend 950PR/Ascend 950DT将支持SIMT编程、SIMD与SIMT混合编程，并通过Layout进一步强化Tensor编程能力；
176- 持续丰富语言扩展层C API(含SIMD、SIMT)和基础API的关键特性介绍，并基于融合编译与 <<<>>>调用完善样例；176- 持续丰富语言扩展层C API(含SIMD、SIMT)和基础API的关键特性介绍，并基于融合编译与 <<<>>>调用完善样例；
177 177 
178## 📝相关信息178## 📝相关信息
Mdocs/api/context/核函数定义.md+1-1
@@ -1,7 +1,7 @@
1# 核函数定义<a name="ZH-CN_TOPIC_0000002484718282"></a>1# 核函数定义<a name="ZH-CN_TOPIC_0000002484718282"></a>
2 2 
3> [!NOTE]说明 3> [!NOTE]说明
4->当前仅展示SIMD与SIMT混合编程场景的核函数定义，纯SIMT场景的核函数定义将在后续版本中发布。4+>当前仅展示SIMD与SIMT混合编程场景的核函数定义，SIMT场景的核函数定义将在后续版本中发布。
5 5 
6- **[asc\_vf\_call](asc_vf_call.md)** 6- **[asc\_vf\_call](asc_vf_call.md)**
7 7 
Mdocs/api/context/简介-140.md+1-1
@@ -1,7 +1,7 @@
1# 简介<a name="ZH-CN_TOPIC_0000002066575561"></a>1# 简介<a name="ZH-CN_TOPIC_0000002066575561"></a>
2 2 
3> [!NOTE]说明 3> [!NOTE]说明
4->**SIMT API当前仅支持SIMD与SIMT混合编程，纯SIMT编程场景待正式发布支持，请关注后续版本**。4+>**SIMT API当前仅支持SIMD与SIMT混合编程，SIMT编程场景待正式发布支持，请关注后续版本**。
5 5 
6SIMT API是面向AI处理器的并行计算编程接口，基于SIMT编程模型实现高效的数据并行计算，具体API分类请参考[表1](#table1452018407354)。在SIMT编程中，每32个线程被分为一组，这一组线程称为Warp。Warp level接口负责处理单个Warp内32个线程的数据操作。Warp中的每个线程称为Lane，每个线程在Warp内的编号为LaneId，LaneId的范围是\[0, 31\]。6SIMT API是面向AI处理器的并行计算编程接口，基于SIMT编程模型实现高效的数据并行计算，具体API分类请参考[表1](#table1452018407354)。在SIMT编程中，每32个线程被分为一组，这一组线程称为Warp。Warp level接口负责处理单个Warp内32个线程的数据操作。Warp中的每个线程称为Lane，每个线程在Warp内的编号为LaneId，LaneId的范围是\[0, 31\]。
7 7 
Dexamples/00_introduction/01_add/README.md+0-17
@@ -1,17 +0,0 @@
1-# Add样例介绍
2- 
3-## 概述
4- 
5-基于Ascend C的Add算子的<<<>>>直调方法，支持main函数和kernel函数在同一个cpp文件中实现。
6- 
7-## 算子开发样例
8- 
9-| 目录名称 | 功能描述 |
10-| ------------------------------------------------------------ | ---------------------------------------------------- |
11-| [basic_api_memory_allocator_add](./basic_api_memory_allocator_add) | 样例基于静态Tensor方式编程实现Add样例，展示了LocalMemAllocator进行线性内存分配 |
12-| [basic_api_tque_add](./basic_api_tque_add) | 样例以Add算子为例，采用TQue内存管理机制实现数据搬运与计算任务的协同调度 |
13-| [c_api_async_add](./c_api_async_add) | 本样例采用C_API接口编写Add算子样例，基于异步搬运、计算接口实现 |
14-| [c_api_delicacy_async_add](./c_api_delicacy_async_add) | 本样例采用C_API接口编写Add算子样例，基于异步搬运、计算接口和手动添加的同步指令实现 |
15-| [c_api_sync_add](./c_api_sync_add) | 本样例采用C_API接口编写Add算子样例，基于同步搬运、计算接口实现 |
16-| [reg_compute_add](./reg_compute_add) | 样例基于微指令API实现Add样例，展示了通过微指令API直接对芯片中涉及Vector计算的寄存器进行操作 |
17-| [simt_add](./simt_add) | 样例基于纯SIMT编程方式实现Add样例，展示了SIMT单指令多线程的编程方式完成加法计算 |
Dexamples/00_introduction/04_simple_operator/README.md+0-17
@@ -1,17 +0,0 @@
1-# SimpleOperator样例介绍
2- 
3-## 概述
4- 
5-样例介绍了5个基于Ascend C的算子的核函数直调样例，涵盖AddN、Broadcast、Gather、Sub以及向量Add等典型算子，展示了动态Tensor、纯SIMT编程、临时缓冲区使用等关键技术，充分体现了Ascend C在高性能算子开发中的灵活性与高效性。
6- 
7-## 算子开发样例
8- 
9-| 目录名称 | 功能描述 |
10-| ------------------------------------------------------------ | ---------------------------------------------------- |
11-| [add_broadcast](./add_broadcast) | 本样例介绍Add算子的核函数直调方法，多核&tiling场景下增加输入Broadcast |
12-| [add_dynamic](./add_dynamic) | 本样例演示基于动态Tensor编程模型的AddN算子实现，该实现采用ListTensorDesc结构处理多输入参数，结合TQue内存管理机制实现数据搬运与计算任务的协同调度 |
13-| [broadcast](./broadcast) | 本样例展示了一个支持多种数据类型（如bfloat，int8，float，half等）和多种形状（如(32, 1024)，(8, 1023)等）的输入张量执行逐元素加法 |
14-| [pure_simt_gather](./pure_simt_gather) | 样例基于Ascend C纯SIMT编程方式实现Gather算子，从输入张量中采集指定的m行数据，展示离散内存访问类算子的开发方法 |
15-| [sub](./sub) | 本样例演示了如何通过自定义核函数实现高性能的逐元素减法（Sub）运算。算子核心功能是完成两个形状相同的输入张量x与y的逐元素相减 |
16-| [tmp_buffer](./tmp_buffer) | 本样例展示了一个支持bfloat16_t数据类型的向量加法（Add）算子，并重点演示了在算子计算过程中使用临时缓冲区（TmpBuf）进行数据转换的典型方法 |
17-| [vector_add](./vector_add) | 本样例介绍Add算子的核函数直调方法，算子支持单核运行 |
Dexamples/00_introduction/README.md+0-16
@@ -1,16 +0,0 @@
1-# Introduce样例介绍
2- 
3-## 概述
4- 
5-基于Ascend C的简单的示例，通过Ascend C编程语言实现了自定义算子，分别给出对应的<<<>>>直调实现，适合初学者
6- 
7-## 算子开发样例
8- 
9-| 目录名称 | 功能描述 |
10-| ------------------------------------------------------------ | ---------------------------------------------------- |
11-| [00_helloworld](./00_helloworld) | 样例介绍了基于Ascend C的HelloWorld算子的核函数直调方法，分别从NPU、AICPU测运行核验证算子核函数，展示核函数从调用到执行的整体流程 |
12-| [01_add](./01_add) | 本样例介绍了基于Ascend C的Add自定义Vector算子的核函数直调方法，实现两个输入张量的逐元素相加，支持main函数和kernel函数在同一个cpp文件中实现 |
13-| [02_matmul](./02_matmul) | 本样例介绍了基于Ascend C的Matmul算子的核函数直调方法，可最大化利用AI处理器的并行计算能力，显著提升算子的执行效率，使用与高性能推理与训练场景 |
14-| [03_matmulleakyrelu](./03_matmulleakyrelu) | 本样例介绍了基于Ascend C的MatmulLeakyRelu自定义算子的核函数直调方法，能够完成矩阵乘加与LeakyReLU激活的融合计算，该方式将关键计算步骤在硬件层面高效协同执行，显著降低内存访问开销与计算延时 |
15-| [04_simple_operator](./04_simple_operator) | 样例介绍了5个基于Ascend C的算子的核函数直调样例，涵盖AddN、Broadcast、Gather、Sub以及向量Add等典型算子，展示了动态Tensor、纯SIMT编程、临时缓冲区使用等关键技术，充分体现了Ascend C在高性能算子开发中的灵活性与高效性 |
16-| [04_reg_compute](./04_reg_compute) | 本样例介绍了基于Ascend C的Add算子核函数直调方法（RegBase场景），通过C_API实现两个输入张量的逐元素相加，展示了片上存储和寄存器层级的向量计算流程。|
Mexamples/01_simd_cpp_api/00_introduction/03_matrix_vector_fusion/simt_gather_and_simd_adds/README.md+1-1
@@ -1,4 +1,4 @@
1-# SIMT与SIMD新同构编程实现gather & adds算子样例1+# SIMT与SIMD混合编程实现gather & adds算子样例
2 2 
3## 概述3## 概述
4本样例以实现gather和adds功能的算子为例，展示了SIMT和SIMD混合编程的算子开发方式。算子中使用SIMT编程方式实现离散内存访问操作gather，使用SIMD编程方式实现连续内存访问操作adds。4本样例以实现gather和adds功能的算子为例，展示了SIMT和SIMD混合编程的算子开发方式。算子中使用SIMT编程方式实现离散内存访问操作gather，使用SIMD编程方式实现连续内存访问操作adds。
Dexamples/02_features/03_simt/simt_gather/CMakeLists.txt+0-53
@@ -1,53 +0,0 @@
1-# ----------------------------------------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------------------------------------
11- 
12-cmake_minimum_required(VERSION 3.16)
13- 
14-project(kernel_samples)
15- 
16-set(ASC_SRC_FILE gather.asc)
17- 
18-if("${RUN_MODE}" STREQUAL "cpu")
19- find_package(tikicpulib REQUIRED)
20- set_source_files_properties(${ASC_SRC_FILE} PROPERTIES LANGUAGE CXX)
21-else()
22- find_package(ASC REQUIRED)
23- enable_language(ASC CXX)
24-endif()
25- 
26-add_executable(demo
27- ${ASC_SRC_FILE}
28-)
29- 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36-target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>
38- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},sim>:-g -O0>>
39- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},cpu>:-x c++>>
40- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},cpu>:-g -O2 -std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0 -Wall -Werror>>
41-)
42- 
43-target_link_directories(demo PRIVATE
44- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},sim>:$ENV{ASCEND_HOME_PATH}/tools/simulator/${SOC_VERSION}/lib>>
45-)
46- 
47-target_link_libraries(demo PRIVATE
48- m
49- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},sim>:runtime_camodel>>
50- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},sim>:npu_drv>>
51- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},cpu>:cpudebug_acl_stub>>
52- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},cpu>:tikicpulib::ascend950pr_9599>>
53-)
Dexamples/02_features/03_simt/simt_gather/README.md+0-185
@@ -1,185 +0,0 @@
1-# 新同构编程实现gather算子样例
2- 
3-## 概述
4-样例基于Ascend新同构编程模型SIMT方式实现gather算子，从100000×128的输入张量中采集12288行数据，展示离散内存访问类算子的开发方法。
5- 
6-## 支持的产品
7-- Ascend 950PR/Ascend 950DT
8- 
9-## 目录结构
10-```
11-├── simt_gather
12-│ ├── CMakeLists.txt # cmake编译文件
13-│ ├── gather.asc # Ascend C算子实现gather调用样例
14-| └── README.md
15-```
16- 
17-## 算子描述
18- 
19-- 算子功能：
20- gather算子实现了从形状为100000 * 128的二维向量中获取指定索引的12288行数据的功能。算子输出output第i行数据计算公式为：
21-
22- ```
23- output[i] = input[index[i]]
24- ```
25- 
26-- 算子规格：
27- <table>
28- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center">gather</td></tr>
29- </tr>
30- <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>
31- <tr><td align="center">input</td><td align="center">100000 * 128</td><td align="center">float</td><td align="center">ND</td></tr>
32- <tr><td align="center">index</td><td align="center">12288</td><td align="center">uint32_t</td><td align="center">ND</td></tr>
33- </tr>
34- </tr>
35- <tr><td rowspan="1" align="center">算子输出</td><td align="center">output</td><td align="center">12288 * 128</td><td align="center">float</td><td align="center">ND</td></tr>
36- </tr>
37- <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">gather_kernel</td></tr>
38- </table>
39- 
40-- 数据切分：
41- * 核数：48核
42- * 每核线程数：256线程
43- * 单线程处理：1行（128列）
44- * 总处理能力：48×256=12288行（覆盖索引长度）
45- 
46-- 算子实现：
47- gather算子的实现流程只有1个步骤：simt_gather。
48-
49- simt_gather负责从输入input（Global Memory）中获取指定索引的数据。基于上述数据切分，首先计算线程应处理数据的索引，然后通过赋值操作将一行数据存储到Global Memory上。
50- ```
51- int32_t out_row = blockIdx.x * blockDim.x + threadIdx.x;
52- ...
53- uint32_t in_row = index[out_row];
54- ...
55- gather_output[row[out_row]] = input[row[in_row]];
56- ```
57- 
58-- 调用实现：
59- - CPU调测模式使用ICPU_RUN_KF CPU调测宏调用核函数。
60- - NPU模式使用内核调用符<<<>>>调用核函数。
61- 
62- 应用程序通过ASCENDC_CPU_DEBUG宏区分代码逻辑运行于CPU模式还是NPU模式。
63- 
64-## 编译运行
65-在本样例根目录下执行如下步骤，编译并执行算子。
66-- 配置环境变量
67- 请根据当前环境上CANN开发套件包的[安装方式](../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
68- - 默认路径，root用户安装CANN软件包
69- ```bash
70- source /usr/local/Ascend/cann/set_env.sh
71- ```
72- 
73- - 默认路径，非root用户安装CANN软件包
74- ```bash
75- source $HOME/Ascend/cann/set_env.sh
76- ```
77- 
78- - 指定路径install_path，安装CANN软件包
79- ```bash
80- source ${install_path}/cann/set_env.sh
81- ```
82-
83-- 样例执行
84- ```bash
85- mkdir -p build && cd build; # 创建并进入build目录
86- cmake ..; make -j; # 编译工程
87- ./demo # 执行样例
88- ```
89- 执行结果如下，说明精度对比成功。
90- ```
91- [Success] Case accuracy is verification passed.
92- ```
93- 
94-## CPU Debug调测
95-CPU Debug功能支持对CPU执行过程中的运行状态进行调试，主要通过GDB工具实现。GDB调试支持设置断点、查看寄存器和内存状态、单步执行、查看调用栈等常用调试操作，并支持多线程程序的调试。
96- 
97-- 样例执行
98- ```bash
99- mkdir -p build && cd build; # 创建并进入build目录
100- cmake -DRUN_MODE=cpu ..; make -j; # 编译工程
101- ./demo # 执行样例
102- ```
103- 执行结果如下，说明精度对比成功。
104- ```
105- [Success] Case accuracy is verification passed.
106- ```
107- 
108-- 进入GDB模式调试
109- 在上述指令中"./demo"前加入"gdb --args"，再次执行指令即可进入GDB模式。调试时需要选择子进程进行调试。
110- ```bash
111- mkdir -p build && cd build;
112- cmake -DRUN_MODE=cpu ..; make -j;
113- gdb --args ./demo
114- ```
115- 
116-## 性能调优
117-算子调优工具支持上板调优和仿真调优两种模式，可分别获取算子在实际硬件/仿真环境下的性能数据，用于定位性能瓶颈、优化算子实现。
118-### 上板调优
119-基于编译生成的可执行文件，直接在NPU硬件上采集算子性能数据，数据为算子预热后的真实运行指标。
120- 
121-**操作步骤**
122- 
123-**1.执行调优命令**
124- 
125-基于编译得到的demo文件，运行算子调优工具。
126-```bash
127-msprof op ./demo
128-```
129-**2.查看性能数据**
130- 
131-在当前目录下会生成OPPPROF_为前缀的文件夹，目录结构及文件说明如下：
132-```bash
133-OPPROF_xxxxxxxx_XXXXXX
134-├── dump # 原始性能数据（无需关注）
135-├── OpBasicInfo.csv # 算子基础数据
136-├── ArithmeticUtilization.csv # cube及vector类型指令的cycle占比数据
137-├── ResourceConflictRatio.csv # 资源冲突占比数据
138-├── ... (开启的aic-metrics)
139-└── visualize_data.bin # 算子可视化文件（可通过MindStudio Insight加载，直观查看算子性能）
140-```
141-可直接查看.csv文件获取算子block级性能数据或通过MindStudio Insight工具打开`visualize_data.bin`文件可视化查看性能数据。
142- 
143-### 仿真调优
144-在无NPU硬件环境的场景下，通过编译仿真版算子可执行文件，结合模拟器获取算子仿真性能数据，适用于开发阶段的快速调优。
145- 
146-**操作步骤**
147- 
148-**1.仿真算子编译**
149- 
150-```bash
151-# 创建并进入build目录
152-mkdir -p build && cd build
153-# 替换${SOC_VERSION}为实际NPU型号，可通过npu-smi info命令进行查询，如Ascend910_957c。
154-cmake -DRUN_MODE=sim -DSOC_VERSION=${SOC_VERSION} ..
155-# 多线程编译
156-make -j$(nproc)
157-```
158-**2.配置运行时依赖**
159- 
160-添加运行时依赖库路径（需替换{SOC_VERSION}为实际NPU型号）：
161-```bash
162-export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/tools/simulator/${SOC_VERSION}/lib/:$LD_LIBRARY_PATH
163-```
164-**3.执行仿真调优命令**
165- 
166-```bash
167-msprof op simulator ./demo
168-```
169-**4.查看仿真性能数据**
170- 
171-当前目录下会生成OPPROF_前缀的文件夹，目录结构如下：
172-```bash
173-OPPROF_xxxxxxxx_XXXXXX
174-├── dump # 原始性能数据，无需关注
175-└── simulation # 仿真性能数据分析结果
176- ├── core0.veccore0 # 算子block级子核
177- ├── core0.veccore1_code_exe.csv # 代码行耗时
178- ├── core0.veccore1_instr_exe.csv # 程序代码指令详细信息
179- └── trace.json # 算子block级子核流水图
180- ├── ...
181- ├── visualize_data.bin # 算子可视化文件（可通过MindStudio Insight加载，直观查看算子性能）
182- └── trace.json # 算子所有核的流水图
183-```
184-**补充说明**
185-更多性能指标的详细说明及调优方案，可参考《算子开发工具》手册。
Dexamples/02_features/03_simt/simt_gather/gather.asc+0-190
@@ -1,190 +0,0 @@
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 gather.asc
14- * \brief
15- */
16- 
17-#include <chrono>
18-#include <iostream>
19-#include <iterator>
20-#include <random>
21-#include <vector>
22-#include "acl/acl.h"
23-#include "kernel_operator.h"
24-#include "simt_api/asc_simt.h"
25-#if defined(ASCENDC_CPU_DEBUG)
26-#include "tikicpulib.h"
27-#endif
28- 
29-namespace {
30- constexpr uint32_t THREAD_COUNT = 256;
31-}
32- 
33-__simt_vf__ __launch_bounds__(THREAD_COUNT) inline void simt_gather(
34- __gm__ float* input,
35- __gm__ uint32_t* index,
36- __gm__ float* gather_output,
37- uint32_t in_shape0,
38- uint32_t in_shape1,
39- uint32_t index_total_length)
40-{
41- // Calculate global thread ID -> maps to the row index of output tensor
42- uint32_t out_row = blockIdx.x * blockDim.x + threadIdx.x;
43- 
44- // Calculate global thread ID -> maps to the row index of output tensor
45- if (out_row >= index_total_length) {
46- return;
47- }
48- 
49- // Single thread processes entire row (all columns) - enables coalesced memory access
50- uint32_t in_row = index[out_row];
51- for (uint32_t col = 0; col < in_shape1; col++) {
52- uint32_t input_idx = in_row * in_shape1 + col;
53- uint32_t output_idx = out_row * in_shape1 + col;
54- gather_output[output_idx] = input[input_idx];
55- }
56-}
57- 
58-__global__ __vector__ void gather_kernel(__gm__ float* input, __gm__ uint32_t* index, __gm__ float* output,
59- uint32_t in_shape0, uint32_t in_shape1, uint32_t index_total_length)
60-{
61- asc_vf_call<simt_gather>(dim3(THREAD_COUNT), input, index,
62- output,
63- in_shape0,
64- in_shape1,
65- index_total_length);
66-}
67- 
68-std::vector<float> gather(std::vector<float>& input, const uint32_t* in_shape, std::vector<uint32_t>& index)
69-{
70- uint32_t input_total_length = input.size();
71- size_t input_total_byte_size = input_total_length * sizeof(float);
72- 
73- uint32_t index_total_length = index.size();
74- size_t index_total_byte_size = index_total_length * sizeof(uint32_t);
75- 
76- uint32_t output_total_length = index.size() * in_shape[1];
77- size_t output_total_byte_size = output_total_length * sizeof(float);
78- uint32_t num_blocks = (index_total_length + THREAD_COUNT - 1) / THREAD_COUNT;
79- 
80- int32_t device_id = 0;
81- aclrtStream stream = nullptr;
82- 
83- uint8_t* input_host = reinterpret_cast<uint8_t *>(input.data());
84- uint8_t* index_host = reinterpret_cast<uint8_t *>(index.data());
85- uint8_t* output_host = nullptr;
86- float* input_device = nullptr;
87- uint32_t* index_device = nullptr;
88- float* output_device = nullptr;
89- 
90- aclInit(nullptr);
91- aclrtSetDevice(device_id);
92- aclrtCreateStream(&stream);
93- aclrtMallocHost((void **)(&output_host), output_total_byte_size);
94- aclrtMalloc((void **)&input_device, input_total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
95- aclrtMalloc((void **)&index_device, index_total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
96- aclrtMalloc((void **)&output_device, output_total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
97- 
98- aclrtMemcpy(input_device, input_total_byte_size, input_host, input_total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE);
99- aclrtMemcpy(index_device, index_total_byte_size, index_host, index_total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE);
100- 
101-#if defined(ASCENDC_CPU_DEBUG)
102- AscendC::SetKernelMode(KernelMode::AIV_MODE);
103- ICPU_RUN_KF(gather_kernel, num_blocks, input_device, index_device, output_device, in_shape[0], in_shape[1], index_total_length);
104-#else
105- uint32_t dyn_ubuf_size = 0; // No need to alloc dynamic memory.
106- gather_kernel<<<num_blocks, dyn_ubuf_size, stream>>>(input_device, index_device, output_device,
107- in_shape[0], in_shape[1], index_total_length);
108-#endif
109- 
110- aclrtSynchronizeStream(stream);
111- 
112- aclrtMemcpy(output_host, output_total_byte_size, output_device, output_total_byte_size, ACL_MEMCPY_DEVICE_TO_HOST);
113-
114- std::vector<float> output((float *)output_host, (float *)(output_host + output_total_byte_size));
115- 
116- aclrtFree(input_device);
117- aclrtFree(index_device);
118- aclrtFree(output_device);
119- aclrtFreeHost(output_host);
120- 
121- aclrtDestroyStream(stream);
122- aclrtResetDevice(device_id);
123- aclFinalize();
124- 
125- return output;
126-}
127- 
128-uint32_t verify_result(std::vector<float>& output, std::vector<float>& golden)
129-{
130- auto print_tensor = [](std::vector<float>& tensor, const char* name) {
131- constexpr size_t max_print_size = 20;
132- std::cout << name << ": ";
133- std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), max_print_size),
134- std::ostream_iterator<float>(std::cout, " "));
135- if (tensor.size() > max_print_size) {
136- std::cout << "...";
137- }
138- std::cout << std::endl;
139- };
140- print_tensor(output, "Output");
141- print_tensor(golden, "Golden");
142- if (std::equal(output.begin(), output.end(), golden.begin())) {
143- std::cout << "[Success] Case accuracy is verification passed." << std::endl;
144- return 0;
145- } else {
146- std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
147- return 1;
148- }
149- return 0;
150-}
151- 
152-float random_float(float min, float max)
153-{
154- static std::mt19937 engine(std::chrono::system_clock::now().time_since_epoch().count());
155- std::uniform_real_distribution<float> dist(min, max);
156- return dist(engine);
157-}
158- 
159-uint32_t random_uint32(uint32_t min, uint32_t max)
160-{
161- static std::mt19937 engine(std::chrono::system_clock::now().time_since_epoch().count());
162- std::uniform_int_distribution<uint32_t> dist(min, max);
163- return dist(engine);
164-}
165- 
166-int32_t main(int32_t argc, char* argv[])
167-{
168- constexpr uint32_t in_shape[2] = {100000, 128};
169- constexpr uint32_t input_total_length = in_shape[0] * in_shape[1];
170- std::vector<float> input(input_total_length);
171- for (uint32_t i = 0; i < input_total_length; i++) {
172- input[i] = random_float(0.0f, 100.0f);
173- }
174- 
175- constexpr uint32_t index_total_length = 12288;
176- std::vector<uint32_t> index(index_total_length);
177- for (uint32_t i = 0; i < index_total_length; i++) {
178- index[i] = random_uint32(0, in_shape[0] - 1);
179- }
180- 
181- std::vector<float> golden(index_total_length * in_shape[1]);
182- for (uint32_t i = 0; i < index_total_length; i++) {
183- for (uint32_t j = 0; j < in_shape[1]; j++) {
184- golden[i * in_shape[1] + j] = input[index[i] * in_shape[1] + j];
185- }
186- }
187- 
188- std::vector<float> output = gather(input, in_shape, index);
189- return verify_result(output, golden);
190-}
Dexamples/02_features/03_simt/simt_gather_and_simd_adds/CMakeLists.txt+0-53
@@ -1,53 +0,0 @@
1-# ----------------------------------------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------------------------------------
11- 
12-cmake_minimum_required(VERSION 3.16)
13- 
14-project(kernel_samples)
15- 
16-set(ASC_SRC_FILE gather_and_adds.asc)
17- 
18-if("${RUN_MODE}" STREQUAL "cpu")
19- find_package(tikicpulib REQUIRED)
20- set_source_files_properties(${ASC_SRC_FILE} PROPERTIES LANGUAGE CXX)
21-else()
22- find_package(ASC REQUIRED)
23- enable_language(ASC CXX)
24-endif()
25- 
26-add_executable(demo
27- ${ASC_SRC_FILE}
28-)
29- 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36-target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>
38- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},sim>:-g -O0>>
39- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},cpu>:-x c++>>
40- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},cpu>:-g -O2 -std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0 -Wall -Werror>>
41-)
42- 
43-target_link_directories(demo PRIVATE
44- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},sim>:$ENV{ASCEND_HOME_PATH}/tools/simulator/${SOC_VERSION}/lib>>
45-)
46- 
47-target_link_libraries(demo PRIVATE
48- m
49- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},sim>:runtime_camodel>>
50- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},sim>:npu_drv>>
51- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},cpu>:cpudebug_acl_stub>>
52- $<BUILD_INTERFACE:$<$<STREQUAL:${RUN_MODE},cpu>:tikicpulib::ascend950pr_9599>>
53-)
Dexamples/02_features/03_simt/simt_gather_and_simd_adds/README.md+0-204
@@ -1,204 +0,0 @@
1-# SIMT与SIMD新同构编程实现gather & adds算子样例
2- 
3-## 概述
4-本样例以实现gather和adds功能的算子为例，展示了SIMT和SIMD混合编程的算子开发方式。算子中使用SIMT编程方式实现离散内存访问操作gather，使用SIMD编程方式实现连续内存访问操作adds。
5- 
6- 
7-## 支持的产品
8-- Ascend 950PR/Ascend 950DT
9-## 目录结构
10-```
11-├── simt_gather_and_simd_adds
12-│ ├── CMakeLists.txt # cmake编译文件
13-│ ├── gather_and_adds.asc # Ascend C算子实现 & 调用样例
14-| └── README.md
15-```
16- 
17-## 算子描述
18- 
19-- 算子功能：
20- gather & adds算子实现了从长度为10万的一维向量中获取指定索引的8192个数据，将获取到的数据分别加1，返回相加结果的功能。算子输出output第i个数据计算公式为：
21-
22- ```
23- output[i] = input[index[i]] + 1
24- ```
25- 
26-- 算子规格：
27- <table>
28- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center">gather & adds</td></tr>
29- </tr>
30- <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>
31- <tr><td align="center">input</td><td align="center">100000</td><td align="center">float</td><td align="center">ND</td></tr>
32- <tr><td align="center">index</td><td align="center">8192</td><td align="center">uint32_t</td><td align="center">ND</td></tr>
33- </tr>
34- </tr>
35- <tr><td rowspan="1" align="center">算子输出</td><td align="center">output</td><td align="center">8192</td><td align="center">float</td><td align="center">ND</td></tr>
36- </tr>
37- <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">gather_and_adds_kernel</td></tr>
38- </table>
39- 
40-- 基础知识：
41- 在SIMT编程中，可以直接读取和使用Global Memory上的数据。在Vector Core中，SIMT单元和SIMD单元共享片上存储，因此可以使用片上存储完成SIMT和SIMD的混合编程。本例中，首先使用SIMT编程方式获取指定索引位置的输入数据，存储到UB，再使用SIMD编程方式利用UB完成adds运算。
42- 
43- 具体来说，SIMT编程常通过组织线程的层次结构来实现数据的切分，函数中使用threadIdx等内置变量计算线程应处理的数据索引，完成索引对应数据的计算，函数被简化为标量计算。SIMD编程提供基于寄存器（Regbase）的编程API，可以直接操作Vector Core中的寄存器。单次API处理的数据量上限为寄存器大小，通过AscendC::GetVecLen函数获取该值。在计算过程中，需要多次调用微指令API完成单核数据的处理。
44- 
45-- 数据切分：
46- 本例中算子输入index的形状为8192，可设置核数为8，每个核处理数据量为1024，设置线程数THREAD_COUNT为1024，每个线程处理1个数据，因此单个核只需调用1次simt_gather函数即可完成gather运算。
47- 
48- > ⚠️ **注意** 当单核处理数据量大于设置的线程数时，需要切分数据到多个线程块，可使用asc_vf_call多次调用simt_gather函数启动多个线程块完成获取指定索引数据的操作。
49- 
50- 基于上述数据拆分，在simd_adds函数中，处理1024个数据的加1操作。
51- 
52- > ⚠️ **注意** simd_adds中加1运算实际可以直接在simt_gather函数中快速实现，本例目的仅仅是通过一个简单用例展示SIMT和SIMD两种编程模式的混合编程方式，不是该算子最佳实践。
53- 
54-- 算子实现：
55- gather & adds算子的实现流程主要分为3个步骤：simt_gather，simd_adds和DataCopy。
56-
57- simt_gather负责从输入input（Global Memory）中获取指定索引的数据。基于上述数据切分，首先计算线程应处理数据的索引，然后通过赋值操作将数据存储到Local Memory上。
58- ```
59- int idx = blockIdx.x * blockDim.x + threadIdx.x;
60- ...
61- uint32_t gather_idx = index[idx];
62- ...
63- gather_output[threadIdx.x] = input[gather_idx];
64- ```
65- 
66- simd_adds负责将Local Memory中数据做加1操作。使用Reg::LoadAlign将数据从Local Memory搬运到寄存器上，调用Reg::Adds完成加1运算并输出到目标寄存器，最后调用Reg::StoreAlign将数据从寄存器搬运到Local Memory。重复上述操作即可完成1024个数据的加1运算。
67- ```
68- for (uint16_t i = 0; i < repeat_times; i++) {
69- AscendC::Reg::LoadAlign(src_reg0, input + i * one_repeat_size);
70- AscendC::Reg::Adds(dst_reg0, src_reg0, ADDS_ADDEND, mask_reg);
71- AscendC::Reg::StoreAlign(output + i * one_repeat_size, dst_reg0, mask_reg);
72- }
73- ```
74- 
75- DataCopy负责将输出数据从Local Memory搬运至Global Memory上。
76- 
77-- 调用实现：
78- - CPU调测模式使用ICPU_RUN_KF CPU调测宏调用核函数。
79- - NPU模式使用内核调用符<<<>>>调用核函数。
80- 
81- 应用程序通过ASCENDC_CPU_DEBUG宏区分代码逻辑运行于CPU模式还是NPU模式。
82- 
83-## 编译运行
84-在本样例根目录下执行如下步骤，编译并执行算子。
85-- 配置环境变量
86- 请根据当前环境上CANN开发套件包的[安装方式](../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
87- - 默认路径，root用户安装CANN软件包
88- ```bash
89- source /usr/local/Ascend/cann/set_env.sh
90- ```
91- 
92- - 默认路径，非root用户安装CANN软件包
93- ```bash
94- source $HOME/Ascend/cann/set_env.sh
95- ```
96- 
97- - 指定路径install_path，安装CANN软件包
98- ```bash
99- source ${install_path}/cann/set_env.sh
100- ```
101-
102-- 样例执行
103- ```bash
104- mkdir -p build && cd build; # 创建并进入build目录
105- cmake ..; make -j; # 编译工程
106- ./demo # 执行样例
107- ```
108- 执行结果如下，说明精度对比成功。
109- ```
110- [Success] Case accuracy is verification passed.
111- ```
112- 
113-## CPU Debug调测
114-CPU Debug功能支持对CPU执行过程中的运行状态进行调试，主要通过GDB工具实现。GDB调试支持设置断点、查看寄存器和内存状态、单步执行、查看调用栈等常用调试操作，并支持多线程程序的调试。
115- 
116-- 样例执行
117- ```bash
118- mkdir -p build && cd build; # 创建并进入build目录
119- cmake -DRUN_MODE=cpu ..; make -j; # 编译工程
120- ./demo # 执行样例
121- ```
122- 执行结果如下，说明精度对比成功。
123- ```
124- [Success] Case accuracy is verification passed.
125- ```
126- 
127-- 进入GDB模式调试
128- 在上述指令中"./demo"前加入"gdb --args"，再次执行指令即可进入GDB模式。调试时需要选择子进程进行调试。
129- ```bash
130- mkdir -p build && cd build;
131- cmake -DRUN_MODE=cpu ..; make -j;
132- gdb --args ./demo
133- ```
134- 
135-## 性能调优
136-算子调优工具支持上板调优和仿真调优两种模式，可分别获取算子在实际硬件/仿真环境下的性能数据，用于定位性能瓶颈、优化算子实现。
137-### 上板调优
138-基于编译生成的可执行文件，直接在NPU硬件上采集算子性能数据，数据为算子预热后的真实运行指标。
139- 
140-**操作步骤**
141- 
142-**1.执行调优命令**
143- 
144-基于编译得到的demo文件，运行算子调优工具。
145-```bash
146-msprof op ./demo
147-```
148-**2.查看性能数据**
149- 
150-在当前目录下会生成OPPPROF_为前缀的文件夹，目录结构及文件说明如下：
151-```bash
152-OPPROF_xxxxxxxx_XXXXXX
153-├── dump # 原始性能数据（无需关注）
154-├── OpBasicInfo.csv # 算子基础数据
155-├── ArithmeticUtilization.csv # cube及vector类型指令的cycle占比数据
156-├── ResourceConflictRatio.csv # 资源冲突占比数据
157-├── ... (开启的aic-metrics)
158-└── visualize_data.bin # 算子可视化文件（可通过MindStudio Insight加载，直观查看算子性能）
159-```
160-可直接查看.csv文件获取算子block级性能数据或通过MindStudio Insight工具打开`visualize_data.bin`文件可视化查看性能数据。
161- 
162-### 仿真调优
163-在无NPU硬件环境的场景下，通过编译仿真版算子可执行文件，结合模拟器获取算子仿真性能数据，适用于开发阶段的快速调优。
164- 
165-**操作步骤**
166- 
167-**1.仿真算子编译**
168- 
169-```bash
170-# 创建并进入build目录
171-mkdir -p build && cd build
172-# 替换${SOC_VERSION}为实际NPU型号，可通过npu-smi info命令进行查询，如Ascend910_957c。
173-cmake -DRUN_MODE=sim -DSOC_VERSION=${SOC_VERSION} ..
174-# 多线程编译
175-make -j$(nproc)
176-```
177-**2.配置运行时依赖**
178- 
179-添加运行时依赖库路径（需替换{SOC_VERSION}为实际NPU型号）：
180-```bash
181-export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/tools/simulator/${SOC_VERSION}/lib/:$LD_LIBRARY_PATH
182-```
183-**3.执行仿真调优命令**
184- 
185-```bash
186-msprof op simulator ./demo
187-```
188-**4.查看仿真性能数据**
189- 
190-当前目录下会生成OPPROF_前缀的文件夹，目录结构如下：
191-```bash
192-OPPROF_xxxxxxxx_XXXXXX
193-├── dump # 原始性能数据，无需关注
194-└── simulation # 仿真性能数据分析结果
195- ├── core0.veccore0 # 算子block级子核
196- ├── core0.veccore1_code_exe.csv # 代码行耗时
197- ├── core0.veccore1_instr_exe.csv # 程序代码指令详细信息
198- └── trace.json # 算子block级子核流水图
199- ├── ...
200- ├── visualize_data.bin # 算子可视化文件（可通过MindStudio Insight加载，直观查看算子性能）
201- └── trace.json # 算子所有核的流水图
202-```
203-**补充说明**
204-更多性能指标的详细说明及调优方案，可参考《算子开发工具》手册。
Dexamples/02_features/03_simt/simt_gather_and_simd_adds/gather_and_adds.asc+0-232
@@ -1,232 +0,0 @@
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 gather_and_adds.asc
14- * \brief
15- */
16- 
17-#include <chrono>
18-#include <iostream>
19-#include <iterator>
20-#include <random>
21-#include <vector>
22-#include "acl/acl.h"
23-#include "kernel_operator.h"
24-#include "simt_api/asc_simt.h"
25-#if defined(ASCENDC_CPU_DEBUG)
26-#include "tikicpulib.h"
27-#endif
28- 
29-namespace {
30- constexpr float ADDS_ADDEND = 1.0f;
31- constexpr uint32_t THREAD_COUNT = 1024;
32-}
33- 
34-__simt_vf__ __launch_bounds__(THREAD_COUNT) inline void simt_gather(
35- __gm__ float* input,
36- __gm__ uint32_t* index,
37- __ubuf__ float* gather_output,
38- uint32_t input_total_length,
39- uint32_t index_total_length,
40- uint32_t output_total_length)
41-{
42- if (threadIdx.x >= output_total_length) {
43- return;
44- }
45- 
46- // calculate thread ID of current thread.
47- uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
48- if (idx >= index_total_length) {
49- return;
50- }
51- 
52- // calculate index of the number we need in input.
53- uint32_t gather_idx = index[idx];
54- if (gather_idx >= input_total_length) {
55- return;
56- }
57- 
58- gather_output[threadIdx.x] = input[gather_idx];
59-}
60- 
61-__simd_vf__ inline void simd_adds(__ubuf__ float* output, __ubuf__ float* input,
62- uint32_t count, uint32_t one_repeat_size, uint16_t repeat_times)
63-{
64- AscendC::Reg::RegTensor<float> src_reg0;
65- AscendC::Reg::RegTensor<float> dst_reg0;
66- // asc_update_mask() will be supported later.
67- // init MaskReg with the count of all numbers.
68- AscendC::Reg::MaskReg mask_reg;
69- 
70- for (uint16_t i = 0; i < repeat_times; i++) {
71- mask_reg = AscendC::Reg::UpdateMask<float>(count);
72- // asc_load, asc_adds and asc_store will be supported later.
73- // load data from UB to RegTensor.
74- AscendC::Reg::LoadAlign(src_reg0, input + i * one_repeat_size);
75- AscendC::Reg::Adds(dst_reg0, src_reg0, ADDS_ADDEND, mask_reg);
76- // store data from RegTensor to UB.
77- AscendC::Reg::StoreAlign(output + i * one_repeat_size, dst_reg0, mask_reg);
78- }
79-}
80- 
81-__global__ __vector__ void gather_and_adds_kernel(__gm__ float* input, __gm__ uint32_t* index, __gm__ float* output,
82- uint32_t input_total_length, uint32_t index_total_length)
83-{
84- 
85- AscendC::LocalMemAllocator<AscendC::Hardware::UB> ub_allocator;
86- 
87- // 1. use SIMT to gather numbers from input.
88- uint32_t index_total_length_per_block = index_total_length / AscendC::GetBlockNum();
89- AscendC::LocalTensor<float> gather_output = ub_allocator.Alloc<float>(index_total_length_per_block);
90- asc_vf_call<simt_gather>(dim3(THREAD_COUNT), input, index,
91- (__ubuf__ float *)gather_output.GetPhyAddr(),
92- input_total_length,
93- index_total_length,
94- index_total_length_per_block);
95- 
96- // 2. use SIMD Micro API to do addition.
97- AscendC::LocalTensor<float> adds_output = ub_allocator.Alloc<float>(index_total_length_per_block);
98- // asc_get_vf_len() will be supported later, use GetVecLen.
99- constexpr uint32_t one_repeat_size = AscendC::GetVecLen() / sizeof(float);
100- uint16_t repeat_times = (index_total_length_per_block + one_repeat_size - 1) / one_repeat_size;
101- 
102- asc_vf_call<simd_adds>((__ubuf__ float *)adds_output.GetPhyAddr(),
103- (__ubuf__ float *)gather_output.GetPhyAddr(), index_total_length_per_block, one_repeat_size, repeat_times);
104- 
105- AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(0);
106- AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(0);
107- 
108- // 3. copy data to global memory.
109- AscendC::GlobalTensor<float> output_global_tensor;
110- output_global_tensor.SetGlobalBuffer(output + index_total_length_per_block * AscendC::GetBlockIdx());
111- AscendC::DataCopy(output_global_tensor, adds_output, index_total_length_per_block);
112-}
113- 
114-std::vector<float> gather_and_adds(std::vector<float>& input, std::vector<uint32_t>& index)
115-{
116- constexpr uint32_t numBlocks = 8;
117- 
118- uint32_t input_total_length = input.size();
119- size_t input_total_byte_size = input_total_length * sizeof(float);
120- 
121- uint32_t index_total_length = index.size();
122- uint32_t index_total_byte_size = index_total_length * sizeof(uint32_t);
123- 
124- int32_t device_id = 0;
125- aclrtStream stream = nullptr;
126- 
127- uint8_t* input_host = reinterpret_cast<uint8_t *>(input.data());
128- uint8_t* index_host = reinterpret_cast<uint8_t *>(index.data());
129- uint8_t* output_host = nullptr;
130- float* input_device = nullptr;
131- uint32_t* index_device = nullptr;
132- float* output_device = nullptr;
133- 
134- aclInit(nullptr);
135- aclrtSetDevice(device_id);
136- aclrtCreateStream(&stream);
137- 
138- aclrtMallocHost((void **)(&output_host), index_total_byte_size);
139- aclrtMalloc((void **)&input_device, input_total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
140- aclrtMalloc((void **)&index_device, index_total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
141- aclrtMalloc((void **)&output_device, index_total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
142- 
143- aclrtMemcpy(input_device, input_total_byte_size, input_host, input_total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE);
144- aclrtMemcpy(index_device, index_total_byte_size, index_host, index_total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE);
145- 
146-#if defined(ASCENDC_CPU_DEBUG)
147- AscendC::SetKernelMode(KernelMode::AIV_MODE);
148- ICPU_RUN_KF(gather_and_adds_kernel, numBlocks, input_device, index_device, output_device, input_total_length, index_total_length);
149-#else
150- constexpr uint32_t dynUBufSize = 2048;
151- // launch kernel with <<<>>>, dynUBufSize is UB byte size we will allocate in kernel.
152- gather_and_adds_kernel<<<numBlocks, dynUBufSize, stream>>>(input_device, index_device, output_device, input_total_length, index_total_length);
153-#endif
154- 
155- aclrtSynchronizeStream(stream);
156- 
157- aclrtMemcpy(output_host, index_total_byte_size, output_device, index_total_byte_size, ACL_MEMCPY_DEVICE_TO_HOST);
158- std::vector<float> output((float *)output_host, (float *)(output_host + index_total_byte_size));
159- 
160- aclrtFree(input_device);
161- aclrtFree(index_device);
162- aclrtFree(output_device);
163- aclrtFreeHost(output_host);
164- 
165- aclrtDestroyStream(stream);
166- aclrtResetDevice(device_id);
167- aclFinalize();
168- 
169- return output;
170-}
171- 
172-uint32_t verify_result(std::vector<float>& output, std::vector<float>& golden)
173-{
174- auto print_tensor = [](std::vector<float>& tensor, const char* name) {
175- constexpr size_t max_print_size = 20;
176- std::cout << name << ": ";
177- std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), max_print_size),
178- std::ostream_iterator<float>(std::cout, " "));
179- if (tensor.size() > max_print_size) {
180- std::cout << "...";
181- }
182- std::cout << std::endl;
183- };
184- print_tensor(output, "Output");
185- print_tensor(golden, "Golden");
186- if (std::equal(output.begin(), output.end(), golden.begin())) {
187- std::cout << "[Success] Case accuracy is verification passed." << std::endl;
188- return 0;
189- } else {
190- std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
191- return 1;
192- }
193- return 0;
194-}
195- 
196-float random_float(float min, float max)
197-{
198- static std::mt19937 engine(std::chrono::system_clock::now().time_since_epoch().count());
199- std::uniform_real_distribution<float> dist(min, max);
200- return dist(engine);
201-}
202- 
203-uint32_t random_uint32(uint32_t min, uint32_t max)
204-{
205- static std::mt19937 engine(std::chrono::system_clock::now().time_since_epoch().count());
206- std::uniform_int_distribution<uint32_t> dist(min, max);
207- return dist(engine);
208-}
209- 
210-int32_t main(int32_t argc, char* argv[])
211-{
212- constexpr uint32_t input_total_length = 100000;
213- std::vector<float> input(input_total_length);
214- for (uint32_t i = 0; i < input_total_length; i++) {
215- input[i] = random_float(0.0f, 100.0f);
216- }
217- 
218- constexpr uint32_t index_total_length = 8 * 1024;
219- std::vector<uint32_t> index(index_total_length);
220- for (uint32_t i = 0; i < index_total_length; i++) {
221- index[i] = random_uint32(0, input_total_length - 1);
222- }
223- 
224- std::vector<float> golden(index_total_length);
225- for (uint32_t i = 0; i < index_total_length; i++) {
226- golden[i] = input[index[i]] + ADDS_ADDEND;
227- }
228- 
229- std::vector<float> output = gather_and_adds(input, index);
230- 
231- return verify_result(output, golden);
232-}
Dexamples/02_features/README.md+0-26
@@ -1,26 +0,0 @@
1-# Features样例介绍
2- 
3-## 概述
4- 
5-基于Ascend C特性样例，介绍了Aclnn（ge入图）工程、LocalMemoryAllocator、Barrier单独内存申请分配等特性
6- 
7-## 算子开发样例
8- 
9-| 目录名称 | 功能描述 |
10-| --------- | --------- |
11-| [00_framework_launch](./00_framework_launch/) | 本样例以Add算子为样例，展示了Tiling模板编程。Add算子实现了两个数据相加，返回相加结果的功能。本样例使用自定义算子工程，编译并部署自定义算子包到自定义算子库中，并调用执行自定义算子 |
12-| [01_triple_chevron_notation](./01_triple_chevron_notation/) | 本样例展示了如何使用pybind11与torch.library注册自定义算子，并通过`<<<>>>`内核调用符调用核函数 |
13-| [02_c_api](./02_c_api) | 本样例展示了使用C_API构建Add算子样例的编译流程 |
14-| [03_simt](./03_simt/) | 本样例展示了SIMT算子实现，包括基于SIMT编程模型、SIMD与SIMT混合编程模型两种模式的实现 |
15-| [04_reg_compute](./04_reg_compute/) | 通过Reg矢量计算API实现自定义算子，分别给出对应的<<<>>>直调实现 |
16-| [06_static_tensor_programming](./06_static_tensor_programming) | 本样例介绍基于静态Tensor方式编程的场景下Add算子的实现方法，并提供核函数直调方法 |
17-| [07_data_movement](./07_data_movement) | 本样例路径以copy、data_copy_pad、data_copy_ub2l1等算子为示例，展示了数据搬运接口的使用。|
18-| [08_tiling](./08_tiling) | 本样例实现了一个支持多种数据类型的向量加法算子，其核心在于采用精细的多核并行数据切分（Tiling）策略，通过17个参数动态管理任务分配，处理核间负载不均及尾核尾块场景，以实现高效能计算。|
19-| [09_unalign](./09_unalign/) | 本样例路径以abs、reduce_min、whole_reduce_sum等算子为示例，展示了数据非32字节对齐场景中的处理方式，包括数据搬入，计算和搬出的处理 |
20-| [10_memory_management](./10_memory_management/) | 本路径下包含了与资源管理相关的多个API的样例。每个样例均基于Ascend C的<<<>>>直调方法，支持main函数和kernel函数在同一个cpp文件中实现。 |
21-| [11_synchronous_control](./11_synchronous_control/) | 本路径下包含了与同步控制相关的多个API的样例。每个样例均基于Ascend C的<<<>>>直调方法，支持main函数和kernel函数在同一个cpp文件中实现。 |
22-| [12_system_variable_access](./12_system_variable_access/) | 本路径下包含了与系统变量访问相关的多个API的样例。每个样例均基于Ascend C的<<<>>>直调方法，支持main函数和kernel函数在同一个cpp文件中实现。 |
23-| [13_atomic_operations](./13_atomic_operations/) | 本路径下包含了与原子操作相关的多个API的样例。每个样例均基于Ascend C的<<<>>>直调方法，支持main函数和kernel函数在同一个cpp文件中实现。 |
24-| [14_cube_group_management](./14_cube_group_management/) | 本路径下包含了与Cube分组管理相关的多个API的样例。每个样例均基于Ascend C的<<<>>>直调方法，支持main函数和kernel函数在同一个cpp文件中实现。
25-| [15_utility_function](./15_utility_function/) | 本路径下包含了与工具函数相关的多个API的样例。每个样例均基于Ascend C的<<<>>>直调方法，支持main函数和kernel函数在同一个cpp文件中实现。 |
26-| [16_scalar_computation](./16_scalar_computation/) | 本路径下包含了与标量计算相关的多个API的样例。每个样例均基于Ascend C的<<<>>>直调方法，支持main函数和kernel函数在同一个cpp文件中实现。 |
Rexamples/00_introduction/01_add/simt_add/README.md→examples/03_simt_api/00_introduction/01_add/README.md+3-3
@@ -1,8 +1,8 @@
1-# 纯SIMT编程模式实现Add算子样例1+# SIMT编程模式实现Add算子样例
2 2 
3## 概述3## 概述
4 4 
5-样例基于Ascend C纯SIMT编程方式实现Add算子，实现两个输入张量逐元素相加得到输出张量的功能，展示纯SIMT编程的基本流程。5+样例基于Ascend C SIMT编程方式实现Add算子，实现两个输入张量逐元素相加得到输出张量的功能，展示SIMT编程的基本流程。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -11,7 +11,7 @@
11## 目录结构11## 目录结构
12 12 
13```13```
14-├── simt_add14+├── 01_add
15│ ├── add.asc # SIMT实现add调用样例15│ ├── add.asc # SIMT实现add调用样例
16| └── README.md16| └── README.md
17```17```
Rexamples/00_introduction/01_add/simt_add/add.asc→examples/03_simt_api/00_introduction/01_add/add.asc+0-0
文件重命名但无更改。
Rexamples/02_features/03_simt/README.md→examples/03_simt_api/00_introduction/README.md+4-5
@@ -1,12 +1,11 @@
1-# Simt样例介绍1+# 入门样例介绍
2 2 
3## 概述3## 概述
4-基于SIMT实现的算子实现样例。4+ 
5+基于Ascend C的简单SIMT算子实现示例，展示SIMT编程基本语法、编译及基于<<<>>>的调用方法。
5 6 
6## 算子开发样例7## 算子开发样例
7 8 
8| 目录名称 | 功能描述 |9| 目录名称 | 功能描述 |
9| ------------------------------------------------------------ | ---------------------------------------------------- |10| ------------------------------------------------------------ | ---------------------------------------------------- |
10-| [simt_gather](./simt_gather) | 本样例以Gather算子为示例，展示了在SIMD与SIMT混合编程模式下实现一个仅包含SIMT VF函数的算子 |11+| [01_add](./01_add) | 样例基于SIMT编程方式实现Add样例，展示了SIMT单指令多线程的编程方式完成加法计算 |
11-| [simt_gather_and_simd_adds](./simt_gather_and_simd_adds) | 本样例以gather & adds算子为示例，展示了在SIMD与SIMT混合编程模式下实现包含SIMT VF函数和SIMD VF函数的算子 |
12-| [pure_simt_gather](./pure_simt_gather) | 本样例以gather算子为示例，展示了在纯SIMT场景下实现动态算子功能 |
Aexamples/03_simt_api/01_utilities/00_printf/README.md+3-0
@@ -0,0 +1,3 @@
1+# printf样例介绍
2+ 
3+待开发。
Aexamples/03_simt_api/01_utilities/01_assert/README.md+3-0
@@ -0,0 +1,3 @@
1+# assert样例介绍
2+ 
3+待开发。
A
Aai_xin3月24日

这些样例readme建议与样例代码一起增加，目前不要加一个空的内容

likedislike
Aexamples/03_simt_api/01_utilities/README.md+9-0
@@ -0,0 +1,9 @@
1+# Utilities样例介绍
2+## 概述
3+基于Ascend C SIMT编程的简单样例，通过printf、clock等API介绍上板打印、时钟打点等维测方法，适用于调试阶段。
4+ 
5+## 算子开发样例
6+| 目录名称 | 功能描述 |
7+| -------------------------- | -------------------- |
8+| [00_printf](./00_printf) | 本样例展示上板打印功能使用方法。 |
9+| [01_asserk](./01_assert/) | 本样例展示上板assert断言的功能。 |
Aexamples/03_simt_api/02_features/00_resource_management/README.md+12-0
@@ -0,0 +1,12 @@
1+# 编程基础样例介绍
2+ 
3+## 概述
4+ 
5+基于Ascend C SIMT编程的基础样例，展示基于SIMT编程的算子开发用户界面、开发流程等。
6+ 
7+## 算子开发样例
8+ 
9+| 目录名称 | 功能描述 |
10+| --------- | --------- |
11+| [basic_gather](./basic_gather/) | 样例基于Ascend C SIMT编程方式实现Gather算子，展示简化场景离散内存访问类算子的开发方法。 |
12+| [general_gather](./general_gather/) | 样例基于Ascend C SIMT编程方式实现泛化场景的Gather算子，展示SIMT算子开发的端到端流程。 |
Rexamples/00_introduction/04_simple_operator/pure_simt_gather/README.md→examples/03_simt_api/02_features/00_resource_management/basic_gather/README.md+4-4
@@ -1,8 +1,8 @@
1-# 纯SIMT编程模式实现Gather算子样例1+# SIMT编程模式实现Gather算子样例
2 2 
3## 概述3## 概述
4 4 
5-样例基于Ascend C纯SIMT编程方式实现Gather算子，从输入张量中采集指定的m行数据，展示离散内存访问类算子的开发方法。5+样例基于Ascend C SIMT编程方式实现简单场景（固定shape的）Gather算子，从输入张量中采集指定的m行数据，展示简化场景离散内存访问类算子的开发方法。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -11,7 +11,7 @@
11## 目录结构11## 目录结构
12 12 
13```13```
14-├── pure_simt_gather14+├── basic_gather
15│ ├── gather.asc # SIMT实现gather调用样例15│ ├── gather.asc # SIMT实现gather调用样例
16| └── README.md16| └── README.md
17```17```
@@ -55,7 +55,7 @@
55 55 
56在本样例根目录下执行如下步骤，编译并执行算子。56在本样例根目录下执行如下步骤，编译并执行算子。
57- 配置环境变量 57- 配置环境变量
58- 请根据当前环境上CANN开发套件包的[安装方式](../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。58+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
59 - 默认路径，root用户安装CANN软件包59 - 默认路径，root用户安装CANN软件包
60 ```bash60 ```bash
61 source /usr/local/Ascend/cann/set_env.sh61 source /usr/local/Ascend/cann/set_env.sh
Rexamples/00_introduction/04_simple_operator/pure_simt_gather/gather.asc→examples/03_simt_api/02_features/00_resource_management/basic_gather/gather.asc+0-0
文件重命名但无更改。
Rexamples/02_features/03_simt/pure_simt_gather/README.md→examples/03_simt_api/02_features/00_resource_management/general_gather/README.md+4-4
@@ -1,14 +1,14 @@
1-# 纯SIMT编程模式实现Gather算子样例1+# SIMT编程模式实现Gather算子样例
2 2 
3## 概述3## 概述
4-样例基于Ascend C纯SIMT编程方式实现Gather算子，从输入张量中采集指定的m行数据，展示离散内存访问类算子的开发方法。4+样例基于Ascend C SIMT编程方式实现支持泛化shape的Gather算子，从输入张量中采集指定的m行数据，展示泛化场景离散内存访问类算子的开发方法。
5 5 
6## 支持的产品6## 支持的产品
7- Ascend 950PR/Ascend 950DT7- Ascend 950PR/Ascend 950DT
8 8 
9## 目录结构9## 目录结构
10```10```
11-├── pure_simt_gather11+├── general_gather
12│ ├── gather.asc # SIMT实现gather调用样例12│ ├── gather.asc # SIMT实现gather调用样例
13| └── README.md13| └── README.md
14```14```
@@ -51,7 +51,7 @@
51## 编译运行51## 编译运行
52在本样例根目录下执行如下步骤，编译并执行算子。52在本样例根目录下执行如下步骤，编译并执行算子。
53- 配置环境变量 53- 配置环境变量
54- 请根据当前环境上CANN开发套件包的[安装方式](../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。54+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
55 - 默认路径，root用户安装CANN软件包55 - 默认路径，root用户安装CANN软件包
56 ```bash56 ```bash
57 source /usr/local/Ascend/cann/set_env.sh57 source /usr/local/Ascend/cann/set_env.sh
Rexamples/02_features/03_simt/pure_simt_gather/gather.asc→examples/03_simt_api/02_features/00_resource_management/general_gather/gather.asc+0-0
文件重命名但无更改。
Aexamples/03_simt_api/02_features/README.md+11-0
@@ -0,0 +1,11 @@
1+# Features样例介绍
2+ 
3+## 概述
4+ 
5+基于Ascend C SIMT特性样例，展示SIMT编程的各类特性。
6+ 
7+## 算子开发样例
8+ 
9+| 目录名称 | 功能描述 |
10+| --------- | --------- |
11+| [00_resource_management](./00_resource_management/) | 展示SIMT资源管理相关的特性样例。 |
Aexamples/03_simt_api/README.md+17-0
@@ -0,0 +1,17 @@
1+# 样例运行验证
2+ 
3+开发者基于Ascend C SIMT编程实现自定义算子后，可通过单算子调用的方式验证算子功能。本代码仓提供部分算子实现及其调用样例，具体如下。
4+ 
5+## 算子开发样例
6+| 目录名称 | 功能描述 |
7+| ------------------------------------- | ------------------------------------- |
8+| [00_introduction](./00_introduction) | 基于Ascend C SIMT编程的简单的示例，通过Ascend C编程语言实现了自定义算子，分别给出对应的<<<>>>直调实现，适合初学者 |
9+| [01_utilities](./01_utilities) | 基于Ascend C SIMT的调试调优样例，通过printf、assert、clock等API介绍上板打印、打点、CPU孪生调试等系统工具使用方法，适用于调试阶段 |
10+| [02_features](./02_features) | 展示Ascend C SIMT编程的重点特性的样例，包含基础编程、内存资源管理、线程同步、原子操作等特性。 |
11+ 
12+## npu-arch编译选项说明
13+ 
14+开发者需根据实际的执行环境，修改具体样例目录下CMakeLists.txt文件中的--npu-arch编译选项，参考下表中的对应关系，修改为环境对应的npu-arch参数值。
15+| 产品型号 | npu-arch参数 |
16+| ---- | ---- |
17+| Ascend 950PR/Ascend 950DT | --npu-arch=dav-3510 |
Mexamples/README.md+1-3
@@ -5,11 +5,9 @@
5## 算子开发样例5## 算子开发样例
6| 目录名称 | 功能描述 |6| 目录名称 | 功能描述 |
7| ------------------------------------------------------------ | ---------------------------------------------------- |7| ------------------------------------------------------------ | ---------------------------------------------------- |
8-| [00_introduction](./00_introduction) | 基于Ascend C的简单的示例，通过Ascend C编程语言实现了自定义算子，分别给出对应的<<<>>>直调实现 |
9| [01_simd_cpp_api](./01_simd_cpp_api) | 基于Ascend C的SIMD API样例，通过<<<>>>直调的实现方式，介绍了SIMD API的使用方法 |8| [01_simd_cpp_api](./01_simd_cpp_api) | 基于Ascend C的SIMD API样例，通过<<<>>>直调的实现方式，介绍了SIMD API的使用方法 |
10-| [02_features](./02_features) | 基于Ascend C特性样例，介绍了Aclnn（ge入图）工程、LocalMemoryAllocator、Barrier单独内存申请分配等特性 |
11| [02_simd_c_api](./02_simd_c_api) | 基于Ascend C的C_API样例，通过C_API的实现方式，介绍了C_API的使用方法 |9| [02_simd_c_api](./02_simd_c_api) | 基于Ascend C的C_API样例，通过C_API的实现方式，介绍了C_API的使用方法 |
12- 10+| [03_simt_api](./03_simt_api) | 基于Ascend C SIMT编程的算子样例，通过<<<>>>直调的实现方式，介绍了SIMT的使用方法 |
13## npu-arch编译选项说明11## npu-arch编译选项说明
14 12 
15开发者需根据实际的执行环境，修改具体样例目录下CMakeLists.txt文件中的--npu-arch编译选项，参考下表中的对应关系，修改为环境对应的npu-arch参数值。13开发者需根据实际的执行环境，修改具体样例目录下CMakeLists.txt文件中的--npu-arch编译选项，参考下表中的对应关系，修改为环境对应的npu-arch参数值。