Star743
1034
代码
代码
Issues55
Pull Requests198
Actions
项目讨论
Wiki
分析
Star743
1034
已合并
新增simt的快速入门样例add #193
wulinyu创建于 1月30日
新增simt的快速入门样例add #193
已合并
wulinyu创建于 1月30日
从wulinyu/asc-devkitmaster合入到cann/asc-devkitmaster
共 3 个文件变更+241-1
Mexamples/00_introduction/01_add/README.md+2-1
@@ -13,4 +13,5 @@
13| [c_api_async_add](./c_api_async_add) | 本样例采用C_API接口编写Add算子样例，基于异步搬运、计算接口实现 |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算子样例，基于异步搬运、计算接口和手动添加的同步指令实现 |14| [c_api_delicacy_async_add](./c_api_delicacy_async_add) | 本样例采用C_API接口编写Add算子样例，基于异步搬运、计算接口和手动添加的同步指令实现 |
15| [c_api_delicacy_async_add](./c_api_delicacy_async_add) | 本样例采用C_API接口编写Add算子样例，基于同步搬运、计算接口实现 |15| [c_api_delicacy_async_add](./c_api_delicacy_async_add) | 本样例采用C_API接口编写Add算子样例，基于同步搬运、计算接口实现 |
16-| [micro_api_add](./micro_api_add) | 样例基于微指令API实现Add样例，展示了通过微指令API直接对芯片中涉及Vector计算的寄存器进行操作 |16+| [micro_api_add](./micro_api_add) | 样例基于微指令API实现Add样例，展示了通过微指令API直接对芯片中涉及Vector计算的寄存器进行操作 |
17+| [simt_add](./simt_add) | 样例基于纯SIMT编程方式实现Add样例，展示了SIMT单指令多线程的编程方式完成加法计算 |
Aexamples/00_introduction/01_add/simt_add/README.md+114-0
@@ -0,0 +1,114 @@
1+# 纯SIMT编程模式实现Add算子样例
2+ 
3+## 概述
A
Aai_xin2月2日

标题前后加空行

likedislike
A
ai_xin
2月2日 评论：
4+ 
5+样例基于Ascend C纯SIMT编程方式实现Add算子，实现两个输入张量逐元素相加得到输出张量的功能，展示纯SIMT编程的基本流程。
6+ 
7+## 支持的产品
8+ 
9+- Ascend 950PR/Ascend 950DT
10+ 
11+## 目录结构
12+ 
13+```
14+├── aimt_add
15+│ ├── add.asc # SIMT实现add调用样例
16+| └── README.md
17+```
18+ 
19+## 算子描述
20+ 
21+- 算子功能：
22+ 本算子实现了形状为48 * 256的两个张量x，y相加得到算子输出z。第i个元素的计算公式为：
23+
24+ ```
25+ z[i] = x[i] + y[i]
26+ ```
27+ 
28+- 算子规格：
29+ <table>
30+ <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center">add</td></tr>
31+ </tr>
32+ <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>
33+ <tr><td align="center">x</td><td align="center">48 * 256</td><td align="center">float</td><td align="center">ND</td></tr>
34+ <tr><td align="center">y</td><td align="center">48 * 256</td><td align="center">float</td><td align="center">ND</td></tr>
35+ </tr>
36+ </tr>
37+ <tr><td rowspan="1" align="center">算子输出</td><td align="center">z</td><td align="center">48 * 256</td><td align="center">float</td><td align="center">ND</td></tr>
38+ </tr>
39+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">add_custom</td></tr>
40+ </table>
41+ 
42+- 数据切分：
43+ * 核数：48核
44+ * 每核线程数：256线程
45+ * 单线程处理：1个元素
46+ * 总处理能力：48×256=12288
47+ 
48+- 算子实现：
49+ 算子的实现流程为从输入x（Global Memory上的指针）中获取指定索引的数据。基于上述数据切分，首先计算线程应处理数据的索引，然后通过加法运算符计算得到输出值。
50+ 
51+- 调用实现：
52+ 使用内核调用符<<<>>>调用核函数。
53+ 
54+## 编译运行
55+ 
56+在本样例根目录下执行如下步骤，编译并执行算子。
57+- 配置环境变量
58+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
59+ - 默认路径，root用户安装CANN软件包
60+ ```bash
61+ source /usr/local/Ascend/cann/set_env.sh
62+ ```
63+ 
64+ - 默认路径，非root用户安装CANN软件包
65+ ```bash
66+ source $HOME/Ascend/cann/set_env.sh
67+ ```
68+ 
69+ - 指定路径install_path，安装CANN软件包
70+ ```bash
71+ source ${install_path}/cann/set_env.sh
72+ ```
73+
74+- 样例执行
75+ ```bash
76+ CANN_PATH=$(printenv ASCEND_HOME_PATH) # 获取CANN包安装路径
77+ OUTPUT="demo" #用户自定义编译生成的二进制文件名称
78+ 
79+ bisheng \
80+ -x dpp --cce-aicore-arch=dav-c310-vec \
81+ -std=c++17 \
82+ add.asc \
83+ -I${CANN_PATH}/include \
84+ -I${CANN_PATH}/include/ascendc/host_api \
85+ -I${CANN_PATH}/compiler/ascendc/include/highlevel_api \
86+ -I${CANN_PATH}/compiler/tikcpp/tikcfw \
87+ -I${CANN_PATH}/compiler/tikcpp/tikcfw/lib \
88+ -I${CANN_PATH}/compiler/tikcpp/tikcfw/lib/matmul \
89+ -I${CANN_PATH}/compiler/tikcpp/tikcfw/impl \
90+ -I${CANN_PATH}/compiler/tikcpp/tikcfw/interface \
91+ -L${CANN_PATH}/lib64 \
92+ -lascendc_runtime \
93+ -lascendcl \
94+ -lruntime \
95+ -lregister \
96+ -lerror_manager \
97+ -lprofapi \
98+ -lascendalog \
99+ -lmmpa \
100+ -lascend_dump \
101+ -ltiling_api \
102+ -lplatform \
103+ -ldl \
104+ -lc_sec \
105+ -lstdc++ \
106+ -o ${OUTPUT}
107+ 
108+ ./${OUTPUT} # 执行样例
109+ 
110+ ```
111+ 执行结果如下，说明精度对比成功。
A
Aai_xin2月2日

