Star173
257
代码
代码
Issues64
Pull Requests88
Actions
项目讨论
Wiki
分析
Star173
257
已合并
feat(rot): 新增 arch35 平台 aclblasSrot 算子接口 #233
陈思创建于 6月30日
feat(rot): 新增 arch35 平台 aclblasSrot 算子接口 #233
已合并
陈思创建于 6月30日
从chensi79/ops-blasaclblasSrot合入到cann/ops-blasmaster
共 17 个文件变更+1429-36
Mblas/asum/README.md+12-8
@@ -115,6 +115,13 @@ private:
115 bool deviceSet_ = false;115 bool deviceSet_ = false;
116};116};
117 117 
118+struct AclMemDeleter {
119+ void operator()(void* p) const { aclrtFree(p); }
120+};
121+struct BlasHandleDeleter {
122+ void operator()(aclblasHandle_t h) const { aclblasDestroy(h); }
123+};
124+ 
118int aclblasSasumTest(AclContext& ctx)125int aclblasSasumTest(AclContext& ctx)
119{126{
120 aclrtStream stream = ctx.Stream();127 aclrtStream stream = ctx.Stream();
@@ -124,9 +131,9 @@ int aclblasSasumTest(AclContext& ctx)
124 auto blasRet = aclblasCreate(&rawHandle);131 auto blasRet = aclblasCreate(&rawHandle);
125 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);132 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);
126 return blasRet);133 return blasRet);
127- std::unique_ptr<void, aclblasStatus_t (*)(void*)> handlePtr(rawHandle, aclblasDestroy);134+ std::unique_ptr<std::remove_pointer<aclblasHandle_t>::type, BlasHandleDeleter> handlePtr(rawHandle);
128 135 
129- blasRet = aclblasSetStream(static_cast<aclblasHandle_t>(handlePtr.get()), stream);136+ blasRet = aclblasSetStream(handlePtr.get(), stream);
130 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);137 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);
131 return blasRet);138 return blasRet);
132 139 
@@ -140,21 +147,18 @@ int aclblasSasumTest(AclContext& ctx)
140 void* rawMemX = nullptr;147 void* rawMemX = nullptr;
141 auto aclRet = aclrtMalloc(&rawMemX, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);148 auto aclRet = aclrtMalloc(&rawMemX, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);
142 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);149 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);
143- std::unique_ptr<void, aclError (*)(void*)> xDevicePtr(rawMemX, aclrtFree);150+ std::unique_ptr<float, AclMemDeleter> xDevicePtr(static_cast<float*>(rawMemX));
144 151 
145 void* rawMemResult = nullptr;152 void* rawMemResult = nullptr;
146 aclRet = aclrtMalloc(&rawMemResult, sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST);153 aclRet = aclrtMalloc(&rawMemResult, sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST);
147 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for result failed. ERROR: %d\n", aclRet); return aclRet);154 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for result failed. ERROR: %d\n", aclRet); return aclRet);
148- std::unique_ptr<void, aclError (*)(void*)> resultDevicePtr(rawMemResult, aclrtFree);155+ std::unique_ptr<float, AclMemDeleter> resultDevicePtr(static_cast<float*>(rawMemResult));
149 156 
150 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);157 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);
151 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);158 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);
152 159 
153 // 4. 调用 aclblasSasum160 // 4. 调用 aclblasSasum
154- blasRet = aclblasSasum(161+ blasRet = aclblasSasum(handlePtr.get(), n, xDevicePtr.get(), incx, resultDevicePtr.get());
155- static_cast<aclblasHandle_t>(handlePtr.get()), n,
156- static_cast<const float*>(xDevicePtr.get()), incx,
157- static_cast<float*>(resultDevicePtr.get()));
158 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSasum failed. ERROR: %d\n", blasRet);162 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSasum failed. ERROR: %d\n", blasRet);
159 return blasRet);163 return blasRet);
160 164 
Mblas/nrm2/README.md+12-8
@@ -144,6 +144,13 @@ private:
144 bool deviceSet_ = false;144 bool deviceSet_ = false;
145};145};
146 146 
147+struct AclMemDeleter {
148+ void operator()(void* p) const { aclrtFree(p); }
149+};
150+struct BlasHandleDeleter {
151+ void operator()(aclblasHandle_t h) const { aclblasDestroy(h); }
152+};
153+ 
147int aclblasSnrm2Test(AclContext& ctx)154int aclblasSnrm2Test(AclContext& ctx)
148{155{
149 aclrtStream stream = ctx.Stream();156 aclrtStream stream = ctx.Stream();
@@ -153,9 +160,9 @@ int aclblasSnrm2Test(AclContext& ctx)
153 auto blasRet = aclblasCreate(&rawHandle);160 auto blasRet = aclblasCreate(&rawHandle);
154 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);161 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);
155 return blasRet);162 return blasRet);
156- std::unique_ptr<void, aclblasStatus_t (*)(void*)> handlePtr(rawHandle, aclblasDestroy);163+ std::unique_ptr<std::remove_pointer<aclblasHandle_t>::type, BlasHandleDeleter> handlePtr(rawHandle);
157 164 
158- blasRet = aclblasSetStream(static_cast<aclblasHandle_t>(handlePtr.get()), stream);165+ blasRet = aclblasSetStream(handlePtr.get(), stream);
159 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);166 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);
160 return blasRet);167 return blasRet);
161 168 
@@ -169,21 +176,18 @@ int aclblasSnrm2Test(AclContext& ctx)
169 void* rawMemX = nullptr;176 void* rawMemX = nullptr;
170 auto aclRet = aclrtMalloc(&rawMemX, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);177 auto aclRet = aclrtMalloc(&rawMemX, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);
171 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);178 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);
172- std::unique_ptr<void, aclError (*)(void*)> xDevicePtr(rawMemX, aclrtFree);179+ std::unique_ptr<float, AclMemDeleter> xDevicePtr(static_cast<float*>(rawMemX));
173 180 
174 void* rawMemResult = nullptr;181 void* rawMemResult = nullptr;
175 aclRet = aclrtMalloc(&rawMemResult, sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST);182 aclRet = aclrtMalloc(&rawMemResult, sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST);
176 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for result failed. ERROR: %d\n", aclRet); return aclRet);183 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for result failed. ERROR: %d\n", aclRet); return aclRet);
177- std::unique_ptr<void, aclError (*)(void*)> resultDevicePtr(rawMemResult, aclrtFree);184+ std::unique_ptr<float, AclMemDeleter> resultDevicePtr(static_cast<float*>(rawMemResult));
178 185 
179 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);186 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);
180 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);187 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);
181 188 
182 // 4. 调用 aclblasSnrm2189 // 4. 调用 aclblasSnrm2
183- blasRet = aclblasSnrm2(190+ blasRet = aclblasSnrm2(handlePtr.get(), n, xDevicePtr.get(), incx, resultDevicePtr.get());
184- static_cast<aclblasHandle_t>(handlePtr.get()), n,
185- static_cast<const float*>(xDevicePtr.get()), incx,
186- static_cast<float*>(resultDevicePtr.get()));
187 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSnrm2 failed. ERROR: %d\n", blasRet);191 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSnrm2 failed. ERROR: %d\n", blasRet);
188 return blasRet);192 return blasRet);
189 193 
Mblas/rot/README.md+210-3
@@ -2,7 +2,7 @@
2 2 
3## 算子概述3## 算子概述
4 4 
5-向量旋转算子，实现对两个向量的平面旋转（Givens 旋转），常用于 QR 分解、求解线性方程组和特征值计算等数值算法中。5+向量旋转算子，实现对两个等长向量的平面旋转（Givens 旋转），常用于 QR 分解、求解线性方程组和特征值计算等数值算法中。rot 算子族支持实数（FP32）与复数（FP32 complex）两种数据类型的 Givens 旋转。
6 6 
7数学表达式：7数学表达式：
8 8 
@@ -15,7 +15,8 @@ y[i] = c * y[i] - s * x[i] (使用原始 x[i])
15 15 
16| 接口名 | 功能简述 |16| 接口名 | 功能简述 |
17|--------|---------|17|--------|---------|
18-| aclblasCsrot | 复数向量平面旋转 |18+| aclblasCsrot | 复数（FP32 complex）向量平面旋转 |
19+| aclblasSrot | 单精度（FP32）实数向量平面旋转 |
19 20 
20## 算子执行接口21## 算子执行接口
21 22 
@@ -48,4 +49,210 @@ aclblasStatus_t aclblasCsrot(aclblasHandle_t handle, const int64_t n, uint8_t* x
48 49 
49#### 约束说明50#### 约束说明
50 51 
51-- n >= 052+- n >= 0
53+ 
54+### aclblasSrot
55+ 
56+#### 产品支持情况
57+ 
58+- Ascend 950PR / Ascend 950DT：支持
59+- Atlas A3 训练系列产品 / Atlas A3 推理系列产品：不支持
60+- Atlas A2 训练系列产品 / Atlas A2 推理系列产品：不支持
61+ 
62+#### 函数原型
63+ 
64+```cpp
65+aclblasStatus_t aclblasSrot(
66+ aclblasHandle_t handle, int n, float* x, int incx, float* y, int incy, const float* c, const float* s);
67+```
68+ 
69+#### 参数说明
70+ 
71+| 参数名 | 输入/输出 | 参数类型 | 说明 |
72+|--------|----------|---------|------|
73+| handle | 输入 | aclblasHandle_t | ops-blas 库上下文句柄，携带 stream，Host 内存 |
74+| n | 输入 | int | 参与旋转的元素个数，Host 内存 |
75+| x | 输入/输出 | float*（FP32） | FP32 向量，in-place 修改，按 incx 步长访问，Device 内存 |
76+| incx | 输入 | int | 向量 x 中连续元素之间的步长，可为正/负/零，Host 内存 |
77+| y | 输入/输出 | float*（FP32） | FP32 向量，in-place 修改，按 incy 步长访问，Device 内存 |
78+| incy | 输入 | int | 向量 y 中连续元素之间的步长，可为正/负/零，Host 内存 |
79+| c | 输入 | const float* | 指向 Givens 旋转参数 cos(θ) 的标量指针，Host 内存或 Device 内存，不可为 nullptr；算子运行时通过 aclrtPointerGetAttributes 自动判定其内存位置 |
80+| s | 输入 | const float* | 指向 Givens 旋转参数 sin(θ) 的标量指针，Host 内存或 Device 内存，不可为 nullptr；算子运行时通过 aclrtPointerGetAttributes 自动判定其内存位置 |
81+ 
82+#### 约束说明
83+ 
84+**Ascend 950PR / Ascend 950DT（arch35）：**
85+ 
86+- handle 不能为 nullptr，否则返回 ACLBLAS_STATUS_HANDLE_IS_NULLPTR
87+- n 为整数；n <= 0 时为 no-op（直接返回 ACLBLAS_STATUS_SUCCESS，不修改 x、y，对齐参考 BLAS 的 `IF (N.LE.0) RETURN` 语义）
88+- n > 0 时 x、y 不能为 nullptr，否则返回 ACLBLAS_STATUS_INVALID_VALUE
89+- incx、incy 为整数，可为正、负、零，均不拦截：
90+ - incx == 1 且 incy == 1 时走连续路径
91+ - 其余 stride 组合（含正非 1、负、零及其混合）走 stride 路径；零 stride 时按参考实现语义对同一元素反复旋转，负 stride 时从向量尾端起算沿负方向步进
92+- 当 (n-1) * stride 的乘积超出 int32 表示范围时，返回 ACLBLAS_STATUS_INVALID_VALUE，避免 kernel 侧地址偏移溢出
93+- c、s 可为 host 内存指针或 device 内存指针，指向 cos/sin 标量值，由算子运行时通过 aclrtPointerGetAttributes 自动判定内存位置；c 或 s 为 nullptr 时返回 ACLBLAS_STATUS_INVALID_VALUE
94+- c、s 指向的标量值可取任意 FP32 值（包括 c == 1 且 s == 0 的单位旋转，不做短路，正常执行旋转公式）
95+- 精度标准：FP32 单标杆（MARE ≤ 10·2⁻¹³，MERE ≤ 2⁻¹³）
96+ 
97+**Atlas A2 / Atlas A3 系列产品：**
98+ 
99+- 不支持
100+ 
101+#### 调用示例
102+ 
103+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../docs/zh/develop/compile_and_run_example.md)。
104+ 
105+```cpp
106+#include <cstdio>
107+#include <memory>
108+#include <vector>
109+ 
110+#include "acl/acl.h"
111+#include "cann_ops_blas.h"
112+ 
113+#define CHECK_RET(cond, return_expr) \
114+ do { \
115+ if (!(cond)) { \
116+ return_expr; \
117+ } \
118+ } while (0)
119+ 
120+#define LOG_PRINT(message, ...) \
121+ do { \
122+ printf(message, ##__VA_ARGS__); \
123+ } while (0)
124+ 
125+class AclContext {
126+public:
127+ explicit AclContext(int32_t deviceId) : deviceId_(deviceId) {}
128+ 
129+ ~AclContext()
130+ {
131+ if (stream_ != nullptr) {
132+ aclrtDestroyStream(stream_);
133+ stream_ = nullptr;
134+ }
135+ if (deviceSet_) {
136+ aclrtResetDevice(deviceId_);
137+ deviceSet_ = false;
138+ }
139+ if (aclInited_) {
140+ aclFinalize();
141+ aclInited_ = false;
142+ }
143+ }
144+ 
145+ int Init()
146+ {
147+ auto ret = aclInit(nullptr);
148+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
149+ aclInited_ = true;
150+ 
151+ ret = aclrtSetDevice(deviceId_);
152+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
153+ deviceSet_ = true;
154+ 
155+ ret = aclrtCreateStream(&stream_);
156+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
157+ return ACL_SUCCESS;
158+ }
159+ 
160+ aclrtStream Stream() const { return stream_; }
161+ 
162+private:
163+ int32_t deviceId_;
164+ aclrtStream stream_ = nullptr;
165+ bool aclInited_ = false;
166+ bool deviceSet_ = false;
167+};
168+ 
169+struct AclMemDeleter {
170+ void operator()(void* p) const { aclrtFree(p); }
171+};
172+struct BlasHandleDeleter {
173+ void operator()(aclblasHandle_t h) const { aclblasDestroy(h); }
174+};
175+ 
176+int aclblasSrotTest(AclContext& ctx)
177+{
178+ aclrtStream stream = ctx.Stream();
179+ 
180+ // 1. 创建 ops-blas 句柄
181+ aclblasHandle_t rawHandle = nullptr;
182+ auto blasRet = aclblasCreate(&rawHandle);
183+ CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);
184+ return blasRet);
185+ std::unique_ptr<std::remove_pointer<aclblasHandle_t>::type, BlasHandleDeleter> handlePtr(rawHandle);
186+ 
187+ blasRet = aclblasSetStream(handlePtr.get(), stream);
188+ CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);
189+ return blasRet);
190+ 
191+ // 2. 准备 Host 数据
192+ int n = 4;
193+ int incx = 1;
194+ int incy = 1;
195+ float c = 0.8f;
196+ float s = 0.6f;
197+ std::vector<float> xHostData = {1.0f, 2.0f, 3.0f, 4.0f};
198+ std::vector<float> yHostData = {5.0f, 6.0f, 7.0f, 8.0f};
199+ size_t xBytes = n * sizeof(float);
200+ size_t yBytes = n * sizeof(float);
201+ 
202+ // 3. 申请 Device 内存并拷贝数据
203+ void* rawXMem = nullptr;
204+ auto aclRet = aclrtMalloc(&rawXMem, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);
205+ CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);
206+ std::unique_ptr<float, AclMemDeleter> xDevicePtr(static_cast<float*>(rawXMem));
207+ 
208+ void* rawYMem = nullptr;
209+ aclRet = aclrtMalloc(&rawYMem, yBytes, ACL_MEM_MALLOC_HUGE_FIRST);
210+ CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for y failed. ERROR: %d\n", aclRet); return aclRet);
211+ std::unique_ptr<float, AclMemDeleter> yDevicePtr(static_cast<float*>(rawYMem));
212+ 
213+ aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);
214+ CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);
215+ 
216+ aclRet = aclrtMemcpy(yDevicePtr.get(), yBytes, yHostData.data(), yBytes, ACL_MEMCPY_HOST_TO_DEVICE);
217+ CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for y failed. ERROR: %d\n", aclRet); return aclRet);
218+ 
219+ // 4. 调用 aclblasSrot（c、s 为 Host 指针，原地旋转）
220+ blasRet = aclblasSrot(handlePtr.get(), n, xDevicePtr.get(), incx, yDevicePtr.get(), incy, &c, &s);
221+ CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSrot failed. ERROR: %d\n", blasRet);
222+ return blasRet);
223+ 
224+ // 5. 同步等待任务执行结束
225+ aclRet = aclrtSynchronizeStream(stream);
226+ CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", aclRet); return aclRet);
227+ 
228+ // 6. 将结果从 Device 拷贝回 Host 并打印
229+ std::vector<float> resultX(n, 0);
230+ aclRet = aclrtMemcpy(resultX.data(), xBytes, xDevicePtr.get(), xBytes, ACL_MEMCPY_DEVICE_TO_HOST);
231+ CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("copy x from device to host failed. ERROR: %d\n", aclRet);
232+ return aclRet);
233+ for (int i = 0; i < n; i++) {
234+ LOG_PRINT("x[%d] is: %f\n", i, resultX[i]);
235+ }
236+ 
237+ std::vector<float> resultY(n, 0);
238+ aclRet = aclrtMemcpy(resultY.data(), yBytes, yDevicePtr.get(), yBytes, ACL_MEMCPY_DEVICE_TO_HOST);
239+ CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("copy y from device to host failed. ERROR: %d\n", aclRet);
240+ return aclRet);
241+ for (int i = 0; i < n; i++) {
242+ LOG_PRINT("y[%d] is: %f\n", i, resultY[i]);
243+ }
244+ 
245+ return ACL_SUCCESS;
246+}
247+ 
248+int main()
249+{
250+ AclContext ctx(0);
251+ auto ret = ctx.Init();
252+ CHECK_RET(ret == ACL_SUCCESS, return ret);
253+ 
254+ ret = aclblasSrotTest(ctx);
255+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclblasSrotTest failed. ERROR: %d\n", ret); return ret);
256+ return 0;
257+}
258+```
Ablas/rot/arch35/srot_host.cpp+256-0
@@ -0,0 +1,256 @@
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 srot_host.cpp
13+ * \brief Host-side API for BLAS Level-1 Givens plane rotation: aclblasSrot.
14+ * Arch35 (ascend950) implementation.
15+ *
16+ * Applies the Givens rotation to vectors x and y in place:
17+ * x[i] = c*x[i] + s*y[i]
18+ * y[i] = c*y[i] - s*x[i]
19+ *
20+ * Two compute paths selected by contiguity of the strides:
21+ * - incx==1 && incy==1 : contiguous path, SIMD membase kernel
22+ * - otherwise : strided path, SIMT kernel (grid-stride loop)
23+ *
24+ * netlib srot.f boundary alignment: n<=0 short-circuits to SUCCESS;
25+ * zero / negative strides and the identity rotation (c==1 && s==0) are
26+ * intentionally not intercepted, matching the reference behavior.
27+ */
28+ 
29+#include <cstdint>
30+#include <climits>
31+#include <algorithm>
32+#include "acl/acl.h"
33+#include "log/log.h"
34+#include "cann_ops_blas.h"
35+#include "srot_kernel.h"
36+#include "common/helper/aclblas_handle_internal.h"
37+#include "common/helper/host_utils.h"
38+#include "common/helper/kernel_constant.h"
39+ 
40+static aclblasStatus_t ValidateSrotParams(const float* x, const float* y, const float* c, const float* s)
41+{
42+ if (x == nullptr || y == nullptr) {
43+ OP_LOGE("aclblasSrot", "x/y must not be nullptr");
44+ return ACLBLAS_STATUS_INVALID_VALUE;
45+ }
46+ if (c == nullptr || s == nullptr) {
47+ OP_LOGE("aclblasSrot", "c/s must not be nullptr");
48+ return ACLBLAS_STATUS_INVALID_VALUE;
49+ }
50+ return ACLBLAS_STATUS_SUCCESS;
51+}
52+ 
53+// ==========================================================================
54+// Pointer location query: determine whether a c/s scalar pointer lives on the
55+// host or on the device (NPU HBM). Used to select the host-dereference path vs
56+// the GM-forwarding path. Mirrors the aclrtPointerGetAttributes usage in rotg
57+// and scalex. On query failure the call is rejected (consistent with scalex).
58+// ==========================================================================
59+static aclblasStatus_t SrotCheckPtrLocation(const void* ptr, bool* isDevice)
60+{
61+ aclrtPtrAttributes ptrAttr{};
62+ aclError aclRet = aclrtPointerGetAttributes(ptr, &ptrAttr);
63+ if (aclRet != ACL_SUCCESS) {
64+ OP_LOGE("aclblasSrot", "aclrtPointerGetAttributes failed, ret=%d", aclRet);
65+ return ACLBLAS_STATUS_INVALID_VALUE;
66+ }
67+ *isDevice = (ptrAttr.location.type == ACL_MEM_LOCATION_TYPE_DEVICE);
68+ return ACLBLAS_STATUS_SUCCESS;
69+}
70+ 
71+// ==========================================================================
72+// Tiling computation for the contiguous path (SIMD membase).
73+// The contiguous kernel uses a TQue three-stage pipeline with five UB tiles:
74+// inQueueX / inQueueY (VECIN, MTE2 input)
75+// outNewX / outNewY (VECOUT, Vector output)
76+// workBuf (VECCALC, pure-Vector scratch for round(c*y))
77+// Each tile is tileSize elements, so UB must hold bufferCount * tileSize floats.
78+// ==========================================================================
79+static void CalSrotTilingContiguous(uint32_t n, uint32_t numBlocks, SrotTilingData& tiling)
80+{
81+ // Even split (no alignment rounding): perCoreN = n/numBlocks base elements per
82+ // core, remainder = n%numBlocks front cores each get one extra element. Unlike
83+ // the old align-down scheme, no core starves when n/numBlocks < ELEMENTS_PER_BLOCK:
84+ // the load stays balanced down to n < numBlocks. The contiguous kernel copies with
85+ // pure DataCopyPad, whose GM side requires only 1-byte alignment, so the
86+ // non-block-aligned per-core start offset and element count are both safe.
87+ tiling.perCoreN = n / numBlocks;
88+ tiling.remainder = n % numBlocks;
89+ 
90+ // Five UB tiles: inQueueX + inQueueY + outNewX + outNewY + workBuf.
91+ // tileSize is the per-copy UB cap; kept block-aligned for UB buffer allocation.
92+ constexpr uint32_t alignUnit = SROT_ELEMENTS_PER_BLOCK;
93+ constexpr uint32_t bufferCount = 5;
94+ uint32_t maxElements = UB_SIZE / (bufferCount * sizeof(float)); // UB_SIZE from kernel_constant.h
95+ tiling.tileSize = (maxElements / alignUnit) * alignUnit;
96+}
97+ 
98+// ==========================================================================
99+// Tiling computation for the strided path (SIMT).
100+// The kernel handles negative strides internally (negX/negY flags + (n-1-i)*absInc),
101+// so it consumes only incx/incy from the tiling. The (n-1)*stride product is still
102+// computed here in int64_t and range-checked before narrowing: if it exceeds the
103+// int32_t range, the kernel-side int32_t offset would overflow (risk R5), so the
104+// call is rejected with an explicit error instead of launching a broken kernel.
105+// The narrowed last offsets are returned via outLastX/outLastY for host-side
106+// diagnostics (OP_LOGD) only; they are NOT consumed by the kernel.
107+//
108+// nthreads: SIMT threads per block. The grid-stride loop only needs as many
109+// threads as the per-core element count justifies, so this is sized from
110+// ceilDiv(n,numBlocks) — rounded UP to SIMT_MIN_THREAD_NUM (the hardware's thread
111+// scheduling granularity) and capped at SIMT_MAX_THREAD_NUM — instead of always
112+// launching SIMT_MAX_THREAD_NUM threads. A small n (e.g. n=10, 1 block) thus
113+// launches 128 threads rather than 2048, eliminating the wasted-thread startup
114+// cost while large n still saturates the full 2048 per block.
115+// ==========================================================================
116+static aclblasStatus_t CalSrotTilingStrided(uint32_t n, int incx, int incy, uint32_t numBlocks,
117+ SrotTilingData& tiling, int32_t* outLastX, int32_t* outLastY)
118+{
119+ tiling.incx = incx;
120+ tiling.incy = incy;
121+ tiling.nthreads = std::min(
122+ CeilAlign<uint32_t>(CeilDiv<uint32_t>(n, numBlocks), SIMT_MIN_THREAD_NUM),
123+ SIMT_MAX_THREAD_NUM);
124+ int64_t lastX = static_cast<int64_t>(n - 1) * static_cast<int64_t>(incx);
125+ int64_t lastY = static_cast<int64_t>(n - 1) * static_cast<int64_t>(incy);
126+ if (lastX > INT32_MAX || lastX < INT32_MIN || lastY > INT32_MAX || lastY < INT32_MIN) {
127+ OP_LOGE("aclblasSrot", "strided offset overflow: n=%u incx=%d incy=%d lastX=%lld lastY=%lld",
128+ static_cast<unsigned>(n), incx, incy, static_cast<long long>(lastX), static_cast<long long>(lastY));
129+ return ACLBLAS_STATUS_INVALID_VALUE;
130+ }
131+ *outLastX = static_cast<int32_t>(lastX);
132+ *outLastY = static_cast<int32_t>(lastY);
133+ return ACLBLAS_STATUS_SUCCESS;
134+}
135+ 
136+// ==========================================================================
137+// Resolve the c/s source: query each pointer's location, fill the corresponding
138+// TilingData fields, and produce the float* pointers to forward to the kernel.
139+// host pointer -> dereference once, store the scalar, forward nullptr
140+// device pointer-> store 0.0f placeholder, forward the pointer verbatim
141+// The two pointers are independent (c on host + s on device is legal). Device
142+// pointers are forwarded as-is with no dereference, no D2H copy, no stream sync.
143+// ==========================================================================
144+static aclblasStatus_t PrepareSrotCsSource(const float* c, const float* s, SrotTilingData& tiling, float** cPtr,
145+ float** sPtr)
146+{
147+ bool cIsDevice = false;
148+ bool sIsDevice = false;
149+ aclblasStatus_t cLocSt = SrotCheckPtrLocation(c, &cIsDevice);
150+ if (cLocSt != ACLBLAS_STATUS_SUCCESS) {
151+ return cLocSt;
152+ }
153+ aclblasStatus_t sLocSt = SrotCheckPtrLocation(s, &sIsDevice);
154+ if (sLocSt != ACLBLAS_STATUS_SUCCESS) {
155+ return sLocSt;
156+ }
157+ 
158+ tiling.cIsDevice = cIsDevice ? 1u : 0u;
159+ tiling.sIsDevice = sIsDevice ? 1u : 0u;
160+ tiling.cosValue = cIsDevice ? 0.0f : (*c);
161+ tiling.sinValue = sIsDevice ? 0.0f : (*s);
162+ // Device path forwards the pointer verbatim; host path forwards nullptr (kernel
163+ // reads the tiling scalar instead).
164+ *cPtr = cIsDevice ? const_cast<float*>(c) : nullptr;
165+ *sPtr = sIsDevice ? const_cast<float*>(s) : nullptr;
166+ return ACLBLAS_STATUS_SUCCESS;
167+}
168+ 
169+// ==========================================================================
170+// Launch — tiling computation + asynchronous kernel launch
171+// ==========================================================================
172+static aclblasStatus_t LaunchSrotKernel(aclblasHandle_t handle, int n, float* x, int incx, float* y, int incy,
173+ const float* c, const float* s)
174+{
175+ auto* h = reinterpret_cast<_aclblas_handle*>(handle);
176+ aclrtStream useStream = h->stream;
177+ 
178+ uint32_t aivCoreNum = GetAivCoreCount();
179+ if (aivCoreNum == 0) {
180+ OP_LOGE("aclblasSrot", "GetAivCoreCount failed");
181+ return ACLBLAS_STATUS_INTERNAL_ERROR;
182+ }
183+ 
184+ uint32_t totalN = static_cast<uint32_t>(n);
185+ bool contiguous = (incx == 1 && incy == 1);
186+ 
187+ SrotTilingData tiling{};
188+ tiling.totalN = totalN;
189+ 
190+ // Resolve c/s source (host scalar vs device GM pointer); fills tiling.cosValue/
191+ // sinValue/cIsDevice/sIsDevice and yields the pointers to forward.
192+ float* cPtr = nullptr;
193+ float* sPtr = nullptr;
194+ aclblasStatus_t csSt = PrepareSrotCsSource(c, s, tiling, &cPtr, &sPtr);
195+ if (csSt != ACLBLAS_STATUS_SUCCESS) {
196+ return csSt;
197+ }
198+ 
199+ // lastX/lastY: strided-path tail offset ((n-1)*stride), host-side diagnostic only.
200+ // Contiguous path leaves them at 0 (no stride offset); the kernel never reads them.
201+ int32_t lastX = 0;
202+ int32_t lastY = 0;
203+ uint32_t numBlocks;
204+ if (contiguous) {
205+ tiling.tilingKey = 0;
206+ numBlocks = (totalN < aivCoreNum) ? totalN : aivCoreNum;
207+ CalSrotTilingContiguous(totalN, numBlocks, tiling);
208+ } else {
209+ tiling.tilingKey = 1;
210+ uint32_t blocksByThreads = CeilDiv<uint32_t>(totalN, SIMT_MIN_THREAD_NUM);
211+ numBlocks = (blocksByThreads < aivCoreNum) ? blocksByThreads : aivCoreNum;
212+ if (numBlocks == 0) {
213+ numBlocks = 1;
214+ }
215+ aclblasStatus_t stridedSt = CalSrotTilingStrided(totalN, incx, incy, numBlocks, tiling, &lastX, &lastY);
216+ if (stridedSt != ACLBLAS_STATUS_SUCCESS) {
217+ return stridedSt;
218+ }
219+ }
220+ 
221+ OP_LOGD("aclblasSrot",
222+ "tiling: key=%u totalN=%u incx=%d incy=%d lastX=%d lastY=%d c=%f s=%f cIsDevice=%u sIsDevice=%u "
223+ "perCoreN=%u remainder=%u tileSize=%u numBlocks=%u nthreads=%u",
224+ tiling.tilingKey, tiling.totalN, incx, incy, lastX, lastY, tiling.cosValue,
225+ tiling.sinValue, tiling.cIsDevice, tiling.sIsDevice, tiling.perCoreN, tiling.remainder, tiling.tileSize,
226+ numBlocks, tiling.nthreads);
227+ OP_LOGI("aclblasSrot", "launching kernel: key=%u blocks=%u cIsDevice=%u sIsDevice=%u", tiling.tilingKey,
228+ numBlocks, tiling.cIsDevice, tiling.sIsDevice);
229+ 
230+ // In-place operation: no workspace needed. Pass nullptr; no aclrtMalloc here.
231+ srot_kernel_do(x, y, cPtr, sPtr, numBlocks, tiling, useStream);
232+ 
233+ return ACLBLAS_STATUS_SUCCESS;
234+}
235+ 
236+// ==========================================================================
237+// Public API entry — dispatch only (validate then launch)
238+// ==========================================================================
239+extern "C" aclblasStatus_t aclblasSrot(aclblasHandle_t handle, int n, float* x, int incx, float* y, int incy,
240+ const float* c, const float* s)
241+{
242+ if (handle == nullptr) {
243+ OP_LOGE("aclblasSrot", "handle is nullptr");
244+ return ACLBLAS_STATUS_HANDLE_IS_NULLPTR;
245+ }
246+ if (n <= 0) {
247+ return ACLBLAS_STATUS_SUCCESS;
248+ }
249+ 
250+ aclblasStatus_t st = ValidateSrotParams(x, y, c, s);
251+ if (st != ACLBLAS_STATUS_SUCCESS) {
252+ return st;
253+ }
254+ 
255+ return LaunchSrotKernel(handle, n, x, incx, y, incy, c, s);
256+}
Ablas/rot/arch35/srot_kernel.cpp+335-0
@@ -0,0 +1,335 @@
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+// Two separate kernel entries selected by tilingKey in srot_kernel_do:
13+// tilingKey == 0: srot_aiv_kernel — contiguous path (incx==1 && incy==1) -> SIMD membase
14+// DataCopy/DataCopyPad GM<->UB, paired vector compute in UB
15+// tilingKey == 1: srot_simt_kernel — strided path (any other incx/incy) -> SIMT
16+// multi-core grid-stride loop on GM (numBlocks cores in parallel)
17+// In-place overwrite avoidance:
18+// - contiguous: original x/y are held in the VECIN input queues during the
19+// whole Compute stage (only freed at the end of Compute), so both new_x and
20+// new_y can still read the originals. The new values are written to separate
21+// VECOUT queues (outNewX/outNewY) and copied back to the original GM in
22+// CopyOut. No staging buffer is overwritten before its consumer finishes.
23+// - strided: each thread loads xi/yi into registers, computes both new values,
24+// then writes back; the originals survive in registers until both stores issue.
25+ 
26+#include <cstdint>
27+#include "acl/acl.h"
28+#include "kernel_operator.h"
29+#include "simt_api/asc_simt.h"
30+#include "common/helper/kernel_constant.h"
31+#include "srot_kernel.h"
32+ 
33+using namespace AscendC;
34+ 
35+// ==========================================================================
36+// Contiguous path: SIMD membase operator class
37+// ==========================================================================
38+class SrotAIV {
39+public:
40+ __aicore__ inline SrotAIV() {}
41+ __aicore__ inline void Init(__gm__ float* x, __gm__ float* y, __gm__ float* cPtr, __gm__ float* sPtr,
42+ const SrotTilingData& tiling, TPipe* pipe);
43+ __aicore__ inline void Process();
44+ 
45+private:
46+ __aicore__ inline void CopyIn(uint32_t curOffset, uint32_t dataCount);
47+ __aicore__ inline void Compute(uint32_t dataCount);
48+ __aicore__ inline void CopyOut(uint32_t curOffset, uint32_t dataCount);
49+ 
50+ GlobalTensor<float> xGM_;
51+ GlobalTensor<float> yGM_;
52+ TQue<TPosition::VECIN, 1> inQueueX_; // MTE2 input for x (EnQue signals MTE2 done)
53+ TQue<TPosition::VECIN, 1> inQueueY_; // MTE2 input for y (EnQue signals MTE2 done)
54+ TQue<TPosition::VECOUT, 1> outNewX_; // Vector output for new x (EnQue signals Vector done)
55+ TQue<TPosition::VECOUT, 1> outNewY_; // Vector output for new y (EnQue signals Vector done)
56+ TBuf<TPosition::VECCALC> workBuf_; // pure-Vector scratch (round(c*y)), no MTE traffic
57+ SrotTilingData tiling_;
58+ float cosValue_;
59+ float sinValue_;
60+ uint32_t blockIdx_;
61+ uint32_t myOffset_;
62+ uint32_t myCount_;
63+};
64+ 
65+__aicore__ inline void SrotAIV::Init(__gm__ float* x, __gm__ float* y, __gm__ float* cPtr, __gm__ float* sPtr,
66+ const SrotTilingData& tiling, TPipe* pipe)
67+{
68+ blockIdx_ = GetBlockIdx();
69+ tiling_ = tiling;
70+ // c/s source: device pointer -> read scalar from GM; host pointer -> use tiling scalar.
71+ // Device path avoids any host dereference / D2H sync; host path keeps scalar broadcast.
72+ cosValue_ = tiling.cIsDevice ? *cPtr : tiling.cosValue;
73+ sinValue_ = tiling.sIsDevice ? *sPtr : tiling.sinValue;
74+ 
75+ // Even split: base perCoreN elements, front `remainder` cores each get one extra.
76+ // Offset = blockIdx*perCoreN + (extra ones already issued by earlier cores), so the
77+ // per-core ranges butt exactly with no gap/overlap regardless of alignment.
78+ uint32_t perCore = tiling_.perCoreN;
79+ myOffset_ = blockIdx_ * perCore;
80+ if (blockIdx_ < tiling_.remainder) {
81+ myOffset_ += blockIdx_;
82+ perCore += 1;
83+ } else {
84+ myOffset_ += tiling_.remainder;
85+ }
86+ myCount_ = perCore;
87+ 
88+ xGM_.SetGlobalBuffer(x, tiling_.totalN);
89+ yGM_.SetGlobalBuffer(y, tiling_.totalN);
90+ 
91+ uint32_t bufSize = tiling_.tileSize * sizeof(float);
92+ pipe->InitBuffer(inQueueX_, 1, bufSize);
93+ pipe->InitBuffer(inQueueY_, 1, bufSize);
94+ pipe->InitBuffer(outNewX_, 1, bufSize);
95+ pipe->InitBuffer(outNewY_, 1, bufSize);
96+ pipe->InitBuffer(workBuf_, bufSize);
97+}
98+ 
99+// ==========================================================================
100+// CopyIn: MTE2 GM -> UB. EnQue at the end is the MTE2 -> Vector sync point.
101+// ==========================================================================
102+__aicore__ inline void SrotAIV::CopyIn(uint32_t curOffset, uint32_t dataCount)
103+{
104+ // Pure DataCopyPad: the GM side needs only 1-byte alignment, so the per-core start
105+ // offset (non-block-aligned under even split) and any element count are both fine.
106+ // isPad=false lets the framework auto-pad the tail in UB; the Vector stage consumes
107+ // only `dataCount` elements, so the dummy tail is never read.
108+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(dataCount * sizeof(float)), 0, 0, 0};
109+ DataCopyPadExtParams<float> padParams{false, 0, 0, 0.0f};
110+ 
111+ LocalTensor<float> xLocal = inQueueX_.AllocTensor<float>();
112+ LocalTensor<float> yLocal = inQueueY_.AllocTensor<float>();
113+ DataCopyPad(xLocal, xGM_[curOffset], copyParams, padParams);
114+ DataCopyPad(yLocal, yGM_[curOffset], copyParams, padParams);
115+ // EnQue marks both tiles ready and signals Vector that MTE2 has finished.
116+ inQueueX_.EnQue(xLocal);
117+ inQueueY_.EnQue(yLocal);
118+}
119+ 
120+// ==========================================================================
121+// Compute: Vector stage. DeQue at the start waits for MTE2; EnQue at the end
122+// signals MTE3. Original x/y are held for the whole stage (freed only at the
123+// bottom) so both new_x and new_y can read the originals -> in-place safe.
124+// ==========================================================================
125+__aicore__ inline void SrotAIV::Compute(uint32_t dataCount)
126+{
127+ // DeQue blocks until the MTE2 copies from CopyIn have landed in UB.
128+ LocalTensor<float> xLocal = inQueueX_.DeQue<float>();
129+ LocalTensor<float> yLocal = inQueueY_.DeQue<float>();
130+ LocalTensor<float> newX = outNewX_.AllocTensor<float>();
131+ LocalTensor<float> newY = outNewY_.AllocTensor<float>();
132+ LocalTensor<float> work = workBuf_.Get<float>();
133+ int32_t count = static_cast<int32_t>(dataCount);
134+ 
135+ // Givens rotation, two separate legs using the ORIGINAL x and y:
136+ // new_x = c*x + s*y
137+ // new_y = c*y - s*x
138+ // Both legs read xLocal/yLocal, which stay live until the end of Compute.
139+ 
140+ // ----- new_x = c*x + s*y (Axpy FMA: tmp = c*x; tmp += s*y) -----
141+ Muls(newX, xLocal, cosValue_, count); // newX = round(c*x)
142+ Axpy(newX, yLocal, sinValue_, count); // newX = round(c*x) + s*y (FMA, single rounding on the addend)
143+ 
144+ // ----- new_y = c*y - s*x (per-step rounding, Muls + Muls + Add) -----
145+ // Use per-step rounding (Muls + Muls + Add) instead of Axpy's FMA: when c==s
146+ // (45° rotation) and x≈y the result cancels toward 0, and FMA's unrounded c*y
147+ // leg diverges from the OpenBLAS reference (which rounds each step), inflating
148+ // the relative error past the FP32 gate. workBuf holds the rounded c*y; the two
149+ // rounded legs are then added.
150+ Muls(newY, xLocal, -sinValue_, count); // newY = round(-s*x)
151+ Muls(work, yLocal, cosValue_, count); // work = round(c*y)
152+ Add(newY, newY, work, count); // newY = round(c*y - s*x)
153+ 
154+ // EnQue signals MTE3 that Vector is done with these outputs.
155+ outNewX_.EnQue(newX);
156+ outNewY_.EnQue(newY);
157+ // Original x/y released only after both new values are computed -> in-place safe.
158+ inQueueX_.FreeTensor(xLocal);
159+ inQueueY_.FreeTensor(yLocal);
160+}
161+ 
162+// ==========================================================================
163+// CopyOut: MTE3 UB -> GM. DeQue at the start waits for Vector; in-place write
164+// back to the original xGM_/yGM_.
165+// ==========================================================================
166+__aicore__ inline void SrotAIV::CopyOut(uint32_t curOffset, uint32_t dataCount)
167+{
168+ // Pure DataCopyPad (UB->GM): writes exactly dataCount floats to GM and auto-strips
169+ // the in-UB dummy tail. GM side needs only 1-byte alignment.
170+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(dataCount * sizeof(float)), 0, 0, 0};
171+ 
172+ // DeQue blocks until the Vector compute on these outputs has finished.
173+ LocalTensor<float> newX = outNewX_.DeQue<float>();
174+ LocalTensor<float> newY = outNewY_.DeQue<float>();
175+ DataCopyPad(xGM_[curOffset], newX, copyParams);
176+ DataCopyPad(yGM_[curOffset], newY, copyParams);
177+ outNewX_.FreeTensor(newX);
178+ outNewY_.FreeTensor(newY);
179+}
180+ 
181+__aicore__ inline void SrotAIV::Process()
182+{
183+ if (myCount_ == 0) {
184+ return;
185+ }
186+ 
187+ uint32_t tileLoop = myCount_ / tiling_.tileSize;
188+ uint32_t tileTail = myCount_ % tiling_.tileSize;
189+ uint32_t curOffset = myOffset_;
190+ 
191+ for (uint32_t i = 0; i < tileLoop; i++) {
192+ CopyIn(curOffset, tiling_.tileSize);
193+ Compute(tiling_.tileSize);
194+ CopyOut(curOffset, tiling_.tileSize);
195+ curOffset += tiling_.tileSize;
196+ }
197+ if (tileTail > 0) {
198+ CopyIn(curOffset, tileTail);
199+ Compute(tileTail);
200+ CopyOut(curOffset, tileTail);
201+ }
202+}
203+ 
204+// ==========================================================================
205+// Strided path: SIMT compute function
206+//
207+// Multi-core grid-stride loop: the host launches the SIMT kernel on numBlocks
208+// cores (>= 1). Every thread in every block walks the N elements with a
209+// grid-stride of (block_num * blockDim.x), starting at
210+// (blockIdx.x * blockDim.x + threadIdx.x). blockDim.x = tiling.nthreads, sized by
211+// the host to the per-core element count (rounded up to SIMT_MIN_THREAD_NUM,
212+// capped at SIMT_MAX_THREAD_NUM) so a small n no longer launches the full
213+// SIMT_MAX_THREAD_NUM only to retire most threads. LAUNCH_BOUND stays at the max
214+// for register allocation; the runtime launch may use any blockDim.x <= it.
215+// ==========================================================================
216+ 
217+// Zero-stride serial fallback (R2): netlib srot.f with incx==0 holds IX constant across the
218+// whole loop, so each of the N iterations reuses the element just written by the previous
219+// iteration (serial accumulation). A multi-thread grid-stride loop would read the same
220+// original value concurrently and race on the single address, diverging from netlib.
221+// When zeroIncX or zeroIncY is set, only block0.thread0 runs the N iterations in strict
222+// netlib order (IX/IY recomputed exactly as the reference: start anchor + i*inc, inc
223+// possibly 0), matching the serial accumulation bit-for-bit. All other threads/blocks
224+// return immediately so the reused address is touched by exactly one program-order stream.
225+// The non-zero stride stays on the parallel grid-stride path; the zero-stride side is
226+// simply re-read/re-written each iteration.
227+__simt_vf__ __aicore__ LAUNCH_BOUND(SIMT_MAX_THREAD_NUM)
228+inline void srot_simt_compute(uint32_t n, uint32_t absIncX, uint32_t absIncY, int32_t negX, int32_t negY,
229+ int32_t zeroIncX, int32_t zeroIncY, float c, float s, __gm__ float* xGm,
230+ __gm__ float* yGm)
231+{
232+ // Zero stride (either side) forces single-thread serial execution to reproduce netlib's
233+ // serial accumulation over the reused address. absInc is set to 1 for the zero side so the
234+ // address math stays valid (the index is forced to element 0 by jx/jy below).
235+ // Multi-core guard: with numBlocks cores in flight, the race would span blocks too, so the
236+ // guard must retire every thread except block0.thread0.
237+ if ((zeroIncX != 0) || (zeroIncY != 0)) {
238+ if ((blockIdx.x != 0) || (threadIdx.x != 0)) {
239+ return;
240+ }
241+ for (uint32_t i = 0; i < n; i++) {
242+ uint32_t jx = (zeroIncX != 0) ? 0U : ((negX != 0) ? (n - 1 - i) : i);
243+ uint32_t jy = (zeroIncY != 0) ? 0U : ((negY != 0) ? (n - 1 - i) : i);
244+ uint32_t xIdx = jx * absIncX;
245+ uint32_t yIdx = jy * absIncY;
246+ float xi = xGm[xIdx];
247+ float yi = yGm[yIdx];
248+ float newX = c * xi + s * yi;
249+ float newY = c * yi - s * xi;
250+ xGm[xIdx] = newX;
251+ yGm[yIdx] = newY;
252+ }
253+ return;
254+ }
255+ 
256+ // Grid-stride loop: thread (blockIdx.x*blockDim.x + threadIdx.x) starts at its global
257+ // index and strides by the total number of threads in the grid (block_num*blockDim.x).
258+ // Each thread handles disjoint elements across the whole N range; with numBlocks cores
259+ // this is genuine multi-core parallelism (vs the old single-block threadIdx.x loop).
260+ uint32_t tidGlobal = threadIdx.x + blockIdx.x * blockDim.x;
261+ uint32_t gridStride = block_num * blockDim.x;
262+ for (uint32_t i = tidGlobal; i < n; i += gridStride) {
263+ uint32_t jx = (negX != 0) ? (n - 1 - i) : i;
264+ uint32_t jy = (negY != 0) ? (n - 1 - i) : i;
265+ uint32_t xIdx = jx * absIncX;
266+ uint32_t yIdx = jy * absIncY;
267+ float xi = xGm[xIdx];
268+ float yi = yGm[yIdx];
269+ float newX = c * xi + s * yi;
270+ float newY = c * yi - s * xi;
271+ xGm[xIdx] = newX;
272+ yGm[yIdx] = newY;
273+ }
274+}
275+ 
276+// ==========================================================================
277+// Kernel entries — two separate binaries so SIMD (TPipe/DataCopy) and SIMT
278+// (asc_vf_call) code never coexist in one kernel image.
279+// ==========================================================================
280+ 
281+// Contiguous path: SIMD membase (incx==1 && incy==1)
282+extern "C" __global__ __aicore__ void srot_aiv_kernel(__gm__ float* x, __gm__ float* y, __gm__ float* cPtr,
283+ __gm__ float* sPtr, SrotTilingData tiling)
284+{
285+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
286+ TPipe pipe;
287+ SrotAIV op;
288+ op.Init(x, y, cPtr, sPtr, tiling, &pipe);
289+ op.Process();
290+}
291+ 
292+// Strided path: SIMT (any other incx/incy combination)
293+extern "C" __global__ __aicore__ void srot_simt_kernel(__gm__ float* x, __gm__ float* y, __gm__ float* cPtr,
294+ __gm__ float* sPtr, SrotTilingData tiling)
295+{
296+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
297+ 
298+ __gm__ float* xGm = x;
299+ __gm__ float* yGm = y;
300+ uint32_t n = tiling.totalN;
301+ int32_t incx = tiling.incx;
302+ int32_t incy = tiling.incy;
303+ int32_t zeroIncX = (incx == 0) ? 1 : 0;
304+ int32_t zeroIncY = (incy == 0) ? 1 : 0;
305+ // |incx|/|incy| computed via int64 to avoid signed-negation UB when incx == INT32_MIN.
306+ uint32_t absIncX = static_cast<uint32_t>(incx < 0 ? -static_cast<int64_t>(incx) : static_cast<int64_t>(incx));
307+ uint32_t absIncY = static_cast<uint32_t>(incy < 0 ? -static_cast<int64_t>(incy) : static_cast<int64_t>(incy));
308+ if (absIncX == 0) {
309+ absIncX = 1; // inc==0 reuses element 0; a single stride covers it
310+ }
311+ if (absIncY == 0) {
312+ absIncY = 1;
313+ }
314+ int32_t negX = (incx < 0) ? 1 : 0;
315+ int32_t negY = (incy < 0) ? 1 : 0;
316+ // c/s source: device pointer -> read scalar from GM; host pointer -> use tiling scalar.
317+ float c = tiling.cIsDevice ? *cPtr : tiling.cosValue;
318+ float s = tiling.sIsDevice ? *sPtr : tiling.sinValue;
319+ asc_vf_call<srot_simt_compute>(dim3{tiling.nthreads, 1, 1}, n, absIncX, absIncY, negX, negY, zeroIncX,
320+ zeroIncY, c, s, xGm, yGm);
321+}
322+ 
323+// ==========================================================================
324+// Host-side kernel launcher
325+// ==========================================================================
326+void srot_kernel_do(float* x, float* y, float* cPtr, float* sPtr, uint32_t numBlocks,
327+ const SrotTilingData& tiling, void* stream)
328+{
329+ auto aclStream = static_cast<aclrtStream>(stream);
330+ if (tiling.tilingKey == 0) {
331+ srot_aiv_kernel<<<numBlocks, nullptr, aclStream>>>(x, y, cPtr, sPtr, tiling);
332+ } else {
333+ srot_simt_kernel<<<numBlocks, nullptr, aclStream>>>(x, y, cPtr, sPtr, tiling);
334+ }
335+}
Ablas/rot/arch35/srot_kernel.h+18-0
@@ -0,0 +1,18 @@
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+#pragma once
12+ 
13+#include <cstdint>
14+#include "srot_tiling_data.h"
15+ 
16+// srot_kernel_do: host-side kernel launcher.
17+void srot_kernel_do(float* x, float* y, float* cPtr, float* sPtr,
18+ uint32_t numBlocks, const SrotTilingData& tiling, void* stream);
Ablas/rot/arch35/srot_tiling_data.h+51-0
@@ -0,0 +1,51 @@
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+#pragma once
12+ 
13+#include <cstdint>
14+ 
15+// SrotTilingData carries the runtime parameters for both code paths of aclblasSrot:
16+// tilingKey == 0: contiguous path (SIMD membase), uses perCoreN/remainder/tileSize
17+// tilingKey == 1: strided path (SIMT), uses incx/incy
18+// int32_t is required for incx/incy to carry negative strides.
19+//
20+// c/s pointer location: c and s may each independently be a host pointer or a
21+// device pointer (queried on the host via aclrtPointerGetAttributes). When a
22+// pointer is on the host, the host dereferences it once and forwards the scalar
23+// via cosValue/sinValue (scalar broadcast in the kernel). When a pointer is on
24+// the device, the host forwards the device pointer verbatim as a GM_ADDR to the
25+// kernel (no dereference, no D2H copy, no stream sync) and fills the scalar
26+// field with a 0.0f placeholder; the kernel then reads the value from GM.
27+// cIsDevice/sIsDevice select which source each path uses. The two are independent
28+// (c on host + s on device is legal), unlike rotg which requires all-or-nothing.
29+struct SrotTilingData {
30+ // ===== common fields (both paths) =====
31+ uint32_t tilingKey; // 0 = contiguous (SIMD), 1 = strided (SIMT)
32+ uint32_t totalN; // number of elements to rotate
33+ float cosValue; // Givens cosine c (host value; 0.0f placeholder when cIsDevice)
34+ float sinValue; // Givens sine s (host value; 0.0f placeholder when sIsDevice)
35+ uint32_t cIsDevice; // 1 = c is a device pointer (kernel reads cPtr from GM), 0 = host scalar
36+ uint32_t sIsDevice; // 1 = s is a device pointer (kernel reads sPtr from GM), 0 = host scalar
37+ 
38+ // ===== contiguous path fields (valid when tilingKey == 0) =====
39+ uint32_t perCoreN; // elements per core, aligned down to ELEMENTS_PER_BLOCK
40+ uint32_t remainder; // extra elements assigned to the last core
41+ uint32_t tileSize; // UB tile size in elements, aligned to ELEMENTS_PER_BLOCK
42+ 
43+ // ===== strided path fields (valid when tilingKey == 1) =====
44+ int32_t incx; // x stride (positive / negative / zero)
45+ int32_t incy; // y stride (positive / negative / zero)
46+ uint32_t nthreads; // SIMT threads/block: ceilDiv(n,numBlocks) rounded up to
47+ // SIMT_MIN_THREAD_NUM, capped at SIMT_MAX_THREAD_NUM
48+};
49+ 
50+// elements per 32-byte block for FP32 (= 8)
51+constexpr uint32_t SROT_ELEMENTS_PER_BLOCK = 32 / sizeof(float);
Mblas/scal/README.md+11-6
@@ -126,6 +126,13 @@ private:
126 bool deviceSet_ = false;126 bool deviceSet_ = false;
127};127};
128 128 
129+struct AclMemDeleter {
130+ void operator()(void* p) const { aclrtFree(p); }
131+};
132+struct BlasHandleDeleter {
133+ void operator()(aclblasHandle_t h) const { aclblasDestroy(h); }
134+};
135+ 
129int aclblasSscalTest(AclContext& ctx)136int aclblasSscalTest(AclContext& ctx)
130{137{
131 aclrtStream stream = ctx.Stream();138 aclrtStream stream = ctx.Stream();
@@ -135,9 +142,9 @@ int aclblasSscalTest(AclContext& ctx)
135 auto blasRet = aclblasCreate(&rawHandle);142 auto blasRet = aclblasCreate(&rawHandle);
136 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);143 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);
137 return blasRet);144 return blasRet);
138- std::unique_ptr<void, aclblasStatus_t (*)(void*)> handlePtr(rawHandle, aclblasDestroy);145+ std::unique_ptr<std::remove_pointer<aclblasHandle_t>::type, BlasHandleDeleter> handlePtr(rawHandle);
139 146 
140- blasRet = aclblasSetStream(static_cast<aclblasHandle_t>(handlePtr.get()), stream);147+ blasRet = aclblasSetStream(handlePtr.get(), stream);
141 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);148 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);
142 return blasRet);149 return blasRet);
143 150 
@@ -152,15 +159,13 @@ int aclblasSscalTest(AclContext& ctx)
152 void* rawMem = nullptr;159 void* rawMem = nullptr;
153 auto aclRet = aclrtMalloc(&rawMem, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);160 auto aclRet = aclrtMalloc(&rawMem, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);
154 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);161 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);
155- std::unique_ptr<void, aclError (*)(void*)> xDevicePtr(rawMem, aclrtFree);162+ std::unique_ptr<float, AclMemDeleter> xDevicePtr(static_cast<float*>(rawMem));
156 163 
157 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);164 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);
158 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);165 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);
159 166 
160 // 4. 调用 aclblasSscal（alpha 为 Host 指针，原地缩放）167 // 4. 调用 aclblasSscal（alpha 为 Host 指针，原地缩放）
161- blasRet = aclblasSscal(168+ blasRet = aclblasSscal(handlePtr.get(), n, &alpha, xDevicePtr.get(), incx);
162- static_cast<aclblasHandle_t>(handlePtr.get()), n, &alpha,
163- static_cast<float*>(xDevicePtr.get()), incx);
164 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSscal failed. ERROR: %d\n", blasRet);169 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSscal failed. ERROR: %d\n", blasRet);
165 return blasRet);170 return blasRet);
166 171 
Mblas/scalex/README.md+11-6
@@ -122,6 +122,13 @@ private:
122 bool deviceSet_ = false;122 bool deviceSet_ = false;
123};123};
124 124 
125+struct AclMemDeleter {
126+ void operator()(void* p) const { aclrtFree(p); }
127+};
128+struct BlasHandleDeleter {
129+ void operator()(aclblasHandle_t h) const { aclblasDestroy(h); }
130+};
131+ 
125int aclblasScalexTest(AclContext& ctx)132int aclblasScalexTest(AclContext& ctx)
126{133{
127 aclrtStream stream = ctx.Stream();134 aclrtStream stream = ctx.Stream();
@@ -131,9 +138,9 @@ int aclblasScalexTest(AclContext& ctx)
131 auto blasRet = aclblasCreate(&rawHandle);138 auto blasRet = aclblasCreate(&rawHandle);
132 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);139 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);
133 return blasRet);140 return blasRet);
134- std::unique_ptr<void, aclblasStatus_t (*)(void*)> handlePtr(rawHandle, aclblasDestroy);141+ std::unique_ptr<std::remove_pointer<aclblasHandle_t>::type, BlasHandleDeleter> handlePtr(rawHandle);
135 142 
136- blasRet = aclblasSetStream(static_cast<aclblasHandle_t>(handlePtr.get()), stream);143+ blasRet = aclblasSetStream(handlePtr.get(), stream);
137 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);144 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);
138 return blasRet);145 return blasRet);
139 146 
@@ -151,15 +158,13 @@ int aclblasScalexTest(AclContext& ctx)
151 void* rawMem = nullptr;158 void* rawMem = nullptr;
152 auto aclRet = aclrtMalloc(&rawMem, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);159 auto aclRet = aclrtMalloc(&rawMem, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);
153 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);160 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);
154- std::unique_ptr<void, aclError (*)(void*)> xDevicePtr(rawMem, aclrtFree);161+ std::unique_ptr<void, AclMemDeleter> xDevicePtr(rawMem);
155 162 
156 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);163 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);
157 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);164 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);
158 165 
159 // 4. 调用 aclblasScalex（原地缩放）166 // 4. 调用 aclblasScalex（原地缩放）
160- blasRet = aclblasScalex(167+ blasRet = aclblasScalex(handlePtr.get(), n, &alpha, alphaType, xDevicePtr.get(), xType, incx, executionType);
161- static_cast<aclblasHandle_t>(handlePtr.get()), n, &alpha, alphaType,
162- xDevicePtr.get(), xType, incx, executionType);
163 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasScalex failed. ERROR: %d\n", blasRet);168 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasScalex failed. ERROR: %d\n", blasRet);
164 return blasRet);169 return blasRet);
165 170 
Mdocs/zh/develop/compile_and_run_example.md+11-5
@@ -131,6 +131,13 @@
131 bool deviceSet_ = false;131 bool deviceSet_ = false;
132 };132 };
133 133
134+ struct AclMemDeleter {
135+ void operator()(void* p) const { aclrtFree(p); }
136+ };
137+ struct BlasHandleDeleter {
138+ void operator()(aclblasHandle_t h) const { aclblasDestroy(h); }
139+ };
140+
134 int aclblasSscalTest(AclContext& ctx)141 int aclblasSscalTest(AclContext& ctx)
135 {142 {
136 aclrtStream stream = ctx.Stream();143 aclrtStream stream = ctx.Stream();
@@ -140,9 +147,9 @@
140 auto blasRet = aclblasCreate(&rawHandle);147 auto blasRet = aclblasCreate(&rawHandle);
141 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);148 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", blasRet);
142 return blasRet);149 return blasRet);
143- std::unique_ptr<void, aclblasStatus_t (*)(void*)> handlePtr(rawHandle, aclblasDestroy);150+ std::unique_ptr<std::remove_pointer<aclblasHandle_t>::type, BlasHandleDeleter> handlePtr(rawHandle);
144 151
145- blasRet = aclblasSetStream(static_cast<aclblasHandle_t>(handlePtr.get()), stream);152+ blasRet = aclblasSetStream(handlePtr.get(), stream);
146 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);153 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", blasRet);
147 return blasRet);154 return blasRet);
148 155
@@ -157,14 +164,13 @@
157 void* rawMem = nullptr;164 void* rawMem = nullptr;
158 auto aclRet = aclrtMalloc(&rawMem, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);165 auto aclRet = aclrtMalloc(&rawMem, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);
159 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);166 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", aclRet); return aclRet);
160- std::unique_ptr<void, aclError (*)(void*)> xDevicePtr(rawMem, aclrtFree);167+ std::unique_ptr<float, AclMemDeleter> xDevicePtr(static_cast<float*>(rawMem));
161 168
162 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);169 aclRet = aclrtMemcpy(xDevicePtr.get(), xBytes, xHostData.data(), xBytes, ACL_MEMCPY_HOST_TO_DEVICE);
163 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);170 CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy for x failed. ERROR: %d\n", aclRet); return aclRet);
164 171
165 // 4. 调用 aclblasSscal（alpha 为 Host 指针）172 // 4. 调用 aclblasSscal（alpha 为 Host 指针）
166- blasRet = aclblasSscal(173+ blasRet = aclblasSscal(handlePtr.get(), n, &alpha, xDevicePtr.get(), incx);
167- static_cast<aclblasHandle_t>(handlePtr.get()), n, &alpha, static_cast<float*>(xDevicePtr.get()), incx);
168 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSscal failed. ERROR: %d\n", blasRet);174 CHECK_RET(blasRet == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSscal failed. ERROR: %d\n", blasRet);
169 return blasRet);175 return blasRet);
170 176
Minclude/cann_ops_blas.h+5-0
@@ -354,6 +354,11 @@ aclblasStatus_t aclblasStpttr(
354aclblasStatus_t aclblasStrttp(354aclblasStatus_t aclblasStrttp(
355 aclblasHandle_t handle, aclblasFillMode_t uplo, int n, const float* A, int lda, float* AP);355 aclblasHandle_t handle, aclblasFillMode_t uplo, int n, const float* A, int lda, float* AP);
356 356 
357+aclblasStatus_t aclblasSrotm(aclblasHandle_t handle, int n, float* x, int incx, float* y, int incy, const float* param);
wangzitao
wangzitaowangzitao7月2日

这个和364行重复定义了。

likedislike
陈思
陈思
7月2日 评论：
358+ 
359+aclblasStatus_t aclblasSrot(
360+ aclblasHandle_t handle, int n, float* x, int incx, float* y, int incy, const float* c, const float* s);
361+ 
357aclblasStatus_t aclblasSrotg(aclblasHandle_t handle, float* a, float* b, float* c, float* s);362aclblasStatus_t aclblasSrotg(aclblasHandle_t handle, float* a, float* b, float* c, float* s);
358 363 
359aclblasStatus_t aclblasSrotm(364aclblasStatus_t aclblasSrotm(
Atest/rot/srot/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+ops_blas_add_gtest_tests(${OPS_BLAS} srot_test)
Atest/rot/srot/arch35/srot_npu_wrapper.h+174-0
@@ -0,0 +1,174 @@
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+#pragma once
12+ 
13+#include <cstdint>
14+#include <string>
15+ 
16+#include "acl/acl.h"
17+#include "cann_ops_blas.h"
18+ 
19+// NPU wrapper — c / s accepted by value, plus an optional csPtrMode that controls where the
20+// c / s pointers are materialized before being passed to the operator:
21+// csPtrMode == "host" : operator invoked with &c / &s (host stack scalars, v1 behavior).
22+// csPtrMode == "device" : c and s each allocated on device (aclrtMalloc + H2D copy), operator
23+// invoked with the device pointers.
24+// csPtrMode == "mixed" : c stays on host (&c), s allocated on device.
25+// The operator auto-determines each pointer's location via aclrtPointerGetAttributes, so the
26+// wrapper only needs to construct the pointer at the desired side; the API signature
27+// (const float* c, const float* s) is unchanged.
28+//
29+// Fast paths (passed through to the operator without device allocation):
30+// - handle == nullptr : operator returns ACLBLAS_STATUS_HANDLE_IS_NULLPTR
31+// - n <= 0 : operator returns ACLBLAS_STATUS_SUCCESS (no-op)
32+// - x == nullptr or y == nullptr (n > 0): operator returns ACLBLAS_STATUS_INVALID_VALUE
33+//
34+// Normal path: malloc device buffers covering the stride access range
35+// (element span = (n-1)*|inc|+1 for both x and y), H2D copy x/y, optionally H2D copy the
36+// device-side c / s scalar(s), call the operator, synchronize, D2H copy back x/y
37+// (in-place modification), then free.
38+//
39+// Every ACL call is checked; on any failure already-allocated device memory is freed via
40+// the freeAll lambda before returning a structured error code (no leak).
41+ 
42+// Element span covered by stride access: (n-1)*|inc| + 1 elements (>= n when |inc|>=1).
43+static inline size_t SrotBufElems(int n, int inc)
44+{
45+ if (n <= 0)
46+ return 0;
47+ int absInc = (inc < 0) ? -inc : inc;
48+ if (absInc == 0)
49+ absInc = 1; // inc==0 reuses element 0; a single element span suffices.
50+ return static_cast<size_t>(n - 1) * static_cast<size_t>(absInc) + 1;
51+}
52+ 
53+// Resolve csPtrMode ("host"/"device"/"mixed") into per-scalar on-device flags.
54+static inline void SrotResolveCsPtrMode(const std::string& csPtrMode, bool& cOnDevice, bool& sOnDevice)
55+{
56+ cOnDevice = (csPtrMode == "device");
57+ sOnDevice = (csPtrMode == "device") || (csPtrMode == "mixed");
58+}
59+ 
60+// Allocate a device scalar (sizeof(float)) and H2D-copy the host value in. On success
61+// dScalar is set and scalarArg points to it; on failure dScalar stays nullptr.
62+// Returns ACLBLAS_STATUS_ALLOC_FAILED / ACLBLAS_STATUS_INTERNAL_ERROR on ACL failure.
63+static inline aclblasStatus_t SrotPrepareDeviceScalar(float hostVal, void*& dScalar, const float*& scalarArg)
64+{
65+ scalarArg = nullptr; // filled on success
66+ if (aclrtMalloc(&dScalar, sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) {
67+ return ACLBLAS_STATUS_ALLOC_FAILED;
68+ }
69+ if (aclrtMemcpy(dScalar, sizeof(float), &hostVal, sizeof(float), ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) {
70+ aclrtFree(dScalar);
71+ dScalar = nullptr;
72+ return ACLBLAS_STATUS_INTERNAL_ERROR;
73+ }
74+ scalarArg = static_cast<const float*>(dScalar);
75+ return ACLBLAS_STATUS_SUCCESS;
76+}
77+ 
78+inline aclblasStatus_t aclblasSrot_npu(
79+ aclblasHandle_t handle, int n, float* x, int incx, float* y, int incy, const float c, const float s,
80+ const std::string& csPtrMode = "host")
81+{
82+ // Fast path: no device buffers needed (handle null / n<=0 / nullptr x or y). c / s are
83+ // passed as host pointers here regardless of csPtrMode, since the operator short-circuits
84+ // before touching them in these cases.
85+ if (handle == nullptr || n <= 0 || x == nullptr || y == nullptr) {
86+ return aclblasSrot(handle, n, x, incx, y, incy, &c, &s);
87+ }
88+ 
89+ bool cOnDevice = false;
90+ bool sOnDevice = false;
91+ SrotResolveCsPtrMode(csPtrMode, cOnDevice, sOnDevice);
92+ 
93+ const size_t xElems = SrotBufElems(n, incx);
94+ const size_t yElems = SrotBufElems(n, incy);
95+ const size_t xBytes = xElems * sizeof(float);
96+ const size_t yBytes = yElems * sizeof(float);
97+ 
98+ void* dX = nullptr;
99+ void* dY = nullptr;
100+ void* dC = nullptr;
101+ void* dS = nullptr;
102+ 
103+ auto freeAll = [&]() {
104+ if (dX)
105+ aclrtFree(dX);
106+ if (dY)
107+ aclrtFree(dY);
108+ if (dC)
109+ aclrtFree(dC);
110+ if (dS)
111+ aclrtFree(dS);
112+ };
113+ 
114+ if (aclrtMalloc(&dX, xBytes, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) {
115+ freeAll();
116+ return ACLBLAS_STATUS_ALLOC_FAILED;
117+ }
118+ if (aclrtMalloc(&dY, yBytes, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) {
119+ freeAll();
120+ return ACLBLAS_STATUS_ALLOC_FAILED;
121+ }
122+ 
123+ if (aclrtMemcpy(dX, xBytes, x, xBytes, ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) {
124+ freeAll();
125+ return ACLBLAS_STATUS_INTERNAL_ERROR;
126+ }
127+ if (aclrtMemcpy(dY, yBytes, y, yBytes, ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) {
128+ freeAll();
129+ return ACLBLAS_STATUS_INTERNAL_ERROR;
130+ }
131+ 
132+ // Construct c / s pointers on the side indicated by csPtrMode. The operator auto-detects
133+ // each pointer's location, so the wrapper just materializes it where requested.
134+ const float* cArg = &c;
135+ const float* sArg = &s;
136+ if (cOnDevice) {
137+ aclblasStatus_t st = SrotPrepareDeviceScalar(c, dC, cArg);
138+ if (st != ACLBLAS_STATUS_SUCCESS) {
139+ freeAll();
140+ return st;
141+ }
142+ }
143+ if (sOnDevice) {
144+ aclblasStatus_t st = SrotPrepareDeviceScalar(s, dS, sArg);
145+ if (st != ACLBLAS_STATUS_SUCCESS) {
146+ freeAll();
147+ return st;
148+ }
149+ }
150+ 
151+ aclblasStatus_t ret = aclblasSrot(handle, n, static_cast<float*>(dX), incx, static_cast<float*>(dY), incy,
152+ cArg, sArg);
153+ if (ret != ACLBLAS_STATUS_SUCCESS) {
154+ freeAll();
155+ return ret;
156+ }
157+ 
158+ if (aclrtSynchronizeDevice() != ACL_SUCCESS) {
159+ freeAll();
160+ return ACLBLAS_STATUS_EXECUTION_FAILED;
161+ }
162+ 
163+ if (aclrtMemcpy(x, xBytes, dX, xBytes, ACL_MEMCPY_DEVICE_TO_HOST) != ACL_SUCCESS) {
164+ freeAll();
165+ return ACLBLAS_STATUS_INTERNAL_ERROR;
166+ }
167+ if (aclrtMemcpy(y, yBytes, dY, yBytes, ACL_MEMCPY_DEVICE_TO_HOST) != ACL_SUCCESS) {
168+ freeAll();
169+ return ACLBLAS_STATUS_INTERNAL_ERROR;
170+ }
171+ 
172+ freeAll();
173+ return ret;
174+}
Atest/rot/srot/arch35/srot_test.cpp+157-0
@@ -0,0 +1,157 @@
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 <cstdint>
12+#include <algorithm>
13+#include <string>
14+#include <vector>
15+ 
16+#include "fill.h"
17+#include "verify.h"
18+#include "blas_test.h"
19+#include "csv_loader.h"
20+#include "srot_param.h"
21+#include "srot_golden.h"
22+#include "srot_npu_wrapper.h"
23+ 
24+class SrotArch35Test : public BlasTest<SrotParam> {};
25+ 
26+// Null handle test — separate TEST_F (not in CSV).
27+TEST_F(SrotArch35Test, NullHandle)
28+{
29+ float x[5] = {1, 2, 3, 4, 5};
30+ float y[5] = {6, 7, 8, 9, 10};
31+ aclblasStatus_t ret = aclblasSrot_npu(nullptr, 5, x, 1, y, 1, 0.6f, 0.8f);
32+ EXPECT_EQ(ret, ACLBLAS_STATUS_HANDLE_IS_NULLPTR);
33+}
34+ 
35+INSTANTIATE_TEST_SUITE_P(
36+ Srot, SrotArch35Test,
37+ ::testing::ValuesIn(GetCasesFromCsv<SrotParam>(ReplaceFileExtension2Csv(__FILE__))),
38+ PrintCaseInfoString<SrotParam>);
39+ 
40+// Helper: build a host buffer matching the stride access pattern.
41+// Continuous (|inc|==1) uses makeBlasArray (size = n); strided uses makeBlasStrided
42+// (size = (n-1)*|inc|+1, elements laid out at the stride positions).
43+//
44+// inc==0 special case: netlib srot.f / cblas_srot repeatedly access element 0 (serial
45+// accumulation, only x[0]/y[0] is touched). The NPU kernel matches this netlib
46+// semantics (empirically verified: only element 0 is rotated in-place N times).
47+// However the NPU wrapper's element span (SrotBufElems) coerces absInc 0 -> 1 before
48+// sizing the device buffer, yielding a span of n elements. We therefore size the host
49+// buffer to n (via makeBlasArray) whenever inc==0 so the H2D copy does not read past
50+// the host buffer end, while the golden (cblas_srot with inc==0) only mutates element 0.
51+static inline std::vector<float> MakeSrotBuffer(int n, int inc, const BlasFillMode& fill, uint32_t seed)
52+{
53+ int absInc = std::abs(inc);
54+ if (absInc <= 1) {
55+ // |inc|==1 (contiguous) AND inc==0 (device span coerces to n elements).
56+ return makeBlasArray(static_cast<int64_t>(n), fill, seed);
57+ }
58+ return makeBlasStrided(n, inc, fill, seed);
59+}
60+ 
61+TEST_P(SrotArch35Test, CsvDriven)
62+{
63+ const auto& p = GetParam();
64+ 
65+ // Step 1: generate host data (x/y are both input and output — in-place).
66+ std::vector<float> xHost = MakeSrotBuffer(p.n, p.incx, p.x, p.randomSeed);
67+ std::vector<float> yHost = MakeSrotBuffer(p.n, p.incy, p.y, p.randomSeed);
68+ 
69+ float* xPtr = xHost.empty() ? nullptr : xHost.data();
70+ float* yPtr = yHost.empty() ? nullptr : yHost.data();
71+ 
72+ // Step 2: execute on NPU. csPtrMode ("host"/"device"/"mixed") tells the wrapper where to
73+ // materialize the c / s scalar pointers; the operator auto-detects each pointer's location.
74+ aclblasStatus_t ret = aclblasSrot_npu(SrotArch35Test::handle_, p.n, xPtr, p.incx, yPtr, p.incy, p.c, p.s,
75+ p.csPtrMode);
76+ 
77+ // Step 3: check return code against expected.
78+ EXPECT_EQ(static_cast<int>(ret), static_cast<int>(p.expectResult));
79+ if (p.expectResult != ACLBLAS_STATUS_SUCCESS)
80+ return;
81+ 
82+ // n <= 0 is a no-op (x/y untouched) — nothing to verify numerically.
83+ if (p.n <= 0)
84+ return;
85+ 
86+ // Step 4: compute CPU golden on a fresh copy of the original data.
87+ std::vector<float> goldenX = MakeSrotBuffer(p.n, p.incx, p.x, p.randomSeed);
88+ std::vector<float> goldenY = MakeSrotBuffer(p.n, p.incy, p.y, p.randomSeed);
89+ aclblasSrot_cpu(SrotArch35Test::handle_, p.n, goldenX.data(), p.incx, goldenY.data(), p.incy, &p.c, &p.s);
90+ // Step 5: verify both x and y. Precision mode is chosen by data characteristics:
91+ // - NaN inputs (TC_L1_20): rotation propagates NaN; verify each element is NaN
92+ // directly (std::isnan), bypassing Verifier (NaN rel-err is undefined).
93+ // - INF inputs with identity rotation c==1&&s==0 (TC_L1_19): result must equal
94+ // the input bit-for-bit; verify EXACT.
95+ // - Otherwise: FP32 single benchmark MERE/MARE (CSV-driven thresholds).
96+ int absIncX = std::max(1, std::abs(p.incx));
97+ int absIncY = std::max(1, std::abs(p.incy));
98+ 
99+ auto containsNaN = [](const float* buf, int n, int inc) -> bool {
100+ int step = std::max(1, std::abs(inc));
101+ for (int i = 0; i < n; i++) {
102+ if (std::isnan(buf[static_cast<int64_t>(i) * step]))
103+ return true;
104+ }
105+ return false;
106+ };
107+ 
108+ if (containsNaN(xPtr, p.n, p.incx) || containsNaN(yPtr, p.n, p.incy)) {
109+ // NaN propagation: every output element must be NaN (c*x+s*y with any NaN -> NaN).
110+ bool allNanX = true;
111+ bool allNanY = true;
112+ for (int i = 0; i < p.n; i++) {
113+ if (!std::isnan(xPtr[static_cast<int64_t>(i) * absIncX]))
114+ allNanX = false;
115+ if (!std::isnan(yPtr[static_cast<int64_t>(i) * absIncY]))
116+ allNanY = false;
117+ }
118+ std::cout << "[" << p.caseName << "] NaN-propagation check: x_allNan=" << allNanX
119+ << " y_allNan=" << allNanY << std::endl;
120+ EXPECT_TRUE(allNanX && allNanY);
121+ return;
122+ }
123+ 
124+ bool identityRotation = (p.c == 1.0f && p.s == 0.0f);
125+ auto containsInf = [](const float* buf, int n, int inc) -> bool {
126+ int step = std::max(1, std::abs(inc));
127+ for (int i = 0; i < n; i++) {
128+ if (std::isinf(buf[static_cast<int64_t>(i) * step]))
129+ return true;
130+ }
131+ return false;
132+ };
133+ 
134+ if (identityRotation && (containsInf(xPtr, p.n, p.incx) || containsInf(yPtr, p.n, p.incy))) {
135+ // Identity rotation (c==1, s==0) must not short-circuit: x' = 1*x + 0*y = x,
136+ // y' = 1*y - 0*x = y. Verify bit-exact equality against the golden (which is
137+ // the original input, since cblas_srot with c=1 s=0 also returns input unchanged).
138+ VerifyConfig cfg;
139+ cfg.mode = PrecisionMode::EXACT;
140+ EXPECT_TRUE(Verifier::verifyVector(xPtr, goldenX.data(), static_cast<size_t>(p.n),
141+ static_cast<int64_t>(absIncX), cfg, p.caseName + "_x"));
142+ EXPECT_TRUE(Verifier::verifyVector(yPtr, goldenY.data(), static_cast<size_t>(p.n),
143+ static_cast<int64_t>(absIncY), cfg, p.caseName + "_y"));
144+ return;
145+ }
146+ 
147+ // Default: MERE/MARE (FP32 single benchmark), CSV columns drive per-case thresholds.
148+ VerifyConfig cfg;
149+ cfg.mode = PrecisionMode::MERE_MARE;
150+ cfg.mereThreshold = (p.mereThreshold > 0.0) ? p.mereThreshold : 1.220703125e-4; // 2^-13
151+ cfg.mareMultiplier = (p.mareMultiplier > 0.0) ? p.mareMultiplier : 10.0;
152+ 
153+ EXPECT_TRUE(Verifier::verifyVector(xPtr, goldenX.data(), static_cast<size_t>(p.n),
154+ static_cast<int64_t>(absIncX), cfg, p.caseName + "_x"));
155+ EXPECT_TRUE(Verifier::verifyVector(yPtr, goldenY.data(), static_cast<size_t>(p.n),
156+ static_cast<int64_t>(absIncY), cfg, p.caseName + "_y"));
157+}
Atest/rot/srot/arch35/srot_test.csv+51-0
@@ -0,0 +1,51 @@
1+ 
2+case_name,description,n,x,incx,y,incy,c,s,mere_threshold,mare_multiplier,expect_result,random_seed,cs_ptr_mode
3+# --- cs_ptr_mode baseline: existing functional cases on host (v1 regression) ---
4+TC_L0_02,x_nullptr_n64,64,NULLPTR,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_INVALID_VALUE,0,host
5+TC_L0_03,y_nullptr_n64,64,RANDOM_NORM_1,1,NULLPTR,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_INVALID_VALUE,0,host
6+TC_L0_04,contig_n64,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,1,host
7+TC_L0_05,contig_n1024,1024,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,2,host
8+TC_L0_06,contig_cos_sin_45deg,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.70710678,0.70710678,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,3,host
9+TC_L0_07,stride_incx2_incy3_n100,100,RANDOM_NORM_1,2,RANDOM_NORM_1,3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,4,host
10+TC_L0_08,mix_incx1_incy2_n100,100,RANDOM_NORM_1,1,RANDOM_NORM_1,2,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,5,host
11+TC_L0_09,mix_incx3_incy1_n100,100,RANDOM_NORM_1,3,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,6,host
12+TC_L0_10,contig_n33_unaligned_tail,33,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,7,host
13+TC_L1_01,n0_noop,0,INDEX,1,INDEX,1,0.6,0.8,,,ACLBLAS_STATUS_SUCCESS,0,host
14+TC_L1_02,nneg_noop,-5,INDEX,1,INDEX,1,0.6,0.8,,,ACLBLAS_STATUS_SUCCESS,0,host
15+TC_L1_03,n1_contig,1,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,8,host
16+TC_L1_04,n1_stride,1,RANDOM_NORM_1,2,RANDOM_NORM_1,3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,9,host
17+TC_L1_05,zero_stride_x,10,RANDOM_NORM_1,0,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,10,host
18+TC_L1_06,zero_stride_y,10,RANDOM_NORM_1,1,RANDOM_NORM_1,0,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,11,host
19+TC_L1_07,zero_stride_both,10,RANDOM_NORM_1,0,RANDOM_NORM_1,0,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,12,host
20+TC_L1_08,neg_stride_x,64,RANDOM_NORM_1,-1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,13,host
21+TC_L1_09,neg_stride_y,64,RANDOM_NORM_1,1,RANDOM_NORM_1,-2,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,14,host
22+TC_L1_10,neg_stride_both,64,RANDOM_NORM_1,-2,RANDOM_NORM_1,-3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,15,host
23+TC_L1_11,mix_pos_neg_stride,64,RANDOM_NORM_1,2,RANDOM_NORM_1,-3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,16,host
24+TC_L1_12,mix_neg_pos_stride,64,RANDOM_NORM_1,-2,RANDOM_NORM_1,3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,17,host
25+TC_L1_13,identity_contig_c1s0,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,1.0,0.0,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,18,host
26+TC_L1_14,pure_swap_c0s1,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.0,1.0,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,19,host
27+TC_L1_15,cs_both_neg,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,-0.6,-0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,20,host
28+TC_L1_16,c_neg_s_pos,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,-0.8,0.6,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,21,host
29+TC_L1_17,identity_stride_c1s0,64,RANDOM_NORM_1,2,RANDOM_NORM_1,3,1.0,0.0,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,22,host
30+TC_L1_18,zero_input,64,VALUE_NORM_0,1,VALUE_NORM_0,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,23,host
31+TC_L1_19,inf_input_identity,8,VALUE_NORM_INF,1,VALUE_NORM_INF,1,1.0,0.0,,,ACLBLAS_STATUS_SUCCESS,24,host
32+TC_L1_20,nan_input,8,VALUE_NORM_NAN,1,VALUE_NORM_NAN,1,0.6,0.8,,,ACLBLAS_STATUS_SUCCESS,25,host
33+TC_L1_26,large_n1048576_stride,1048576,RANDOM_NORM_1,2,RANDOM_NORM_1,3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,31,host
34+# --- cs device pointer: both c and s on device HBM (aclrtMalloc + H2D) ---
35+TC_L1_27,contig_n64_cs_device,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,32,device
36+TC_L1_28,contig_n1024_cs_device,1024,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,33,device
37+TC_L1_29,stride_incx2_incy3_cs_device,100,RANDOM_NORM_1,2,RANDOM_NORM_1,3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,34,device
38+TC_L1_30,n1_contig_cs_device,1,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,35,device
39+TC_L1_31,neg_stride_both_cs_device,64,RANDOM_NORM_1,-2,RANDOM_NORM_1,-3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,36,device
40+TC_L1_32,cos_sin_45deg_cs_device,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.70710678,0.70710678,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,37,device
41+TC_L1_33,identity_c1s0_cs_device,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,1.0,0.0,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,38,device
42+TC_L1_34,pure_swap_c0s1_cs_device,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.0,1.0,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,39,device
43+TC_L1_35,cs_both_neg_cs_device,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,-0.6,-0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,40,device
44+TC_L1_36,large_n1048576_stride_cs_device,1048576,RANDOM_NORM_1,2,RANDOM_NORM_1,3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,41,device
45+# --- cs mixed pointer: c on host," s on device (independent per-pointer determination) ---"
46+TC_L1_37,contig_n64_cs_mixed,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,42,mixed
47+TC_L1_38,contig_n1024_cs_mixed,1024,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,43,mixed
48+TC_L1_39,stride_incx2_incy3_cs_mixed,100,RANDOM_NORM_1,2,RANDOM_NORM_1,3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,44,mixed
49+TC_L1_40,neg_stride_both_cs_mixed,64,RANDOM_NORM_1,-2,RANDOM_NORM_1,-3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,45,mixed
50+TC_L1_41,cos_sin_45deg_cs_mixed,64,RANDOM_NORM_1,1,RANDOM_NORM_1,1,0.70710678,0.70710678,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,46,mixed
51+TC_L1_42,large_n1048576_stride_cs_mixed,1048576,RANDOM_NORM_1,2,RANDOM_NORM_1,3,0.6,0.8,1.220703125e-4,10,ACLBLAS_STATUS_SUCCESS,47,mixed
Atest/rot/srot/srot_golden.h+48-0
@@ -0,0 +1,48 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#pragma once
12+ 
13+#include <cstdint>
14+ 
15+#include "acl/acl.h"
16+#include "cann_ops_blas.h"
17+#include "cblas_compat.h"
18+ 
19+// CPU golden for aclblasSrot — signature identical to the API (c / s passed by pointer,
20+// aligned with cublasSrot), returns aclblasStatus_t.
21+//
22+// Parameter validation is aligned with the NPU operator AND netlib srot.f (1.2 requirement):
23+// - handle == nullptr -> ACLBLAS_STATUS_HANDLE_IS_NULLPTR
24+// - n <= 0 -> ACLBLAS_STATUS_SUCCESS (no-op, x/y untouched)
25+// - x == nullptr || y == nullptr -> ACLBLAS_STATUS_INVALID_VALUE (n > 0)
26+// - c == nullptr || s == nullptr -> ACLBLAS_STATUS_INVALID_VALUE (n > 0)
27+// - incx == 0 / incy == 0 -> NOT rejected (netlib reuses x[0]/y[0] N times)
28+// - incx < 0 / incy < 0 -> NOT rejected (netlib walks from tail end)
29+// - c == 1 && s == 0 -> NOT short-circuited (netlib runs the rotation)
30+//
31+// After validation, cblas_srot (column-major, OpenBLAS) is used as the reference. cblas_srot
32+// takes c / s by value, so the pointers are dereferenced (*c / *s) at the call site. Its srot
33+// implementation matches netlib srot.f boundary behavior including inc==0 / negative stride.
34+inline aclblasStatus_t aclblasSrot_cpu(
35+ aclblasHandle_t handle, int n, float* x, int incx, float* y, int incy, const float* c, const float* s)
36+{
37+ if (handle == nullptr)
38+ return ACLBLAS_STATUS_HANDLE_IS_NULLPTR;
39+ if (n <= 0)
40+ return ACLBLAS_STATUS_SUCCESS;
41+ if (x == nullptr || y == nullptr)
42+ return ACLBLAS_STATUS_INVALID_VALUE;
43+ if (c == nullptr || s == nullptr)
44+ return ACLBLAS_STATUS_INVALID_VALUE;
45+ 
46+ cblas_srot(n, x, incx, y, incy, *c, *s);
47+ return ACLBLAS_STATUS_SUCCESS;
48+}
Atest/rot/srot/srot_param.h+56-0
@@ -0,0 +1,56 @@
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+#pragma once
12+ 
13+#include <cstdint>
14+#include <string>
15+ 
16+#include "acl/acl.h"
17+#include "cann_ops_blas.h"
18+#include "csv_loader.h"
19+ 
20+// cs_ptr_mode — controls where the c / s scalar pointers live when passed to the operator.
21+// Aligned with the design (v2): the operator queries each pointer's location via
22+// aclrtPointerGetAttributes and treats it independently, so c / s need NOT be on the same side.
23+// "host" : c and s both as host stack scalars (v1 behavior, baseline regression).
24+// "device" : c and s both allocated on device HBM (aclrtMalloc + H2D), passed by device ptr.
25+// "mixed" : c on host, s on device (verifies the independent per-pointer determination).
26+// Default is "host" so pre-existing CSV rows (which omit the column) keep v1 behavior.
27+// The actual parsing helper used at runtime lives in arch35/srot_npu_wrapper.h
28+// (SrotResolveCsPtrMode), which is called from aclblasSrot_npu.
29+ 
30+// Parameter struct for aclblasSrot.
31+// Field order matches the aclblasSrot API signature:
32+// aclblasSrot(handle, n, x, incx, y, incy, c, s)
33+// mere_threshold / mare_multiplier come from BlasTestParamBase (CSV columns),
34+// used in MERE_MARE precision mode (FP32 single benchmark: MERE <= 2^-13, MARE <= 10*2^-13).
35+struct SrotParam : public BlasTestParamBase {
36+ int n = 0;
37+ BlasFillMode x = parseFill("RANDOM_NORM_1");
38+ int incx = 1;
39+ BlasFillMode y = parseFill("RANDOM_NORM_1");
40+ int incy = 1;
41+ float c = 0.6f;
42+ float s = 0.8f;
43+ std::string csPtrMode = "host";
44+ 
45+ SrotParam(const csv_map& m) : BlasTestParamBase(m)
46+ {
47+ n = parseInt(ReadMap(m, "n", "0"));
48+ x = parseFill(ReadMap(m, "x", "RANDOM_NORM_1"));
49+ incx = parseInt(ReadMap(m, "incx", "1"));
50+ y = parseFill(ReadMap(m, "y", "RANDOM_NORM_1"));
51+ incy = parseInt(ReadMap(m, "incy", "1"));
52+ c = parseFloat(ReadMap(m, "c", "0.6"));
53+ s = parseFloat(ReadMap(m, "s", "0.8"));
54+ csPtrMode = ReadMap(m, "cs_ptr_mode", "host");
55+ }
56+};