Star596
205
代码
代码
Issues11
Pull Requests40
Actions
项目讨论
Wiki
分析
Star596
205
已合并
增加scatter story sample #346
huang-jz创建于 7月21日
增加scatter story sample #346
已合并
huang-jz创建于 7月21日
从huang-jz/cann-samples0720-scatter合入到cann/cann-samplesmaster
共 9 个文件变更+873-0
MSamples/2_Performance/CMakeLists.txt+1-0
@@ -17,4 +17,5 @@ add_subdirectory(moe_init_routing_story)
17add_subdirectory(moe_dispatch_and_combine_story)17add_subdirectory(moe_dispatch_and_combine_story)
18add_subdirectory(kv_rms_norm_rope_cache_story)18add_subdirectory(kv_rms_norm_rope_cache_story)
19add_subdirectory(simd_vf_story)19add_subdirectory(simd_vf_story)
20+add_subdirectory(simt_scatter_story)
20add_subdirectory(flash_attn_lite_story)21add_subdirectory(flash_attn_lite_story)
MSamples/2_Performance/README.md+4-0
@@ -22,3 +22,7 @@
22 22 
23### [kv_rms_norm_rope_cache_story](./kv_rms_norm_rope_cache_story)23### [kv_rms_norm_rope_cache_story](./kv_rms_norm_rope_cache_story)
24围绕 Ascend 950 上的 KvRmsNormRopeCache full-load 路径给出 MemBase 与 RegBase 两个 BF16 直调样例，展示 RMSNorm、interleave RoPE 与 Norm cache 更新的融合实现，以及从 MemBase 到 RegBase 的寄存器化优化点。24围绕 Ascend 950 上的 KvRmsNormRopeCache full-load 路径给出 MemBase 与 RegBase 两个 BF16 直调样例，展示 RMSNorm、interleave RoPE 与 Norm cache 更新的融合实现，以及从 MemBase 到 RegBase 的寄存器化优化点。
25+ 
26+### [simt_scatter_story](./simt_scatter_story)
27+以 Scatter 算子为例的 SIMT 递进教学样例（`dav-3510`）：演示 SIMT 直接访问 GM 完成不规则写，以及通过目标地址分组和单写者选择处理重复 index 带来的写冲突。
28+ 
ASamples/2_Performance/simt_scatter_story/CMakeLists.txt+59-0
@@ -0,0 +1,59 @@
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, INCLUDING
8+# 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+set(SUPPORTED_NPU_ARCHS dav-3510)
13+file(RELATIVE_PATH SAMPLE_PATH ${PROJECT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR})
14+if(NOT "${NPU_ARCH}" IN_LIST SUPPORTED_NPU_ARCHS)
15+ message(STATUS "Skip sample ${SAMPLE_PATH}: NPU_ARCH=${NPU_ARCH} is not supported")
16+ return()
17+endif()
18+ 
19+set(CPP_FILES
20+ src/0_direct_unique.asc
21+ src/1_grouped_conflict.asc
22+ src/2_grouped_conflict_2d.asc
23+)
24+ 
25+configure_file(scripts/gen_data.py gen_data.py COPYONLY)
26+ 
27+set(INSTALL_TARGETS "")
28+foreach(CPP_FILE ${CPP_FILES})
29+ get_filename_component(TARGET_NAME ${CPP_FILE} NAME_WE)
30+ set(TARGET_NAME "simt_scatter_${TARGET_NAME}")
31+ add_executable(${TARGET_NAME} ${CPP_FILE})
32+ target_include_directories(${TARGET_NAME} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
33+ target_compile_definitions(${TARGET_NAME} PRIVATE SOURCE_DIR="${CMAKE_CURRENT_SOURCE_DIR}")
34+ target_compile_options(${TARGET_NAME} PRIVATE
35+ "$<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${NPU_ARCH}>"
36+ "$<$<COMPILE_LANGUAGE:ASC>:-O3>"
37+ )
38+ target_link_libraries(${TARGET_NAME}
39+ PRIVATE
40+ m
41+ dl
42+ platform
43+ tiling_api
44+ ascendcl
45+ runtime
46+ stdc++
47+ )
48+ list(APPEND INSTALL_TARGETS ${TARGET_NAME})
49+endforeach()
50+ 
51+add_custom_target(simt_scatter_story DEPENDS ${INSTALL_TARGETS})
52+ 
53+install(TARGETS ${INSTALL_TARGETS}
54+ RUNTIME DESTINATION 2_Performance/simt_scatter_story
55+)
56+ 
57+install(FILES scripts/gen_data.py
58+ DESTINATION 2_Performance/simt_scatter_story
59+)
ASamples/2_Performance/simt_scatter_story/README.md+244-0
@@ -0,0 +1,244 @@
1+# 以 Scatter 算子为例：SIMT 不规则写与写冲突处理指南
2+ 
3+本样例面向 Ascend 950 系列上的 SIMT 编程教学，围绕 Scatter 算子展示两类核心问题：
4+ 
5+- **不规则写**：每个 update 根据 `indices[i]` 写到离散的 GM 地址，地址不连续，难以用传统 SIMD/MTE 连续搬运高效覆盖。
6+- **写冲突**：多个 update 指向同一个输出地址时，不能让多个 SIMT 线程同时无序覆盖同一位置，否则结果不可重复。要求使用最后写入值为准的冲突处理语义。
7+ 
8+样例采用递进式 story 结构，每个 step 都是独立可执行目标，方便用户对照代码和 msprof 结果。
9+ 
10+## 支持范围
11+ 
12+| 项目 | 说明 |
13+|:---|:---|
14+| **硬件 / 架构** | Ascend 950PR / 950DT，`NPU_ARCH=dav-3510` |
15+| **特性线** | SIMT |
16+| **算子** | Scatter，覆盖语义 |
17+| **数据类型** | `int32` |
18+| **输入 / 输出 shape** | `base[4096, 8]`，`updates[N, 8]`，`indices[N]`，`y[4096, 8]` |
19+| **冲突语义** | 重复 index 采用 last-writer-wins，即原始 update 顺序中最后一次写生效 |
20+ 
21+已知限制：
22+ 
23+- 本样例只覆盖覆盖语义 Scatter，不覆盖 ScatterAdd / ScatterMax 等规约语义。
24+- 本样例的冲突处理依赖“按目标地址分组后再写”的模式；真实业务中若输入未分组，可在前处理阶段用 Sort / ArgSort / owner map 等方式生成同等结构。
25+- 性能数据与硬件、CANN 版本、shape、index 分布强相关，请以本地 msprof 为准。
26+ 
27+## 目录说明
28+ 
29+```text
30+simt_scatter_story/
31+├── CMakeLists.txt
32+├── README.md
33+├── include/
34+│ └── sample_common.h # ACL 初始化、数据生成、bin 读取、golden 校验、RunSample 入口
35+├── scripts/
36+│ └── gen_data.py # 生成 unique / conflict 两套输入和 golden
37+└── src/
38+ ├── 0_direct_unique.asc # Step 0：唯一 index，SIMT 直接离散写
39+ ├── 1_grouped_conflict.asc # Step 1：重复 index，分组后单写者处理冲突
40+ └── 2_grouped_conflict_2d.asc# Step 2：二维 SIMT 线程布局，并行写 row 内元素
41+```
42+ 
43+## Scatter 语义
44+ 
45+本样例计算：
46+ 
47+```text
48+y = base
49+for i in range(updateRows):
50+ y[indices[i], :] = updates[i, :]
51+```
52+ 
53+当 `indices` 不重复时，每个 update 写不同输出行，可以直接并行。
54+ 
55+当 `indices` 重复时，多个 update 会写同一行。若多个 SIMT 线程直接执行 `y[indices[i]] = updates[i]`，线程调度顺序会影响最终结果。为获得确定结果，本样例定义 last-writer-wins：
56+ 
57+```text
58+y[dst, :] = updates[last_i, :]
59+last_i = max(i) where indices[i] == dst
60+```
61+ 
62+## Step 0：唯一 index 的 SIMT 直接写
63+ 
64+源文件：`src/0_direct_unique.asc`
65+ 
66+目标：先理解 SIMT 为什么适合 Scatter 的不规则写。
67+ 
68+核心代码：
69+ 
70+```cpp
71+for (int32_t row = rowTid + coreId * rowThreadNum; row < updateRows; row += coreNum * rowThreadNum) {
72+ int32_t dst = indices[row];
73+ int64_t srcOffset = static_cast<int64_t>(row) * innerDim;
74+ int64_t dstOffset = static_cast<int64_t>(dst) * innerDim;
75+ for (int32_t col = colTid; col < innerDim; col += colThreadNum) {
76+ y[dstOffset + col] = updates[srcOffset + col];
77+ }
78+}
79+```
80+ 
81+要点：
82+ 
83+- SIMT 线程可直接访问 GM，`indices[row]` 决定每个线程的离散写地址。
84+- 当 `indices` 唯一时，没有两个线程写同一目标行，直接写是正确的。
85+- 这里用二维线程布局：x 维负责 update row，y 维负责 row 内 `innerDim` 元素。
86+ 
87+运行：
88+ 
89+```bash
90+cmake --build build --target simt_scatter_0_direct_unique
91+./build/Samples/2_Performance/simt_scatter_story/simt_scatter_0_direct_unique
92+```
93+ 
94+预期输出：
95+ 
96+```text
97+[0_direct_unique] step 0 PASSED
98+```
99+ 
100+## Step 1：重复 index 的单写者处理
101+ 
102+源文件：`src/1_grouped_conflict.asc`
103+ 
104+目标：解决重复 index 导致的写冲突。
105+ 
106+如果直接写：
107+ 
108+```cpp
109+y[indices[row], col] = updates[row, col];
110+```
111+ 
112+当两个线程的 `indices[row]` 相同，就会出现多写者覆盖同一 GM 地址。覆盖语义下这不是可交换规约，不能靠线程调度获得稳定结果。
113+ 
114+本样例采用“先分组，后单写者”的处理模式。`gen_data.py` 会把冲突输入按 `(dst, 原始位置)` 排序，因此同一个 `dst` 的 update 连续排列，且组内仍保持原始顺序。kernel 中只让每个目标地址所在分组的最后一行写出：
115+ 
116+```cpp
117+int32_t dst = indices[row];
118+bool isLastInGroup = (row == updateRows - 1) || (dst != indices[row + 1]);
119+if (!isLastInGroup) {
120+ continue;
121+}
122+```
123+ 
124+这样每个目标行最终只被一个 SIMT 线程写，结果就等价于 last-writer-wins。
125+ 
126+运行：
127+ 
128+```bash
129+cmake --build build --target simt_scatter_1_grouped_conflict
130+./build/Samples/2_Performance/simt_scatter_story/simt_scatter_1_grouped_conflict
131+```
132+ 
133+预期输出：
134+ 
135+```text
136+[1_grouped_conflict] step 1 PASSED
137+```
138+ 
139+## Step 2：二维 SIMT 布局处理 row 内并行
140+ 
141+源文件：`src/2_grouped_conflict_2d.asc`
142+ 
143+目标：在 Step 1 的冲突处理基础上，让 row 内元素也由 SIMT y 维线程并行写出。
144+ 
145+Step 1 中，一个线程负责一个目标 row 的全部 `innerDim` 元素：
146+ 
147+```cpp
148+for (int32_t col = 0; col < innerDim; ++col) {
149+ y[dstOffset + col] = updates[srcOffset + col];
150+}
151+```
152+ 
153+Step 2 改为二维线程：
154+ 
155+```cpp
156+for (int32_t col = colTid; col < innerDim; col += colThreadNum) {
157+ y[dstOffset + col] = updates[srcOffset + col];
158+}
159+```
160+ 
161+要点：
162+ 
163+- x 维线程负责不同 update row。
164+- y 维线程负责同一个 row 内不同列。
165+- `isLastInGroup` 仍然在 row 维判断，确保一个目标 row 只有组尾 update 写出。
166+ 
167+运行：
168+ 
169+```bash
170+cmake --build build --target simt_scatter_2_grouped_conflict_2d
171+./build/Samples/2_Performance/simt_scatter_story/simt_scatter_2_grouped_conflict_2d
172+```
173+ 
174+预期输出：
175+ 
176+```text
177+[2_grouped_conflict_2d] step 2 PASSED
178+```
179+ 
180+## 实验运行
181+ 
182+以下命令在 `cann-samples` 仓库根目录执行。
183+ 
184+### 环境准备
185+ 
186+```bash
187+source ${ASCEND_HOME_PATH}/set_env.sh
188+# 或：source /usr/local/Ascend/ascend-toolkit/set_env.sh
189+```
190+ 
191+Python 侧至少需要 `numpy`：
192+ 
193+```bash
194+pip install numpy
195+```
196+ 
197+### 构建
198+ 
199+```bash
200+cmake -S . -B build -DNPU_ARCH=dav-3510
201+cmake --build build --target simt_scatter_story
202+```
203+ 
204+也可以只构建单个 step：
205+ 
206+```bash
207+cmake --build build --target simt_scatter_0_direct_unique
208+cmake --build build --target simt_scatter_1_grouped_conflict
209+cmake --build build --target simt_scatter_2_grouped_conflict_2d
210+```
211+ 
212+### 运行
213+ 
214+```bash
215+./build/Samples/2_Performance/simt_scatter_story/simt_scatter_0_direct_unique
216+./build/Samples/2_Performance/simt_scatter_story/simt_scatter_1_grouped_conflict
217+./build/Samples/2_Performance/simt_scatter_story/simt_scatter_2_grouped_conflict_2d
218+```
219+ 
220+每个可执行文件启动后会自动调用 `gen_data.py --output <exe_dir>` 生成输入和 golden，再执行 kernel 并校验输出。
221+ 
222+## 什么时候使用这种模式
223+ 
224+适合：
225+ 
226+- `indices` 离散，SIMD 连续搬运难以高效覆盖。
227+- update 粒度较小，直接 GM 访问和多线程调度可以隐藏部分访存延迟。
228+- 覆盖语义下存在重复 index，需要确定性结果。
229+ 
230+不适合：
231+ 
232+- `innerDim` 很大且每个目标连续搬运占主导，此时 MTE 搬运加 SIMD 计算可能更合适。
233+- 重复 index 需要加和、最大值等规约语义，应优先考虑原子操作或分块归约，而不是 last-writer-wins。
234+- 输入完全未分组且冲突率极高，前处理成本可能成为主瓶颈，需要结合业务分布评估。
235+ 
236+## 总结
237+ 
238+| Step | 场景 | 关键点 | 源文件 |
239+|:---:|:---|:---|:---|
240+| 0 | 唯一 index | SIMT 直接离散写 GM | `src/0_direct_unique.asc` |
241+| 1 | 重复 index | 目标地址分组，只让组尾写 | `src/1_grouped_conflict.asc` |
242+| 2 | 重复 index + row 内并行 | 二维 SIMT 线程布局 | `src/2_grouped_conflict_2d.asc` |
243+ 
244+处理 SIMT Scatter 的推荐思路是：**先确认写语义 -> 再判断 index 是否唯一 -> 对重复 index 先收敛为单写者或规约者 -> 最后再做 SIMT 线程布局优化**。
ASamples/2_Performance/simt_scatter_story/include/sample_common.h+265-0
@@ -0,0 +1,265 @@
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+#ifndef SIMT_SCATTER_STORY_SAMPLE_COMMON_H_
12+#define SIMT_SCATTER_STORY_SAMPLE_COMMON_H_
13+ 
14+#include <algorithm>
15+#include <cstdint>
16+#include <cstdlib>
17+#include <fstream>
18+#include <iostream>
19+#include <sstream>
20+#include <stdexcept>
21+#include <string>
22+#include <vector>
23+#include <libgen.h>
24+#include <linux/limits.h>
25+#include <unistd.h>
26+ 
27+#include "acl/acl.h"
28+#include "acl/acl_rt.h"
29+#include "kernel_operator.h"
30+ 
31+#ifndef SOURCE_DIR
32+#define SOURCE_DIR "."
33+#endif
34+ 
35+#define CHECK_RET(cond, return_expr) \
36+ do { \
37+ if (!(cond)) { \
38+ return_expr; \
39+ } \
40+ } while (0)
41+ 
42+#define LOG_PRINT(message, ...) \
43+ do { \
44+ printf(message, ##__VA_ARGS__); \
45+ } while (0)
46+ 
47+namespace SimtScatterStorySample {
48+ 
49+using DataType = int32_t;
50+using IndexType = int32_t;
51+ 
52+constexpr int32_t DST_ROWS = 4096;
53+constexpr int32_t INNER_DIM = 8;
54+constexpr int32_t UNIQUE_UPDATES = 4096;
55+constexpr int32_t CONFLICT_UPDATES = 8192;
56+constexpr int32_t OUTPUT_ELEMS = DST_ROWS * INNER_DIM;
57+constexpr uint32_t BLOCKS = 4;
58+constexpr int32_t SIMT_THREAD_NUM = 2048;
59+constexpr int32_t SIMT_X_THREAD_NUM = 256;
60+constexpr int32_t SIMT_Y_THREAD_NUM = 8;
61+constexpr int32_t MAX_ERROR_ELEM_NUM = 20;
62+ 
63+enum class ScatterDataCase {
64+ UNIQUE = 0,
65+ CONFLICT = 1,
66+};
67+ 
68+using LaunchKernelFunc = void (*)(uint32_t blocks, aclrtStream stream, IndexType* indices, DataType* updates,
69+ DataType* y, int32_t updateRows, int32_t dstRows, int32_t innerDim);
70+ 
71+inline std::string GetExeDir()
72+{
73+ char path[PATH_MAX];
74+ ssize_t len = readlink("/proc/self/exe", path, sizeof(path) - 1);
75+ if (len != -1) {
76+ path[len] = '\0';
77+ return std::string(dirname(path));
78+ }
79+ return ".";
80+}
81+ 
82+template <typename T>
83+inline void ReadBin(const std::string& filename, std::vector<T>& data)
84+{
85+ std::ifstream file(filename, std::ios::binary);
86+ if (!file.is_open()) {
87+ throw std::runtime_error("Can not open file: " + filename);
88+ }
89+ file.seekg(0, std::ios::end);
90+ std::streampos fileSize = file.tellg();
91+ file.seekg(0, std::ios::beg);
92+ size_t elemNum = static_cast<size_t>(fileSize) / sizeof(T);
93+ data.resize(elemNum);
94+ if (elemNum > 0) {
95+ file.read(reinterpret_cast<char*>(data.data()), elemNum * sizeof(T));
96+ }
97+}
98+ 
99+inline int InitAcl(int32_t deviceId, aclrtStream* stream)
100+{
101+ auto ret = aclInit(nullptr);
102+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
103+ ret = aclrtSetDevice(deviceId);
104+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
105+ ret = aclrtCreateStream(stream);
106+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
107+ return ACL_SUCCESS;
108+}
109+ 
110+inline std::string FindGenDataScript(const std::string& exeDir)
111+{
112+ std::vector<std::string> candidates = {
113+ exeDir + "/gen_data.py",
114+ exeDir + "/scripts/gen_data.py",
115+ std::string(SOURCE_DIR) + "/scripts/gen_data.py",
116+ };
117+ for (const auto& path : candidates) {
118+ std::ifstream script(path);
119+ if (script.is_open()) {
120+ return path;
121+ }
122+ }
123+ return candidates.back();
124+}
125+ 
126+inline int GenerateData(const std::string& exeDir)
127+{
128+ std::ostringstream cmd;
129+ cmd << "env -u LD_LIBRARY_PATH python3 " << FindGenDataScript(exeDir) << " --output " << exeDir;
130+ int ret = std::system(cmd.str().c_str());
131+ if (ret != 0) {
132+ std::cerr << "Generate data failed, command: " << cmd.str() << std::endl;
133+ }
134+ return ret;
135+}
136+ 
137+inline const char* CasePrefix(ScatterDataCase dataCase)
138+{
139+ return dataCase == ScatterDataCase::UNIQUE ? "unique" : "conflict";
140+}
141+ 
142+inline int32_t CaseUpdateRows(ScatterDataCase dataCase)
143+{
144+ return dataCase == ScatterDataCase::UNIQUE ? UNIQUE_UPDATES : CONFLICT_UPDATES;
145+}
146+ 
147+template <typename T>
148+inline void CheckSize(const std::string& name, const std::vector<T>& data, size_t expected)
149+{
150+ if (data.size() != expected) {
151+ std::ostringstream oss;
152+ oss << name << " size mismatch, expected " << expected << ", actual " << data.size();
153+ throw std::runtime_error(oss.str());
154+ }
155+}
156+ 
157+inline int CompareInt(const std::string& name, const DataType* actual, const std::vector<DataType>& golden)
158+{
159+ int errorCount = 0;
160+ for (size_t i = 0; i < golden.size(); ++i) {
161+ if (actual[i] != golden[i]) {
162+ if (errorCount < MAX_ERROR_ELEM_NUM) {
163+ std::cout << name << " mismatch index " << i << ", expected " << golden[i] << ", actual "
164+ << actual[i] << std::endl;
165+ }
166+ ++errorCount;
167+ }
168+ }
169+ float precision = golden.empty() ? 100.0f
170+ : static_cast<float>(golden.size() - errorCount) / golden.size() * 100.0f;
171+ std::cout << name << " precision " << precision << "%, errors " << errorCount << std::endl;
172+ return errorCount;
173+}
174+ 
175+inline int PrepareInputs(ScatterDataCase dataCase, int32_t& deviceId, aclrtStream& stream, IndexType*& dIndices,
176+ DataType*& dUpdates, DataType*& dOut, int32_t& updateRows, std::vector<DataType>& golden)
177+{
178+ deviceId = 0;
179+ stream = nullptr;
180+ auto ret = InitAcl(deviceId, &stream);
181+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
182+ 
183+ std::string exeDir = GetExeDir();
184+ ret = GenerateData(exeDir);
185+ CHECK_RET(ret == 0, return ret);
186+ 
187+ const std::string prefix = CasePrefix(dataCase);
188+ updateRows = CaseUpdateRows(dataCase);
189+ std::vector<DataType> base;
190+ std::vector<IndexType> indices;
191+ std::vector<DataType> updates;
192+ try {
193+ ReadBin(exeDir + "/input/base.bin", base);
194+ ReadBin(exeDir + "/input/" + prefix + "_indices.bin", indices);
195+ ReadBin(exeDir + "/input/" + prefix + "_updates.bin", updates);
196+ ReadBin(exeDir + "/output/" + prefix + "_golden.bin", golden);
197+ CheckSize("base", base, OUTPUT_ELEMS);
198+ CheckSize("indices", indices, static_cast<size_t>(updateRows));
199+ CheckSize("updates", updates, static_cast<size_t>(updateRows) * INNER_DIM);
200+ CheckSize("golden", golden, OUTPUT_ELEMS);
201+ } catch (const std::exception& e) {
202+ std::cerr << "Read input/golden failed: " << e.what() << std::endl;
203+ return 1;
204+ }
205+ 
206+ const size_t indicesBytes = static_cast<size_t>(updateRows) * sizeof(IndexType);
207+ const size_t updatesBytes = static_cast<size_t>(updateRows) * INNER_DIM * sizeof(DataType);
208+ const size_t outputBytes = static_cast<size_t>(OUTPUT_ELEMS) * sizeof(DataType);
209+ 
210+ dIndices = nullptr;
211+ dUpdates = nullptr;
212+ dOut = nullptr;
213+ ret = aclrtMalloc(reinterpret_cast<void**>(&dIndices), indicesBytes, ACL_MEM_MALLOC_HUGE_FIRST);
214+ CHECK_RET(ret == ACL_SUCCESS, return ret);
215+ ret = aclrtMalloc(reinterpret_cast<void**>(&dUpdates), updatesBytes, ACL_MEM_MALLOC_HUGE_FIRST);
216+ CHECK_RET(ret == ACL_SUCCESS, aclrtFree(dIndices); return ret);
217+ ret = aclrtMalloc(reinterpret_cast<void**>(&dOut), outputBytes, ACL_MEM_MALLOC_HUGE_FIRST);
218+ CHECK_RET(ret == ACL_SUCCESS, aclrtFree(dIndices); aclrtFree(dUpdates); return ret);
219+ 
220+ ret = aclrtMemcpy(dIndices, indicesBytes, indices.data(), indicesBytes, ACL_MEMCPY_HOST_TO_DEVICE);
221+ CHECK_RET(ret == ACL_SUCCESS, aclrtFree(dIndices); aclrtFree(dUpdates); aclrtFree(dOut); return ret);
222+ ret = aclrtMemcpy(dUpdates, updatesBytes, updates.data(), updatesBytes, ACL_MEMCPY_HOST_TO_DEVICE);
223+ CHECK_RET(ret == ACL_SUCCESS, aclrtFree(dIndices); aclrtFree(dUpdates); aclrtFree(dOut); return ret);
224+ ret = aclrtMemcpy(dOut, outputBytes, base.data(), outputBytes, ACL_MEMCPY_HOST_TO_DEVICE);
225+ CHECK_RET(ret == ACL_SUCCESS, aclrtFree(dIndices); aclrtFree(dUpdates); aclrtFree(dOut); return ret);
226+ return ACL_SUCCESS;
227+}
228+ 
229+template <int kStep>
230+inline int RunSample(LaunchKernelFunc launchKernel, const std::string& sampleName, ScatterDataCase dataCase)
231+{
232+ int32_t deviceId = 0;
233+ aclrtStream stream = nullptr;
234+ IndexType* dIndices = nullptr;
235+ DataType* dUpdates = nullptr;
236+ DataType* dOut = nullptr;
237+ int32_t updateRows = 0;
238+ std::vector<DataType> golden;
239+ auto ret = PrepareInputs(dataCase, deviceId, stream, dIndices, dUpdates, dOut, updateRows, golden);
240+ CHECK_RET(ret == ACL_SUCCESS, return ret);
241+ 
242+ launchKernel(BLOCKS, stream, dIndices, dUpdates, dOut, updateRows, DST_ROWS, INNER_DIM);
243+ ret = aclrtSynchronizeStream(stream);
244+ CHECK_RET(ret == ACL_SUCCESS, aclrtFree(dIndices); aclrtFree(dUpdates); aclrtFree(dOut); return ret);
245+ 
246+ std::vector<DataType> hostOut(OUTPUT_ELEMS);
247+ const size_t outputBytes = static_cast<size_t>(OUTPUT_ELEMS) * sizeof(DataType);
248+ ret = aclrtMemcpy(hostOut.data(), outputBytes, dOut, outputBytes, ACL_MEMCPY_DEVICE_TO_HOST);
249+ CHECK_RET(ret == ACL_SUCCESS, aclrtFree(dIndices); aclrtFree(dUpdates); aclrtFree(dOut); return ret);
250+ 
251+ int errors = CompareInt("output", hostOut.data(), golden);
252+ std::cout << "[" << sampleName << "] step " << kStep << (errors == 0 ? " PASSED" : " FAILED") << std::endl;
253+ 
254+ aclrtFree(dIndices);
255+ aclrtFree(dUpdates);
256+ aclrtFree(dOut);
257+ aclrtDestroyStream(stream);
258+ aclrtResetDevice(deviceId);
259+ aclFinalize();
260+ return errors == 0 ? 0 : 1;
261+}
262+ 
263+} // namespace SimtScatterStorySample
264+ 
265+#endif // SIMT_SCATTER_STORY_SAMPLE_COMMON_H_
ASamples/2_Performance/simt_scatter_story/scripts/gen_data.py+95-0
@@ -0,0 +1,95 @@
1+#!/usr/bin/env python3
2+# ----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# 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+# Generate input and golden files for simt_scatter_story.
13+#
14+# Shape:
15+# - base/golden: [4096, 8] int32
16+# - unique updates: [4096, 8], unique destination rows
17+# - conflict updates: [8192, 8], sorted by (destination row, original update position)
18+import argparse
19+import os
20+ 
21+import numpy as np
22+ 
23+ 
24+DST_ROWS = 4096
25+INNER_DIM = 8
26+UNIQUE_UPDATES = 4096
27+CONFLICT_UPDATES = 8192
28+ 
29+ 
30+def make_base():
31+ data = np.arange(DST_ROWS * INNER_DIM, dtype=np.int32)
32+ return (data.reshape(DST_ROWS, INNER_DIM) * 3 - 17).astype(np.int32)
33+ 
34+ 
35+def make_updates(row_count, start):
36+ row_id = np.arange(row_count, dtype=np.int32).reshape(row_count, 1)
37+ col_id = np.arange(INNER_DIM, dtype=np.int32).reshape(1, INNER_DIM)
38+ return (start + row_id * 13 + col_id).astype(np.int32)
39+ 
40+ 
41+def write_case(output_dir, prefix, base, indices, updates):
42+ golden = base.copy()
43+ for row, dst in enumerate(indices):
44+ golden[int(dst), :] = updates[row, :]
45+ 
46+ input_dir = os.path.join(output_dir, "input")
47+ golden_dir = os.path.join(output_dir, "output")
48+ os.makedirs(input_dir, exist_ok=True)
49+ os.makedirs(golden_dir, exist_ok=True)
50+ indices.astype(np.int32).tofile(os.path.join(input_dir, f"{prefix}_indices.bin"))
51+ updates.astype(np.int32).tofile(os.path.join(input_dir, f"{prefix}_updates.bin"))
52+ golden.astype(np.int32).tofile(os.path.join(golden_dir, f"{prefix}_golden.bin"))
53+ 
54+ 
55+def main():
56+ parser = argparse.ArgumentParser()
57+ parser.add_argument("--output", default=".", help="directory to write input/ and output/")
58+ parser.add_argument("--seed", type=int, default=42, help="numpy random seed (default: 42)")
59+ args = parser.parse_args()
60+ 
61+ rng = np.random.default_rng(args.seed)
62+ base = make_base()
63+ 
64+ input_dir = os.path.join(args.output, "input")
65+ os.makedirs(input_dir, exist_ok=True)
66+ base.tofile(os.path.join(input_dir, "base.bin"))
67+ 
68+ unique_indices = rng.permutation(DST_ROWS).astype(np.int32)[:UNIQUE_UPDATES]
69+ unique_updates = make_updates(UNIQUE_UPDATES, 100000)
70+ write_case(args.output, "unique", base, unique_indices, unique_updates)
71+ 
72+ raw_indices = rng.integers(0, DST_ROWS, size=CONFLICT_UPDATES, dtype=np.int32)
73+ hot_count = CONFLICT_UPDATES // 2
74+ raw_indices[:hot_count] = rng.integers(0, 512, size=hot_count, dtype=np.int32)
75+ raw_updates = make_updates(CONFLICT_UPDATES, 200000)
76+ 
77+ original_pos = np.arange(CONFLICT_UPDATES, dtype=np.int32)
78+ order = np.lexsort((original_pos, raw_indices))
79+ conflict_indices = raw_indices[order]
80+ conflict_updates = raw_updates[order]
81+ 
82+ conflict_golden = base.copy()
83+ for row, dst in enumerate(raw_indices):
84+ conflict_golden[int(dst), :] = raw_updates[row, :]
85+ 
86+ conflict_indices.astype(np.int32).tofile(os.path.join(input_dir, "conflict_indices.bin"))
87+ conflict_updates.astype(np.int32).tofile(os.path.join(input_dir, "conflict_updates.bin"))
88+ 
89+ golden_dir = os.path.join(args.output, "output")
90+ os.makedirs(golden_dir, exist_ok=True)
91+ conflict_golden.astype(np.int32).tofile(os.path.join(golden_dir, "conflict_golden.bin"))
92+ 
93+ 
94+if __name__ == "__main__":
95+ main()
ASamples/2_Performance/simt_scatter_story/src/0_direct_unique.asc+68-0
@@ -0,0 +1,68 @@
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+ * Step 0: SIMT direct scatter for unique destination indices.
11+ */
12+ 
13+#include "kernel_operator.h"
14+#include "sample_common.h"
15+#include "simt_api/asc_simt.h"
16+ 
17+namespace {
18+ 
19+using namespace SimtScatterStorySample;
20+ 
21+__simt_vf__ LAUNCH_BOUND(SIMT_THREAD_NUM) inline void ScatterDirectUniqueSimt(
22+ int32_t coreId, int32_t coreNum, int32_t updateRows, int32_t dstRows, int32_t innerDim,
23+ __gm__ IndexType* indices, __gm__ DataType* updates, __gm__ DataType* y)
24+{
25+ auto rowTid = static_cast<int32_t>(AscendC::Simt::GetThreadIdx<0>());
26+ auto rowThreadNum = static_cast<int32_t>(AscendC::Simt::GetThreadNum<0>());
27+ auto colTid = static_cast<int32_t>(AscendC::Simt::GetThreadIdx<1>());
28+ auto colThreadNum = static_cast<int32_t>(AscendC::Simt::GetThreadNum<1>());
29+ 
30+ for (int32_t row = rowTid + coreId * rowThreadNum; row < updateRows; row += coreNum * rowThreadNum) {
31+ int32_t dst = indices[row];
32+ if (dst < 0 || dst >= dstRows) {
33+ continue;
34+ }
35+ int64_t srcOffset = static_cast<int64_t>(row) * innerDim;
36+ int64_t dstOffset = static_cast<int64_t>(dst) * innerDim;
37+ for (int32_t col = colTid; col < innerDim; col += colThreadNum) {
38+ y[dstOffset + col] = updates[srcOffset + col];
39+ }
40+ }
41+}
42+ 
43+__global__ __aicore__ __vector__ void simt_scatter_direct_unique_kernel(
44+ __gm__ IndexType* indices, __gm__ DataType* updates, __gm__ DataType* y, int32_t updateRows, int32_t dstRows,
45+ int32_t innerDim)
46+{
47+ AscendC::InitSocState();
48+ int32_t coreId = static_cast<int32_t>(AscendC::GetBlockIdx());
49+ int32_t coreNum = static_cast<int32_t>(AscendC::GetBlockNum());
50+ AscendC::Simt::VF_CALL<ScatterDirectUniqueSimt>(
51+ AscendC::Simt::Dim3{SIMT_X_THREAD_NUM, SIMT_Y_THREAD_NUM, 1}, coreId, coreNum, updateRows, dstRows, innerDim, indices,
52+ updates, y);
53+ AscendC::PipeBarrier<PIPE_ALL>();
54+}
55+ 
56+void LaunchScatterKernel(uint32_t blocks, aclrtStream stream, IndexType* indices, DataType* updates, DataType* y,
57+ int32_t updateRows, int32_t dstRows, int32_t innerDim)
58+{
59+ simt_scatter_direct_unique_kernel<<<blocks, nullptr, stream>>>(indices, updates, y, updateRows, dstRows, innerDim);
60+}
61+ 
62+} // namespace
63+ 
64+int main()
65+{
66+ return SimtScatterStorySample::RunSample<0>(
67+ LaunchScatterKernel, "0_direct_unique", SimtScatterStorySample::ScatterDataCase::UNIQUE);
68+}
ASamples/2_Performance/simt_scatter_story/src/1_grouped_conflict.asc+67-0
@@ -0,0 +1,67 @@
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+ * Step 1: conflict-safe grouped scatter. Only the last row in each destination group writes.
11+ */
12+ 
13+#include "kernel_operator.h"
14+#include "sample_common.h"
15+#include "simt_api/asc_simt.h"
16+ 
17+namespace {
18+ 
19+using namespace SimtScatterStorySample;
20+ 
21+__simt_vf__ LAUNCH_BOUND(SIMT_THREAD_NUM) inline void ScatterGroupedConflictSimt(
22+ int32_t coreId, int32_t coreNum, int32_t updateRows, int32_t dstRows, int32_t innerDim,
23+ __gm__ IndexType* indices, __gm__ DataType* updates, __gm__ DataType* y)
24+{
25+ auto tid = static_cast<int32_t>(AscendC::Simt::GetThreadIdx());
26+ auto threadNum = static_cast<int32_t>(AscendC::Simt::GetThreadNum());
27+ 
28+ for (int32_t row = tid + coreId * threadNum; row < updateRows; row += coreNum * threadNum) {
29+ int32_t dst = indices[row];
30+ bool isLastInGroup = (row == updateRows - 1) || (dst != indices[row + 1]);
31+ if (!isLastInGroup || dst < 0 || dst >= dstRows) {
32+ continue;
33+ }
34+ int64_t srcOffset = static_cast<int64_t>(row) * innerDim;
35+ int64_t dstOffset = static_cast<int64_t>(dst) * innerDim;
36+ for (int32_t col = 0; col < innerDim; ++col) {
37+ y[dstOffset + col] = updates[srcOffset + col];
38+ }
39+ }
40+}
41+ 
42+__global__ __aicore__ __vector__ void simt_scatter_grouped_conflict_kernel(
43+ __gm__ IndexType* indices, __gm__ DataType* updates, __gm__ DataType* y, int32_t updateRows, int32_t dstRows,
44+ int32_t innerDim)
45+{
46+ AscendC::InitSocState();
47+ int32_t coreId = static_cast<int32_t>(AscendC::GetBlockIdx());
48+ int32_t coreNum = static_cast<int32_t>(AscendC::GetBlockNum());
49+ AscendC::Simt::VF_CALL<ScatterGroupedConflictSimt>(
50+ AscendC::Simt::Dim3{SIMT_THREAD_NUM, 1, 1}, coreId, coreNum, updateRows, dstRows, innerDim, indices, updates, y);
51+ AscendC::PipeBarrier<PIPE_ALL>();
52+}
53+ 
54+void LaunchScatterKernel(uint32_t blocks, aclrtStream stream, IndexType* indices, DataType* updates, DataType* y,
55+ int32_t updateRows, int32_t dstRows, int32_t innerDim)
56+{
57+ simt_scatter_grouped_conflict_kernel<<<blocks, nullptr, stream>>>(indices, updates, y, updateRows, dstRows,
58+ innerDim);
59+}
60+ 
61+} // namespace
62+ 
63+int main()
64+{
65+ return SimtScatterStorySample::RunSample<1>(
66+ LaunchScatterKernel, "1_grouped_conflict", SimtScatterStorySample::ScatterDataCase::CONFLICT);
67+}
ASamples/2_Performance/simt_scatter_story/src/2_grouped_conflict_2d.asc+70-0
@@ -0,0 +1,70 @@
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+ * Step 2: conflict-safe grouped scatter with a 2-D SIMT layout.
11+ */
12+ 
13+#include "kernel_operator.h"
14+#include "sample_common.h"
15+#include "simt_api/asc_simt.h"
16+ 
17+namespace {
18+ 
19+using namespace SimtScatterStorySample;
20+ 
21+__simt_vf__ LAUNCH_BOUND(SIMT_THREAD_NUM) inline void ScatterGroupedConflict2dSimt(
22+ int32_t coreId, int32_t coreNum, int32_t updateRows, int32_t dstRows, int32_t innerDim,
23+ __gm__ IndexType* indices, __gm__ DataType* updates, __gm__ DataType* y)
24+{
25+ auto rowTid = static_cast<int32_t>(AscendC::Simt::GetThreadIdx<0>());
26+ auto rowThreadNum = static_cast<int32_t>(AscendC::Simt::GetThreadNum<0>());
27+ auto colTid = static_cast<int32_t>(AscendC::Simt::GetThreadIdx<1>());
28+ auto colThreadNum = static_cast<int32_t>(AscendC::Simt::GetThreadNum<1>());
29+ 
30+ for (int32_t row = rowTid + coreId * rowThreadNum; row < updateRows; row += coreNum * rowThreadNum) {
31+ int32_t dst = indices[row];
32+ bool isLastInGroup = (row == updateRows - 1) || (dst != indices[row + 1]);
33+ if (!isLastInGroup || dst < 0 || dst >= dstRows) {
34+ continue;
35+ }
36+ int64_t srcOffset = static_cast<int64_t>(row) * innerDim;
37+ int64_t dstOffset = static_cast<int64_t>(dst) * innerDim;
38+ for (int32_t col = colTid; col < innerDim; col += colThreadNum) {
39+ y[dstOffset + col] = updates[srcOffset + col];
40+ }
41+ }
42+}
43+ 
44+__global__ __aicore__ __vector__ void simt_scatter_grouped_conflict_2d_kernel(
45+ __gm__ IndexType* indices, __gm__ DataType* updates, __gm__ DataType* y, int32_t updateRows, int32_t dstRows,
46+ int32_t innerDim)
47+{
48+ AscendC::InitSocState();
49+ int32_t coreId = static_cast<int32_t>(AscendC::GetBlockIdx());
50+ int32_t coreNum = static_cast<int32_t>(AscendC::GetBlockNum());
51+ AscendC::Simt::VF_CALL<ScatterGroupedConflict2dSimt>(
52+ AscendC::Simt::Dim3{SIMT_X_THREAD_NUM, SIMT_Y_THREAD_NUM, 1}, coreId, coreNum, updateRows, dstRows, innerDim, indices,
53+ updates, y);
54+ AscendC::PipeBarrier<PIPE_ALL>();
55+}
56+ 
57+void LaunchScatterKernel(uint32_t blocks, aclrtStream stream, IndexType* indices, DataType* updates, DataType* y,
58+ int32_t updateRows, int32_t dstRows, int32_t innerDim)
59+{
60+ simt_scatter_grouped_conflict_2d_kernel<<<blocks, nullptr, stream>>>(indices, updates, y, updateRows, dstRows,
61+ innerDim);
62+}
63+ 
64+} // namespace
65+ 
66+int main()
67+{
68+ return SimtScatterStorySample::RunSample<2>(
69+ LaunchScatterKernel, "2_grouped_conflict_2d", SimtScatterStorySample::ScatterDataCase::CONFLICT);
70+}