上层readme未修改

likedislike
112+ ```
113+ [Success] Case accuracy is verification passed.
114+ ```
Aexamples/00_introduction/01_add/simt_add/add.asc+125-0
@@ -0,0 +1,125 @@
1+/**
2+* Copyright (c) 2025 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 gather.asc
13+ * \brief
14+ */
15+ 
16+#include <iostream>
17+#include <iterator>
18+#include <vector>
19+#include "acl/acl.h"
20+#include "tiling/platform/platform_ascendc.h"
21+ 
22+ 
23+__global__ void add_custom(float* x, float* y, float* z, uint64_t total_length)
24+{
25+ // Calculate global thread ID
26+ int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
27+ 
28+ // Maps to the row index of output tensor
29+ if (idx >= total_length) {
30+ return;
31+ }
32+ z[idx] = x[idx] + y[idx];
33+}
34+ 
35+std::vector<float> add(std::vector<float>& x, std::vector<float>& y)
36+{
37+ size_t total_byte_size =x.size() * sizeof(float);
38+ int32_t device_id = 0;
39+ aclrtStream stream = nullptr;
40+ 
41+ uint8_t* x_host = reinterpret_cast<uint8_t *>(x.data());
42+ uint8_t* y_host = reinterpret_cast<uint8_t *>(y.data());
43+ uint8_t* z_host = nullptr;
44+ float* x_device = nullptr;
45+ float* y_device = nullptr;
46+ float* z_device = nullptr;
47+ // Init
48+ aclInit(nullptr);
49+ aclrtSetDevice(device_id);
50+ aclrtCreateStream(&stream);
51+ // Malloc memory in host and device
52+ aclrtMallocHost((void **)(&z_host), total_byte_size);
53+ aclrtMalloc((void **)&x_device, total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
54+ aclrtMalloc((void **)&y_device, total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
55+ aclrtMalloc((void **)&z_device, total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
56+ aclrtMemcpy(x_device, total_byte_size, x_host, total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE);
57+ aclrtMemcpy(y_device, total_byte_size, y_host, total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE);
58+ // Calc splite params
59+ uint32_t block_num = 48;
60+ uint32_t thread_num_per_block = 256;
61+ uint32_t dyn_ubuf_size = 0; // No need to alloc dynamic memory.
62+ // Call kernel funtion with <<<...>>>
63+ add_custom<<<block_num, thread_num_per_block, dyn_ubuf_size, stream>>>(x_device, y_device, z_device, x.size());
64+ aclrtSynchronizeStream(stream);
65+ // Copy result from device to host
66+ aclrtMemcpy(z_host, total_byte_size, z_device, total_byte_size, ACL_MEMCPY_DEVICE_TO_HOST);
67+ std::vector<float> output((float *)z_host, (float *)(z_host + total_byte_size));
68+ // Free memory
69+ aclrtFree(x_device);
70+ aclrtFree(y_device);
71+ aclrtFree(z_device);
72+ aclrtFreeHost(z_host);
73+ // DeInt
74+ aclrtDestroyStream(stream);
75+ aclrtResetDevice(device_id);
76+ aclFinalize();
77+ return output;
78+}
79+ 
80+uint32_t verify_result(std::vector<float>& output, std::vector<float>& golden)
81+{
82+ auto print_tensor = [](std::vector<float>& tensor, const char* name) {
83+ constexpr size_t max_print_size = 20;
84+ std::cout << name << ": ";
85+ std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), max_print_size),
86+ std::ostream_iterator<float>(std::cout, " "));
87+ if (tensor.size() > max_print_size) {
88+ std::cout << "...";
89+ }
90+ std::cout << std::endl;
91+ };
92+ print_tensor(output, "Output");
93+ print_tensor(golden, "Golden");
94+ if (std::equal(output.begin(), output.end(), golden.begin())) {
95+ std::cout << "[Success] Case accuracy is verification passed." << std::endl;
96+ return 0;
97+ } else {
98+ std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
99+ return 1;
100+ }
101+ return 0;
102+}
103+ 
104+int32_t main(int32_t argc, char* argv[])
105+{
106+ constexpr uint32_t in_shape = 48 * 256;
107+ std::vector<float> x(in_shape);
108+ for (uint32_t i = 0; i < in_shape; i++) {
109+ x[i] = i * 1.1f;
110+ }
111+ 
112+ std::vector<float> y(in_shape);
113+ for (uint32_t i = 0; i < in_shape; i++) {
114+ y[i] = i + 3.4f;
115+ }
116+ 
117+ std::vector<float> golden(in_shape);
118+ for (uint32_t i = 0; i < in_shape; i++) {
119+ golden[i] = x[i] + y[i];
120+ }
121+ 
122+ std::vector<float> output = add(x, y);
123+ 
124+ return verify_result(output, golden);
125+}