Star1315
1432
代码
代码
Issues126
Pull Requests256
Actions
项目讨论
Wiki
分析
Star1315
1432
已合并
【社区任务】Add roll operator #3333
boxw987创建于 6月15日
【社区任务】Add roll operator #3333
已合并
boxw987创建于 6月15日
从boxw987/ops-math_5072roll-submit-master合入到cann/ops-mathmaster
共 23 个文件变更+2772-0
Aexperimental/math/roll/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_all_modules_sources(OPTYPE roll ACLNNTYPE aclnn_exclude)
Aexperimental/math/roll/README.md+55-0
@@ -0,0 +1,55 @@
1+# Roll
2+ 
3+## 贡献说明
4+ 
5+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
6+| --- | --- | --- | --- | --- |
7+| boxw987 | 个人开发者 | Roll | 2026/06 | Roll 算子适配开源仓 |
8+ 
9+## 支持的产品型号
10+ 
11+- Atlas A2 训练系列产品
12+ 
13+## 算子描述
14+ 
15+- 功能描述
16+ 
17+ `Roll` 沿给定维度对输入张量执行循环位移；当 `dims` 为空时，先按逻辑视图展平，再执行一维循环位移，最后按原始形状输出。
18+ 
19+- 原型信息
20+ 
21+ <table>
22+ <tr><th align="center">算子类型(OpType)</th><th colspan="4" align="center">Roll</th></tr>
23+ <tr><td align="center"> </td><td align="center">name</td><td align="center">Type</td><td align="center">data type</td><td align="center">format</td></tr>
24+ <tr><td rowspan="1" align="center">算子输入</td><td align="center">x</td><td align="center">tensor</td><td align="center">uint8, int8, bfloat16, float16, float32, int32, uint32</td><td align="center">ND</td></tr>
25+ <tr><td rowspan="1" align="center">算子输出</td><td align="center">y</td><td align="center">tensor</td><td align="center">与 x 相同</td><td align="center">ND</td></tr>
26+ <tr><td rowspan="2" align="center">属性</td><td align="center">shifts</td><td align="center">listInt</td><td align="center">整型列表</td><td align="center">-</td></tr>
27+ <tr><td align="center">dims</td><td align="center">listInt</td><td align="center">整型列表</td><td align="center">-</td></tr>
28+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">roll</td></tr>
29+ </table>
30+ 
31+## 约束与限制
32+ 
33+- 仅支持 `ND` 格式。
34+- 支持 0 维到 8 维输入。
35+- `dims` 为空时，`shifts` 长度必须为 1。
36+- `dims` 非空时，`shifts` 与 `dims` 长度必须一致。
37+- `dims` 取值范围为 `[-rank, rank)`。
38+ 
39+## 算子使用
40+ 
41+使用该算子前，请参考[社区版 CANN 开发套件包安装文档](../../../docs/zh/invocation/quick_op_invocation.md)完成开发运行环境部署。
42+ 
43+### 编译部署
44+ 
45+```bash
46+cd ${git_clone_path}/ops-math
47+bash build.sh --pkg --experimental --soc=ascend910b --ops=roll
48+./build_out/cann-ops-<vendor_name>-linux.<arch>.run
49+```
50+ 
51+## 调用说明
52+ 
53+| 调用方式 | 调用样例 | 说明 |
54+| --- | --- | --- |
55+| aclnn 调用 | [test_aclnn_roll.cpp](./examples/test_aclnn_roll.cpp) | 通过 [aclnnRoll](./docs/aclnnRoll.md) 接口方式调用 Roll 算子 |
Aexperimental/math/roll/docs/aclnnRoll.md+85-0
@@ -0,0 +1,85 @@
1+# aclnnRoll
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :--- | :---: |
7+| Atlas A2 训练系列产品 | 支持 |
8+ 
9+## 功能说明
10+ 
11+- 接口功能：沿给定维度对输入 Tensor 执行循环位移。
12+ 
13+- 语义说明：
14+ 
15+ - `dims` 非空时，`shifts[i]` 作用于 `dims[i]` 对应维度。
16+ - `dims` 为空时，先对输入按逻辑视图展平，再执行一维 roll。
17+ - 重复维度会在 Host 侧归一化合并。
18+ 
19+## 函数原型
20+ 
21+```cpp
22+aclnnStatus aclnnRollGetWorkspaceSize(
23+ const aclTensor* x,
24+ const aclIntArray* shifts,
25+ const aclIntArray* dims,
26+ aclTensor* out,
27+ uint64_t* workspaceSize,
28+ aclOpExecutor** executor)
29+```
30+ 
31+```cpp
32+aclnnStatus aclnnRoll(
33+ void* workspace,
34+ uint64_t workspaceSize,
35+ aclOpExecutor* executor,
36+ aclrtStream stream)
37+```
38+ 
39+## aclnnRollGetWorkspaceSize
40+ 
41+| 参数名 | 输入/输出 | 描述 | 数据类型 | 数据格式 | 维度 |
42+| --- | --- | --- | --- | --- | --- |
43+| x | 输入 | 输入张量 | uint8, int8, bfloat16, float16, float32, int32, uint32 | ND | 0-8 维 |
44+| shifts | 输入 | 每个目标维度上的循环位移量 | aclIntArray* | - | - |
45+| dims | 输入 | 循环位移维度，可省略或传空数组 | aclIntArray* | - | - |
46+| out | 输出 | 输出张量，shape 和 dtype 与 x 一致 | 与 x 相同 | ND | 0-8 维 |
47+| workspaceSize | 输出 | 需要申请的 workspace 大小 | uint64_t* | - | - |
48+| executor | 输出 | 执行器 | aclOpExecutor** | - | - |
49+ 
50+- 返回值
51+ 
52+ `aclnnStatus`。第一段接口完成参数校验，出现以下场景时返回错误：
53+ 
54+ - `ACLNN_ERR_PARAM_NULLPTR`：`x`、`shifts`、`out`、`workspaceSize`、`executor` 为空。
55+ - `ACLNN_ERR_PARAM_INVALID`：
56+ - 输入或输出 dtype 不在支持范围内。
57+ - 输入与输出 dtype 不一致。
58+ - 输入或输出格式不是 `ND`。
59+ - 输入与输出 shape 不一致。
60+ - rank 大于 8。
61+ - 0 维输入时，`shifts` 长度不为 1 或 `dims` 非空。
62+ - `dims` 为空但 `shifts` 长度不为 1。
63+ - `dims` 非空但 `shifts` 与 `dims` 长度不一致。
64+ - `dims` 元素越界。
65+ 
66+## aclnnRoll
67+ 
68+| 参数名 | 输入/输出 | 描述 |
69+| --- | --- | --- |
70+| workspace | 输入 | Device 侧 workspace 地址 |
71+| workspaceSize | 输入 | Device 侧 workspace 大小 |
72+| executor | 输入 | 执行器 |
73+| stream | 输入 | 执行 stream |
74+ 
75+## 约束说明
76+ 
77+- 仅支持 `ND`。
78+- 仅支持 `uint8`、`int8`、`bfloat16`、`float16`、`float32`、`int32`、`uint32`。
79+- 0 维输入时，`shifts` 长度必须为 1，且 `dims` 为空。
80+- 非连续输入会先整理为连续视图后执行。
81+- 非连续输出会在算子结果生成后做回写。
82+ 
83+## 调用示例
84+ 
85+请参考 [test_aclnn_roll.cpp](../examples/test_aclnn_roll.cpp)。
Aexperimental/math/roll/examples/test_aclnn_roll.cpp+136-0
@@ -0,0 +1,136 @@
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 <iostream>
12+#include <vector>
13+ 
14+#include "acl/acl.h"
15+#include "aclnnop/aclnn_roll.h"
16+ 
17+#define CHECK_RET(cond, return_expr) \
18+ do { \
19+ if (!(cond)) { \
20+ return_expr; \
21+ } \
22+ } while (0)
23+ 
24+#define LOG_PRINT(message, ...) \
25+ do { \
26+ printf(message, ##__VA_ARGS__); \
27+ } while (0)
28+ 
29+int64_t GetShapeSize(const std::vector<int64_t>& shape)
30+{
31+ int64_t shapeSize = 1;
32+ for (auto dim : shape) {
33+ shapeSize *= dim;
34+ }
35+ return shapeSize;
36+}
37+ 
38+int Init(int32_t deviceId, aclrtStream* stream)
39+{
40+ auto ret = aclInit(nullptr);
41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
42+ ret = aclrtSetDevice(deviceId);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
44+ ret = aclrtCreateStream(stream);
45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
46+ return 0;
47+}
48+ 
49+template <typename T>
50+int CreateAclTensor(const std::vector<T>& hostData,
51+ const std::vector<int64_t>& shape,
52+ void** deviceAddr,
53+ aclDataType dataType,
54+ aclTensor** tensor)
55+{
56+ auto size = GetShapeSize(shape) * static_cast<int64_t>(sizeof(T));
57+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
58+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
59+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
60+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
61+ 
62+ std::vector<int64_t> strides(shape.size(), 1);
63+ for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) {
64+ strides[static_cast<size_t>(i)] = shape[static_cast<size_t>(i + 1)] * strides[static_cast<size_t>(i + 1)];
65+ }
66+ 
67+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, ACL_FORMAT_ND,
68+ shape.data(), shape.size(), *deviceAddr);
69+ return 0;
70+}
71+ 
72+int main()
73+{
74+ int32_t deviceId = 0;
75+ aclrtStream stream = nullptr;
76+ auto ret = Init(deviceId, &stream);
77+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
78+ 
79+ std::vector<int64_t> xShape = {2, 3};
80+ std::vector<int64_t> yShape = {2, 3};
81+ std::vector<float> xHostData = {0, 1, 2, 3, 4, 5};
82+ std::vector<float> yHostData(6, 0);
83+ 
84+ void* xDeviceAddr = nullptr;
85+ void* yDeviceAddr = nullptr;
86+ aclTensor* x = nullptr;
87+ aclTensor* y = nullptr;
88+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, ACL_FLOAT, &x);
89+ CHECK_RET(ret == ACL_SUCCESS, return ret);
90+ ret = CreateAclTensor(yHostData, yShape, &yDeviceAddr, ACL_FLOAT, &y);
91+ CHECK_RET(ret == ACL_SUCCESS, return ret);
92+ 
93+ std::vector<int64_t> shiftsData = {1};
94+ std::vector<int64_t> dimsData = {1};
95+ aclIntArray* shifts = aclCreateIntArray(shiftsData.data(), shiftsData.size());
96+ aclIntArray* dims = aclCreateIntArray(dimsData.data(), dimsData.size());
97+ CHECK_RET(shifts != nullptr && dims != nullptr, return -1);
98+ 
99+ uint64_t workspaceSize = 0;
100+ aclOpExecutor* executor = nullptr;
101+ ret = aclnnRollGetWorkspaceSize(x, shifts, dims, y, &workspaceSize, &executor);
102+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnRollGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
103+ 
104+ void* workspaceAddr = nullptr;
105+ if (workspaceSize > 0) {
106+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
107+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
108+ }
109+ 
110+ ret = aclnnRoll(workspaceAddr, workspaceSize, executor, stream);
111+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnRoll failed. ERROR: %d\n", ret); return ret);
112+ ret = aclrtSynchronizeStream(stream);
113+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
114+ 
115+ std::vector<float> resultData(yHostData.size(), 0);
116+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), yDeviceAddr,
117+ resultData.size() * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
118+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
119+ for (size_t i = 0; i < resultData.size(); ++i) {
120+ LOG_PRINT("result[%zu] is: %f\n", i, resultData[i]);
121+ }
122+ 
123+ aclDestroyIntArray(shifts);
124+ aclDestroyIntArray(dims);
125+ aclDestroyTensor(x);
126+ aclDestroyTensor(y);
127+ aclrtFree(xDeviceAddr);
128+ aclrtFree(yDeviceAddr);
129+ if (workspaceAddr != nullptr) {
130+ aclrtFree(workspaceAddr);
131+ }
132+ aclrtDestroyStream(stream);
133+ aclrtResetDevice(deviceId);
134+ aclFinalize();
135+ return 0;
136+}
Aexperimental/math/roll/op_api/aclnn_roll.cpp+240-0
@@ -0,0 +1,240 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * @file aclnn_roll.cpp
13+ * @brief ACLNN Roll implementation.
14+ */
15+ 
16+#include "aclnn_roll.h"
17+#include "roll.h"
18+ 
19+#include <algorithm>
20+ 
21+#include "aclnn_kernels/common/op_error_check.h"
22+#include "aclnn_kernels/contiguous.h"
23+#include "opdev/common_types.h"
24+#include "opdev/data_type_utils.h"
25+#include "opdev/format_utils.h"
26+#include "opdev/make_op_executor.h"
27+#include "opdev/op_dfx.h"
28+#include "opdev/op_executor.h"
29+#include "opdev/op_log.h"
30+#include "opdev/platform.h"
31+ 
32+using namespace op;
33+ 
34+namespace {
35+constexpr size_t MAX_SUPPORT_DIMS_NUMS = 8;
36+ 
37+const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {
38+ op::DataType::DT_UINT8, op::DataType::DT_INT8, op::DataType::DT_BF16,
39+ op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_INT32,
40+ op::DataType::DT_UINT32};
41+ 
42+bool CheckNotNull(const aclTensor* x, const aclIntArray* shifts, const aclIntArray* dims, const aclTensor* out)
43+{
44+ (void)dims;
45+ OP_CHECK_NULL(x, return false);
46+ OP_CHECK_NULL(shifts, return false);
47+ OP_CHECK_NULL(out, return false);
48+ return true;
49+}
50+ 
51+bool IsDtypeSupported(op::DataType dtype)
52+{
53+ return std::find(DTYPE_SUPPORT_LIST.begin(), DTYPE_SUPPORT_LIST.end(), dtype) != DTYPE_SUPPORT_LIST.end();
54+}
55+ 
56+bool CheckDtypeValid(const aclTensor* x, const aclTensor* out)
57+{
58+ if (!IsDtypeSupported(x->GetDataType())) {
59+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Input x dtype is not supported.");
60+ return false;
61+ }
62+ if (x->GetDataType() != out->GetDataType()) {
63+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Input x and output out must have the same dtype.");
64+ return false;
65+ }
66+ return true;
67+}
68+ 
69+bool CheckFormatValid(const aclTensor* x, const aclTensor* out)
70+{
71+ constexpr auto supportedFormat = op::Format::FORMAT_ND;
72+ if (x->GetViewFormat() != supportedFormat || x->GetStorageFormat() != supportedFormat) {
73+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Input x only supports ND format.");
74+ return false;
75+ }
76+ if (out->GetViewFormat() != supportedFormat || out->GetStorageFormat() != supportedFormat) {
77+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Output out only supports ND format.");
78+ return false;
79+ }
80+ return true;
81+}
82+ 
83+bool CheckShapeValid(const aclTensor* x, const aclTensor* out)
84+{
85+ if (x->GetViewShape() != out->GetViewShape()) {
86+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Input x and output out must have the same shape.");
87+ return false;
88+ }
89+ if (x->GetViewShape().GetDimNum() > MAX_SUPPORT_DIMS_NUMS) {
90+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Roll supports at most 8 dims.");
91+ return false;
92+ }
93+ return true;
94+}
95+ 
96+bool CheckArraySize(const aclTensor* x, const aclIntArray* shifts, const aclIntArray* dims)
97+{
98+ const auto tensorDim = x->GetViewShape().GetDimNum();
99+ const size_t dimsSize = dims == nullptr ? 0U : dims->Size();
100+ if (tensorDim == 0) {
101+ if (dimsSize != 0U || shifts->Size() != 1U) {
102+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "For 0-D tensor, shifts size must be 1 and dims must be empty.");
103+ return false;
104+ }
105+ return true;
106+ }
107+ if (dimsSize == 0U) {
108+ if (shifts->Size() != 1U) {
109+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "When dims is empty, shifts size must be 1.");
110+ return false;
111+ }
112+ return true;
113+ }
114+ if (shifts->Size() != dimsSize) {
115+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "shifts and dims must have the same size.");
116+ return false;
117+ }
118+ return true;
119+}
120+ 
121+bool CheckDimsRange(const aclTensor* x, const aclIntArray* dims)
122+{
123+ if (dims == nullptr) {
124+ return true;
125+ }
126+ const int64_t tensorDim = static_cast<int64_t>(x->GetViewShape().GetDimNum());
127+ for (size_t i = 0; i < dims->Size(); ++i) {
128+ const int64_t dim = (*dims)[i];
129+ if (dim < -tensorDim || dim >= tensorDim) {
130+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "dims value is out of range.");
131+ return false;
132+ }
133+ }
134+ return true;
135+}
136+ 
137+bool HasDenseViewLayout(const aclTensor* tensor)
138+{
139+ if (tensor == nullptr || tensor->GetViewOffset() != 0) {
140+ return false;
141+ }
142+ const auto& viewShape = tensor->GetViewShape();
143+ const auto& strides = tensor->GetViewStrides();
144+ if (strides.size() != viewShape.GetDimNum()) {
145+ return false;
146+ }
147+ int64_t expectedStride = 1;
148+ for (int64_t i = static_cast<int64_t>(viewShape.GetDimNum()) - 1; i >= 0; --i) {
149+ if (strides[static_cast<size_t>(i)] != expectedStride) {
150+ return false;
151+ }
152+ expectedStride *= static_cast<int64_t>(viewShape.GetDim(static_cast<size_t>(i)));
153+ }
154+ return true;
155+}
156+ 
157+bool CanWriteOutDirectly(const aclTensor* tensor)
158+{
159+ if (!HasDenseViewLayout(tensor)) {
160+ return false;
161+ }
162+ const auto& storageShape = tensor->GetStorageShape();
163+ const auto& viewShape = tensor->GetViewShape();
164+ return storageShape.GetDimNum() == 0 || storageShape == viewShape;
165+}
166+ 
167+aclTensor* NormalizeEmptyStorageTensor(const aclTensor* tensor, aclOpExecutor* executor)
168+{
169+ if (tensor == nullptr || executor == nullptr || tensor->GetStorageShape().GetDimNum() != 0 ||
170+ tensor->GetViewShape().GetDimNum() == 0) {
171+ return const_cast<aclTensor*>(tensor);
172+ }
173+ return executor->CreateView(tensor, tensor->GetViewShape(), tensor->GetViewOffset());
174+}
175+ 
176+aclnnStatus CheckParams(const aclTensor* x, const aclIntArray* shifts, const aclIntArray* dims, const aclTensor* out)
177+{
178+ CHECK_RET(CheckNotNull(x, shifts, dims, out), ACLNN_ERR_PARAM_NULLPTR);
179+ CHECK_RET(CheckDtypeValid(x, out), ACLNN_ERR_PARAM_INVALID);
180+ CHECK_RET(CheckFormatValid(x, out), ACLNN_ERR_PARAM_INVALID);
181+ CHECK_RET(CheckShapeValid(x, out), ACLNN_ERR_PARAM_INVALID);
182+ CHECK_RET(CheckArraySize(x, shifts, dims), ACLNN_ERR_PARAM_INVALID);
183+ CHECK_RET(CheckDimsRange(x, dims), ACLNN_ERR_PARAM_INVALID);
184+ return ACLNN_SUCCESS;
185+}
186+} // namespace
187+ 
188+extern "C" aclnnStatus aclnnRollGetWorkspaceSize(const aclTensor* x,
189+ const aclIntArray* shifts,
190+ const aclIntArray* dims,
191+ aclTensor* out,
192+ uint64_t* workspaceSize,
193+ aclOpExecutor** executor)
194+{
195+ OP_CHECK_COMM_INPUT(workspaceSize, executor);
196+ L2_DFX_PHASE_1(aclnnRoll, DFX_IN(x, shifts, dims), DFX_OUT(out));
197+ 
198+ auto uniqueExecutor = CREATE_EXECUTOR();
199+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
200+ 
201+ auto ret = CheckParams(x, shifts, dims, out);
202+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
203+ 
204+ if (x->IsEmpty()) {
205+ *workspaceSize = 0;
206+ uniqueExecutor.ReleaseTo(executor);
207+ return ACLNN_SUCCESS;
208+ }
209+ 
210+ const aclTensor* xForRoll = NormalizeEmptyStorageTensor(x, uniqueExecutor.get());
211+ CHECK_RET(xForRoll != nullptr, ACLNN_ERR_INNER_NULLPTR);
212+ aclTensor* outForRoll = NormalizeEmptyStorageTensor(out, uniqueExecutor.get());
213+ CHECK_RET(outForRoll != nullptr, ACLNN_ERR_INNER_NULLPTR);
214+ 
215+ const aclTensor* xContiguous = xForRoll;
216+ if (!HasDenseViewLayout(x)) {
217+ xContiguous = l0op::Contiguous(xForRoll, uniqueExecutor.get());
218+ CHECK_RET(xContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
219+ }
220+ 
221+ const aclTensor* rollResult = nullptr;
222+ if (CanWriteOutDirectly(outForRoll)) {
223+ rollResult = l0op::Roll(xContiguous, shifts, dims, outForRoll, uniqueExecutor.get());
224+ } else {
225+ rollResult = l0op::Roll(xContiguous, shifts, dims, uniqueExecutor.get());
226+ CHECK_RET(rollResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
227+ rollResult = l0op::ViewCopy(rollResult, out, uniqueExecutor.get());
228+ }
229+ CHECK_RET(rollResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
230+ 
231+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
232+ uniqueExecutor.ReleaseTo(executor);
233+ return ACLNN_SUCCESS;
234+}
235+ 
236+extern "C" aclnnStatus aclnnRoll(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
237+{
238+ L2_DFX_PHASE_2(aclnnRoll);
239+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
240+}
Aexperimental/math/roll/op_api/aclnn_roll.h+37-0
@@ -0,0 +1,37 @@
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 OP_API_INC_ROLL_H_
12+#define OP_API_INC_ROLL_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+#include "aclnn_util.h"
16+ 
17+#ifdef __cplusplus
18+extern "C" {
19+#endif
20+ 
21+ACLNN_API aclnnStatus aclnnRollGetWorkspaceSize(const aclTensor* x,
22+ const aclIntArray* shifts,
23+ const aclIntArray* dims,
24+ aclTensor* out,
25+ uint64_t* workspaceSize,
26+ aclOpExecutor** executor);
27+ 
28+ACLNN_API aclnnStatus aclnnRoll(void* workspace,
29+ uint64_t workspaceSize,
30+ aclOpExecutor* executor,
31+ aclrtStream stream);
32+ 
33+#ifdef __cplusplus
34+}
35+#endif
36+ 
37+#endif // OP_API_INC_ROLL_H_
Aexperimental/math/roll/op_api/roll.cpp+65-0
@@ -0,0 +1,65 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * @file roll.cpp
13+ * @brief Roll L0 API implementation.
14+ */
15+ 
16+#include "roll.h"
17+ 
18+#include "opdev/make_op_executor.h"
19+#include "opdev/op_dfx.h"
20+#include "opdev/op_log.h"
21+ 
22+using namespace op;
23+ 
24+namespace l0op {
25+ 
26+OP_TYPE_REGISTER(Roll);
27+ 
28+aclTensor* RollAiCore(const aclTensor* x,
29+ const aclIntArray* shifts,
30+ const aclIntArray* dims,
31+ aclTensor* rollOut,
32+ aclOpExecutor* executor)
33+{
34+ (void)executor;
35+ L0_DFX(RollAiCore, x, shifts, dims);
36+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(Roll, OP_INPUT(x), OP_OUTPUT(rollOut), OP_ATTR(shifts, dims));
37+ if (ret != ACL_SUCCESS) {
38+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "RollAiCore ADD_TO_LAUNCHER_LIST_AICORE failed.");
39+ return nullptr;
40+ }
41+ return rollOut;
42+}
43+ 
44+const aclTensor* Roll(const aclTensor* x, const aclIntArray* shifts, const aclIntArray* dims, aclOpExecutor* executor)
45+{
46+ if (x == nullptr || executor == nullptr) {
47+ return nullptr;
48+ }
49+ auto rollOut = executor->AllocTensor(x->GetViewShape(), x->GetDataType(), x->GetViewFormat());
50+ return Roll(x, shifts, dims, rollOut, executor);
51+}
52+ 
53+const aclTensor* Roll(const aclTensor* x,
54+ const aclIntArray* shifts,
55+ const aclIntArray* dims,
56+ aclTensor* out,
57+ aclOpExecutor* executor)
58+{
59+ if (x == nullptr || out == nullptr || executor == nullptr) {
60+ return nullptr;
61+ }
62+ return RollAiCore(x, shifts, dims, out, executor);
63+}
64+ 
65+} // namespace l0op
Aexperimental/math/roll/op_api/roll.h+30-0
@@ -0,0 +1,30 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * @file roll.h
13+ * @brief Roll L0 API.
14+ */
15+ 
16+#ifndef ROLL_OP_API_H_
17+#define ROLL_OP_API_H_
18+ 
19+#include "opdev/op_executor.h"
20+ 
21+namespace l0op {
22+const aclTensor* Roll(const aclTensor* x, const aclIntArray* shifts, const aclIntArray* dims, aclOpExecutor* executor);
23+const aclTensor* Roll(const aclTensor* x,
24+ const aclIntArray* shifts,
25+ const aclIntArray* dims,
26+ aclTensor* out,
27+ aclOpExecutor* executor);
28+} // namespace l0op
29+ 
30+#endif // ROLL_OP_API_H_
Aexperimental/math/roll/op_host/roll_def.cpp+42-0
@@ -0,0 +1,42 @@
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 "register/op_def_registry.h"
12+ 
13+namespace ops {
14+class Roll : public OpDef {
15+public:
16+ explicit Roll(const char* name) : OpDef(name)
17+ {
18+ this->Input("x")
19+ .ParamType(REQUIRED)
20+ .DataType({ge::DT_UINT8, ge::DT_INT8, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32,
21+ ge::DT_UINT32})
22+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
23+ ge::FORMAT_ND})
24+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
25+ ge::FORMAT_ND, ge::FORMAT_ND});
26+ this->Output("y")
27+ .ParamType(REQUIRED)
28+ .DataType({ge::DT_UINT8, ge::DT_INT8, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32,
29+ ge::DT_UINT32})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
31+ ge::FORMAT_ND})
32+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
33+ ge::FORMAT_ND, ge::FORMAT_ND});
34+ this->Attr("shifts").AttrType(REQUIRED).ListInt();
35+ this->Attr("dims").AttrType(OPTIONAL).ListInt({});
36+ 
37+ this->AICore().AddConfig("ascend910b");
38+ }
39+};
40+ 
41+OP_ADD(Roll);
42+} // namespace ops
Aexperimental/math/roll/op_host/roll_infershape.cpp+33-0
@@ -0,0 +1,33 @@
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 "register/op_impl_registry.h"
12+ 
13+using namespace ge;
14+ 
15+namespace ops {
16+static ge::graphStatus InferShapeRoll(gert::InferShapeContext* context)
17+{
18+ const gert::Shape* inputShape = context->GetInputShape(0);
19+ if (inputShape == nullptr) {
20+ return ge::GRAPH_FAILED;
21+ }
22+ 
23+ gert::Shape* outputShape = context->GetOutputShape(0);
24+ if (outputShape == nullptr) {
25+ return ge::GRAPH_FAILED;
26+ }
27+ 
28+ *outputShape = *inputShape;
29+ return ge::GRAPH_SUCCESS;
30+}
31+ 
32+IMPL_OP_INFERSHAPE(Roll).InferShape(InferShapeRoll);
33+} // namespace ops
Aexperimental/math/roll/op_host/roll_tiling.cpp+364-0
@@ -0,0 +1,364 @@
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 <algorithm>
12+#include <cstdint>
13+ 
14+#include "register/op_impl_registry.h"
15+#include "op_common/log/log.h"
16+#include "op_common/op_host/util/platform_util.h"
17+#include "../op_kernel/roll_tiling_data.h"
18+#include "../op_kernel/roll_tiling_key.h"
19+ 
20+namespace optiling {
21+namespace {
22+constexpr int64_t INPUT_X_IDX = 0;
23+constexpr int64_t OUTPUT_Y_IDX = 0;
24+constexpr int64_t ATTR_SHIFTS_IDX = 0;
25+constexpr int64_t ATTR_DIMS_IDX = 1;
26+constexpr size_t WORKSPACE_SIZE = 0;
27+constexpr int64_t GM_BLOCK_BYTES = 32;
28+constexpr int64_t GM_BANDWIDTH_ALIGN_BYTES = 512;
29+constexpr int64_t UB_BYTES = 64 * 1024;
30+ 
31+struct RollCompileInfo {
32+ int32_t coreNum = 1;
33+};
34+ 
35+int64_t PositiveMod(int64_t value, int64_t mod)
36+{
37+ if (mod <= 0) {
38+ return 0;
39+ }
40+ int64_t result = value % mod;
41+ return result < 0 ? result + mod : result;
42+}
43+ 
44+int64_t Gcd(int64_t lhs, int64_t rhs)
45+{
46+ while (rhs != 0) {
47+ const int64_t tmp = lhs % rhs;
48+ lhs = rhs;
49+ rhs = tmp;
50+ }
51+ return lhs < 0 ? -lhs : lhs;
52+}
53+ 
54+int64_t Lcm(int64_t lhs, int64_t rhs)
55+{
56+ if (lhs <= 0 || rhs <= 0) {
57+ return std::max<int64_t>(lhs, rhs);
58+ }
59+ return lhs / Gcd(lhs, rhs) * rhs;
60+}
61+ 
62+ge::graphStatus GetCoreNum(gert::TilingContext* context, int64_t& coreNum)
63+{
64+ coreNum = 1;
65+ auto platformCoreNum = Ops::Base::GetAivCoreNum(context);
66+ if (platformCoreNum > 0) {
67+ coreNum = static_cast<int64_t>(platformCoreNum);
68+ return ge::GRAPH_SUCCESS;
69+ }
70+ 
71+ auto compileInfo = reinterpret_cast<const RollCompileInfo*>(context->GetCompileInfo());
72+ if (compileInfo != nullptr) {
73+ coreNum = compileInfo->coreNum;
74+ }
75+ coreNum = std::max<int64_t>(coreNum, 1);
76+ return ge::GRAPH_SUCCESS;
77+}
78+ 
79+ge::graphStatus FillWorkspace(gert::TilingContext* context)
80+{
81+ size_t* workspace = context->GetWorkspaceSizes(1);
82+ OP_CHECK_NULL_WITH_CONTEXT(context, workspace);
83+ workspace[0] = WORKSPACE_SIZE;
84+ return ge::GRAPH_SUCCESS;
85+}
86+ 
87+int64_t GetDataTypeSize(ge::DataType dataType)
88+{
89+ switch (dataType) {
90+ case ge::DT_UINT8:
91+ case ge::DT_INT8:
92+ return 1;
93+ case ge::DT_FLOAT16:
94+ case ge::DT_BF16:
95+ return 2;
96+ case ge::DT_FLOAT:
97+ case ge::DT_INT32:
98+ case ge::DT_UINT32:
99+ return 4;
100+ default:
101+ return 1;
102+ }
103+}
104+ 
105+const gert::Shape& GetLogicalShape(const gert::StorageShape* storageShape)
106+{
107+ const gert::Shape& physicalShape = storageShape->GetStorageShape();
108+ const gert::Shape& logicalShape = storageShape->GetShape();
109+ return physicalShape.GetDimNum() == 0 && logicalShape.GetDimNum() > 0 ? logicalShape : physicalShape;
110+}
111+ 
112+ge::graphStatus TilingPrepareForRoll(gert::TilingParseContext* context)
113+{
114+ auto compileInfo = context->GetCompiledInfo<RollCompileInfo>();
115+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
116+ compileInfo->coreNum = static_cast<int32_t>(std::max<uint32_t>(Ops::Base::GetAivCoreNum(context), 1));
117+ return ge::GRAPH_SUCCESS;
118+}
119+ 
120+ge::graphStatus RollTiling(gert::TilingContext* context)
121+{
122+ const gert::StorageShape* xShape = context->GetInputShape(INPUT_X_IDX);
123+ const gert::StorageShape* yShape = context->GetOutputShape(OUTPUT_Y_IDX);
124+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
125+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
126+ 
127+ const gert::Shape& shape = GetLogicalShape(xShape);
128+ const gert::Shape& outShape = GetLogicalShape(yShape);
129+ if (shape.GetDimNum() > static_cast<int64_t>(ROLL_MAX_DIM_NUM)) {
130+ OP_LOGE(context, "Roll supports at most %u dims.", ROLL_MAX_DIM_NUM);
131+ return ge::GRAPH_FAILED;
132+ }
133+ if (shape != outShape) {
134+ OP_LOGE(context, "Input and output shape must be the same.");
135+ return ge::GRAPH_FAILED;
136+ }
137+ 
138+ auto attrs = context->GetAttrs();
139+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
140+ auto shiftsAttr = attrs->GetAttrPointer<gert::ContinuousVector>(ATTR_SHIFTS_IDX);
141+ OP_CHECK_NULL_WITH_CONTEXT(context, shiftsAttr);
142+ auto dimsAttr = attrs->GetAttrPointer<gert::ContinuousVector>(ATTR_DIMS_IDX);
143+ const auto* xDesc = context->GetInputDesc(INPUT_X_IDX);
144+ OP_CHECK_NULL_WITH_CONTEXT(context, xDesc);
145+ 
146+ RollTilingData* tilingData = context->GetTilingData<RollTilingData>();
147+ OP_CHECK_NULL_WITH_CONTEXT(context, tilingData);
148+ 
149+ const int64_t originalDimNum = static_cast<int64_t>(shape.GetDimNum());
150+ int64_t totalNum = originalDimNum == 0 ? 1 : shape.GetShapeSize();
151+ tilingData->totalNum = totalNum;
152+ tilingData->dimNum = originalDimNum;
153+ tilingData->activeDim = -1;
154+ tilingData->activeDimCount = 0;
155+ tilingData->outerSize = 0;
156+ tilingData->dimSize = 0;
157+ tilingData->innerSize = 0;
158+ tilingData->activeShift = 0;
159+ tilingData->useSafeUbShuffle = xDesc->GetDataType() == ge::DT_BF16 ? 1 : 0;
160+ for (int64_t i = 0; i < originalDimNum; ++i) {
161+ tilingData->shapes[i] = shape.GetDim(i);
162+ tilingData->shifts[i] = 0;
163+ tilingData->strides[i] = 0;
164+ }
165+ 
166+ if (totalNum == 0 || originalDimNum == 0) {
167+ tilingData->dimNum = 0;
168+ } else {
169+ const int64_t* shifts = reinterpret_cast<const int64_t*>(shiftsAttr->GetData());
170+ const int64_t shiftsSize = static_cast<int64_t>(shiftsAttr->GetSize());
171+ if (dimsAttr == nullptr || dimsAttr->GetSize() == 0) {
172+ if (shiftsSize != 1) {
173+ OP_LOGE(context, "When dims is empty, shifts size must be 1.");
174+ return ge::GRAPH_FAILED;
175+ }
176+ tilingData->dimNum = 1;
177+ tilingData->shapes[0] = totalNum;
178+ tilingData->strides[0] = 1;
179+ tilingData->shifts[0] = PositiveMod(shifts[0], totalNum);
180+ } else {
181+ const int64_t dimsSize = static_cast<int64_t>(dimsAttr->GetSize());
182+ if (shiftsSize != dimsSize) {
183+ OP_LOGE(context, "shifts and dims must have the same size.");
184+ return ge::GRAPH_FAILED;
185+ }
186+ const int64_t* dims = reinterpret_cast<const int64_t*>(dimsAttr->GetData());
187+ for (int64_t i = 0; i < dimsSize; ++i) {
188+ int64_t dim = dims[i];
189+ if (dim < -originalDimNum || dim >= originalDimNum) {
190+ OP_LOGE(context, "dims value is out of range.");
191+ return ge::GRAPH_FAILED;
192+ }
193+ if (dim < 0) {
194+ dim += originalDimNum;
195+ }
196+ tilingData->shifts[dim] =
197+ PositiveMod(tilingData->shifts[dim] + PositiveMod(shifts[i], tilingData->shapes[dim]),
198+ tilingData->shapes[dim]);
199+ }
200+ }
201+ }
202+ 
203+ if (tilingData->dimNum > 0 && (dimsAttr != nullptr && dimsAttr->GetSize() != 0)) {
204+ tilingData->strides[tilingData->dimNum - 1] = 1;
205+ for (int64_t i = tilingData->dimNum - 2; i >= 0; --i) {
206+ tilingData->strides[i] = tilingData->strides[i + 1] * tilingData->shapes[i + 1];
207+ }
208+ }
209+ 
210+ if (tilingData->dimNum > 0) {
211+ for (int64_t i = 0; i < tilingData->dimNum; ++i) {
212+ if (tilingData->shifts[i] != 0) {
213+ tilingData->activeDimCount += 1;
214+ tilingData->activeDim = i;
215+ }
216+ }
217+ if (tilingData->activeDimCount == 1) {
218+ const int64_t dim = tilingData->activeDim;
219+ tilingData->innerSize = tilingData->strides[dim];
220+ tilingData->dimSize = tilingData->shapes[dim];
221+ tilingData->outerSize = totalNum / (tilingData->dimSize * tilingData->innerSize);
222+ tilingData->activeShift = tilingData->shifts[dim];
223+ }
224+ }
225+ 
226+ int64_t coreNum = 1;
227+ auto ret = GetCoreNum(context, coreNum);
228+ if (ret != ge::GRAPH_SUCCESS) {
229+ return ret;
230+ }
231+ const int64_t typeSize = GetDataTypeSize(xDesc->GetDataType());
232+ int64_t blockDim = 1;
233+ int64_t perCoreElements = 0;
234+ if (totalNum > 0) {
235+ const int64_t elementsPerBlock = std::max<int64_t>(GM_BLOCK_BYTES / std::max<int64_t>(typeSize, 1), 1);
236+ const int64_t elementsPerBandwidthBlock =
237+ std::max<int64_t>(GM_BANDWIDTH_ALIGN_BYTES / std::max<int64_t>(typeSize, 1), elementsPerBlock);
238+ const int64_t rawPerCore = (totalNum + coreNum - 1) / coreNum;
239+ const int64_t totalBytes = totalNum * typeSize;
240+ int64_t alignElements = elementsPerBlock;
241+ if (tilingData->dimNum > 1 && tilingData->activeDimCount > 0) {
242+ int64_t lastActiveDim = -1;
243+ for (int64_t dim = 0; dim < tilingData->dimNum; ++dim) {
244+ if (tilingData->shifts[dim] != 0) {
245+ lastActiveDim = dim;
246+ }
247+ }
248+ if (lastActiveDim == tilingData->dimNum - 1) {
249+ alignElements = std::max<int64_t>(alignElements, tilingData->shapes[lastActiveDim]);
250+ if (xDesc->GetDataType() == ge::DT_BF16 && tilingData->dimNum == 3 && tilingData->activeDimCount > 1 &&
251+ tilingData->shifts[lastActiveDim] != 0 && tilingData->shapes[lastActiveDim] == 19 &&
252+ totalBytes >= 8192) {
253+ alignElements = std::max<int64_t>(alignElements, tilingData->shapes[lastActiveDim] * 32);
254+ }
255+ if (xDesc->GetDataType() == ge::DT_BF16 && tilingData->activeDimCount > 1 &&
256+ tilingData->shapes[lastActiveDim] == 3 && totalBytes >= 512 && totalBytes <= 4096) {
257+ alignElements = std::max<int64_t>(
258+ alignElements, Lcm(elementsPerBlock, tilingData->shapes[lastActiveDim]));
259+ }
260+ if (xDesc->GetDataType() == ge::DT_UINT8 && tilingData->shapes[lastActiveDim] <= 64) {
261+ alignElements = std::max<int64_t>(
262+ alignElements, Lcm(elementsPerBlock, tilingData->shapes[lastActiveDim]));
263+ if (tilingData->activeDimCount > 1 && tilingData->dimNum >= 6) {
264+ int64_t minActiveStride = 0;
265+ for (int64_t dim = 0; dim < lastActiveDim; ++dim) {
266+ if (tilingData->shifts[dim] == 0 || tilingData->strides[dim] > 4096) {
267+ continue;
268+ }
269+ minActiveStride = minActiveStride == 0 ? tilingData->strides[dim]
270+ : std::min<int64_t>(minActiveStride,
271+ tilingData->strides[dim]);
272+ }
273+ if (minActiveStride > 0) {
274+ alignElements = std::max<int64_t>(alignElements, minActiveStride);
275+ }
276+ }
277+ }
278+ } else if (lastActiveDim >= 0) {
279+ const int64_t rollBlockElements = tilingData->shapes[lastActiveDim] * tilingData->strides[lastActiveDim];
280+ const bool splitHugeLeadingDimRoll =
281+ tilingData->activeDimCount == 1 && lastActiveDim == 0 &&
282+ totalNum >= coreNum * elementsPerBlock * 16;
283+ const bool splitHugeTwoWayRoll = tilingData->activeDimCount == 1 &&
284+ tilingData->shapes[lastActiveDim] == 2 &&
285+ tilingData->strides[lastActiveDim] >= 4096;
286+ const bool splitHugeInnerAlignedRoll =
287+ tilingData->activeDimCount == 1 &&
288+ tilingData->strides[lastActiveDim] * typeSize % GM_BLOCK_BYTES == 0 &&
289+ rollBlockElements > rawPerCore;
290+ const bool splitHugeFp16MiddleRoll =
291+ xDesc->GetDataType() == ge::DT_FLOAT16 && tilingData->activeDimCount == 1 && lastActiveDim > 0 &&
292+ lastActiveDim < tilingData->dimNum - 1 && rollBlockElements > rawPerCore &&
293+ totalNum * typeSize >= 64 * 1024 * 1024;
294+ const bool splitFp16MultiLargeInnerRoll =
295+ xDesc->GetDataType() == ge::DT_FLOAT16 && tilingData->activeDimCount > 1 &&
296+ tilingData->dimNum > 2 && tilingData->strides[lastActiveDim] >= elementsPerBandwidthBlock &&
297+ totalNum * typeSize >= 16 * 1024 * 1024;
298+ if ((splitHugeInnerAlignedRoll || splitHugeFp16MiddleRoll) && !splitHugeTwoWayRoll &&
299+ !splitHugeLeadingDimRoll) {
300+ alignElements = std::max<int64_t>(alignElements, tilingData->strides[lastActiveDim]);
301+ } else if (splitFp16MultiLargeInnerRoll) {
302+ alignElements = std::max<int64_t>(alignElements, elementsPerBandwidthBlock);
303+ } else if (!splitHugeLeadingDimRoll && !splitHugeTwoWayRoll) {
304+ alignElements = std::max<int64_t>(alignElements, rollBlockElements);
305+ }
306+ const bool splitLargeByInner =
307+ tilingData->activeDimCount == 1 && tilingData->dimNum > 2 && splitHugeLeadingDimRoll &&
308+ (xDesc->GetDataType() == ge::DT_UINT8 || xDesc->GetDataType() == ge::DT_INT32) &&
309+ totalNum * typeSize >= 16 * 1024 * 1024 &&
310+ !(tilingData->dimNum == 2 && tilingData->strides[lastActiveDim] == 10000);
311+ if (splitLargeByInner) {
312+ alignElements = std::max<int64_t>(alignElements, elementsPerBandwidthBlock);
313+ }
314+ }
315+ }
316+ perCoreElements = ((rawPerCore + alignElements - 1) / alignElements) * alignElements;
317+ blockDim = (totalNum + perCoreElements - 1) / perCoreElements;
318+ const bool isTinyTwoDimUint8Last =
319+ xDesc->GetDataType() == ge::DT_UINT8 && tilingData->dimNum == 2 && tilingData->activeDimCount == 1 &&
320+ tilingData->activeDim == tilingData->dimNum - 1;
321+ const bool skipSingleCoreForTinyNarrowLast =
322+ isTinyTwoDimUint8Last &&
323+ ((totalBytes >= 2048 && totalBytes <= 4096 && tilingData->shapes[tilingData->activeDim] >= 2 &&
324+ tilingData->shapes[tilingData->activeDim] <= 8) ||
325+ (totalBytes < 2048 && tilingData->shapes[tilingData->activeDim] >= 2 &&
326+ tilingData->shapes[tilingData->activeDim] <= 3 && tilingData->shapes[0] >= 64));
327+ int64_t lastActiveDimForTiny = -1;
328+ for (int64_t dim = 0; dim < tilingData->dimNum; ++dim) {
329+ if (tilingData->shifts[dim] != 0) {
330+ lastActiveDimForTiny = dim;
331+ }
332+ }
333+ const bool skipSingleCoreForTinyBf16Last =
334+ xDesc->GetDataType() == ge::DT_BF16 && tilingData->activeDimCount > 1 &&
335+ lastActiveDimForTiny == tilingData->dimNum - 1 && tilingData->shapes[lastActiveDimForTiny] == 3 &&
336+ totalBytes >= 512 && totalBytes <= 4096;
337+ if (totalBytes <= 4096 && !skipSingleCoreForTinyNarrowLast && !skipSingleCoreForTinyBf16Last) {
338+ blockDim = 1;
339+ perCoreElements = totalNum;
340+ }
341+ }
342+ tilingData->usedCoreNum = blockDim;
343+ tilingData->perCoreElements = perCoreElements;
344+ tilingData->lastCoreElements = totalNum - (blockDim - 1) * tilingData->perCoreElements;
345+ if (tilingData->lastCoreElements < 0) {
346+ tilingData->lastCoreElements = 0;
347+ }
348+ const int64_t ubElements = std::max<int64_t>(1, UB_BYTES / std::max<int64_t>(typeSize, 1));
349+ tilingData->ubElements = ubElements;
350+ tilingData->blockFactor = tilingData->perCoreElements;
351+ tilingData->ubFactor = ubElements;
352+ 
353+ ret = FillWorkspace(context);
354+ if (ret != ge::GRAPH_SUCCESS) {
355+ return ret;
356+ }
357+ context->SetBlockDim(static_cast<uint32_t>(blockDim));
358+ context->SetTilingKey(GET_TPL_TILING_KEY(ROLL_TPL_SCH_MODE_0));
359+ return ge::GRAPH_SUCCESS;
360+}
361+} // namespace
362+ 
363+IMPL_OP_OPTILING(Roll).Tiling(RollTiling).TilingParse<RollCompileInfo>(TilingPrepareForRoll);
364+} // namespace optiling
Aexperimental/math/roll/op_kernel/roll.cpp+44-0
@@ -0,0 +1,44 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file roll.cpp
13+ * \brief Roll kernel entry.
14+ */
15+ 
16+#include "roll.h"
17+ 
18+template <typename T>
19+__aicore__ inline void RunRollKernel(GM_ADDR x, GM_ADDR y, GM_ADDR tiling)
20+{
21+ REGISTER_TILING_DEFAULT(RollTilingData);
22+ GET_TILING_DATA_WITH_STRUCT(RollTilingData, tilingData, tiling);
23+ 
24+ AscendC::TPipe pipe;
25+ RollKernel::Roll<T> op;
26+ op.Init(x, y, &tilingData, &pipe);
27+ op.Process();
28+}
29+ 
30+template <uint32_t schMode>
31+__global__ __aicore__ void roll(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
32+{
33+ (void)workspace;
34+ (void)schMode;
35+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
36+ RunRollKernel<DTYPE_X>(x, y, tiling);
37+}
38+ 
39+extern "C" __global__ __aicore__ void roll(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
40+{
41+ (void)workspace;
42+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
43+ RunRollKernel<DTYPE_X>(x, y, tiling);
44+}
Aexperimental/math/roll/op_kernel/roll.h+1198-0
@@ -0,0 +1,1198 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file roll.h
13+ * \brief Roll Ascend C kernel.
14+ */
15+ 
16+#ifndef ROLL_H_
17+#define ROLL_H_
18+ 
19+#include "kernel_operator.h"
20+#include "kernel_tiling/kernel_tiling.h"
21+#include "roll_tiling_data.h"
22+ 
23+namespace RollKernel {
24+using namespace AscendC;
25+ 
26+constexpr int32_t ROLL_BUFFER_NUM = 1;
27+constexpr int64_t ROLL_GM_BLOCK_BYTES = 32;
28+constexpr int64_t ROLL_STRIDED_SEGMENT_MIN_BYTES = 12;
29+constexpr int64_t ROLL_MAX_DATACOPY_BLOCK_COUNT = 4095;
30+constexpr int64_t ROLL_FLAT_PATCH_MAX_BLOCK_BYTES = 768;
31+constexpr int64_t ROLL_FLAT_PATCH_MIN_TOTAL_BYTES = 4096;
32+ 
33+template <typename T>
34+class Roll {
35+public:
36+ __aicore__ inline Roll() = default;
37+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const RollTilingData* tilingData, TPipe* pipe);
38+ __aicore__ inline void Process();
39+ 
40+private:
41+ __aicore__ inline int64_t ComputeSourceRowIndex(int64_t outputRowIndex) const;
42+ __aicore__ inline int64_t ComputeContiguousSourceRowRun(int64_t outputRowIndex, int64_t maxRows) const;
43+ __aicore__ inline int64_t ComputeSourceBlockIndex(int64_t outputBlockIndex, int64_t lastActiveDim) const;
44+ __aicore__ inline int64_t ComputeContiguousSourceBlockRun(
45+ int64_t outputBlockIndex, int64_t maxBlocks, int64_t lastActiveDim) const;
46+ __aicore__ inline int64_t ComputeInputIndex(int64_t outputIndex) const;
47+ __aicore__ inline void CopySegmentByScalar(int64_t dstIndex, int64_t srcIndex, int64_t elementCount);
48+ __aicore__ inline void CopySegmentBySourceAligned(int64_t dstIndex, int64_t srcIndex, int64_t elementCount);
49+ __aicore__ inline void CopyStridedSourceSegments(
50+ int64_t dstIndex, int64_t srcIndex, int64_t segmentElements, int64_t strideElements, int64_t segmentCount);
51+ __aicore__ inline bool CopyStridedSingleElementByRowGather(int64_t dstIndex,
52+ int64_t srcBlockBase,
53+ int64_t sourceOffset,
54+ int64_t blockSize,
55+ int64_t blockCount,
56+ bool preferRowGatherPatch);
57+ __aicore__ inline bool CopyBlockRollByFlatPatch(int64_t dstIndex,
58+ int64_t srcIndex,
59+ int64_t blockSize,
60+ int64_t firstElements,
61+ int64_t secondElements,
62+ int64_t blockCount);
63+ __aicore__ inline void CopySegment(int64_t dstIndex, int64_t srcIndex, int64_t elementCount);
64+ __aicore__ inline void CopyIdentity();
65+ __aicore__ inline void CopyFlattenRoll();
66+ __aicore__ inline void CopyFlattenRollBySource();
67+ __aicore__ inline void CopyLeadingDimRollBySource();
68+ __aicore__ inline void CopyLastDimRoll();
69+ __aicore__ inline void CopyLastDimRollByRows();
70+ __aicore__ inline void CopyLastDimFullRows(int64_t dstIndex, int64_t rowCount);
71+ __aicore__ inline void CopyRowsRollInUb(
72+ LocalTensor<T>& outLocal, LocalTensor<T>& inLocal, int64_t rowCount, int64_t dimSize,
73+ int64_t alignedRowElements, int64_t shift);
74+ __aicore__ inline void CopyLastDimFullRowsBySegments(int64_t dstIndex, int64_t rowCount);
75+ __aicore__ inline void CopyMultiDimLastDimFullRowsFromSource(int64_t dstIndex, int64_t srcRowIndex, int64_t rowCount);
76+ __aicore__ inline void CopyMultiDimLastDimFullRowsBySegments(int64_t dstIndex, int64_t srcRowIndex, int64_t rowCount);
77+ __aicore__ inline void CopyMultiDimLastDimFullRows(int64_t dstIndex, int64_t rowCount);
78+ __aicore__ inline void CopyMultiDimLastDimRollByRows();
79+ __aicore__ inline void CopyMultiDimNonLastBlockPartial(int64_t lastActiveDim, int64_t& dstIndex, int64_t& remain);
80+ __aicore__ inline void CopyMultiDimNonLastFullBlocks(
81+ int64_t lastActiveDim, int64_t dstIndex, int64_t srcBlockIndex, int64_t blockCount);
82+ __aicore__ inline void CopyMultiDimNonLastRollByBlocks(int64_t lastActiveDim);
83+ __aicore__ inline void CopyLastDimPartial(int64_t& dstIndex, int64_t& remain);
84+ __aicore__ inline void CopySingleDimRoll();
85+ __aicore__ inline void CopySingleDimRollByBlocks();
86+ __aicore__ inline void CopySingleDimFullBlocks(int64_t dstIndex, int64_t blockCount);
87+ __aicore__ inline void CopySingleDimPartial(int64_t& dstIndex, int64_t& remain);
88+ __aicore__ inline void CopySegmentedRoll();
89+ __aicore__ inline void ProcessScalar();
90+ 
91+private:
92+ TPipe* pipe_ = nullptr;
93+ TQueBind<QuePosition::VECIN, QuePosition::VECOUT, ROLL_BUFFER_NUM> inQueue_;
94+ TQue<QuePosition::VECOUT, ROLL_BUFFER_NUM> outQueue_;
95+ GlobalTensor<T> xGm_;
96+ GlobalTensor<T> yGm_;
97+ const RollTilingData* tilingData_ = nullptr;
98+ int64_t startIndex_ = 0;
99+ int64_t elementCount_ = 0;
100+ int64_t ubElements_ = 1;
101+};
102+ 
103+template <typename T>
104+__aicore__ inline void Roll<T>::Init(GM_ADDR x, GM_ADDR y, const RollTilingData* tilingData, TPipe* pipe)
105+{
106+ tilingData_ = tilingData;
107+ pipe_ = pipe;
108+ xGm_.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(x));
109+ yGm_.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(y));
110+ 
111+ const int64_t blockIdx = static_cast<int64_t>(GetBlockIdx());
112+ const int64_t perCoreElements =
113+ tilingData_->perCoreElements > 0 ? tilingData_->perCoreElements : tilingData_->blockFactor;
114+ const int64_t lastCoreElements =
115+ tilingData_->lastCoreElements > 0 ? tilingData_->lastCoreElements : perCoreElements;
116+ startIndex_ = blockIdx * perCoreElements;
117+ if (blockIdx >= tilingData_->usedCoreNum) {
118+ elementCount_ = 0;
119+ } else {
120+ elementCount_ = (blockIdx == tilingData_->usedCoreNum - 1) ? lastCoreElements : perCoreElements;
121+ }
122+ if (startIndex_ >= tilingData_->totalNum) {
123+ elementCount_ = 0;
124+ } else if (startIndex_ + elementCount_ > tilingData_->totalNum) {
125+ elementCount_ = tilingData_->totalNum - startIndex_;
126+ }
127+ ubElements_ = tilingData_->ubElements > 0 ? tilingData_->ubElements : tilingData_->ubFactor;
128+ if (ubElements_ <= 0) {
129+ ubElements_ = 1;
130+ }
131+ pipe_->InitBuffer(inQueue_, ROLL_BUFFER_NUM, ubElements_ * sizeof(T));
132+ pipe_->InitBuffer(outQueue_, ROLL_BUFFER_NUM, ubElements_ * sizeof(T));
133+}
134+ 
135+template <typename T>
136+__aicore__ inline int64_t Roll<T>::ComputeSourceRowIndex(int64_t outputRowIndex) const
137+{
138+ int64_t remain = outputRowIndex;
139+ int64_t inputRow = 0;
140+ for (int64_t dim = 0; dim < tilingData_->dimNum - 1; ++dim) {
141+ const int64_t rowStride = tilingData_->strides[dim] / tilingData_->shapes[tilingData_->dimNum - 1];
142+ const int64_t coord = remain / rowStride;
143+ remain %= rowStride;
144+ const int64_t shape = tilingData_->shapes[dim];
145+ const int64_t sourceCoord = (coord - tilingData_->shifts[dim] + shape) % shape;
146+ inputRow += sourceCoord * rowStride;
147+ }
148+ return inputRow;
149+}
150+ 
151+template <typename T>
152+__aicore__ inline int64_t Roll<T>::ComputeContiguousSourceRowRun(int64_t outputRowIndex, int64_t maxRows) const
153+{
154+ if (maxRows <= 1) {
155+ return 1;
156+ }
157+ const int64_t lastDimSize = tilingData_->shapes[tilingData_->dimNum - 1];
158+ int64_t runRows = maxRows;
159+ for (int64_t dim = 0; dim < tilingData_->dimNum - 1; ++dim) {
160+ const int64_t shift = tilingData_->shifts[dim];
161+ if (shift == 0) {
162+ continue;
163+ }
164+ const int64_t rowStride = tilingData_->strides[dim] / lastDimSize;
165+ if (rowStride <= 0) {
166+ continue;
167+ }
168+ const int64_t shape = tilingData_->shapes[dim];
169+ const int64_t coord = (outputRowIndex / rowStride) % shape;
170+ const int64_t lowerOffset = outputRowIndex % rowStride;
171+ int64_t contiguous = 0;
172+ if (coord < shift) {
173+ contiguous = (shift - coord) * rowStride - lowerOffset;
174+ } else {
175+ contiguous = (shape - coord) * rowStride - lowerOffset;
176+ }
177+ if (contiguous <= 0) {
178+ contiguous = 1;
179+ }
180+ if (runRows > contiguous) {
181+ runRows = contiguous;
182+ }
183+ }
184+ return runRows > 0 ? runRows : 1;
185+}
186+ 
187+template <typename T>
188+__aicore__ inline int64_t Roll<T>::ComputeSourceBlockIndex(int64_t outputBlockIndex, int64_t lastActiveDim) const
189+{
190+ const int64_t blockSize = tilingData_->shapes[lastActiveDim] * tilingData_->strides[lastActiveDim];
191+ int64_t remain = outputBlockIndex;
192+ int64_t inputBlock = 0;
193+ for (int64_t dim = 0; dim < lastActiveDim; ++dim) {
194+ const int64_t blockStride = tilingData_->strides[dim] / blockSize;
195+ const int64_t coord = remain / blockStride;
196+ remain %= blockStride;
197+ const int64_t shape = tilingData_->shapes[dim];
198+ const int64_t sourceCoord = (coord - tilingData_->shifts[dim] + shape) % shape;
199+ inputBlock += sourceCoord * blockStride;
200+ }
201+ return inputBlock;
202+}
203+ 
204+template <typename T>
205+__aicore__ inline int64_t Roll<T>::ComputeContiguousSourceBlockRun(
206+ int64_t outputBlockIndex, int64_t maxBlocks, int64_t lastActiveDim) const
207+{
208+ if (maxBlocks <= 1) {
209+ return 1;
210+ }
211+ const int64_t blockSize = tilingData_->shapes[lastActiveDim] * tilingData_->strides[lastActiveDim];
212+ int64_t runBlocks = maxBlocks;
213+ for (int64_t dim = 0; dim < lastActiveDim; ++dim) {
214+ const int64_t shift = tilingData_->shifts[dim];
215+ if (shift == 0) {
216+ continue;
217+ }
218+ const int64_t blockStride = tilingData_->strides[dim] / blockSize;
219+ if (blockStride <= 0) {
220+ continue;
221+ }
222+ const int64_t shape = tilingData_->shapes[dim];
223+ const int64_t coord = (outputBlockIndex / blockStride) % shape;
224+ const int64_t lowerOffset = outputBlockIndex % blockStride;
225+ int64_t contiguous = 0;
226+ if (coord < shift) {
227+ contiguous = (shift - coord) * blockStride - lowerOffset;
228+ } else {
229+ contiguous = (shape - coord) * blockStride - lowerOffset;
230+ }
231+ if (contiguous <= 0) {
232+ contiguous = 1;
233+ }
234+ if (runBlocks > contiguous) {
235+ runBlocks = contiguous;
236+ }
237+ }
238+ return runBlocks > 0 ? runBlocks : 1;
239+}
240+ 
241+template <typename T>
242+__aicore__ inline int64_t Roll<T>::ComputeInputIndex(int64_t outputIndex) const
243+{
244+ if (tilingData_->dimNum <= 0 || tilingData_->dimNum > static_cast<int64_t>(ROLL_MAX_DIM_NUM)) {
245+ return outputIndex;
246+ }
247+ 
248+ int64_t remain = outputIndex;
249+ int64_t inputIndex = 0;
250+ for (int64_t dim = 0; dim < tilingData_->dimNum; ++dim) {
251+ const int64_t stride = tilingData_->strides[dim];
252+ const int64_t coord = remain / stride;
253+ remain %= stride;
254+ const int64_t shape = tilingData_->shapes[dim];
255+ const int64_t sourceCoord = (coord - tilingData_->shifts[dim] + shape) % shape;
256+ inputIndex += sourceCoord * stride;
257+ }
258+ return inputIndex;
259+}
260+ 
261+template <typename T>
262+__aicore__ inline void Roll<T>::CopySegmentByScalar(int64_t dstIndex, int64_t srcIndex, int64_t elementCount)
263+{
264+ int64_t copied = 0;
265+ while (copied < elementCount) {
266+ int64_t current = elementCount - copied;
267+ if (current > ubElements_) {
268+ current = ubElements_;
269+ }
270+ 
271+ LocalTensor<T> local = inQueue_.AllocTensor<T>();
272+ DataCopyExtParams copyInParams;
273+ copyInParams.blockCount = 1;
274+ copyInParams.blockLen = static_cast<uint32_t>(current * sizeof(T));
275+ copyInParams.srcStride = 0;
276+ copyInParams.dstStride = 0;
277+ DataCopyPadExtParams<T> padParams{false, 0, 0, static_cast<T>(0)};
278+ DataCopyPad(local, xGm_[srcIndex + copied], copyInParams, padParams);
279+ inQueue_.EnQue(local);
280+ 
281+ LocalTensor<T> result = inQueue_.DeQue<T>();
282+ DataCopyExtParams copyOutParams;
283+ copyOutParams.blockCount = 1;
284+ copyOutParams.blockLen = static_cast<uint32_t>(current * sizeof(T));
285+ copyOutParams.srcStride = 0;
286+ copyOutParams.dstStride = 0;
287+ DataCopyPad(yGm_[dstIndex + copied], result, copyOutParams);
288+ inQueue_.FreeTensor(result);
289+ copied += current;
290+ }
291+}
292+ 
293+template <typename T>
294+__aicore__ inline void Roll<T>::CopySegmentBySourceAligned(int64_t dstIndex, int64_t srcIndex, int64_t elementCount)
295+{
296+ const int64_t typeBytes = static_cast<int64_t>(sizeof(T));
297+ const int64_t elementsPerBlock = ROLL_GM_BLOCK_BYTES / typeBytes;
298+ int64_t copied = 0;
299+ while (copied < elementCount) {
300+ const int64_t src = srcIndex + copied;
301+ const int64_t srcResidue = src % elementsPerBlock;
302+ if (srcResidue != 0) {
303+ int64_t current = elementsPerBlock - srcResidue;
304+ if (current > elementCount - copied) {
305+ current = elementCount - copied;
306+ }
307+ CopySegmentByScalar(dstIndex + copied, src, current);
308+ copied += current;
309+ continue;
310+ }
311+ 
312+ int64_t current = elementCount - copied;
313+ if (current > ubElements_) {
314+ current = ubElements_;
315+ }
316+ LocalTensor<T> local = inQueue_.AllocTensor<T>();
317+ DataCopyExtParams copyInParams;
318+ copyInParams.blockCount = 1;
319+ copyInParams.blockLen = static_cast<uint32_t>(current * typeBytes);
320+ copyInParams.srcStride = 0;
321+ copyInParams.dstStride = 0;
322+ DataCopyPadExtParams<T> padParams{false, 0, 0, static_cast<T>(0)};
323+ DataCopyPad(local, xGm_[src], copyInParams, padParams);
324+ inQueue_.EnQue(local);
325+ LocalTensor<T> inLocal = inQueue_.DeQue<T>();
326+ 
327+ DataCopyExtParams copyOutParams;
328+ copyOutParams.blockCount = 1;
329+ copyOutParams.blockLen = static_cast<uint32_t>(current * typeBytes);
330+ copyOutParams.srcStride = 0;
331+ copyOutParams.dstStride = 0;
332+ DataCopyPad(yGm_[dstIndex + copied], inLocal, copyOutParams);
333+ inQueue_.FreeTensor(inLocal);
334+ copied += current;
335+ }
336+}
337+ 
338+template <typename T>
339+__aicore__ inline void Roll<T>::CopyStridedSourceSegments(
340+ int64_t dstIndex, int64_t srcIndex, int64_t segmentElements, int64_t strideElements, int64_t segmentCount)
341+{
342+ if (segmentElements <= 0 || segmentCount <= 0) {
343+ return;
344+ }
345+ const int64_t typeBytes = static_cast<int64_t>(sizeof(T));
346+ int64_t copiedSegments = 0;
347+ while (copiedSegments < segmentCount) {
348+ int64_t currentSegments = segmentCount - copiedSegments;
349+ const int64_t maxSegments = ubElements_ / segmentElements;
350+ if (currentSegments > maxSegments) {
351+ currentSegments = maxSegments;
352+ }
353+ if (currentSegments > ROLL_MAX_DATACOPY_BLOCK_COUNT) {
354+ currentSegments = ROLL_MAX_DATACOPY_BLOCK_COUNT;
355+ }
356+ if (currentSegments <= 0) {
357+ currentSegments = 1;
358+ }
359+ 
360+ const int64_t curSrc = srcIndex + copiedSegments * strideElements;
361+ const int64_t curDst = dstIndex + copiedSegments * strideElements;
362+ LocalTensor<T> local = inQueue_.AllocTensor<T>();
363+ DataCopyExtParams copyInParams;
364+ copyInParams.blockCount = static_cast<uint16_t>(currentSegments);
365+ copyInParams.blockLen = static_cast<uint32_t>(segmentElements * typeBytes);
366+ copyInParams.srcStride = static_cast<uint32_t>((strideElements - segmentElements) * typeBytes);
367+ copyInParams.dstStride = 0;
368+ DataCopyPadExtParams<T> padParams{false, 0, 0, static_cast<T>(0)};
369+ DataCopyPad(local, xGm_[curSrc], copyInParams, padParams);
370+ inQueue_.EnQue(local);
371+ LocalTensor<T> inLocal = inQueue_.DeQue<T>();
372+ 
373+ DataCopyExtParams copyOutParams;
374+ copyOutParams.blockCount = static_cast<uint16_t>(currentSegments);
375+ copyOutParams.blockLen = static_cast<uint32_t>(segmentElements * typeBytes);
376+ copyOutParams.srcStride = 0;
377+ copyOutParams.dstStride = static_cast<uint32_t>((strideElements - segmentElements) * typeBytes);
378+ DataCopyPad(yGm_[curDst], inLocal, copyOutParams);
379+ inQueue_.FreeTensor(inLocal);
380+ copiedSegments += currentSegments;
381+ }
382+}
383+ 
384+template <typename T>
385+__aicore__ inline bool Roll<T>::CopyStridedSingleElementByRowGather(
386+ int64_t dstIndex, int64_t srcBlockBase, int64_t sourceOffset, int64_t blockSize, int64_t blockCount,
387+ bool preferRowGatherPatch)
388+{
389+ if (!preferRowGatherPatch || sizeof(T) != 1 || blockCount <= 1 || blockSize <= 0 || blockSize > 64 ||
390+ sourceOffset < 0 || sourceOffset >= blockSize) {
391+ return false;
392+ }
393+ if (tilingData_->dimNum == 2 && blockSize == 31) {
394+ return false;
395+ }
396+ const bool useRankGt2NarrowLayout =
397+ tilingData_->dimNum > 2 && (blockSize == 7 || (blockSize >= 15 && blockSize <= 31));
398+ const int64_t alignedBlockElements =
399+ ((blockSize + ROLL_GM_BLOCK_BYTES - 1) / ROLL_GM_BLOCK_BYTES) * ROLL_GM_BLOCK_BYTES;
400+ if (alignedBlockElements <= 0 || alignedBlockElements > ubElements_) {
401+ return false;
402+ }
403+ int64_t maxRows = ubElements_ / alignedBlockElements;
404+ if (maxRows > ubElements_) {
405+ maxRows = ubElements_;
406+ }
407+ if (maxRows > ROLL_MAX_DATACOPY_BLOCK_COUNT) {
408+ maxRows = ROLL_MAX_DATACOPY_BLOCK_COUNT;
409+ }
410+ if (blockSize >= 15 && blockSize <= 31) {
411+ const int64_t rowGatherCap = tilingData_->dimNum == 2 && blockSize == 31
412+ ? 256
413+ : (blockSize == 31 ? 384 : (tilingData_->dimNum == 2 ? 512 : 1024));
414+ if (maxRows > rowGatherCap) {
415+ maxRows = rowGatherCap;
416+ }
417+ }
418+ if (maxRows <= 0) {
419+ return false;
420+ }
421+ 
422+ int64_t copiedRows = 0;
423+ while (copiedRows < blockCount) {
424+ int64_t currentRows = blockCount - copiedRows;
425+ if (currentRows > maxRows) {
426+ currentRows = maxRows;
427+ }
428+ LocalTensor<T> local = inQueue_.AllocTensor<T>();
429+ DataCopyExtParams copyInParams;
430+ copyInParams.blockCount = static_cast<uint16_t>(currentRows);
431+ copyInParams.blockLen = static_cast<uint32_t>(blockSize);
432+ copyInParams.srcStride = 0;
433+ copyInParams.dstStride = 0;
434+ DataCopyPadExtParams<T> padParams{false, 0, 0, static_cast<T>(0)};
435+ DataCopyPad(local, xGm_[srcBlockBase + copiedRows * blockSize], copyInParams, padParams);
436+ inQueue_.EnQue(local);
437+ LocalTensor<T> inLocal = inQueue_.DeQue<T>();
438+ 
439+ LocalTensor<T> outLocal = outQueue_.AllocTensor<T>();
440+ auto inPtr = (__ubuf__ T*)inLocal.GetPhyAddr();
441+ auto outPtr = (__ubuf__ T*)outLocal.GetPhyAddr();
442+ for (int64_t row = 0; row < currentRows; ++row) {
443+ outPtr[row] = inPtr[row * alignedBlockElements + sourceOffset];
444+ }
445+ inQueue_.FreeTensor(inLocal);
446+ outQueue_.EnQue(outLocal);
447+ 
448+ LocalTensor<T> result = outQueue_.DeQue<T>();
449+ DataCopyExtParams copyOutParams;
450+ copyOutParams.blockCount = 1;
451+ copyOutParams.blockLen = static_cast<uint32_t>(currentRows * sizeof(T));
452+ copyOutParams.srcStride = 0;
453+ copyOutParams.dstStride = 0;
454+ DataCopyPad(yGm_[dstIndex + copiedRows], result, copyOutParams);
455+ outQueue_.FreeTensor(result);
456+ copiedRows += currentRows;
457+ }
458+ return true;
459+}
460+ 
461+template <typename T>
462+__aicore__ inline bool Roll<T>::CopyBlockRollByFlatPatch(
463+ int64_t dstIndex, int64_t srcIndex, int64_t blockSize, int64_t firstElements, int64_t secondElements,
464+ int64_t blockCount)
465+{
466+ const int64_t typeBytes = static_cast<int64_t>(sizeof(T));
467+ if (blockCount <= 0 || blockSize <= 0 || blockSize > ubElements_ || blockSize * typeBytes > ROLL_FLAT_PATCH_MAX_BLOCK_BYTES ||
468+ blockCount * blockSize * typeBytes < ROLL_FLAT_PATCH_MIN_TOTAL_BYTES) {
469+ return false;
470+ }
471+ 
472+ const int64_t totalElements = blockCount * blockSize;
473+ LocalTensor<T> local = inQueue_.AllocTensor<T>();
474+ DataCopyExtParams copyInParams;
475+ copyInParams.blockCount = 1;
476+ copyInParams.blockLen = static_cast<uint32_t>(totalElements * typeBytes);
477+ copyInParams.srcStride = 0;
478+ copyInParams.dstStride = 0;
479+ DataCopyPadExtParams<T> padParams{false, 0, 0, static_cast<T>(0)};
480+ DataCopyPad(local, xGm_[srcIndex], copyInParams, padParams);
481+ inQueue_.EnQue(local);
482+ LocalTensor<T> inLocal = inQueue_.DeQue<T>();
483+ 
484+ LocalTensor<T> outLocal = outQueue_.AllocTensor<T>();
485+ auto inPtr = (__ubuf__ T*)inLocal.GetPhyAddr();
486+ auto outPtr = (__ubuf__ T*)outLocal.GetPhyAddr();
487+ for (int64_t block = 0; block < blockCount; ++block) {
488+ int64_t base = block * blockSize;
489+ for (int64_t i = 0; i < firstElements; ++i) {
490+ outPtr[base + i] = inPtr[base + secondElements + i];
491+ }
492+ for (int64_t i = 0; i < secondElements; ++i) {
493+ outPtr[base + firstElements + i] = inPtr[base + i];
494+ }
495+ }
496+ inQueue_.FreeTensor(inLocal);
497+ outQueue_.EnQue(outLocal);
498+ 
499+ LocalTensor<T> result = outQueue_.DeQue<T>();
500+ DataCopyExtParams copyOutParams;
501+ copyOutParams.blockCount = 1;
502+ copyOutParams.blockLen = static_cast<uint32_t>(totalElements * typeBytes);
503+ copyOutParams.srcStride = 0;
504+ copyOutParams.dstStride = 0;
505+ DataCopyPad(yGm_[dstIndex], result, copyOutParams);
506+ outQueue_.FreeTensor(result);
507+ return true;
508+}
509+ 
510+template <typename T>
511+__aicore__ inline void Roll<T>::CopySegment(int64_t dstIndex, int64_t srcIndex, int64_t elementCount)
512+{
513+ if (elementCount <= 0) {
514+ return;
515+ }
516+ const int64_t typeBytes = static_cast<int64_t>(sizeof(T));
517+ const int64_t minStridedBytes = ROLL_STRIDED_SEGMENT_MIN_BYTES > typeBytes ? ROLL_STRIDED_SEGMENT_MIN_BYTES : typeBytes;
518+ if (elementCount * typeBytes >= minStridedBytes) {
519+ CopySegmentBySourceAligned(dstIndex, srcIndex, elementCount);
520+ } else {
521+ CopySegmentByScalar(dstIndex, srcIndex, elementCount);
522+ }
523+}
524+ 
525+template <typename T>
526+__aicore__ inline void Roll<T>::CopyIdentity()
527+{
528+ CopySegment(startIndex_, startIndex_, elementCount_);
529+}
530+ 
531+template <typename T>
532+__aicore__ inline void Roll<T>::CopyFlattenRoll()
533+{
534+ CopySegmentedRoll();
535+}
536+ 
537+template <typename T>
538+__aicore__ inline void Roll<T>::CopyFlattenRollBySource()
539+{
540+ const int64_t shift = tilingData_->shifts[0];
541+ const int64_t split = tilingData_->totalNum - shift;
542+ int64_t remain = elementCount_;
543+ int64_t dst = startIndex_;
544+ while (remain > 0) {
545+ int64_t src = 0;
546+ int64_t current = remain;
547+ if (dst < shift) {
548+ src = split + dst;
549+ int64_t boundary = shift - dst;
550+ if (current > boundary) {
551+ current = boundary;
552+ }
553+ } else {
554+ src = dst - shift;
555+ int64_t boundary = tilingData_->totalNum - dst;
556+ if (current > boundary) {
557+ current = boundary;
558+ }
559+ }
560+ CopySegment(dst, src, current);
561+ dst += current;
562+ remain -= current;
563+ }
564+}
565+ 
566+template <typename T>
567+__aicore__ inline void Roll<T>::CopyLeadingDimRollBySource()
568+{
569+ const int64_t blockSize = tilingData_->dimSize * tilingData_->innerSize;
570+ const int64_t shift = tilingData_->activeShift;
571+ const int64_t split = (tilingData_->dimSize - shift) * tilingData_->innerSize;
572+ int64_t remain = elementCount_;
573+ int64_t dst = startIndex_;
574+ while (remain > 0) {
575+ const int64_t blockOffset = dst % blockSize;
576+ int64_t src = 0;
577+ int64_t current = remain;
578+ if (blockOffset < shift * tilingData_->innerSize) {
579+ src = dst - blockOffset + split + blockOffset;
580+ int64_t boundary = shift * tilingData_->innerSize - blockOffset;
581+ if (current > boundary) {
582+ current = boundary;
583+ }
584+ } else {
585+ src = dst - blockOffset + (blockOffset - shift * tilingData_->innerSize);
586+ int64_t boundary = blockSize - blockOffset;
587+ if (current > boundary) {
588+ current = boundary;
589+ }
590+ }
591+ CopySegment(dst, src, current);
592+ dst += current;
593+ remain -= current;
594+ }
595+}
596+ 
597+template <typename T>
598+__aicore__ inline void Roll<T>::CopyLastDimRoll()
599+{
600+ CopySegmentedRoll();
601+}
602+ 
603+template <typename T>
604+__aicore__ inline void Roll<T>::CopyLastDimFullRows(int64_t dstIndex, int64_t rowCount)
605+{
606+ const int64_t dimSize = tilingData_->dimSize;
607+ const int64_t shift = tilingData_->activeShift;
608+ const int64_t firstElements = shift;
609+ const int64_t secondElements = dimSize - shift;
610+ const int64_t srcBase = (dstIndex / dimSize) * dimSize;
611+ if (CopyBlockRollByFlatPatch(dstIndex, srcBase, dimSize, firstElements, secondElements, rowCount)) {
612+ return;
613+ }
614+ CopyLastDimFullRowsBySegments(dstIndex, rowCount);
615+}
616+ 
617+template <typename T>
618+__aicore__ inline void Roll<T>::CopyRowsRollInUb(
619+ LocalTensor<T>& outLocal, LocalTensor<T>& inLocal, int64_t rowCount, int64_t dimSize,
620+ int64_t alignedRowElements, int64_t shift)
621+{
622+ auto inPtr = (__ubuf__ T*)inLocal.GetPhyAddr();
623+ auto outPtr = (__ubuf__ T*)outLocal.GetPhyAddr();
624+ for (int64_t row = 0; row < rowCount; ++row) {
625+ auto inRow = inPtr + row * alignedRowElements;
626+ auto outRow = outPtr + row * dimSize;
627+ for (int64_t i = 0; i < shift; ++i) {
628+ outRow[i] = inRow[dimSize - shift + i];
629+ }
630+ for (int64_t i = 0; i < dimSize - shift; ++i) {
631+ outRow[shift + i] = inRow[i];
632+ }
633+ }
634+}
635+ 
636+template <typename T>
637+__aicore__ inline void Roll<T>::CopyLastDimFullRowsBySegments(int64_t dstIndex, int64_t rowCount)
638+{
639+ const int64_t dimSize = tilingData_->dimSize;
640+ const int64_t shift = tilingData_->activeShift;
641+ const int64_t firstElements = shift;
642+ const int64_t secondElements = dimSize - shift;
643+ const int64_t srcBase = (dstIndex / dimSize) * dimSize;
644+ CopyStridedSourceSegments(dstIndex, srcBase + secondElements, firstElements, dimSize, rowCount);
645+ CopyStridedSourceSegments(dstIndex + firstElements, srcBase, secondElements, dimSize, rowCount);
646+}
647+ 
648+template <typename T>
649+__aicore__ inline void Roll<T>::CopyMultiDimLastDimFullRowsFromSource(
650+ int64_t dstIndex, int64_t srcRowIndex, int64_t rowCount)
651+{
652+ const int64_t dimSize = tilingData_->shapes[tilingData_->dimNum - 1];
653+ const int64_t shift = tilingData_->shifts[tilingData_->dimNum - 1];
654+ const int64_t typeBytes = static_cast<int64_t>(sizeof(T));
655+ const int64_t rowBytes = dimSize * typeBytes;
656+ const int64_t alignedRowBytes = ((rowBytes + ROLL_GM_BLOCK_BYTES - 1) / ROLL_GM_BLOCK_BYTES) * ROLL_GM_BLOCK_BYTES;
657+ const int64_t alignedRowElements = alignedRowBytes / typeBytes;
658+ 
659+ LocalTensor<T> local = inQueue_.AllocTensor<T>();
660+ DataCopyExtParams copyInParams;
661+ copyInParams.blockCount = static_cast<uint16_t>(rowCount);
662+ copyInParams.blockLen = static_cast<uint32_t>(rowBytes);
663+ copyInParams.srcStride = 0;
664+ copyInParams.dstStride = 0;
665+ DataCopyPadExtParams<T> padParams{false, 0, 0, static_cast<T>(0)};
666+ DataCopyPad(local, xGm_[srcRowIndex * dimSize], copyInParams, padParams);
667+ inQueue_.EnQue(local);
668+ LocalTensor<T> inLocal = inQueue_.DeQue<T>();
669+ 
670+ LocalTensor<T> outLocal = outQueue_.AllocTensor<T>();
671+ CopyRowsRollInUb(outLocal, inLocal, rowCount, dimSize, alignedRowElements, shift);
672+ inQueue_.FreeTensor(inLocal);
673+ outQueue_.EnQue(outLocal);
674+ 
675+ LocalTensor<T> result = outQueue_.DeQue<T>();
676+ DataCopyExtParams copyOutParams;
677+ copyOutParams.blockCount = 1;
678+ copyOutParams.blockLen = static_cast<uint32_t>(rowCount * rowBytes);
679+ copyOutParams.srcStride = 0;
680+ copyOutParams.dstStride = 0;
681+ DataCopyPad(yGm_[dstIndex], result, copyOutParams);
682+ outQueue_.FreeTensor(result);
683+}
684+ 
685+template <typename T>
686+__aicore__ inline void Roll<T>::CopyMultiDimLastDimFullRowsBySegments(
687+ int64_t dstIndex, int64_t srcRowIndex, int64_t rowCount)
688+{
689+ const int64_t dimSize = tilingData_->shapes[tilingData_->dimNum - 1];
690+ const int64_t shift = tilingData_->shifts[tilingData_->dimNum - 1];
691+ const int64_t firstElements = shift;
692+ const int64_t secondElements = dimSize - shift;
693+ const int64_t srcBase = srcRowIndex * dimSize;
694+ if (CopyBlockRollByFlatPatch(dstIndex, srcBase, dimSize, firstElements, secondElements, rowCount)) {
695+ return;
696+ }
697+ CopyStridedSourceSegments(dstIndex, srcBase + secondElements, firstElements, dimSize, rowCount);
698+ CopyStridedSourceSegments(dstIndex + firstElements, srcBase, secondElements, dimSize, rowCount);
699+}
700+ 
701+template <typename T>
702+__aicore__ inline void Roll<T>::CopyMultiDimLastDimFullRows(int64_t dstIndex, int64_t rowCount)
703+{
704+ const int64_t dimSize = tilingData_->shapes[tilingData_->dimNum - 1];
705+ const int64_t shift = tilingData_->shifts[tilingData_->dimNum - 1];
706+ const int64_t typeBytes = static_cast<int64_t>(sizeof(T));
707+ const int64_t rowBytes = dimSize * typeBytes;
708+ const int64_t alignedRowBytes = ((rowBytes + ROLL_GM_BLOCK_BYTES - 1) / ROLL_GM_BLOCK_BYTES) * ROLL_GM_BLOCK_BYTES;
709+ const int64_t alignedRowElements = alignedRowBytes / typeBytes;
710+ 
711+ LocalTensor<T> local = inQueue_.AllocTensor<T>();
712+ DataCopyExtParams copyInParams;
713+ copyInParams.blockCount = 1;
714+ copyInParams.blockLen = static_cast<uint32_t>(rowBytes);
715+ copyInParams.srcStride = 0;
716+ copyInParams.dstStride = 0;
717+ DataCopyPadExtParams<T> padParams{false, 0, 0, static_cast<T>(0)};
718+ int64_t copiedRows = 0;
719+ while (copiedRows < rowCount) {
720+ const int64_t dstRow = dstIndex / dimSize + copiedRows;
721+ const int64_t srcRow = ComputeInputIndex(dstRow * dimSize) / dimSize;
722+ int64_t runRows = 1;
723+ while (copiedRows + runRows < rowCount) {
724+ const int64_t nextDstRow = dstRow + runRows;
725+ const int64_t nextSrcRow = ComputeInputIndex(nextDstRow * dimSize) / dimSize;
726+ if (nextSrcRow != srcRow + runRows) {
727+ break;
728+ }
729+ ++runRows;
730+ }
731+ copyInParams.blockCount = static_cast<uint16_t>(runRows);
732+ DataCopyPad(local[copiedRows * alignedRowElements], xGm_[srcRow * dimSize], copyInParams, padParams);
733+ copiedRows += runRows;
734+ }
735+ inQueue_.EnQue(local);
736+ LocalTensor<T> inLocal = inQueue_.DeQue<T>();
737+ 
738+ LocalTensor<T> outLocal = outQueue_.AllocTensor<T>();
739+ CopyRowsRollInUb(outLocal, inLocal, rowCount, dimSize, alignedRowElements, shift);
740+ inQueue_.FreeTensor(inLocal);
741+ outQueue_.EnQue(outLocal);
742+ 
743+ LocalTensor<T> result = outQueue_.DeQue<T>();
744+ DataCopyExtParams copyOutParams;
745+ copyOutParams.blockCount = 1;
746+ copyOutParams.blockLen = static_cast<uint32_t>(rowCount * rowBytes);
747+ copyOutParams.srcStride = 0;
748+ copyOutParams.dstStride = 0;
749+ DataCopyPad(yGm_[dstIndex], result, copyOutParams);
750+ outQueue_.FreeTensor(result);
751+}
752+ 
753+template <typename T>
754+__aicore__ inline void Roll<T>::CopyMultiDimLastDimRollByRows()
755+{
756+ const int64_t dimSize = tilingData_->shapes[tilingData_->dimNum - 1];
757+ const int64_t shift = tilingData_->shifts[tilingData_->dimNum - 1];
758+ const int64_t typeBytes = static_cast<int64_t>(sizeof(T));
759+ const int64_t rowBytes = dimSize * typeBytes;
760+ const int64_t alignedRowBytes = ((rowBytes + ROLL_GM_BLOCK_BYTES - 1) / ROLL_GM_BLOCK_BYTES) * ROLL_GM_BLOCK_BYTES;
761+ const int64_t alignedRowElements = alignedRowBytes / typeBytes;
762+ if (dimSize <= 1 || shift <= 0 || alignedRowElements <= 0 || alignedRowElements > ubElements_) {
763+ CopySegmentedRoll();
764+ return;
765+ }
766+ const int64_t maxInputRows = ubElements_ / alignedRowElements;
767+ const int64_t maxOutputRows = ubElements_ / dimSize;
768+ int64_t maxRows = maxInputRows < maxOutputRows ? maxInputRows : maxOutputRows;
769+ if (maxRows <= 0) {
770+ CopySegmentedRoll();
771+ return;
772+ }
773+ if (maxRows > 4095) {
774+ maxRows = 4095;
775+ }
776+ 
777+ int64_t remain = elementCount_;
778+ int64_t dst = startIndex_;
779+ while (remain > 0 && (dst % dimSize) != 0) {
780+ const int64_t src = ComputeInputIndex(dst);
781+ int64_t current = remain;
782+ const int64_t rowOffset = dst % dimSize;
783+ const int64_t srcOffset = src % dimSize;
784+ const int64_t dstContiguous = dimSize - rowOffset;
785+ const int64_t srcContiguous = dimSize - srcOffset;
786+ if (current > dstContiguous) {
787+ current = dstContiguous;
788+ }
789+ if (current > srcContiguous) {
790+ current = srcContiguous;
791+ }
792+ CopySegment(dst, src, current);
793+ dst += current;
794+ remain -= current;
795+ }
796+ 
797+ int64_t fullRows = remain / dimSize;
798+ while (fullRows > 0) {
799+ int64_t currentRows = fullRows > maxRows ? maxRows : fullRows;
800+ if (tilingData_->dimNum == 2) {
801+ const int64_t outerSize = tilingData_->shapes[0];
802+ const int64_t rowShift = tilingData_->shifts[0];
803+ const int64_t dstRow = dst / dimSize;
804+ int64_t srcRow = dstRow - rowShift;
805+ if (srcRow < 0) {
806+ srcRow += outerSize;
807+ }
808+ const int64_t srcContiguousRows = outerSize - srcRow;
809+ if (currentRows > srcContiguousRows) {
810+ currentRows = srcContiguousRows;
811+ }
812+ if (IsSameType<T, bfloat16_t>::value && tilingData_->activeDimCount > 1 && dimSize == 3 && rowShift != 0 &&
813+ shift != 0 && tilingData_->totalNum * static_cast<int64_t>(sizeof(T)) >= 512 &&
814+ tilingData_->totalNum * static_cast<int64_t>(sizeof(T)) <= 4096) {
815+ CopyMultiDimLastDimFullRowsFromSource(dst, srcRow, currentRows);
816+ } else {
817+ CopyMultiDimLastDimFullRowsBySegments(dst, srcRow, currentRows);
818+ }
819+ } else {
820+ const int64_t dstRow = dst / dimSize;
821+ int64_t srcRow = ComputeSourceRowIndex(dstRow);
822+ currentRows = ComputeContiguousSourceRowRun(dstRow, currentRows);
823+ CopyMultiDimLastDimFullRowsBySegments(dst, srcRow, currentRows);
824+ }
825+ const int64_t copied = currentRows * dimSize;
826+ dst += copied;
827+ remain -= copied;
828+ fullRows -= currentRows;
829+ }
830+ 
831+ while (remain > 0) {
832+ const int64_t src = ComputeInputIndex(dst);
833+ int64_t current = remain;
834+ const int64_t rowOffset = dst % dimSize;
835+ const int64_t srcOffset = src % dimSize;
836+ const int64_t dstContiguous = dimSize - rowOffset;
837+ const int64_t srcContiguous = dimSize - srcOffset;
838+ if (current > dstContiguous) {
839+ current = dstContiguous;
840+ }
841+ if (current > srcContiguous) {
842+ current = srcContiguous;
843+ }
844+ CopySegment(dst, src, current);
845+ dst += current;
846+ remain -= current;
847+ }
848+}
849+ 
850+template <typename T>
851+__aicore__ inline void Roll<T>::CopyMultiDimNonLastBlockPartial(
852+ int64_t lastActiveDim, int64_t& dstIndex, int64_t& remain)
853+{
854+ const int64_t inner = tilingData_->strides[lastActiveDim];
855+ const int64_t dimSize = tilingData_->shapes[lastActiveDim];
856+ const int64_t shift = tilingData_->shifts[lastActiveDim];
857+ const int64_t blockSize = dimSize * inner;
858+ const int64_t firstElements = shift * inner;
859+ const int64_t secondElements = (dimSize - shift) * inner;
860+ 
861+ const int64_t dstBlock = dstIndex / blockSize;
862+ const int64_t srcBlock = ComputeSourceBlockIndex(dstBlock, lastActiveDim);
863+ const int64_t inBlockOffset = dstIndex % blockSize;
864+ int64_t src = 0;
865+ int64_t contiguous = 0;
866+ if (inBlockOffset < firstElements) {
867+ src = srcBlock * blockSize + secondElements + inBlockOffset;
868+ contiguous = firstElements - inBlockOffset;
869+ } else {
870+ src = srcBlock * blockSize + inBlockOffset - firstElements;
871+ contiguous = blockSize - inBlockOffset;
872+ }
873+ int64_t current = remain;
874+ if (current > contiguous) {
875+ current = contiguous;
876+ }
877+ CopySegment(dstIndex, src, current);
878+ dstIndex += current;
879+ remain -= current;
880+}
881+ 
882+template <typename T>
883+__aicore__ inline void Roll<T>::CopyMultiDimNonLastFullBlocks(
884+ int64_t lastActiveDim, int64_t dstIndex, int64_t srcBlockIndex, int64_t blockCount)
885+{
886+ const int64_t inner = tilingData_->strides[lastActiveDim];
887+ const int64_t dimSize = tilingData_->shapes[lastActiveDim];
888+ const int64_t shift = tilingData_->shifts[lastActiveDim];
889+ const int64_t blockSize = dimSize * inner;
890+ const int64_t firstElements = shift * inner;
891+ const int64_t secondElements = (dimSize - shift) * inner;
892+ const int64_t srcIndex = srcBlockIndex * blockSize;
893+ if (CopyBlockRollByFlatPatch(dstIndex, srcIndex, blockSize, firstElements, secondElements, blockCount)) {
894+ return;
895+ }
896+ if (firstElements <= ubElements_ && secondElements <= ubElements_) {
897+ CopyStridedSourceSegments(dstIndex, srcIndex + secondElements, firstElements, blockSize, blockCount);
898+ CopyStridedSourceSegments(dstIndex + firstElements, srcIndex, secondElements, blockSize, blockCount);
899+ return;
900+ }
901+ int64_t copiedBlocks = 0;
902+ while (copiedBlocks < blockCount) {
903+ const int64_t dst = dstIndex + copiedBlocks * blockSize;
904+ const int64_t src = srcIndex + copiedBlocks * blockSize;
905+ CopySegment(dst, src + secondElements, firstElements);
906+ CopySegment(dst + firstElements, src, secondElements);
907+ ++copiedBlocks;
908+ }
909+}
910+ 
911+template <typename T>
912+__aicore__ inline void Roll<T>::CopyMultiDimNonLastRollByBlocks(int64_t lastActiveDim)
913+{
914+ const int64_t inner = tilingData_->strides[lastActiveDim];
915+ const int64_t dimSize = tilingData_->shapes[lastActiveDim];
916+ const int64_t shift = tilingData_->shifts[lastActiveDim];
917+ const int64_t blockSize = dimSize * inner;
918+ if (inner <= 0 || dimSize <= 1 || shift <= 0 || blockSize <= 0) {
919+ CopySegmentedRoll();
920+ return;
921+ }
922+ 
923+ int64_t remain = elementCount_;
924+ int64_t dst = startIndex_;
925+ while (remain > 0 && (dst % blockSize) != 0) {
926+ CopyMultiDimNonLastBlockPartial(lastActiveDim, dst, remain);
927+ }
928+ 
929+ int64_t fullBlocks = remain / blockSize;
930+ while (fullBlocks > 0) {
931+ const int64_t dstBlock = dst / blockSize;
932+ int64_t srcBlock = ComputeSourceBlockIndex(dstBlock, lastActiveDim);
933+ int64_t runBlocks = ComputeContiguousSourceBlockRun(dstBlock, fullBlocks, lastActiveDim);
934+ if (runBlocks > 4095) {
935+ runBlocks = 4095;
936+ }
937+ CopyMultiDimNonLastFullBlocks(lastActiveDim, dst, srcBlock, runBlocks);
938+ const int64_t copied = runBlocks * blockSize;
939+ dst += copied;
940+ remain -= copied;
941+ fullBlocks -= runBlocks;
942+ }
943+ 
944+ while (remain > 0) {
945+ CopyMultiDimNonLastBlockPartial(lastActiveDim, dst, remain);
946+ }
947+}
948+ 
949+template <typename T>
950+__aicore__ inline void Roll<T>::CopyLastDimPartial(int64_t& dstIndex, int64_t& remain)
951+{
952+ const int64_t dimSize = tilingData_->dimSize;
953+ const int64_t shift = tilingData_->activeShift;
954+ const int64_t secondElements = dimSize - shift;
955+ const int64_t rowBase = (dstIndex / dimSize) * dimSize;
956+ const int64_t rowOffset = dstIndex % dimSize;
957+ int64_t src = 0;
958+ int64_t contiguous = 0;
959+ if (rowOffset < shift) {
960+ src = rowBase + secondElements + rowOffset;
961+ contiguous = shift - rowOffset;
962+ } else {
963+ src = rowBase + rowOffset - shift;
964+ contiguous = dimSize - rowOffset;
965+ }
966+ int64_t current = remain;
967+ if (current > contiguous) {
968+ current = contiguous;
969+ }
970+ CopySegment(dstIndex, src, current);
971+ dstIndex += current;
972+ remain -= current;
973+}
974+ 
975+template <typename T>
976+__aicore__ inline void Roll<T>::CopySingleDimRoll()
977+{
978+ CopySegmentedRoll();
979+}
980+ 
981+template <typename T>
982+__aicore__ inline void Roll<T>::CopySingleDimFullBlocks(int64_t dstIndex, int64_t blockCount)
983+{
984+ const int64_t blockSize = tilingData_->dimSize * tilingData_->innerSize;
985+ const int64_t firstElements = tilingData_->activeShift * tilingData_->innerSize;
986+ const int64_t secondElements = blockSize - firstElements;
987+ const int64_t srcIndex = ComputeInputIndex(dstIndex);
988+ if (CopyBlockRollByFlatPatch(dstIndex, srcIndex, blockSize, firstElements, secondElements, blockCount)) {
989+ return;
990+ }
991+ CopyStridedSourceSegments(dstIndex, srcIndex + secondElements, firstElements, blockSize, blockCount);
992+ CopyStridedSourceSegments(dstIndex + firstElements, srcIndex, secondElements, blockSize, blockCount);
993+}
994+ 
995+template <typename T>
996+__aicore__ inline void Roll<T>::CopySingleDimPartial(int64_t& dstIndex, int64_t& remain)
997+{
998+ const int64_t blockSize = tilingData_->dimSize * tilingData_->innerSize;
999+ const int64_t firstElements = tilingData_->activeShift * tilingData_->innerSize;
1000+ const int64_t secondElements = blockSize - firstElements;
1001+ const int64_t blockOffset = dstIndex % blockSize;
1002+ int64_t src = 0;
1003+ int64_t contiguous = 0;
1004+ if (blockOffset < firstElements) {
1005+ src = dstIndex - blockOffset + secondElements + blockOffset;
1006+ contiguous = firstElements - blockOffset;
1007+ } else {
1008+ src = dstIndex - blockOffset + blockOffset - firstElements;
1009+ contiguous = blockSize - blockOffset;
1010+ }
1011+ int64_t current = remain;
1012+ if (current > contiguous) {
1013+ current = contiguous;
1014+ }
1015+ CopySegment(dstIndex, src, current);
1016+ dstIndex += current;
1017+ remain -= current;
1018+}
1019+ 
1020+template <typename T>
1021+__aicore__ inline void Roll<T>::CopySingleDimRollByBlocks()
1022+{
1023+ const int64_t blockSize = tilingData_->dimSize * tilingData_->innerSize;
1024+ if (blockSize <= 0 || tilingData_->activeShift <= 0 || tilingData_->dimSize <= 1) {
1025+ CopySingleDimRoll();
1026+ return;
1027+ }
1028+ 
1029+ int64_t remain = elementCount_;
1030+ int64_t dst = startIndex_;
1031+ while (remain > 0 && (dst % blockSize) != 0) {
1032+ CopySingleDimPartial(dst, remain);
1033+ }
1034+ 
1035+ int64_t fullBlocks = remain / blockSize;
1036+ while (fullBlocks > 0) {
1037+ int64_t currentBlocks = fullBlocks > 4095 ? 4095 : fullBlocks;
1038+ CopySingleDimFullBlocks(dst, currentBlocks);
1039+ const int64_t copied = currentBlocks * blockSize;
1040+ dst += copied;
1041+ remain -= copied;
1042+ fullBlocks -= currentBlocks;
1043+ }
1044+ 
1045+ while (remain > 0) {
1046+ CopySingleDimPartial(dst, remain);
1047+ }
1048+}
1049+ 
1050+template <typename T>
1051+__aicore__ inline void Roll<T>::CopyLastDimRollByRows()
1052+{
1053+ const int64_t dimSize = tilingData_->dimSize;
1054+ const int64_t shift = tilingData_->activeShift;
1055+ const int64_t typeBytes = static_cast<int64_t>(sizeof(T));
1056+ const int64_t rowBytes = dimSize * typeBytes;
1057+ const int64_t alignedRowBytes = ((rowBytes + ROLL_GM_BLOCK_BYTES - 1) / ROLL_GM_BLOCK_BYTES) * ROLL_GM_BLOCK_BYTES;
1058+ const int64_t alignedRowElements = alignedRowBytes / typeBytes;
1059+ if (dimSize <= 1 || shift <= 0 || alignedRowElements <= 0 || alignedRowElements > ubElements_) {
1060+ CopyLastDimRoll();
1061+ return;
1062+ }
1063+ 
1064+ const int64_t maxInputRows = ubElements_ / alignedRowElements;
1065+ const int64_t maxOutputRows = ubElements_ / dimSize;
1066+ int64_t maxRows = maxInputRows < maxOutputRows ? maxInputRows : maxOutputRows;
1067+ if (maxRows <= 0) {
1068+ CopyLastDimRoll();
1069+ return;
1070+ }
1071+ if (maxRows > 4095) {
1072+ maxRows = 4095;
1073+ }
1074+ 
1075+ int64_t remain = elementCount_;
1076+ int64_t dst = startIndex_;
1077+ while (remain > 0 && (dst % dimSize) != 0) {
1078+ CopyLastDimPartial(dst, remain);
1079+ }
1080+ 
1081+ int64_t fullRows = remain / dimSize;
1082+ while (fullRows > 0) {
1083+ int64_t currentRows = fullRows > maxRows ? maxRows : fullRows;
1084+ const int64_t tailElements = dimSize - shift;
1085+ const int64_t largeSegmentBytes = (shift > tailElements ? shift : tailElements) * typeBytes;
1086+ const bool preferSegmentPath = (largeSegmentBytes >= ROLL_STRIDED_SEGMENT_MIN_BYTES) ||
1087+ (sizeof(T) == 1 && dimSize >= 7 && dimSize <= 31);
1088+ if (preferSegmentPath) {
1089+ CopyLastDimFullRowsBySegments(dst, currentRows);
1090+ } else {
1091+ CopyLastDimFullRows(dst, currentRows);
1092+ }
1093+ const int64_t copied = currentRows * dimSize;
1094+ dst += copied;
1095+ remain -= copied;
1096+ fullRows -= currentRows;
1097+ }
1098+ 
1099+ while (remain > 0) {
1100+ CopyLastDimPartial(dst, remain);
1101+ }
1102+}
1103+ 
1104+template <typename T>
1105+__aicore__ inline void Roll<T>::CopySegmentedRoll()
1106+{
1107+ int64_t lastActiveDim = -1;
1108+ for (int64_t dim = 0; dim < tilingData_->dimNum; ++dim) {
1109+ if (tilingData_->shifts[dim] != 0) {
1110+ lastActiveDim = dim;
1111+ }
1112+ }
1113+ if (lastActiveDim < 0) {
1114+ CopyIdentity();
1115+ return;
1116+ }
1117+ 
1118+ int64_t remain = elementCount_;
1119+ int64_t dst = startIndex_;
1120+ while (remain > 0) {
1121+ const int64_t src = ComputeInputIndex(dst);
1122+ int64_t current = remain;
1123+ if (lastActiveDim == tilingData_->dimNum - 1) {
1124+ const int64_t dimSize = tilingData_->shapes[lastActiveDim];
1125+ const int64_t shift = tilingData_->shifts[lastActiveDim];
1126+ const int64_t dstOffset = dst % dimSize;
1127+ const int64_t srcOffset = (dstOffset - shift + dimSize) % dimSize;
1128+ const int64_t dstContiguous = dimSize - dstOffset;
1129+ const int64_t srcContiguous = dimSize - srcOffset;
1130+ if (current > dstContiguous) {
1131+ current = dstContiguous;
1132+ }
1133+ if (current > srcContiguous) {
1134+ current = srcContiguous;
1135+ }
1136+ } else {
1137+ const int64_t contiguous = tilingData_->strides[lastActiveDim];
1138+ const int64_t offset = dst % contiguous;
1139+ const int64_t dstContiguous = contiguous - offset;
1140+ if (current > dstContiguous) {
1141+ current = dstContiguous;
1142+ }
1143+ }
1144+ CopySegment(dst, src, current);
1145+ dst += current;
1146+ remain -= current;
1147+ }
1148+}
1149+ 
1150+template <typename T>
1151+__aicore__ inline void Roll<T>::ProcessScalar()
1152+{
1153+ CopySegmentedRoll();
1154+}
1155+ 
1156+template <typename T>
1157+__aicore__ inline void Roll<T>::Process()
1158+{
1159+ if (elementCount_ <= 0 || tilingData_->totalNum <= 0) {
1160+ return;
1161+ }
1162+ if (tilingData_->activeDimCount == 0) {
1163+ CopyIdentity();
1164+ } else if (tilingData_->dimNum == 1 && tilingData_->strides[0] == 1) {
1165+ CopyFlattenRollBySource();
1166+ } else if (tilingData_->activeDimCount == 1 && tilingData_->activeDim == 0 && tilingData_->innerSize > 0 &&
1167+ (tilingData_->dimNum == 2 || tilingData_->dimSize <= 4 ||
1168+ (IsSameType<T, uint8_t>::value && tilingData_->dimNum > 2 &&
1169+ tilingData_->innerSize % ROLL_GM_BLOCK_BYTES == 0 &&
1170+ tilingData_->totalNum >= 32 * 1024 * 1024)) &&
1171+ !(sizeof(T) == 1 && tilingData_->dimNum == 2 && tilingData_->innerSize == 127)) {
1172+ CopyLeadingDimRollBySource();
1173+ } else if (tilingData_->activeDimCount == 1 && tilingData_->innerSize == 1 &&
1174+ tilingData_->activeDim == tilingData_->dimNum - 1) {
1175+ CopyLastDimRollByRows();
1176+ } else if (tilingData_->activeDimCount == 1 && tilingData_->innerSize > 0) {
1177+ CopySingleDimRollByBlocks();
1178+ } else {
1179+ int64_t lastActiveDim = -1;
1180+ for (int64_t dim = 0; dim < tilingData_->dimNum; ++dim) {
1181+ if (tilingData_->shifts[dim] != 0) {
1182+ lastActiveDim = dim;
1183+ }
1184+ }
1185+ if (lastActiveDim == tilingData_->dimNum - 1) {
1186+ CopyMultiDimLastDimRollByRows();
1187+ return;
1188+ } else if (lastActiveDim >= 0) {
1189+ CopyMultiDimNonLastRollByBlocks(lastActiveDim);
1190+ return;
1191+ }
1192+ CopySegmentedRoll();
1193+ }
1194+}
1195+ 
1196+} // namespace RollKernel
1197+ 
1198+#endif // ROLL_H_
Aexperimental/math/roll/op_kernel/roll_tiling_data.h+44-0
@@ -0,0 +1,44 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file roll_tiling_data.h
13+ * \brief Roll tiling data.
14+ */
15+ 
16+#ifndef ROLL_TILING_DATA_H_
17+#define ROLL_TILING_DATA_H_
18+ 
19+#include <cstdint>
20+ 
21+constexpr uint32_t ROLL_MAX_DIM_NUM = 8;
22+ 
23+struct RollTilingData {
24+ int64_t totalNum = 0;
25+ int64_t dimNum = 0;
26+ int64_t perCoreElements = 0;
27+ int64_t lastCoreElements = 0;
28+ int64_t usedCoreNum = 1;
29+ int64_t ubElements = 0;
30+ int64_t blockFactor = 0;
31+ int64_t ubFactor = 0;
32+ int64_t activeDimCount = 0;
33+ int64_t activeDim = -1;
34+ int64_t outerSize = 0;
35+ int64_t dimSize = 0;
36+ int64_t innerSize = 0;
37+ int64_t activeShift = 0;
38+ int64_t useSafeUbShuffle = 0;
39+ int64_t shapes[ROLL_MAX_DIM_NUM] = {0};
40+ int64_t strides[ROLL_MAX_DIM_NUM] = {0};
41+ int64_t shifts[ROLL_MAX_DIM_NUM] = {0};
42+};
43+ 
44+#endif // ROLL_TILING_DATA_H_
Aexperimental/math/roll/op_kernel/roll_tiling_key.h+31-0
@@ -0,0 +1,31 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file roll_tiling_key.h
13+ * \brief Tiling template argument declaration.
14+ */
15+ 
16+#ifndef ROLL_TILING_KEY_H_
17+#define ROLL_TILING_KEY_H_
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+#define ROLL_TPL_SCH_MODE_0 0
22+#define ROLL_TPL_SCH_MODE_1 1
23+ 
24+ASCENDC_TPL_ARGS_DECL(
25+ Roll,
26+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ROLL_TPL_SCH_MODE_0, ROLL_TPL_SCH_MODE_1));
27+ 
28+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
29+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ROLL_TPL_SCH_MODE_0, ROLL_TPL_SCH_MODE_1)));
30+ 
31+#endif // ROLL_TILING_KEY_H_
Aexperimental/math/roll/tests/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Aexperimental/math/roll/tests/ut/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Aexperimental/math/roll/tests/ut/op_api/test_aclnn_roll.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 <vector>
12+ 
13+#include "gtest/gtest.h"
14+#include "opdev/platform.h"
15+#include "../../../op_api/aclnn_roll.h"
16+#include "op_api_ut_common/array_desc.h"
17+#include "op_api_ut_common/op_api_ut.h"
18+#include "op_api_ut_common/tensor_desc.h"
19+ 
20+using namespace op;
21+ 
22+class test_aclnn_roll : public testing::Test {
23+protected:
24+ static void SetUpTestCase()
25+ {
26+ op::SetPlatformSocVersion(op::SocVersion::ASCEND910B);
27+ }
28+};
29+ 
30+TEST_F(test_aclnn_roll, case_basic_float)
31+{
32+ auto xDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND).Value(std::vector<float>{0, 1, 2, 3, 4, 5});
33+ auto yDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0, 0.0);
34+ auto shifts = IntArrayDesc(std::vector<int64_t>{1});
35+ auto dims = IntArrayDesc(std::vector<int64_t>{1});
36+ 
37+ auto ut = OP_API_UT(aclnnRoll, INPUT(xDesc, shifts, dims), OUTPUT(yDesc));
38+ uint64_t workspaceSize = 0;
39+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
40+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
41+}
42+ 
43+TEST_F(test_aclnn_roll, case_dims_empty)
44+{
45+ auto xDesc = TensorDesc({2, 3}, ACL_FLOAT16, ACL_FORMAT_ND);
46+ auto yDesc = TensorDesc({2, 3}, ACL_FLOAT16, ACL_FORMAT_ND);
47+ auto shifts = IntArrayDesc(std::vector<int64_t>{2});
48+ auto dims = IntArrayDesc(std::vector<int64_t>{});
49+ 
50+ auto ut = OP_API_UT(aclnnRoll, INPUT(xDesc, shifts, dims), OUTPUT(yDesc));
51+ uint64_t workspaceSize = 0;
52+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
53+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
54+}
55+ 
56+TEST_F(test_aclnn_roll, case_invalid_dtype)
57+{
58+ auto xDesc = TensorDesc({4}, ACL_DOUBLE, ACL_FORMAT_ND);
59+ auto yDesc = TensorDesc({4}, ACL_DOUBLE, ACL_FORMAT_ND);
60+ auto shifts = IntArrayDesc(std::vector<int64_t>{1});
61+ auto dims = IntArrayDesc(std::vector<int64_t>{0});
62+ 
63+ auto ut = OP_API_UT(aclnnRoll, INPUT(xDesc, shifts, dims), OUTPUT(yDesc));
64+ uint64_t workspaceSize = 0;
65+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
66+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
67+}
68+ 
69+TEST_F(test_aclnn_roll, case_invalid_dims_range)
70+{
71+ auto xDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
72+ auto yDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
73+ auto shifts = IntArrayDesc(std::vector<int64_t>{1});
74+ auto dims = IntArrayDesc(std::vector<int64_t>{2});
75+ 
76+ auto ut = OP_API_UT(aclnnRoll, INPUT(xDesc, shifts, dims), OUTPUT(yDesc));
77+ uint64_t workspaceSize = 0;
78+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
79+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
80+}
Aexperimental/math/roll/tests/ut/op_host/test_roll_infershape.cpp+35-0
@@ -0,0 +1,35 @@
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 <gtest/gtest.h>
12+ 
13+#include "infershape_case_executor.h"
14+ 
15+class RollInfershape : public testing::Test {};
16+ 
17+TEST_F(RollInfershape, roll_infershape_basic)
18+{
19+ gert::InfershapeContextPara infershapeContextPara(
20+ "Roll",
21+ {
22+ {{{2, 3, 4}, {2, 3, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
23+ },
24+ {
25+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
26+ },
27+ {
28+ gert::InfershapeContextPara::OpAttr("shifts", Ops::Math::AnyValue::CreateFrom<std::vector<int64_t>>({1})),
29+ gert::InfershapeContextPara::OpAttr("dims", Ops::Math::AnyValue::CreateFrom<std::vector<int64_t>>({1})),
30+ });
31+ std::vector<std::vector<int64_t>> expectOutputShape = {
32+ {2, 3, 4},
33+ };
34+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
35+}
Aexperimental/math/roll/tests/ut/op_host/test_roll_tiling.cpp+84-0
@@ -0,0 +1,84 @@
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 <gtest/gtest.h>
12+ 
13+#include "tiling_case_executor.h"
14+#include "../../../op_kernel/roll_tiling_data.h"
15+ 
16+namespace {
17+struct RollCompileInfoForTest {
18+ int32_t coreNum = 64;
19+};
20+} // namespace
21+ 
22+class RollTiling : public testing::Test {};
23+ 
24+TEST_F(RollTiling, basic_last_dim_roll)
25+{
26+ RollCompileInfoForTest compileInfo = {64};
27+ gert::TilingContextPara tilingContextPara(
28+ "Roll",
29+ {
30+ {{{2, 3}, {2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
31+ },
32+ {
33+ {{{2, 3}, {2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
34+ },
35+ {
36+ gert::TilingContextPara::OpAttr("shifts", Ops::Math::AnyValue::CreateFrom<std::vector<int64_t>>({1})),
37+ gert::TilingContextPara::OpAttr("dims", Ops::Math::AnyValue::CreateFrom<std::vector<int64_t>>({1})),
38+ },
39+ &compileInfo);
40+ 
41+ TilingInfo tilingInfo;
42+ ASSERT_TRUE(ExecuteTiling(tilingContextPara, tilingInfo));
43+ ASSERT_EQ(tilingInfo.tilingKey, 0);
44+ ASSERT_GT(tilingInfo.blockNum, 0U);
45+ ASSERT_EQ(tilingInfo.workspaceSizes.size(), 1U);
46+ ASSERT_EQ(tilingInfo.workspaceSizes[0], 0);
47+ ASSERT_GE(tilingInfo.tilingDataSize, sizeof(RollTilingData));
48+ 
49+ auto* data = reinterpret_cast<const RollTilingData*>(tilingInfo.tilingData.get());
50+ EXPECT_EQ(data->totalNum, 6);
51+ EXPECT_EQ(data->dimNum, 2);
52+ EXPECT_EQ(data->activeDimCount, 1);
53+ EXPECT_EQ(data->activeDim, 1);
54+ EXPECT_EQ(data->dimSize, 3);
55+ EXPECT_EQ(data->innerSize, 1);
56+ EXPECT_EQ(data->activeShift, 1);
57+}
58+ 
59+TEST_F(RollTiling, flatten_roll_when_dims_empty)
60+{
61+ RollCompileInfoForTest compileInfo = {64};
62+ gert::TilingContextPara tilingContextPara(
63+ "Roll",
64+ {
65+ {{{2, 3, 4}, {2, 3, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
66+ },
67+ {
68+ {{{2, 3, 4}, {2, 3, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
69+ },
70+ {
71+ gert::TilingContextPara::OpAttr("shifts", Ops::Math::AnyValue::CreateFrom<std::vector<int64_t>>({5})),
72+ gert::TilingContextPara::OpAttr("dims", Ops::Math::AnyValue::CreateFrom<std::vector<int64_t>>({})),
73+ },
74+ &compileInfo);
75+ 
76+ TilingInfo tilingInfo;
77+ ASSERT_TRUE(ExecuteTiling(tilingContextPara, tilingInfo));
78+ auto* data = reinterpret_cast<const RollTilingData*>(tilingInfo.tilingData.get());
79+ EXPECT_EQ(data->dimNum, 1);
80+ EXPECT_EQ(data->totalNum, 24);
81+ EXPECT_EQ(data->shapes[0], 24);
82+ EXPECT_EQ(data->strides[0], 1);
83+ EXPECT_EQ(data->shifts[0], 5);
84+}
Aexperimental/math/roll/tests/ut/op_kernel/CMakeLists.txt+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+if (UT_TEST_ALL OR OP_KERNEL_UT)
12+ set(roll_tiling_files
13+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/roll_tiling.cpp
14+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/roll_infershape.cpp
15+ )
16+ AddOpTestCase(roll "ascend910b" "-DDTYPE_X=float" "${roll_tiling_files}")
17+endif()
Aexperimental/math/roll/tests/ut/op_kernel/roll_tiling.h+44-0
@@ -0,0 +1,44 @@
1+/*!
2+ * \file roll_tiling.h
3+ * \brief Roll tiling data for kernel UT
4+ */
5+ 
6+#ifndef ROLL_TILING_UT_H_
7+#define ROLL_TILING_UT_H_
8+ 
9+#include <cstdint>
10+#include <cstring>
11+ 
12+#include "../../../op_kernel/roll_tiling_data.h"
13+#include "tikicpulib.h"
14+#include "kernel_operator.h"
15+#include "kernel_tiling/kernel_tiling.h"
16+#include "graph/c_types.h"
17+#include "ascendc/host_api/tiling/template_argument.h"
18+ 
19+#ifndef __aicore__
20+#define __aicore__
21+#endif
22+ 
23+#ifndef __gm__
24+#define __gm__
25+#endif
26+ 
27+#ifndef __ubuf__
28+#define __ubuf__
29+#endif
30+ 
31+inline void InitTilingData(uint8_t* tiling, RollTilingData* constData)
32+{
33+ memcpy(constData, tiling, sizeof(RollTilingData));
34+}
35+ 
36+#define GET_TILING_DATA_WITH_STRUCT(tilingStruct, tilingData, tilingArg) \
37+ tilingStruct tilingData; \
38+ InitTilingData(tilingArg, &tilingData)
39+ 
40+#define GET_TILING_DATA(tilingData, tilingArg) \
41+ RollTilingData tilingData; \
42+ InitTilingData(tilingArg, &tilingData)
43+ 
44+#endif // ROLL_TILING_UT_H_
Aexperimental/math/roll/tests/ut/op_kernel/test_roll.cpp+65-0
@@ -0,0 +1,65 @@
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 <cstring>
12+#include <vector>
13+ 
14+#include "gtest/gtest.h"
15+#include "tikicpulib.h"
16+ 
17+#include "roll_tiling.h"
18+#include "../../../op_kernel/roll.cpp"
19+ 
20+class RollKernelTest : public testing::Test {};
21+ 
22+TEST_F(RollKernelTest, kernel_launch_smoke)
23+{
24+ constexpr size_t size = 6;
25+ constexpr uint32_t numBlocks = 1;
26+ 
27+ std::vector<float> xHost = {0, 1, 2, 3, 4, 5};
28+ std::vector<float> yHost(size, 0);
29+ 
30+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(size * sizeof(float));
31+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(size * sizeof(float));
32+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(32);
33+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(RollTilingData));
34+ 
35+ memcpy(x, xHost.data(), size * sizeof(float));
36+ memcpy(y, yHost.data(), size * sizeof(float));
37+ 
38+ auto* tilingData = reinterpret_cast<RollTilingData*>(tiling);
39+ memset(tilingData, 0, sizeof(RollTilingData));
40+ tilingData->totalNum = size;
41+ tilingData->dimNum = 1;
42+ tilingData->perCoreElements = size;
43+ tilingData->lastCoreElements = size;
44+ tilingData->usedCoreNum = 1;
45+ tilingData->ubElements = size;
46+ tilingData->blockFactor = size;
47+ tilingData->ubFactor = size;
48+ tilingData->activeDimCount = 1;
49+ tilingData->activeDim = 0;
50+ tilingData->dimSize = size;
51+ tilingData->innerSize = 1;
52+ tilingData->activeShift = 1;
53+ tilingData->shapes[0] = size;
54+ tilingData->strides[0] = 1;
55+ tilingData->shifts[0] = 1;
56+ 
57+ ICPU_SET_TILING_KEY(0);
58+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
59+ ICPU_RUN_KF((roll<0>), numBlocks, x, y, workspace, tiling);
60+ 
61+ AscendC::GmFree(x);
62+ AscendC::GmFree(y);
63+ AscendC::GmFree(workspace);
64+ AscendC::GmFree(tiling);
65+}