Star67
10
代码
代码
Issues1
Pull Requests5
Actions
项目讨论
Wiki
分析
Star67
10
已合并
新增 Reformer LSH 桶排序算子 #30
海阔天空创建于 15 天前
新增 Reformer LSH 桶排序算子 #30
已合并
海阔天空创建于 15 天前
从haikuo/mat-chem-sim-predfeature/reformer-lsh-bucket-sort合入到cann/mat-chem-sim-predmaster
共 28 个文件变更+1800-0
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/CMakeLists.txt+49-0
@@ -0,0 +1,49 @@
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+cmake_minimum_required(VERSION 3.16)
12+project(reformer_lsh_bucket_sort LANGUAGES CXX)
13+set(CMAKE_CXX_STANDARD 17)
14+set(CMAKE_CXX_STANDARD_REQUIRED ON)
15+ 
16+if(NOT DEFINED ENV{ASCEND_CANN_PACKAGE_PATH})
17+ if(DEFINED ENV{ASCEND_TOOLKIT_HOME})
18+ set(ASCEND_CANN_PACKAGE_PATH $ENV{ASCEND_TOOLKIT_HOME})
19+ else()
20+ set(ASCEND_CANN_PACKAGE_PATH "/usr/local/Ascend/ascend-toolkit/latest")
21+ endif()
22+else()
23+ set(ASCEND_CANN_PACKAGE_PATH $ENV{ASCEND_CANN_PACKAGE_PATH})
24+endif()
25+set(SOC_VERSION "Ascend910B1" CACHE STRING "SOC version")
26+ 
27+foreach(candidate
28+ "${ASCEND_CANN_PACKAGE_PATH}/tikcpp/ascendc_kernel_cmake"
29+ "${ASCEND_CANN_PACKAGE_PATH}/x86_64-linux/tikcpp/ascendc_kernel_cmake"
30+ "${ASCEND_CANN_PACKAGE_PATH}/aarch64-linux/tikcpp/ascendc_kernel_cmake")
31+ if(EXISTS "${candidate}")
32+ set(ASCENDC_CMAKE_PATH "${candidate}")
33+ break()
34+ endif()
35+endforeach()
36+if(NOT ASCENDC_CMAKE_PATH)
37+ message(FATAL_ERROR "Cannot find ascendc_kernel_cmake")
38+endif()
39+include(${ASCENDC_CMAKE_PATH}/ascendc.cmake)
40+ 
41+ascendc_library(reformer_lsh_bucket_sort_kernel_lib SHARED op_kernel/reformer_lsh_bucket_sort_kernel.cpp)
42+ascendc_include_directories(reformer_lsh_bucket_sort_kernel_lib PRIVATE
43+ ${CMAKE_CURRENT_SOURCE_DIR}/op_kernel ${CMAKE_CURRENT_SOURCE_DIR}/op_host)
44+add_library(reformer_lsh_bucket_sort_op_def SHARED
45+ op_host/reformer_lsh_bucket_sort_host.cpp
46+ op_host/reformer_lsh_bucket_sort_def.cpp)
47+target_include_directories(reformer_lsh_bucket_sort_op_def PRIVATE
48+ ${ASCEND_CANN_PACKAGE_PATH}/include ${CMAKE_CURRENT_SOURCE_DIR}/op_host)
49+install(FILES msopgen/reformer_lsh_bucket_sort_msopgen.json DESTINATION msopgen)
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/README.md+57-0
@@ -0,0 +1,57 @@
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+# ReformerLshBucketSort 算子
12+ 
13+`ReformerLshBucketSort` 为 Reformer LSH Attention 提供稳定桶排序。算子将框架侧两次排序及逆置换构造合并为一个 Ascend C kernel，直接输出排序键、原位置索引和逆索引。
14+ 
15+## 模型接入点
16+ 
17+TSLib 调用链为 `layers/SelfAttention_Family.py::ReformerLayer` -> `reformer_pytorch.reformer_pytorch.LSHAttention.forward`。本算子替换 `sort_key_val(buckets_and_t, ticker)` 及随后逆置换构造，不改变 Reformer 其余计算。
18+ 
19+## 输入输出
20+ 
21+`keys[R,N]` 中每个键编码为 `bucket_id * sequence_length + position`。
22+ 
23+| 名称 | 类型 | Shape | 说明 |
24+|---|---|---|---|
25+| `keys` | int64 | `[R,N]` | 编码后的桶键 |
26+| `sorted_keys` | int64 | `[R,N]` | 稳定排序后的键 |
27+| `sticker` | int64 | `[R,N]` | 排序位置对应的原位置 |
28+| `inverse` | int64 | `[R,N]` | 原位置对应的排序位置 |
29+ 
30+支持非空 rank-2 ND 输入；行数和 `sequence_length` 必须为正且可由 uint32 表示，单行长度不超过 int32 上限；`1 <= total_buckets <= 4096`。键值范围 `[0, sequence_length * total_buckets)` 由调用方保证，adapter 对不支持的 dtype、layout 或 shape 回退到框架实现。
31+ 
32+## 构建与测试
33+ 
34+```bash
35+source /usr/local/Ascend/ascend-toolkit/set_env.sh
36+msopgen gen -i msopgen/reformer_lsh_bucket_sort_msopgen.json -f aclnn \
37+ -c ai_core-ascend910b -out build/msopgen_reformer_lsh_bucket_sort -lan cpp
38+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DSOC_VERSION=Ascend910B1
39+cmake --build build -j2
40+python -m unittest discover -s tests -p "test_*.py" -v
41+python -m compileall -q integration reference tests
42+bash tests/run_packaged_gate.sh /tmp/reformer_lsh_bucket_sort_gate 6
43+```
44+ 
45+已在 Node202、CANN 8.1 RC1、Ascend 910B3 上完成 msopgen、独立构建、ACLNN smoke 和精度回读。当前本地 Python 套件共 7 项；官方 SCA、CLA 和仓库 CI 仅能在 MR 阶段执行。
46+ 
47+## 性能价值
48+ 
49+`B` 表示推理 batch size，`L` 表示输入序列长度。计时均为 NPU resident、同步 wall time，基线为相同语义下最快的框架表达式。
50+ 
51+| 范围 | 框架基线 | 自定义算子 | 延迟降低 | 加速比 |
52+|---|---:|---:|---:|---:|
53+| 算子组件，B32/L336 | 37.4029 ms | 0.5987 ms | 98.40% | 62.472x |
54+| 单算子替换完整模型，B32/L336 | 55.1197 ms | 17.5769 ms | 68.11% | 3.136x |
55+| ETTh1 全测试集 E2E，B32 | 1558.82 ms | 682.56 ms | 56.21% | 2.284x |
56+ 
57+E2E 覆盖 ETTh1 官方测试集全部 2,857 个滑窗、7 个特征、输入长度 96、预测长度 24，共 479,976 个预测值。排序结果、索引及最终预测均逐元素完全一致。详细口径见 [性能说明](docs/benchmark.md) 和 [测试报告](docs/test_report.md)。
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/algorithm.md+45-0
@@ -0,0 +1,45 @@
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+# ReformerLshBucketSort 算法说明
12+ 
13+## 问题定义
14+ 
15+输入 `keys[r,n] = bucket_id * sequence_length + position`。对每一行执行稳定升序排序，并同时输出排序置换及其逆置换：
16+ 
17+```text
18+sorted_keys[r,j] = keys[r,sticker[r,j]]
19+inverse[r,sticker[r,j]] = j
20+```
21+ 
22+稳定性要求相同桶键保持原始先后顺序，这是 Reformer LSH 分桶后正确还原 token 顺序的必要条件。
23+ 
24+## 模型接入
25+ 
26+TSLib 的 `ReformerLayer` 最终调用 `LSHAttention.forward`。框架路径使用 `sort_key_val` 生成排序结果，再额外构造 inverse。本算子一次遍历同时产生三项输出，替换范围只覆盖该排序阶段。
27+ 
28+## NPU 实现
29+ 
30+每个逻辑行分配给一个 vector core。kernel 在 Local Memory 中维护不超过 4096 个桶的计数和前缀位置，然后完成：
31+ 
32+1. 清零桶计数；
33+2. 扫描该行键值并统计各桶数量；
34+3. 计算稳定写入起点；
35+4. 再次扫描输入，写出 `sorted_keys`、`sticker` 和 `inverse`。
36+ 
37+行间互不依赖，无跨核共享状态；所有全局地址偏移使用 uint64 计算。相较两次框架排序及逆索引构造，该方案减少 kernel launch、中间 tensor 和重复遍历。
38+ 
39+## 边界条件
40+ 
41+- `keys` 为非空 rank-2、ND、int64 tensor；
42+- 行数和 `sequence_length` 为正且不超过 uint32 上限，单行长度不超过 int32 上限；
43+- `1 <= total_buckets <= 4096`；
44+- 调用方保证编码键位于合法范围；
45+- 不满足合同时，由框架 adapter 回退到原实现。
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/api_reference.md+41-0
@@ -0,0 +1,41 @@
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+# ReformerLshBucketSort API 参考
12+ 
13+## 算子信息
14+ 
15+| 项目 | 内容 |
16+|---|---|
17+| OpDef | `ReformerLshBucketSort` |
18+| ACLNN API | `aclnnReformerLshBucketSort` |
19+| 输入 | `keys: int64 [R,N]` |
20+| 属性 | `sequence_length: int64`，`total_buckets: int64` |
21+| 输出 | `sorted_keys`、`sticker`、`inverse`，均为 int64 `[R,N]` |
22+| Format | ND |
23+| Workspace | 当前 tiling 实现为 0 Byte |
24+| SoC 注册 | `ascend910b` |
25+ 
26+## 参数约束与错误条件
27+ 
28+Host tiling 会拒绝空 tensor、非 rank-2 输入、无法由 uint32 表示的行数或 `sequence_length`、超过 int32 上限的单行长度，以及不在 `[1,4096]` 内的 `total_buckets`。编码键范围属于调用方前置条件，CPU reference 覆盖其合法与非法用例。
29+ 
30+integration adapter 仅对 NPU resident、contiguous、int64 且满足上述 shape 合同的输入调用 ACLNN；其余情况调用传入的 framework fallback。
31+ 
32+## ACLNN 工程生成
33+ 
34+```bash
35+source /usr/local/Ascend/ascend-toolkit/set_env.sh
36+msopgen gen -i msopgen/reformer_lsh_bucket_sort_msopgen.json -f aclnn \
37+ -c ai_core-ascend910b -out build/msopgen_reformer_lsh_bucket_sort -lan cpp
38+bash tests/run_packaged_gate.sh /tmp/reformer_lsh_bucket_sort_gate 6
39+```
40+ 
41+`run_packaged_gate.sh` 会生成 ACLNN 工程、覆盖本目录的 operator-specific Host/Kernel 源码、构建自定义包并执行 device smoke。
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/benchmark.md+48-0
@@ -0,0 +1,48 @@
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+# ReformerLshBucketSort 性能说明
12+ 
13+## 环境与计时口径
14+ 
15+- 设备：Ascend 910B3，研究环境物理 NPU 6；
16+- 软件：CANN 8.1 RC1，FP32 模型数据与 int64 索引；
17+- Shape：B4/B16/B32，L96/L336，预测长度 24；
18+- 算子和单模型计时：3 次预热，9 次 NPU 同步计时取中位数；
19+- 输入常驻 NPU，不计进程启动及 H2D/D2H；
20+- 性能基线：相同语义下最快的 eager/TorchAir 兼容框架表达式；
21+- CPU reference 用于语义正确性，不作为 NPU 加速比分母；
22+- 自定义路径：msopgen 生成的 `GetWorkspaceSize + aclnn...`，不是研究 launcher。
23+ 
24+`B` 是推理 batch size，`L` 是输入序列长度。
25+ 
26+## 算子组件
27+ 
28+| Shape | 框架基线 | ACLNN | 延迟降低 | 加速比 | 最大误差 |
29+|---|---:|---:|---:|---:|---:|
30+| B4/L96 | 1.4896 ms | 0.1854 ms | 87.55% | 8.034x | 0 |
31+| B16/L96 | 5.5665 ms | 0.1941 ms | 96.51% | 28.680x | 0 |
32+| B32/L96 | 10.7926 ms | 0.2741 ms | 97.46% | 39.378x | 0 |
33+| B4/L336 | 4.7998 ms | 0.2307 ms | 95.19% | 20.802x | 0 |
34+| B16/L336 | 18.6948 ms | 0.3623 ms | 98.06% | 51.604x | 0 |
35+| B32/L336 | 37.4029 ms | 0.5987 ms | 98.40% | 62.472x | 0 |
36+ 
37+六档实测均启用自定义路径。
38+ 
39+## 模型与 checkpoint
40+ 
41+| 范围 | 框架基线 | 单算子替换 | 延迟降低 | 加速比 | 最大预测差 |
42+|---|---:|---:|---:|---:|---:|
43+| Reformer B32/L336 单 batch | 55.1197 ms | 17.5769 ms | 68.11% | 3.136x | 0 |
44+| ETTh1 全测试集，B32 | 1558.82 ms | 682.56 ms | 56.21% | 2.284x | 0 |
45+ 
46+checkpoint 闭环使用 ETTh1 官方测试 split：2,857 个重叠窗口、7 个特征、输入 96、预测 24，共 479,976 个预测值和 90 个推理 batch。两条路径严格加载同一 checkpoint，只独立替换本算子，并各自分配输出存储。
47+ 
48+机器可读证据位于 `docs/evidence/`，分别记录六档原始样本、单模型消融、checkpoint 全量指标和多 stream 结果。
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/evidence/checkpoint_etth1.json+50-0
@@ -0,0 +1,50 @@
1+{
2+ "path": "msopgen_packaged_aclnn_checkpoint_ablation",
3+ "device": "Ascend 910B3",
4+ "model": "Reformer",
5+ "operator": "bucket",
6+ "dataset": "ETTh1 official test split",
7+ "test_windows": 2857,
8+ "sequence_length": 96,
9+ "prediction_length": 24,
10+ "channels": 7,
11+ "checkpoint_sha256": "bc60bb3ef492705eeee68d0cb856708a9a8d7a884fa959a5aec807a6de4b034a",
12+ "dataset_sha256": "f18de3ad269cef59bb07b5438d79bb3042d3be49bdeecf01c1cd6d29695ee066",
13+ "output_policy": "allocate",
14+ "result": {
15+ "batch_size": 32,
16+ "windows": 2857,
17+ "wall_seconds_all_variants_and_transfers": 2.529852165025659,
18+ "metrics": {
19+ "original_eager": {
20+ "inference_ms": 1558.8212383445352,
21+ "batches": 90,
22+ "normalized_mse": 0.7956262284549774,
23+ "normalized_mae": 0.6287206820210395,
24+ "original_scale_mse": 20.725634759928102,
25+ "original_scale_mae": 2.6398817723275485,
26+ "mean_batch_ms": 17.320235981605947,
27+ "windows_per_second_model_only": 1832.795146564803,
28+ "speedup_vs_framework_model_time": 1.0
29+ },
30+ "bucket_packaged": {
31+ "inference_ms": 682.5618179282174,
32+ "batches": 90,
33+ "normalized_mse": 0.7956262284549774,
34+ "normalized_mae": 0.6287206820210395,
35+ "original_scale_mse": 20.725634759928102,
36+ "original_scale_mae": 2.6398817723275485,
37+ "mean_batch_ms": 7.584020199202415,
38+ "windows_per_second_model_only": 4185.701463161042,
39+ "speedup_vs_framework_model_time": 2.2837803073661984
40+ }
41+ },
42+ "prediction_differences_vs_framework": {
43+ "bucket_packaged": {
44+ "max_abs_diff": 0.0,
45+ "mean_abs_diff": 0.0,
46+ "rmse_diff": 0.0
47+ }
48+ }
49+ }
50+}
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/evidence/component_matrix.json+377-0
@@ -0,0 +1,377 @@
1+{
2+ "path": "msopgen_packaged_aclnn",
3+ "device": "Ascend 910B3",
4+ "batches": [
5+ 4,
6+ 16,
7+ 32
8+ ],
9+ "lengths": [
10+ 96,
11+ 336
12+ ],
13+ "warmup": 3,
14+ "repeat": 9,
15+ "results": [
16+ {
17+ "operator": "lsh_bucket_sort",
18+ "shape": [
19+ 16,
20+ 480
21+ ],
22+ "framework_two_sorts": {
23+ "median_ms": 1.4895560452714562,
24+ "mean_ms": 1.496461543461515,
25+ "min_ms": 1.4755059964954853,
26+ "max_ms": 1.5585769433528185,
27+ "samples_ms": [
28+ 1.5585769433528185,
29+ 1.5187360113486648,
30+ 1.492515904828906,
31+ 1.4956759987398982,
32+ 1.4895560452714562,
33+ 1.476484932936728,
34+ 1.478336052969098,
35+ 1.4755059964954853,
36+ 1.4827660052105784
37+ ]
38+ },
39+ "custom": {
40+ "median_ms": 0.18541200552135706,
41+ "mean_ms": 0.23448030050430033,
42+ "min_ms": 0.15564204659312963,
43+ "max_ms": 0.4335750127211213,
44+ "samples_ms": [
45+ 0.1870519481599331,
46+ 0.18541200552135706,
47+ 0.39389391895383596,
48+ 0.15771191101521254,
49+ 0.15564204659312963,
50+ 0.1686519244685769,
51+ 0.15568092931061983,
52+ 0.4335750127211213,
53+ 0.27270300779491663
54+ ]
55+ },
56+ "correctness": {
57+ "max_abs_diff": 0.0,
58+ "per_output_max_abs_diff": [
59+ 0.0,
60+ 0.0,
61+ 0.0
62+ ],
63+ "per_output_mean_abs_diff": [
64+ 0.0,
65+ 0.0,
66+ 0.0
67+ ],
68+ "all_equal": true
69+ },
70+ "speedup": 8.033762652439886,
71+ "case": {
72+ "B": 4,
73+ "L": 96
74+ }
75+ },
76+ {
77+ "operator": "lsh_bucket_sort",
78+ "shape": [
79+ 64,
80+ 480
81+ ],
82+ "framework_two_sorts": {
83+ "median_ms": 5.56647009216249,
84+ "mean_ms": 5.5827693625663715,
85+ "min_ms": 5.514518939889967,
86+ "max_ms": 5.799081060104072,
87+ "samples_ms": [
88+ 5.799081060104072,
89+ 5.569719010964036,
90+ 5.56647009216249,
91+ 5.597949028015137,
92+ 5.57322904933244,
93+ 5.524399108253419,
94+ 5.5378490360453725,
95+ 5.514518939889967,
96+ 5.561708938330412
97+ ]
98+ },
99+ "custom": {
100+ "median_ms": 0.19409193191677332,
101+ "mean_ms": 0.19618767934540907,
102+ "min_ms": 0.19191193860024214,
103+ "max_ms": 0.21160300821065903,
104+ "samples_ms": [
105+ 0.21160300821065903,
106+ 0.19898207392543554,
107+ 0.19512197468429804,
108+ 0.1920120557770133,
109+ 0.19191193860024214,
110+ 0.19239203538745642,
111+ 0.19379204604774714,
112+ 0.19409193191677332,
113+ 0.19578204955905676
114+ ]
115+ },
116+ "correctness": {
117+ "max_abs_diff": 0.0,
118+ "per_output_max_abs_diff": [
119+ 0.0,
120+ 0.0,
121+ 0.0
122+ ],
123+ "per_output_mean_abs_diff": [
124+ 0.0,
125+ 0.0,
126+ 0.0
127+ ],
128+ "all_equal": true
129+ },
130+ "speedup": 28.679554256533415,
131+ "case": {
132+ "B": 16,
133+ "L": 96
134+ }
135+ },
136+ {
137+ "operator": "lsh_bucket_sort",
138+ "shape": [
139+ 128,
140+ 480
141+ ],
142+ "framework_two_sorts": {
143+ "median_ms": 10.792564949952066,
144+ "mean_ms": 10.869090005548465,
145+ "min_ms": 10.75938402209431,
146+ "max_ms": 11.03277807123959,
147+ "samples_ms": [
148+ 11.03277807123959,
149+ 11.031217058189213,
150+ 10.977786965668201,
151+ 10.77522395644337,
152+ 10.75938402209431,
153+ 10.915726074017584,
154+ 10.777115006931126,
155+ 10.760013945400715,
156+ 10.792564949952066
157+ ]
158+ },
159+ "custom": {
160+ "median_ms": 0.274072983302176,
161+ "mean_ms": 0.2743184773458375,
162+ "min_ms": 0.2699530450627208,
163+ "max_ms": 0.28260296676307917,
164+ "samples_ms": [
165+ 0.28260296676307917,
166+ 0.27832307387143373,
167+ 0.2745829988270998,
168+ 0.27491303626447916,
169+ 0.2719020703807473,
170+ 0.2699530450627208,
171+ 0.27045304886996746,
172+ 0.27206307277083397,
173+ 0.274072983302176
174+ ]
175+ },
176+ "correctness": {
177+ "max_abs_diff": 0.0,
178+ "per_output_max_abs_diff": [
179+ 0.0,
180+ 0.0,
181+ 0.0
182+ ],
183+ "per_output_mean_abs_diff": [
184+ 0.0,
185+ 0.0,
186+ 0.0
187+ ],
188+ "all_equal": true
189+ },
190+ "speedup": 39.378434240097455,
191+ "case": {
192+ "B": 32,
193+ "L": 96
194+ }
195+ },
196+ {
197+ "operator": "lsh_bucket_sort",
198+ "shape": [
199+ 16,
200+ 1440
201+ ],
202+ "framework_two_sorts": {
203+ "median_ms": 4.7998009249567986,
204+ "mean_ms": 4.812526684771809,
205+ "min_ms": 4.780791001394391,
206+ "max_ms": 4.876471008174121,
207+ "samples_ms": [
208+ 4.876471008174121,
209+ 4.864162066951394,
210+ 4.808421013876796,
211+ 4.791041021235287,
212+ 4.791671060957015,
213+ 4.780791001394391,
214+ 4.7998009249567986,
215+ 4.793841042555869,
216+ 4.806541022844613
217+ ]
218+ },
219+ "custom": {
220+ "median_ms": 0.2307329559698701,
221+ "mean_ms": 0.2331703176928891,
222+ "min_ms": 0.22663292475044727,
223+ "max_ms": 0.24421291891485453,
224+ "samples_ms": [
225+ 0.24421291891485453,
226+ 0.23657199926674366,
227+ 0.23703300394117832,
228+ 0.23021199740469456,
229+ 0.2307329559698701,
230+ 0.22867205552756786,
231+ 0.2294030273333192,
232+ 0.2350619761273265,
233+ 0.22663292475044727
234+ ]
235+ },
236+ "correctness": {
237+ "max_abs_diff": 0.0,
238+ "per_output_max_abs_diff": [
239+ 0.0,
240+ 0.0,
241+ 0.0
242+ ],
243+ "per_output_mean_abs_diff": [
244+ 0.0,
245+ 0.0,
246+ 0.0
247+ ],
248+ "all_equal": true
249+ },
250+ "speedup": 20.802407288465428,
251+ "case": {
252+ "B": 4,
253+ "L": 336
254+ }
255+ },
256+ {
257+ "operator": "lsh_bucket_sort",
258+ "shape": [
259+ 64,
260+ 1440
261+ ],
262+ "framework_two_sorts": {
263+ "median_ms": 18.694829079322517,
264+ "mean_ms": 18.689772993740107,
265+ "min_ms": 18.616786925122142,
266+ "max_ms": 18.75384000595659,
267+ "samples_ms": [
268+ 18.734049052000046,
269+ 18.66802794393152,
270+ 18.730918993242085,
271+ 18.75384000595659,
272+ 18.616786925122142,
273+ 18.694829079322517,
274+ 18.695898936130106,
275+ 18.672468024306,
276+ 18.64113798364997
277+ ]
278+ },
279+ "custom": {
280+ "median_ms": 0.3622740041464567,
281+ "mean_ms": 0.3645883407443762,
282+ "min_ms": 0.3578329924494028,
283+ "max_ms": 0.3801640123128891,
284+ "samples_ms": [
285+ 0.3801640123128891,
286+ 0.3643740201368928,
287+ 0.3622740041464567,
288+ 0.36325398832559586,
289+ 0.3754140343517065,
290+ 0.35980402026325464,
291+ 0.3593940054997802,
292+ 0.35878398921340704,
293+ 0.3578329924494028
294+ ]
295+ },
296+ "correctness": {
297+ "max_abs_diff": 0.0,
298+ "per_output_max_abs_diff": [
299+ 0.0,
300+ 0.0,
301+ 0.0
302+ ],
303+ "per_output_mean_abs_diff": [
304+ 0.0,
305+ 0.0,
306+ 0.0
307+ ],
308+ "all_equal": true
309+ },
310+ "speedup": 51.604114193533874,
311+ "case": {
312+ "B": 16,
313+ "L": 336
314+ }
315+ },
316+ {
317+ "operator": "lsh_bucket_sort",
318+ "shape": [
319+ 128,
320+ 1440
321+ ],
322+ "framework_two_sorts": {
323+ "median_ms": 37.402927060611546,
324+ "mean_ms": 37.45248010899458,
325+ "min_ms": 37.2820160118863,
326+ "max_ms": 37.80223103240132,
327+ "samples_ms": [
328+ 37.80223103240132,
329+ 37.71809092722833,
330+ 37.2820160118863,
331+ 37.46989800129086,
332+ 37.39768802188337,
333+ 37.30176598764956,
334+ 37.41501795593649,
335+ 37.402927060611546,
336+ 37.28268598206341
337+ ]
338+ },
339+ "custom": {
340+ "median_ms": 0.5987160839140415,
341+ "mean_ms": 0.6009386723033256,
342+ "min_ms": 0.5923260468989611,
343+ "max_ms": 0.6122470367699862,
344+ "samples_ms": [
345+ 0.6122470367699862,
346+ 0.5987160839140415,
347+ 0.6015870021656156,
348+ 0.5937359528616071,
349+ 0.5968759069219232,
350+ 0.6080070743337274,
351+ 0.6080459570512176,
352+ 0.596906989812851,
353+ 0.5923260468989611
354+ ]
355+ },
356+ "correctness": {
357+ "max_abs_diff": 0.0,
358+ "per_output_max_abs_diff": [
359+ 0.0,
360+ 0.0,
361+ 0.0
362+ ],
363+ "per_output_mean_abs_diff": [
364+ 0.0,
365+ 0.0,
366+ 0.0
367+ ],
368+ "all_equal": true
369+ },
370+ "speedup": 62.47189288133695,
371+ "case": {
372+ "B": 32,
373+ "L": 336
374+ }
375+ }
376+ ]
377+}
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/evidence/model_ablation.json+55-0
@@ -0,0 +1,55 @@
1+{
2+ "path": "msopgen_packaged_aclnn_model_ablation",
3+ "device": "Ascend 910B3",
4+ "output_policy": "allocate",
5+ "result": {
6+ "model": "Reformer",
7+ "operator": "bucket",
8+ "shape": {
9+ "B": 32,
10+ "L": 336
11+ },
12+ "baseline": {
13+ "median_ms": 55.11973495595157,
14+ "mean_ms": 55.194207454203735,
15+ "min_ms": 54.9688640749082,
16+ "max_ms": 55.76317198574543,
17+ "samples_ms": [
18+ 55.76317198574543,
19+ 55.07133505307138,
20+ 55.03459507599473,
21+ 55.11973495595157,
22+ 55.119946016930044,
23+ 54.9688640749082,
24+ 55.080504971556365,
25+ 55.23228691890836,
26+ 55.35742803476751
27+ ]
28+ },
29+ "packaged_custom": {
30+ "median_ms": 17.576936981640756,
31+ "mean_ms": 17.792527990726132,
32+ "min_ms": 17.486906028352678,
33+ "max_ms": 19.393846043385565,
34+ "samples_ms": [
35+ 17.533286940306425,
36+ 17.535396036691964,
37+ 17.486906028352678,
38+ 19.393846043385565,
39+ 17.828649026341736,
40+ 17.562886932864785,
41+ 17.615076969377697,
42+ 17.599766957573593,
43+ 17.576936981640756
44+ ],
45+ "max_abs_diff": 0.0,
46+ "mean_abs_diff": 0.0,
47+ "rmse": 0.0,
48+ "relative_l2": 0.0,
49+ "reference_abs_max": 2.249037027359009,
50+ "allclose_rtol_1e-5_atol_1e-6": true,
51+ "allclose_rtol_1e-4_atol_1e-5": true
52+ },
53+ "speedup": 3.1359124182742733
54+ }
55+}
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/evidence/multistream.json+17-0
@@ -0,0 +1,17 @@
1+{
2+ "operator": "bucket",
3+ "iterations_per_stream": 20,
4+ "workspace_size": 0,
5+ "distinct_output_storage": true,
6+ "stream_a_max_abs_diff": [
7+ 0.0,
8+ 0.0,
9+ 0.0
10+ ],
11+ "stream_b_max_abs_diff": [
12+ 0.0,
13+ 0.0,
14+ 0.0
15+ ],
16+ "status": "pass"
17+}
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/evidence/node202_gate_summary.json+32-0
@@ -0,0 +1,32 @@
1+{
2+ "date": "2026-07-28",
3+ "host": "node202",
4+ "device": "Ascend 910B3",
5+ "physical_npu": 6,
6+ "cann": "8.1.RC1.alpha001",
7+ "pytorch": "2.5.0",
8+ "torch_npu": "2.5.1",
9+ "upstream_base": "fef85c9be08c70dfb65c2254d37a3791edf9c625",
10+ "source_scope": "joint-submission staged source snapshot",
11+ "global_opp_install_performed": false,
12+ "operator": "reformer_lsh_bucket_sort",
13+ "reproducible_clean_gate": {
14+ "python_tests": 7,
15+ "msopgen": "pass",
16+ "ascend910b_build": "pass",
17+ "aclnn_smoke": "pass",
18+ "smoke_max_abs_error": 0.0,
19+ "smoke_hot_mean_ms": 0.0363553,
20+ "unsupported_contract_status": 561002,
21+ "multistream": {
22+ "streams": 2,
23+ "iterations_per_stream": 20,
24+ "workspace_size": 0,
25+ "distinct_output_storage": true,
26+ "max_abs_error": 0.0,
27+ "status": "pass"
28+ }
29+ },
30+ "formal_package_gate_passed": true,
31+ "official_mr_sca_ci_cla_pending": true
32+}
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/test_report.md+39-0
@@ -0,0 +1,39 @@
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+# ReformerLshBucketSort 测试报告
12+ 
13+## 已完成验证
14+ 
15+| 检查项 | 结果 |
16+|---|---|
17+| Python CPU reference 与 adapter | PASS，7 项测试 |
18+| msopgen schema 生成 | PASS |
19+| CANN 8.1 RC1 / Ascend910B 独立构建 | PASS |
20+| Ascend 910B3 ACLNN smoke 与回读 | PASS |
21+| 非法合同 Host 拒绝 | PASS，返回非零状态 `561002` |
22+| 双 stream 各 20 次持续调用 | PASS，独立输出、0 Byte workspace |
23+| B4/B16/B32 × L96/L336 组件矩阵 | PASS |
24+| Reformer B32/L336 单算子模型消融 | PASS |
25+| ETTh1 全部 2,857 个测试窗口 | PASS |
26+ 
27+reference 覆盖确定性、随机、最小边界和非法输入；adapter 覆盖 custom、常规 fallback 和 uint32 超界 fallback。组件输出、索引、完整模型预测均逐元素完全一致。
28+ 
29+## 可复现命令
30+ 
31+```bash
32+python -m unittest discover -s tests -p "test_*.py" -v
33+python -m compileall -q integration reference tests
34+bash tests/run_packaged_gate.sh /tmp/reformer_lsh_bucket_sort_gate 6
35+```
36+ 
37+## 外部提交门禁
38+ 
39+已基于 2026-07-28 最新 `upstream/master` 完成同名和同功能去重检查。GitCode SCA、CLA 和官方仓库 CI 只能在由 fork 发起 MR 后执行，当前不宣称已通过。
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/docs/upstream_dedup.md+17-0
@@ -0,0 +1,17 @@
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+# ReformerLshBucketSort 上游去重检查
12+ 
13+检查基线：`cann/mat-chem-sim-pred` 的 upstream/master，commit `fef85c9be08c70dfb65c2254d37a3791edf9c625`，检查日期 2026-07-28。
14+ 
15+结论：当前上游 TimeSeriesForecast 目录没有同时实现 Reformer LSH 稳定桶排序、原位置索引和逆置换构造的算子；与已有算子不存在接口或功能重复。
16+ 
17+由于上游仍会变化，推送前和创建 MR 前均需刷新 upstream/master 并再次检查。
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/examples/test_aclnn_reformer_lsh_bucket_sort.cpp+274-0
@@ -0,0 +1,274 @@
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+#include <dlfcn.h>
12+ 
13+#include <algorithm>
14+#include <chrono>
15+#include <cmath>
16+#include <cstdint>
17+#include <iostream>
18+#include <numeric>
19+#include <vector>
20+ 
21+#include "acl/acl.h"
22+#include "aclnn_reformer_lsh_bucket_sort.h"
23+ 
24+namespace
25+{
26+constexpr size_t kElementCount = 8U;
27+constexpr int kWarmup = 3;
28+constexpr int kRepeat = 10;
29+ 
30+aclTensor* MakeTensor(const std::vector<int64_t>& dims, aclDataType dtype, void* data)
31+{
32+ std::vector<int64_t> strides(dims.size(), 1);
33+ for (int64_t i = static_cast<int64_t>(dims.size()) - 2; i >= 0; --i)
34+ {
35+ strides[static_cast<size_t>(i)] = strides[static_cast<size_t>(i + 1)] * dims[static_cast<size_t>(i + 1)];
36+ }
37+ return aclCreateTensor(dims.data(), dims.size(), dtype, strides.data(), 0, ACL_FORMAT_ND, dims.data(), dims.size(),
38+ data);
39+}
40+ 
41+template <typename T>
42+bool CopyToDevice(void** device, const std::vector<T>& host)
43+{
44+ const size_t bytes = host.size() * sizeof(T);
45+ return aclrtMalloc(device, bytes, ACL_MEM_MALLOC_NORMAL_ONLY) == ACL_SUCCESS &&
46+ aclrtMemcpy(*device, bytes, host.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE) == ACL_SUCCESS;
47+}
48+ 
49+template <typename T>
50+bool AllocateOutput(void** device, size_t elements)
51+{
52+ return aclrtMalloc(device, elements * sizeof(T), ACL_MEM_MALLOC_NORMAL_ONLY) == ACL_SUCCESS;
53+}
54+ 
55+template <typename T>
56+bool CopyToHost(std::vector<T>& host, void* device)
57+{
58+ const size_t bytes = host.size() * sizeof(T);
59+ return aclrtMemcpy(host.data(), bytes, device, bytes, ACL_MEMCPY_DEVICE_TO_HOST) == ACL_SUCCESS;
60+}
61+ 
62+template <typename T>
63+size_t MismatchCount(const std::vector<T>& lhs, const std::vector<T>& rhs)
64+{
65+ size_t count = 0;
66+ for (size_t i = 0; i < lhs.size(); ++i)
67+ {
68+ count += lhs[i] != rhs[i] ? 1U : 0U;
69+ }
70+ return count;
71+}
72+ 
73+class RuntimeContext
74+{
75+ public:
76+ bool Initialize()
77+ {
78+ if (aclInit(nullptr) != ACL_SUCCESS)
79+ {
80+ return false;
81+ }
82+ initialized_ = true;
83+ if (aclrtSetDevice(0) != ACL_SUCCESS)
84+ {
85+ return false;
86+ }
87+ deviceSet_ = true;
88+ allOps_ = dlopen("libascend_all_ops.so", RTLD_NOW | RTLD_GLOBAL);
89+ opApi_ = dlopen("libcust_opapi.so", RTLD_NOW | RTLD_GLOBAL);
90+ if (allOps_ == nullptr || opApi_ == nullptr)
91+ {
92+ std::cerr << "custom operator libraries unavailable" << std::endl;
93+ return false;
94+ }
95+ return aclrtCreateStream(&stream_) == ACL_SUCCESS;
96+ }
97+ 
98+ ~RuntimeContext()
99+ {
100+ if (stream_ != nullptr)
101+ {
102+ aclrtDestroyStream(stream_);
103+ }
104+ if (opApi_ != nullptr)
105+ {
106+ dlclose(opApi_);
107+ }
108+ if (allOps_ != nullptr)
109+ {
110+ dlclose(allOps_);
111+ }
112+ if (deviceSet_)
113+ {
114+ aclrtResetDevice(0);
115+ }
116+ if (initialized_)
117+ {
118+ aclFinalize();
119+ }
120+ }
atomgit-bot
atomgit-botatomgit-bot15 天前

🟡 Medium Priority

changed line 96-100 → 当 CopyToDevice(&dKeys, hKeys) 成功（dKeys 已分配），但后续 AllocateOutput<int64_t>(&dSorted, 8) 或更后面的分配失败时，短路求值使整体条件为 true，直接 return 1，dKeys（以及可能已成功分配的 dSorted、dSticker）泄漏。此外，CopyToDevice 内部也存在同类泄漏：若 aclrtMalloc 成功但 aclrtMemcpy 失败，已分配内存同样泄漏。

建议：在 return 前释放已成功分配的资源。将分配步骤拆分为独立的 if 判断，每个失败分支释放已分配的资源；或使用 RAII 包装 aclrtMalloc/aclrtFree。

likedislike
121+ 
122+ aclrtStream Stream() const { return stream_; }
123+ 
124+ private:
125+ void* allOps_ = nullptr;
126+ void* opApi_ = nullptr;
127+ aclrtStream stream_ = nullptr;
128+ bool initialized_ = false;
129+ bool deviceSet_ = false;
130+};
131+ 
132+class SortFixture
133+{
134+ public:
135+ bool Initialize()
136+ {
137+ if (!CopyToDevice(&dKeys_, keys_) || !AllocateOutput<int64_t>(&dSorted_, kElementCount) ||
138+ !AllocateOutput<int64_t>(&dSticker_, kElementCount) || !AllocateOutput<int64_t>(&dInverse_, kElementCount))
139+ {
140+ return false;
141+ }
142+ keysTensor_ = MakeTensor({1, 8}, ACL_INT64, dKeys_);
143+ sortedTensor_ = MakeTensor({1, 8}, ACL_INT64, dSorted_);
144+ stickerTensor_ = MakeTensor({1, 8}, ACL_INT64, dSticker_);
145+ inverseTensor_ = MakeTensor({1, 8}, ACL_INT64, dInverse_);
146+ return keysTensor_ != nullptr && sortedTensor_ != nullptr && stickerTensor_ != nullptr &&
147+ inverseTensor_ != nullptr;
148+ }
149+ 
150+ ~SortFixture()
151+ {
152+ DestroyTensor(keysTensor_);
153+ DestroyTensor(sortedTensor_);
154+ DestroyTensor(stickerTensor_);
155+ DestroyTensor(inverseTensor_);
156+ FreeDevice(dKeys_);
157+ FreeDevice(dSorted_);
158+ FreeDevice(dSticker_);
159+ FreeDevice(dInverse_);
160+ }
161+ 
162+ aclTensor* Keys() const { return keysTensor_; }
163+ 
164+ aclTensor* Sorted() const { return sortedTensor_; }
165+ 
166+ aclTensor* Sticker() const { return stickerTensor_; }
167+ 
168+ aclTensor* Inverse() const { return inverseTensor_; }
169+ 
170+ bool CountMismatches(size_t* mismatch)
171+ {
172+ if (!CopyToHost(sorted_, dSorted_) || !CopyToHost(sticker_, dSticker_) || !CopyToHost(inverse_, dInverse_))
173+ {
174+ return false;
175+ }
176+ *mismatch = MismatchCount(sorted_, expectedSorted_) + MismatchCount(sticker_, expectedSticker_) +
177+ MismatchCount(inverse_, expectedInverse_);
178+ return true;
179+ }
180+ 
181+ private:
182+ static void DestroyTensor(aclTensor* tensor)
183+ {
184+ if (tensor != nullptr)
185+ {
186+ aclDestroyTensor(tensor);
187+ }
188+ }
189+ 
190+ static void FreeDevice(void* device)
191+ {
192+ if (device != nullptr)
193+ {
194+ aclrtFree(device);
195+ }
196+ }
197+ 
198+ const std::vector<int64_t> keys_{9, 1, 8, 0, 5, 4, 2, 10};
199+ const std::vector<int64_t> expectedSorted_{1, 0, 2, 5, 4, 9, 8, 10};
200+ const std::vector<int64_t> expectedSticker_{1, 3, 6, 4, 5, 0, 2, 7};
201+ const std::vector<int64_t> expectedInverse_{5, 0, 6, 1, 3, 4, 2, 7};
202+ std::vector<int64_t> sorted_ = std::vector<int64_t>(kElementCount);
203+ std::vector<int64_t> sticker_ = std::vector<int64_t>(kElementCount);
204+ std::vector<int64_t> inverse_ = std::vector<int64_t>(kElementCount);
205+ void* dKeys_ = nullptr;
206+ void* dSorted_ = nullptr;
207+ void* dSticker_ = nullptr;
208+ void* dInverse_ = nullptr;
209+ aclTensor* keysTensor_ = nullptr;
210+ aclTensor* sortedTensor_ = nullptr;
211+ aclTensor* stickerTensor_ = nullptr;
212+ aclTensor* inverseTensor_ = nullptr;
213+};
214+ 
215+bool RunBenchmark(const SortFixture& fixture, aclrtStream stream, double* meanMs)
216+{
217+ std::vector<double> hotMs;
218+ for (int iteration = 0; iteration < kWarmup + kRepeat; ++iteration)
219+ {
220+ uint64_t workspaceSize = 0;
221+ aclOpExecutor* executor = nullptr;
222+ const aclnnStatus prepareStatus = aclnnReformerLshBucketSortGetWorkspaceSize(
223+ fixture.Keys(), 4, 3, fixture.Sorted(), fixture.Sticker(), fixture.Inverse(), &workspaceSize, &executor);
224+ if (prepareStatus != 0)
225+ {
226+ std::cerr << "GetWorkspaceSize failed: " << prepareStatus << std::endl;
227+ return false;
228+ }
229+ void* workspace = nullptr;
230+ if (workspaceSize > 0 && aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_NORMAL_ONLY) != ACL_SUCCESS)
231+ {
232+ return false;
233+ }
234+ const auto start = std::chrono::steady_clock::now();
235+ const aclnnStatus launchStatus = aclnnReformerLshBucketSort(workspace, workspaceSize, executor, stream);
236+ const aclError synchronizeStatus = aclrtSynchronizeStream(stream);
237+ const auto stop = std::chrono::steady_clock::now();
238+ if (workspace != nullptr)
239+ {
240+ aclrtFree(workspace);
241+ }
242+ if (launchStatus != 0 || synchronizeStatus != ACL_SUCCESS)
243+ {
244+ return false;
245+ }
246+ if (iteration >= kWarmup)
247+ {
248+ hotMs.push_back(std::chrono::duration<double, std::milli>(stop - start).count());
249+ }
250+ }
251+ *meanMs = std::accumulate(hotMs.begin(), hotMs.end(), 0.0) / hotMs.size();
252+ return true;
253+}
254+ 
255+} // namespace
256+ 
257+int main()
258+{
259+ RuntimeContext runtime;
260+ SortFixture fixture;
261+ if (!runtime.Initialize() || !fixture.Initialize())
262+ {
263+ return 1;
264+ }
265+ double meanMs = 0.0;
266+ size_t mismatch = 0U;
267+ if (!RunBenchmark(fixture, runtime.Stream(), &meanMs) || !fixture.CountMismatches(&mismatch))
268+ {
269+ return 1;
270+ }
271+ std::cout << "mismatch_count=" << mismatch << std::endl;
272+ std::cout << "operator=ReformerLshBucketSort,npu_hot_mean_ms=" << meanMs << std::endl;
273+ return mismatch == 0U ? 0 : 2;
274+}
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/integration/adapter.py+53-0
@@ -0,0 +1,53 @@
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+"""Guarded framework adapter; unsupported cases call the exact fallback."""
12+ 
13+_UINT32_MAX = (1 << 32) - 1
14+_INT32_MAX = (1 << 31) - 1
15+ 
16+ 
17+def _resident_contiguous(*tensors):
18+ return all(
19+ getattr(tensor, "device", None) is not None
20+ and tensor.device.type == "npu"
21+ and tensor.is_contiguous()
22+ for tensor in tensors
23+ )
24+ 
25+ 
26+def supports_custom(keys, sequence_length, total_buckets):
27+ try:
28+ return (
29+ _resident_contiguous(keys)
30+ and len(keys.shape) == 2
31+ and str(keys.dtype) == "torch.int64"
32+ and 0 < keys.shape[0] <= _UINT32_MAX
33+ and 0 < keys.shape[1] <= _INT32_MAX
34+ and 0 < sequence_length <= _UINT32_MAX
35+ and 0 < total_buckets <= 4096
36+ )
37+ except (AttributeError, IndexError, TypeError):
38+ return False
39+ 
40+ 
41+def dispatch(keys, sequence_length, total_buckets, *, custom_call, fallback):
42+ """Use custom_call only for the audited contract; fallback owns all other shapes."""
43+ args = (
44+ keys,
45+ sequence_length,
46+ total_buckets,
47+ )
48+ if (
49+ supports_custom(keys, sequence_length, total_buckets)
50+ and custom_call is not None
51+ ):
52+ return custom_call(*args)
53+ return fallback(*args)
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/msopgen/reformer_lsh_bucket_sort_msopgen.json+62-0
@@ -0,0 +1,62 @@
1+[
2+ {
3+ "op": "ReformerLshBucketSort",
4+ "language": "cpp",
5+ "input_desc": [
6+ {
7+ "name": "keys",
8+ "param_type": "required",
9+ "format": [
10+ "ND"
11+ ],
12+ "type": [
13+ "int64"
14+ ]
15+ }
16+ ],
17+ "output_desc": [
18+ {
19+ "name": "sorted_keys",
20+ "param_type": "required",
21+ "format": [
22+ "ND"
23+ ],
24+ "type": [
25+ "int64"
26+ ]
27+ },
28+ {
29+ "name": "sticker",
30+ "param_type": "required",
31+ "format": [
32+ "ND"
33+ ],
34+ "type": [
35+ "int64"
36+ ]
37+ },
38+ {
39+ "name": "inverse",
40+ "param_type": "required",
41+ "format": [
42+ "ND"
43+ ],
44+ "type": [
45+ "int64"
46+ ]
47+ }
48+ ],
49+ "attr": [
50+ {
51+ "name": "sequence_length",
52+ "param_type": "required",
53+ "type": "int"
54+ },
55+ {
56+ "name": "total_buckets",
57+ "param_type": "required",
58+ "type": "int"
59+ }
60+ ]
61+ }
62+]
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/op_host/reformer_lsh_bucket_sort_def.cpp+61-0
@@ -0,0 +1,61 @@
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+#include "../../common/op_def_utils.h"
12+ 
13+namespace optiling
14+{
15+ge::graphStatus ReformerLshBucketSortTiling(gert::TilingContext* context);
16+}
17+ 
18+namespace ge
19+{
20+graphStatus InferReformerLshBucketSortShape(gert::InferShapeContext* context);
21+graphStatus InferReformerLshBucketSortDataType(gert::InferDataTypeContext* context);
22+} // namespace ge
23+ 
24+namespace ops
25+{
26+class ReformerLshBucketSort : public TimeSeriesOpDef
27+{
28+ public:
29+ explicit ReformerLshBucketSort(const char* name) : TimeSeriesOpDef(name)
30+ {
31+ this->Input("keys")
32+ .ParamType(REQUIRED)
33+ .DataType({ge::DT_INT64})
34+ .Format({ge::FORMAT_ND})
35+ .UnknownShapeFormat({ge::FORMAT_ND});
36+ this->Output("sorted_keys")
37+ .ParamType(REQUIRED)
38+ .DataType({ge::DT_INT64})
39+ .Format({ge::FORMAT_ND})
40+ .UnknownShapeFormat({ge::FORMAT_ND});
41+ this->Output("sticker")
42+ .ParamType(REQUIRED)
43+ .DataType({ge::DT_INT64})
44+ .Format({ge::FORMAT_ND})
45+ .UnknownShapeFormat({ge::FORMAT_ND});
46+ this->Output("inverse")
47+ .ParamType(REQUIRED)
48+ .DataType({ge::DT_INT64})
49+ .Format({ge::FORMAT_ND})
50+ .UnknownShapeFormat({ge::FORMAT_ND});
51+ this->Attr("sequence_length").AttrType(REQUIRED).Int();
52+ this->Attr("total_buckets").AttrType(REQUIRED).Int();
53+ this->SetInferShape(ge::InferReformerLshBucketSortShape)
54+ .SetInferDataType(ge::InferReformerLshBucketSortDataType);
55+ this->AICore().SetTiling(optiling::ReformerLshBucketSortTiling);
56+ this->AICore().AddConfig("ascend910b");
57+ }
58+};
59+ 
60+OP_ADD(ReformerLshBucketSort);
61+} // namespace ops
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/op_host/reformer_lsh_bucket_sort_host.cpp+80-0
@@ -0,0 +1,80 @@
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+#include <limits>
12+ 
13+#include "../../common/op_host_utils.h"
14+#include "reformer_lsh_bucket_sort_tiling.h"
15+ 
16+namespace
17+{
18+bool FitsUint32(int64_t value)
19+{
20+ constexpr int64_t kUint32Max = static_cast<int64_t>(std::numeric_limits<uint32_t>::max());
21+ return value > 0 && value <= kUint32Max;
22+}
23+ 
24+bool FitsCountBuffer(int64_t value) { return value > 0 && value <= std::numeric_limits<int32_t>::max(); }
25+} // namespace
26+ 
27+namespace optiling
28+{
29+ge::graphStatus ReformerLshBucketSortTiling(gert::TilingContext* context)
30+{
31+ const auto& keys = context->GetInputShape(0)->GetStorageShape();
32+ if (keys.GetDimNum() != 2U || !FitsUint32(keys.GetDim(0)) || !FitsCountBuffer(keys.GetDim(1)))
33+ {
34+ return ge::GRAPH_FAILED;
35+ }
36+ const auto* attrs = context->GetAttrs();
37+ if (attrs == nullptr)
38+ {
39+ return ge::GRAPH_FAILED;
40+ }
41+ const int64_t* sequenceLength = attrs->GetAttrPointer<int64_t>(0);
42+ const int64_t* totalBuckets = attrs->GetAttrPointer<int64_t>(1);
43+ if (sequenceLength == nullptr || totalBuckets == nullptr || !FitsUint32(*sequenceLength) || *totalBuckets <= 0 ||
44+ *totalBuckets > 4096)
45+ {
46+ return ge::GRAPH_FAILED;
47+ }
48+ ReformerLshBucketSortTilingData tiling;
49+ tiling.set_rows(static_cast<uint32_t>(keys.GetDim(0)));
50+ tiling.set_total_length(static_cast<uint32_t>(keys.GetDim(1)));
51+ tiling.set_sequence_length(static_cast<uint32_t>(*sequenceLength));
52+ tiling.set_total_buckets(static_cast<uint32_t>(*totalBuckets));
53+ const uint32_t coreCount = timeseries_host::GetVectorCoreCount(context);
54+ context->SetBlockDim(timeseries_host::SelectBlockDim(static_cast<uint32_t>(keys.GetDim(0)), coreCount));
55+ timeseries_host::StoreTilingData(context, tiling);
56+ return ge::GRAPH_SUCCESS;
57+}
58+} // namespace optiling
59+ 
60+namespace ge
61+{
62+graphStatus InferReformerLshBucketSortShape(gert::InferShapeContext* context)
63+{
64+ const gert::Shape* input = context->GetInputShape(0);
65+ for (uint32_t i = 0; i < 3U; ++i)
66+ {
67+ *context->GetOutputShape(i) = *input;
68+ }
69+ return ge::GRAPH_SUCCESS;
70+}
71+ 
72+graphStatus InferReformerLshBucketSortDataType(gert::InferDataTypeContext* context)
73+{
74+ for (uint32_t i = 0; i < 3U; ++i)
75+ {
76+ context->SetOutputDataType(i, context->GetInputDataType(0));
77+ }
78+ return ge::GRAPH_SUCCESS;
79+}
80+} // namespace ge
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/op_host/reformer_lsh_bucket_sort_tiling.h+28-0
@@ -0,0 +1,28 @@
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 REFORMER_LSH_BUCKET_SORT_TILING_H_
12+#define REFORMER_LSH_BUCKET_SORT_TILING_H_
13+ 
14+#include "register/tilingdata_base.h"
15+ 
16+namespace optiling
17+{
18+BEGIN_TILING_DATA_DEF(ReformerLshBucketSortTilingData)
19+TILING_DATA_FIELD_DEF(uint32_t, rows);
20+TILING_DATA_FIELD_DEF(uint32_t, total_length);
21+TILING_DATA_FIELD_DEF(uint32_t, sequence_length);
22+TILING_DATA_FIELD_DEF(uint32_t, total_buckets);
23+END_TILING_DATA_DEF;
24+ 
25+REGISTER_TILING_DATA_CLASS(ReformerLshBucketSort, ReformerLshBucketSortTilingData)
26+} // namespace optiling
27+ 
28+#endif // REFORMER_LSH_BUCKET_SORT_TILING_H_
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/op_kernel/reformer_lsh_bucket_sort_kernel.cpp+110-0
@@ -0,0 +1,110 @@
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+#include "kernel_operator.h"
12+ 
13+using namespace AscendC;
14+ 
15+namespace
16+{
17+ 
18+struct LshBucketSortTiling
19+{
20+ uint32_t rows;
21+ uint32_t total_length;
22+ uint32_t sequence_length;
23+ uint32_t total_buckets;
24+};
25+ 
26+class LshBucketSortKernel
27+{
28+ public:
29+ __aicore__ inline LshBucketSortKernel() = default;
30+ 
31+ __aicore__ inline void Init(GM_ADDR keys, GM_ADDR sortedKeys, GM_ADDR sticker, GM_ADDR inverse, GM_ADDR tiling)
32+ {
33+ const __gm__ LshBucketSortTiling* tilingData = reinterpret_cast<const __gm__ LshBucketSortTiling*>(tiling);
34+ rows_ = tilingData->rows;
35+ totalLength_ = tilingData->total_length;
36+ sequenceLength_ = tilingData->sequence_length;
37+ totalBuckets_ = tilingData->total_buckets;
38+ const uint64_t elements = static_cast<uint64_t>(rows_) * totalLength_;
39+ keysGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(keys), elements);
40+ sortedKeysGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(sortedKeys), elements);
41+ stickerGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(sticker), elements);
42+ inverseGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(inverse), elements);
43+ pipe_.InitBuffer(countBuffer_, ((totalBuckets_ + 7U) / 8U) * 8U * sizeof(int32_t));
44+ }
45+ 
46+ __aicore__ inline void Process()
47+ {
48+ const uint32_t block = GetBlockIdx();
49+ const uint32_t blocks = GetBlockNum();
50+ LocalTensor<int32_t> counts = countBuffer_.Get<int32_t>();
51+ for (uint64_t row = block; row < rows_; row += blocks)
52+ {
53+ for (uint32_t bucket = 0; bucket < totalBuckets_; ++bucket)
54+ {
55+ counts.SetValue(bucket, 0);
56+ }
57+ pipe_barrier(PIPE_ALL);
58+ const uint64_t base = row * totalLength_;
59+ for (uint32_t position = 0; position < totalLength_; ++position)
60+ {
61+ const int64_t key = keysGm_.GetValue(base + position);
62+ const uint32_t bucket = static_cast<uint32_t>(key / sequenceLength_);
63+ counts.SetValue(bucket, counts.GetValue(bucket) + 1);
64+ }
65+ pipe_barrier(PIPE_ALL);
66+ int32_t running = 0;
67+ for (uint32_t bucket = 0; bucket < totalBuckets_; ++bucket)
68+ {
69+ const int32_t count = counts.GetValue(bucket);
70+ counts.SetValue(bucket, running);
71+ running += count;
72+ }
73+ pipe_barrier(PIPE_ALL);
74+ for (uint32_t position = 0; position < totalLength_; ++position)
75+ {
76+ const int64_t key = keysGm_.GetValue(base + position);
77+ const uint32_t bucket = static_cast<uint32_t>(key / sequenceLength_);
78+ const int32_t sortedPosition = counts.GetValue(bucket);
79+ counts.SetValue(bucket, sortedPosition + 1);
80+ const uint64_t output = base + static_cast<uint32_t>(sortedPosition);
81+ sortedKeysGm_.SetValue(output, key);
82+ stickerGm_.SetValue(output, static_cast<int64_t>(position));
83+ inverseGm_.SetValue(base + position, static_cast<int64_t>(sortedPosition));
84+ }
85+ }
86+ }
87+ 
88+ private:
89+ TPipe pipe_;
90+ TBuf<TPosition::VECCALC> countBuffer_;
91+ GlobalTensor<int64_t> keysGm_;
92+ GlobalTensor<int64_t> sortedKeysGm_;
93+ GlobalTensor<int64_t> stickerGm_;
94+ GlobalTensor<int64_t> inverseGm_;
95+ uint32_t rows_ = 0U;
96+ uint32_t totalLength_ = 0U;
97+ uint32_t sequenceLength_ = 0U;
98+ uint32_t totalBuckets_ = 0U;
99+};
100+ 
101+} // namespace
102+ 
103+extern "C" __global__ __aicore__ void reformer_lsh_bucket_sort(GM_ADDR keys, GM_ADDR sortedKeys, GM_ADDR sticker,
104+ GM_ADDR inverse, GM_ADDR workspace, GM_ADDR tiling)
105+{
106+ (void)workspace;
107+ LshBucketSortKernel kernel;
108+ kernel.Init(keys, sortedKeys, sticker, inverse, tiling);
109+ kernel.Process();
110+}
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/reference/reference.py+38-0
@@ -0,0 +1,38 @@
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+"""Dependency-free CPU reference for ReformerLshBucketSort."""
12+ 
13+ 
14+def _matrix(value, name):
15+ if not isinstance(value, (list, tuple)) or not value or not value[0]:
16+ raise ValueError(f"{name} must be a non-empty matrix")
17+ width = len(value[0])
18+ if any(len(row) != width for row in value):
19+ raise ValueError(f"{name} must be rectangular")
20+ return len(value), width
21+ 
22+ 
23+def reference(keys, sequence_length, total_buckets):
24+ _matrix(keys, "keys")
25+ if sequence_length <= 0 or not 1 <= total_buckets <= 4096:
26+ raise ValueError("invalid bucket attributes")
27+ sorted_keys, stickers, inverses = [], [], []
28+ for row in keys:
29+ if any(key < 0 or key // sequence_length >= total_buckets for key in row):
30+ raise ValueError("key outside encoded bucket range")
31+ sticker = sorted(range(len(row)), key=lambda i: row[i] // sequence_length)
32+ inverse = [0] * len(row)
33+ for sorted_position, source_position in enumerate(sticker):
34+ inverse[source_position] = sorted_position
35+ sorted_keys.append([row[i] for i in sticker])
36+ stickers.append(sticker)
37+ inverses.append(inverse)
38+ return sorted_keys, stickers, inverses
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/tests/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
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+add_subdirectory(ut)
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/tests/benchmark_reformer_lsh_bucket_sort_aclnn.cpp+11-0文件内容审核中，请稍后刷新重试
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/tests/run_packaged_gate.sh+61-0文件内容审核中，请稍后刷新重试
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/tests/test_adapter.py+81-0文件内容审核中，请稍后刷新重试
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/tests/test_reference.py+68-0文件内容审核中，请稍后刷新重试
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/tests/ut/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
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+add_subdirectory(op_kernel)
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/tests/ut/op_kernel/CMakeLists.txt+13-0文件内容审核中，请稍后刷新重试
Aprediction/ProcessControl/TimeSeriesForecast/reformer_lsh_bucket_sort/tests/ut/op_kernel/test_reformer_lsh_bucket_sort.cpp+11-0文件内容审核中，请稍后刷新重试