Star110
87
代码
代码
Issues32
Pull Requests14
Actions
项目讨论
Wiki
分析
Star110
87
已开启
Feat: 新增面向 arch22 的 aclblasSspr 接口 #333
guodong54_创建于 9 天前
Feat: 新增面向 arch22 的 aclblasSspr 接口 #333
已开启
guodong54_创建于 9 天前
从guodong54_/ops-blasarch22-sspr-release合入到cann/ops-blasmaster
共 6 个文件变更+843-0
Ablas/spr/arch22/sspr_host.cpp+97-0
@@ -0,0 +1,97 @@
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 <climits>
13+#include <cstdint>
14+#include "acl/acl.h"
15+#include "log/log.h"
16+#include "cann_ops_blas.h"
17+#include "cann_ops_blas_common.h"
18+#include "sspr_tiling_data.h"
19+#include "common/helper/aclblas_handle_internal.h"
20+#include "common/helper/kernel_constant.h"
21+#include "common/helper/host_utils.h"
22+ 
23+// ---- kernel entry (local declaration, replacing aclblas_kernel_do.h) ----
24+void sspr_kernel_do(uint8_t* x, uint8_t* ap, const SsprTilingData& tiling, uint32_t numBlocks, void* stream);
25+ 
26+static aclblasStatus_t ValidateSsprParams(
27+ aclblasFillMode_t uplo, int incx, const float* alpha, const float* x, const float* ap)
28+{
29+ CHECK_RET(
30+ alpha != nullptr, OP_LOGE("aclblasSspr", "alpha must not be nullptr"); return ACLBLAS_STATUS_INVALID_VALUE);
31+ CHECK_RET(
32+ uplo == ACLBLAS_UPPER || uplo == ACLBLAS_LOWER,
33+ OP_LOGE("aclblasSspr", "uplo must be UPPER(121) or LOWER(122), got %d", static_cast<int>(uplo));
34+ return ACLBLAS_STATUS_INVALID_VALUE);
35+ CHECK_RET(incx != 0, OP_LOGE("aclblasSspr", "incx must not be zero"); return ACLBLAS_STATUS_INVALID_VALUE);
36+ CHECK_RET(incx != INT_MIN, OP_LOGE("aclblasSspr", "incx must not be INT_MIN"); return ACLBLAS_STATUS_INVALID_VALUE);
37+ CHECK_RET(x != nullptr, OP_LOGE("aclblasSspr", "x must not be nullptr"); return ACLBLAS_STATUS_INVALID_VALUE);
38+ CHECK_RET(ap != nullptr, OP_LOGE("aclblasSspr", "ap must not be nullptr"); return ACLBLAS_STATUS_INVALID_VALUE);
39+ return ACLBLAS_STATUS_SUCCESS;
40+}
41+ 
42+static SsprTilingData CalSsprTilingData(uint32_t useNumBlocks, int n, aclblasFillMode_t uplo, float alpha, int incx)
43+{
44+ SsprTilingData tilingData{};
45+ uint32_t nU32 = static_cast<uint32_t>(n);
46+ uint32_t colsPerBlk = CeilDiv<uint32_t>(nU32, useNumBlocks);
47+ tilingData.numThreads = std::min(CeilAlign<uint32_t>(colsPerBlk, SIMT_MIN_THREAD_NUM), SIMT_MAX_THREAD_NUM);
48+ tilingData.columnsPerBlock = colsPerBlk;
49+ tilingData.n = nU32;
50+ tilingData.uplo = static_cast<uint32_t>(uplo);
51+ tilingData.alpha = alpha;
52+ tilingData.incx = static_cast<int64_t>(incx);
53+ return tilingData;
54+}
55+ 
56+aclblasStatus_t aclblasSspr(
57+ aclblasHandle_t handle, aclblasFillMode_t uplo, int n, const float* alpha, const float* x, int incx, float* ap)
58+{
59+ auto* h = handle;
60+ CHECK_RET(h != nullptr, OP_LOGE("aclblasSspr", "handle is nullptr"); return ACLBLAS_STATUS_HANDLE_IS_NULLPTR);
61+ 
62+ CHECK_RET(n >= 0, OP_LOGE("aclblasSspr", "n must be >= 0, got %d", n); return ACLBLAS_STATUS_INVALID_VALUE);
63+ if (n == 0) {
64+ return ACLBLAS_STATUS_SUCCESS;
65+ }
66+ 
67+ aclblasStatus_t st = ValidateSsprParams(uplo, incx, alpha, x, ap);
68+ if (st != ACLBLAS_STATUS_SUCCESS) {
69+ return st;
70+ }
71+ 
72+ float alphaVal = *alpha;
73+ if (alphaVal == 0.0f) {
74+ return ACLBLAS_STATUS_SUCCESS;
75+ }
76+ 
77+ aclrtStream stream = h->stream;
78+ 
79+ uint32_t aivCoreNum = GetAivCoreCount();
80+ if (aivCoreNum == 0) {
81+ OP_LOGE("aclblasSspr", "vector core count is 0");
82+ return ACLBLAS_STATUS_EXECUTION_FAILED;
83+ }
84+ uint32_t useNumBlocks = (incx != 1 && n >= 4096)
85+ ? aivCoreNum
86+ : std::min(CeilDiv<uint32_t>(n, SIMT_MIN_THREAD_NUM), aivCoreNum);
87+ 
88+ SsprTilingData tiling = CalSsprTilingData(useNumBlocks, n, uplo, alphaVal, incx);
89+ 
90+ OP_LOGI(
91+ "aclblasSspr", "launching kernel: blocks=%u, cores=%u, ubEligible=%d", useNumBlocks, aivCoreNum,
92+ (incx == 1 && n >= static_cast<int>(UB_THRESHOLD) && static_cast<uint32_t>(n) <= UB_X_FLOATS));
93+ 
94+ sspr_kernel_do((uint8_t*)const_cast<float*>(x), (uint8_t*)ap, tiling, useNumBlocks, stream);
95+ 
96+ return ACLBLAS_STATUS_SUCCESS;
97+}
Ablas/spr/arch22/sspr_kernel.cpp+501-0
@@ -0,0 +1,501 @@
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 "kernel_operator.h"
13+#include "c_api/asc_simd.h"
14+#include "cann_ops_blas_common.h"
15+#include "sspr_tiling_data.h"
16+ 
17+using namespace AscendC;
18+ 
19+namespace {
20+constexpr uint32_t kChunkFloats = 8192;
21+constexpr uint32_t kStridedTileFloats = 512;
22+constexpr uint32_t kMaxPackedStride = 128;
23+constexpr uint32_t kCanonicalXFloats = 32768 + 8;
24+constexpr uint32_t kMaxCanonicalN = 32768;
25+constexpr uint32_t kPackSpanFloats = 4096;
26+constexpr uint32_t kPackOffsetFloats = 4096;
27+constexpr uint32_t kLocalRowFloats = 4096;
28+constexpr uint32_t kLocalColFloats = 8192;
29+ 
30+// Cost of the columns [0, columns), used to split columns across AI vector cores.
31+//
32+// A core's time is not proportional to the number of packed elements it owns. Owning k
33+// columns costs roughly
34+//
35+// sum over the k columns of len + k * p + q * (extra transfers)
36+//
37+// where p is a flat charge every column pays regardless of length (AP base address, scalar
38+// broadcast of x[col], loop and descriptor setup, first DMA round trip) and q is charged
39+// again for each further kChunkFloats piece a long column has to walk. Equalising element
40+// counts ignores both: at n = 32768 the slowest core took 2x the mean, because half the
41+// columns of a triangular matrix are short and a core owning many of them pays p many times
42+// while moving few elements.
43+//
44+// Implemented below as kCostPerColumn * columns + kCostPerElement * elements +
45+// kCostPerTransfer * transfers, which is the same function since every column contributes at
46+// least one transfer:
47+//
48+// kCostPerColumn + len + kCostPerTransfer * ceil(len / kChunkFloats)
49+// = (kCostPerColumn + kCostPerTransfer) + len + kCostPerTransfer * floor((len - 1) / kChunkFloats)
50+// \___________ p ________________/ \_ extra transfers _/
51+//
52+// Note which quantity is actually identifiable. For len <= kChunkFloats the ceiling is 1, so
53+// only the sum p = kCostPerColumn + kCostPerTransfer affects the split; the two constants are
54+// separable only where columns span several transfers. Measurements bear that out: candidates
55+// sharing a sum scored within 0.2 percentage points of each other, far inside the ~2.7% per
56+// case run-to-run spread, so do not read the individual values as measurements of a per-column
57+// and a per-transfer cost. What the data supports is p in the 4096-8192 range (p = 0 scores
58+// 0.71x, i.e. worse than an equal-column split) and a smaller positive extra-transfer charge.
59+//
60+// Selected by timing 44 (n, uplo, incx) configurations over a grid of candidates on 910B3,
61+// then 9 repeats of the top four. A later LOWER-only sweep tested 25 coefficient pairs and
62+// confirmed four finalists at warmup 300 / measurement 1000 over seven sizes. The alternatives
63+// gained at most 2.0% geometric mean but regressed individual sizes by up to 10.2%, so LOWER
64+// retains the same coefficients as UPPER. Each term has an O(1) closed form so the boundary
65+// search below stays O(log n).
66+constexpr uint64_t kCostPerColumn = 2048; // with kCostPerTransfer: p = 6144 per column
67+constexpr uint64_t kCostPerElement = 1; // per packed element streamed
68+constexpr uint64_t kCostPerTransfer = 4096; // per kChunkFloats-sized transfer
69+ 
70+// sum_{len=1..m} ceil(len / grain)
71+__aicore__ inline uint64_t CeilRunSum(uint64_t m, uint64_t grain)
72+{
73+ if (m == 0ULL) return 0ULL;
74+ const uint64_t q = m / grain;
75+ const uint64_t r = m - q * grain;
76+ return grain * q * (q + 1ULL) / 2ULL + r * (q + 1ULL);
77+}
78+ 
79+__aicore__ inline uint64_t ColumnWorkPrefix(uint32_t columns, uint32_t n, uint32_t uplo)
80+{
81+ // Column c has length c+1 (UPPER) or n-c (LOWER); the sums below run over [0, columns).
82+ const uint64_t elements = uplo == ACLBLAS_UPPER
83+ ? static_cast<uint64_t>(columns) * (columns + 1ULL) / 2ULL
84+ : static_cast<uint64_t>(columns) * (2ULL * n - columns + 1ULL) / 2ULL;
85+ const uint64_t transfers = uplo == ACLBLAS_UPPER
86+ ? CeilRunSum(columns, kChunkFloats)
87+ : CeilRunSum(n, kChunkFloats) - CeilRunSum(static_cast<uint64_t>(n) - columns, kChunkFloats);
88+ return kCostPerColumn * static_cast<uint64_t>(columns)
89+ + kCostPerElement * elements
90+ + kCostPerTransfer * transfers;
91+}
92+ 
93+__aicore__ inline uint64_t PartitionTarget(uint64_t total, uint32_t part, uint32_t parts)
94+{
95+ return (total / parts) * part + (total % parts) * part / parts;
96+}
97+ 
98+__aicore__ inline uint32_t WeightedColumnBoundary(
99+ uint32_t part, uint32_t parts, uint32_t n, uint32_t uplo)
100+{
101+ if (part == 0U) return 0U;
102+ if (part >= parts) return n;
103+ const uint64_t target = PartitionTarget(ColumnWorkPrefix(n, n, uplo), part, parts);
104+ uint32_t lo = 0U;
105+ uint32_t hi = n;
106+ while (lo < hi) {
107+ const uint32_t mid = lo + (hi - lo) / 2U;
108+ if (ColumnWorkPrefix(mid, n, uplo) < target) lo = mid + 1U;
109+ else hi = mid;
110+ }
111+ return lo;
112+}
113+ 
114+__aicore__ inline uint64_t UpperColumnBase(uint32_t col)
115+{
116+ return static_cast<uint64_t>(col) * (col + 1ULL) / 2ULL;
117+}
118+ 
119+__aicore__ inline uint64_t LowerColumnBase(uint32_t col, uint32_t n)
120+{
121+ return static_cast<uint64_t>(col) * (2ULL * n - col + 1ULL) / 2ULL;
122+}
123+ 
124+__aicore__ inline uint64_t XPhysicalIndex(uint32_t logical, uint32_t n, int64_t incx)
125+{
126+ const uint64_t absIncx = static_cast<uint64_t>(incx >= 0 ? incx : -incx);
127+ return incx >= 0 ? static_cast<uint64_t>(logical) * absIncx
128+ : static_cast<uint64_t>(n - 1U - logical) * absIncx;
129+}
130+ 
131+__aicore__ inline void ScalarColumns(
132+ GlobalTensor<float>& xGm, GlobalTensor<float>& apGm, const SsprTilingData& tiling,
133+ uint32_t colStart, uint32_t colEnd)
134+{
135+ for (uint32_t col = colStart; col < colEnd; ++col) {
136+ const float axCol = tiling.alpha * xGm.GetValue(XPhysicalIndex(col, tiling.n, tiling.incx));
137+ if (tiling.uplo == ACLBLAS_UPPER) {
138+ const uint64_t base = UpperColumnBase(col);
139+ for (uint32_t row = 0; row <= col; ++row) {
140+ const uint64_t apIndex = base + row;
141+ apGm.SetValue(apIndex, apGm.GetValue(apIndex) +
142+ axCol * xGm.GetValue(XPhysicalIndex(row, tiling.n, tiling.incx)));
143+ }
144+ } else {
145+ const uint64_t base = LowerColumnBase(col, tiling.n);
146+ for (uint32_t row = col; row < tiling.n; ++row) {
147+ const uint64_t apIndex = base + (row - col);
148+ apGm.SetValue(apIndex, apGm.GetValue(apIndex) +
149+ axCol * xGm.GetValue(XPhysicalIndex(row, tiling.n, tiling.incx)));
150+ }
151+ }
152+ }
153+}
154+ 
155+#if __DAV_C220_VEC__
156+__aicore__ inline void CanonicalizeStridedRange(
157+ LocalTensor<float> cache, LocalTensor<float> span, LocalTensor<uint32_t> offsets,
158+ __gm__ float* source, uint32_t n, int64_t incx, uint32_t logicalBase, uint32_t logicalCount)
159+{
160+ const uint32_t increment = static_cast<uint32_t>(incx > 0 ? incx : -incx);
161+ auto ubSpan = reinterpret_cast<__ubuf__ float*>(span.GetPhyAddr());
162+ auto ubCache = reinterpret_cast<__ubuf__ float*>(cache.GetPhyAddr());
163+ 
164+ // Gather() writes to cache[localStart], and localStart advances by logicalTile every
165+ // iteration. A UB vector destination must begin on a 32-byte (8-float) boundary, so
166+ // logicalTile must be a multiple of 8. The span buffer additionally bounds one tile to
167+ // (logicalTile - 1) * increment + 1 <= kPackSpanFloats floats. Both constraints together
168+ // require increment <= kPackSpanFloats / 8; larger strides cannot fill even 8 logical
169+ // elements from one span read and are handled by the scalar path below.
170+ const uint32_t logicalTile = (kPackSpanFloats / increment) & ~7U;
171+ if (logicalTile == 0U) {
172+ // Rare very large stride. Move one element at a time: correctness over throughput.
173+ for (uint32_t i = 0; i < logicalCount; ++i) {
174+ const uint64_t physical = XPhysicalIndex(logicalBase + i, n, incx);
175+ set_flag(PIPE_S, PIPE_MTE2, EVENT_ID1);
176+ wait_flag(PIPE_S, PIPE_MTE2, EVENT_ID1);
177+ copy_gm_to_ubuf_align_b32(ubSpan, source + physical, 0, 1, sizeof(float), 0, 0, 0, 0);
178+ set_flag(PIPE_MTE2, PIPE_S, EVENT_ID1);
179+ wait_flag(PIPE_MTE2, PIPE_S, EVENT_ID1);
180+ ubCache[i] = ubSpan[0];
181+ }
182+ set_flag(PIPE_S, PIPE_V, EVENT_ID1);
183+ wait_flag(PIPE_S, PIPE_V, EVENT_ID1);
184+ return;
185+ }
186+ 
187+ for (uint32_t localStart = 0; localStart < logicalCount; localStart += logicalTile) {
188+ const uint32_t count = (logicalCount - localStart < logicalTile)
189+ ? logicalCount - localStart : logicalTile;
190+ const uint32_t globalStart = logicalBase + localStart;
191+ const uint32_t physicalStart = incx > 0 ? globalStart * increment
192+ : (n - globalStart - count) * increment;
193+ const uint32_t spanCount = (count - 1U) * increment + 1U;
194+ for (uint32_t i = 0; i < count; ++i) {
195+ const uint32_t sourceIndex = incx > 0 ? i * increment : (count - 1U - i) * increment;
196+ offsets.SetValue(i, sourceIndex * sizeof(float));
197+ }
198+ set_flag(PIPE_S, PIPE_MTE2, EVENT_ID1);
199+ wait_flag(PIPE_S, PIPE_MTE2, EVENT_ID1);
200+ copy_gm_to_ubuf_align_b32(
201+ ubSpan, source + physicalStart, 0, 1, spanCount * sizeof(float), 0, 0, 0, 0);
202+ set_flag(PIPE_MTE2, PIPE_V, EVENT_ID1);
203+ wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID1);
204+ Gather(cache[localStart], span, offsets, 0U, count);
205+ PipeBarrier<PIPE_V>();
206+ set_flag(PIPE_V, PIPE_S, EVENT_ID1);
207+ wait_flag(PIPE_V, PIPE_S, EVENT_ID1);
208+ }
209+}
210+__aicore__ inline void CachedStridedColumns(
211+ __gm__ float* gmX, GlobalTensor<float>& apGm, const SsprTilingData& tiling,
212+ uint32_t colStart, uint32_t colEnd)
213+{
214+ constexpr uint32_t kRowTile = kLocalRowFloats;
215+ TPipe pipe;
216+ TBuf<TPosition::VECCALC> xColBuf;
217+ TBuf<TPosition::VECCALC> xRowBuf;
218+ TBuf<TPosition::VECCALC> spanBuf;
219+ TBuf<TPosition::VECCALC> offsetBuf;
220+ TBuf<TPosition::VECCALC> xWorkBuf;
221+ TQue<QuePosition::VECIN, 2> apQueue;
222+ TQue<QuePosition::VECOUT, 2> outQueue;
223+ pipe.InitBuffer(xColBuf, kLocalColFloats * sizeof(float));
224+ pipe.InitBuffer(xRowBuf, kRowTile * sizeof(float));
225+ pipe.InitBuffer(spanBuf, kPackSpanFloats * sizeof(float));
226+ pipe.InitBuffer(offsetBuf, kPackOffsetFloats * sizeof(uint32_t));
227+ pipe.InitBuffer(xWorkBuf, kRowTile * sizeof(float));
228+ pipe.InitBuffer(apQueue, 2, kRowTile * sizeof(float));
229+ pipe.InitBuffer(outQueue, 2, kRowTile * sizeof(float));
230+ LocalTensor<float> xCol = xColBuf.Get<float>();
231+ LocalTensor<float> xRow = xRowBuf.Get<float>();
232+ LocalTensor<float> xWork = xWorkBuf.Get<float>();
233+ LocalTensor<float> span = spanBuf.Get<float>();
234+ LocalTensor<uint32_t> offsets = offsetBuf.Get<uint32_t>();
235+ auto ubXCol = reinterpret_cast<__ubuf__ float*>(xCol.GetPhyAddr());
236+ auto ubOffsets = reinterpret_cast<__ubuf__ uint32_t*>(offsets.GetPhyAddr());
237+ const uint32_t colCount = colEnd - colStart;
238+ if (colCount > kLocalColFloats) return;
239+ CanonicalizeStridedRange(xCol, span, offsets, gmX, tiling.n, tiling.incx, colStart, colCount);
240+ 
241+ for (uint32_t tileRow = 0; tileRow < tiling.n; tileRow += kRowTile) {
242+ const uint32_t tileEnd = (tiling.n - tileRow < kRowTile) ? tiling.n : tileRow + kRowTile;
243+ const uint32_t fullCount = tileEnd - tileRow;
244+ CanonicalizeStridedRange(xRow, span, offsets, gmX, tiling.n, tiling.incx, tileRow, fullCount);
245+ const DataCopyExtParams fullCopy{1, fullCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
246+ const DataCopyPadExtParams<float> pad{true, 0, 0, 0.0f};
247+ 
248+ const uint32_t rectBegin = tiling.uplo == ACLBLAS_UPPER
249+ ? (colStart > tileEnd ? colStart : tileEnd) : colStart;
250+ const uint32_t rectEnd = tiling.uplo == ACLBLAS_UPPER
251+ ? colEnd : (colEnd < tileRow ? colEnd : tileRow);
252+ if (rectBegin < rectEnd) {
253+ const uint32_t firstCol = rectBegin;
254+ const uint64_t firstOffset = tiling.uplo == ACLBLAS_UPPER
255+ ? UpperColumnBase(firstCol) + tileRow
256+ : LowerColumnBase(firstCol, tiling.n) + (tileRow - firstCol);
257+ LocalTensor<float> firstAp = apQueue.AllocTensor<float>();
258+ DataCopyPad(firstAp, apGm[firstOffset], fullCopy, pad);
259+ apQueue.EnQue(firstAp);
260+ for (uint32_t col = rectBegin; col < rectEnd; ++col) {
261+ if (col + 1U < rectEnd) {
262+ const uint32_t nextCol = col + 1U;
263+ const uint64_t nextOffset = tiling.uplo == ACLBLAS_UPPER
264+ ? UpperColumnBase(nextCol) + tileRow
265+ : LowerColumnBase(nextCol, tiling.n) + (tileRow - nextCol);
266+ LocalTensor<float> nextAp = apQueue.AllocTensor<float>();
267+ DataCopyPad(nextAp, apGm[nextOffset], fullCopy, pad);
268+ apQueue.EnQue(nextAp);
269+ }
270+ LocalTensor<float> apIn = apQueue.DeQue<float>();
271+ LocalTensor<float> out = outQueue.AllocTensor<float>();
272+ const float axCol = tiling.alpha * ubXCol[col - colStart];
273+ Muls(out, xRow, axCol, fullCount);
274+ Add(out, out, apIn, fullCount);
275+ outQueue.EnQue(out);
276+ apQueue.FreeTensor(apIn);
277+ out = outQueue.DeQue<float>();
278+ const uint64_t apOffset = tiling.uplo == ACLBLAS_UPPER
279+ ? UpperColumnBase(col) + tileRow
280+ : LowerColumnBase(col, tiling.n) + (tileRow - col);
281+ DataCopyPad(apGm[apOffset], out, fullCopy);
282+ outQueue.FreeTensor(out);
283+ }
284+ }
285+ 
286+ const uint32_t diagBegin = colStart > tileRow ? colStart : tileRow;
287+ const uint32_t diagEnd = colEnd < tileEnd ? colEnd : tileEnd;
288+ if (diagBegin < diagEnd) {
289+ if (tiling.uplo == ACLBLAS_UPPER) {
290+ for (uint32_t col = diagBegin; col < diagEnd; ++col) {
291+ const uint32_t count = col + 1U - tileRow;
292+ const uint64_t apOffset = UpperColumnBase(col) + tileRow;
293+ DataCopyExtParams copy{1, count * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
294+ LocalTensor<float> apIn = apQueue.AllocTensor<float>();
295+ DataCopyPad(apIn, apGm[apOffset], copy, pad);
296+ apQueue.EnQue(apIn);
297+ apIn = apQueue.DeQue<float>();
298+ LocalTensor<float> out = outQueue.AllocTensor<float>();
299+ const float axCol = tiling.alpha * ubXCol[col - colStart];
300+ Muls(out, xRow, axCol, count);
301+ Add(out, out, apIn, count);
302+ outQueue.EnQue(out);
303+ apQueue.FreeTensor(apIn);
304+ out = outQueue.DeQue<float>();
305+ DataCopyPad(apGm[apOffset], out, copy);
306+ outQueue.FreeTensor(out);
307+ }
308+ } else {
309+ for (uint32_t phase = 0; phase < 8U; ++phase) {
310+ for (uint32_t i = 0; i < fullCount; ++i)
311+ ubOffsets[i] = (i + phase) * sizeof(float);
312+ set_flag(PIPE_S, PIPE_V, EVENT_ID2);
313+ wait_flag(PIPE_S, PIPE_V, EVENT_ID2);
314+ uint32_t col = diagBegin + ((phase + 8U - (diagBegin & 7U)) & 7U);
315+ for (; col < diagEnd; col += 8U) {
316+ const uint32_t localStart = col - tileRow;
317+ const uint32_t count = tileEnd - col;
318+ Gather(xWork, xRow[localStart - phase], offsets, 0U, count);
319+ PipeBarrier<PIPE_V>();
320+ const uint64_t apOffset = LowerColumnBase(col, tiling.n);
321+ DataCopyExtParams copy{1, count * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
322+ LocalTensor<float> apIn = apQueue.AllocTensor<float>();
323+ DataCopyPad(apIn, apGm[apOffset], copy, pad);
324+ apQueue.EnQue(apIn);
325+ apIn = apQueue.DeQue<float>();
326+ LocalTensor<float> out = outQueue.AllocTensor<float>();
327+ const float axCol = tiling.alpha * ubXCol[col - colStart];
328+ Muls(out, xWork, axCol, count);
329+ Add(out, out, apIn, count);
330+ outQueue.EnQue(out);
331+ apQueue.FreeTensor(apIn);
332+ out = outQueue.DeQue<float>();
333+ DataCopyPad(apGm[apOffset], out, copy);
334+ outQueue.FreeTensor(out);
335+ }
336+ set_flag(PIPE_V, PIPE_S, EVENT_ID2);
337+ wait_flag(PIPE_V, PIPE_S, EVENT_ID2);
338+ }
339+ }
340+ }
341+ }
342+}
343+ 
344+__aicore__ inline void UnifiedContiguousColumns(
345+ GlobalTensor<float>& xGm, GlobalTensor<float>& apGm, const SsprTilingData& tiling,
346+ uint32_t colStart, uint32_t colEnd)
347+{
348+ constexpr uint32_t kRowTile = 16384U;
349+ constexpr uint32_t kPipelineFloats = 4096U;
350+ TPipe pipe;
351+ TBuf<TPosition::VECCALC> xTileBuf;
352+ TQue<QuePosition::VECIN, 1> xQueue;
353+ TQue<QuePosition::VECIN, 2> apQueue;
354+ TQue<QuePosition::VECOUT, 2> outQueue;
355+ pipe.InitBuffer(xTileBuf, kRowTile * sizeof(float));
356+ pipe.InitBuffer(xQueue, 1, kPipelineFloats * sizeof(float));
357+ pipe.InitBuffer(apQueue, 2, kPipelineFloats * sizeof(float));
358+ pipe.InitBuffer(outQueue, 2, kPipelineFloats * sizeof(float));
359+ LocalTensor<float> xTile = xTileBuf.Get<float>();
360+ const DataCopyPadExtParams<float> pad{true, 0, 0, 0.0f};
361+ 
362+ for (uint32_t tileRow = 0; tileRow < tiling.n; tileRow += kRowTile) {
363+ const uint32_t tileEnd = (tiling.n - tileRow < kRowTile) ? tiling.n : tileRow + kRowTile;
364+ const uint32_t tileCount = tileEnd - tileRow;
365+ const DataCopyExtParams tileCopy{1, tileCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
366+ DataCopyPad(xTile, xGm[tileRow], tileCopy, pad);
367+ PipeBarrier<PIPE_ALL>();
368+ 
369+ for (uint32_t col = colStart; col < colEnd; ++col) {
370+ const uint32_t rowBegin = tiling.uplo == ACLBLAS_UPPER ? 0U : col;
371+ const uint32_t rowEnd = tiling.uplo == ACLBLAS_UPPER ? col + 1U : tiling.n;
372+ uint32_t row = rowBegin > tileRow ? rowBegin : tileRow;
373+ const uint32_t segmentEnd = rowEnd < tileEnd ? rowEnd : tileEnd;
374+ if (row >= segmentEnd) continue;
375+ const uint64_t apBase = tiling.uplo == ACLBLAS_UPPER ? UpperColumnBase(col)
376+ : LowerColumnBase(col, tiling.n);
377+ const float axCol = tiling.alpha * xGm.GetValue(col);
378+ 
379+ // A LOWER diagonal suffix may start at an arbitrary Unified Buffer offset.
380+ // Read only the first partial pipeline segment directly; every later
381+ // 4096-element segment starts at a zero-based aligned offset in xTile.
382+ const uint32_t localStart = row - tileRow;
383+ if ((localStart & (kPipelineFloats - 1U)) != 0U) {
384+ uint32_t alignedRow = (row + kPipelineFloats - 1U) & ~(kPipelineFloats - 1U);
385+ if (alignedRow > segmentEnd) alignedRow = segmentEnd;
386+ const uint32_t count = alignedRow - row;
387+ const uint64_t apOffset = apBase + row - rowBegin;
388+ const DataCopyExtParams copy{1, count * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
389+ LocalTensor<float> xIn = xQueue.AllocTensor<float>();
390+ LocalTensor<float> apIn = apQueue.AllocTensor<float>();
391+ DataCopyPad(xIn, xGm[row], copy, pad);
392+ DataCopyPad(apIn, apGm[apOffset], copy, pad);
393+ xQueue.EnQue(xIn);
394+ apQueue.EnQue(apIn);
395+ xIn = xQueue.DeQue<float>();
396+ apIn = apQueue.DeQue<float>();
397+ LocalTensor<float> out = outQueue.AllocTensor<float>();
398+ Muls(out, xIn, axCol, count);
399+ Add(out, out, apIn, count);
400+ outQueue.EnQue(out);
401+ xQueue.FreeTensor(xIn);
402+ apQueue.FreeTensor(apIn);
403+ out = outQueue.DeQue<float>();
404+ DataCopyPad(apGm[apOffset], out, copy);
405+ outQueue.FreeTensor(out);
406+ row = alignedRow;
407+ }
408+ 
409+ if (row < segmentEnd) {
410+ uint32_t currentCount = (segmentEnd - row < kPipelineFloats)
411+ ? segmentEnd - row : kPipelineFloats;
412+ uint64_t currentOffset = apBase + row - rowBegin;
413+ DataCopyExtParams currentCopy{
414+ 1, currentCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
415+ LocalTensor<float> firstAp = apQueue.AllocTensor<float>();
416+ DataCopyPad(firstAp, apGm[currentOffset], currentCopy, pad);
417+ apQueue.EnQue(firstAp);
418+ 
419+ while (row < segmentEnd) {
420+ const uint32_t nextRow = row + currentCount;
421+ if (nextRow < segmentEnd) {
422+ const uint32_t nextCount = (segmentEnd - nextRow < kPipelineFloats)
423+ ? segmentEnd - nextRow : kPipelineFloats;
424+ const uint64_t nextOffset = apBase + nextRow - rowBegin;
425+ const DataCopyExtParams nextCopy{
426+ 1, nextCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
427+ LocalTensor<float> nextAp = apQueue.AllocTensor<float>();
428+ DataCopyPad(nextAp, apGm[nextOffset], nextCopy, pad);
429+ apQueue.EnQue(nextAp);
430+ }
431+ 
432+ LocalTensor<float> apIn = apQueue.DeQue<float>();
433+ LocalTensor<float> out = outQueue.AllocTensor<float>();
434+ Muls(out, xTile[row - tileRow], axCol, currentCount);
435+ Add(out, out, apIn, currentCount);
436+ outQueue.EnQue(out);
437+ apQueue.FreeTensor(apIn);
438+ out = outQueue.DeQue<float>();
439+ DataCopyPad(apGm[currentOffset], out, currentCopy);
440+ outQueue.FreeTensor(out);
441+ 
442+ row = nextRow;
443+ if (row < segmentEnd) {
444+ currentCount = (segmentEnd - row < kPipelineFloats)
445+ ? segmentEnd - row : kPipelineFloats;
446+ currentOffset = apBase + row - rowBegin;
447+ currentCopy = DataCopyExtParams{
448+ 1, currentCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
449+ }
450+ }
451+ }
452+ }
453+ }
454+}
455+ 
456+#endif
457+}
458+ 
459+extern "C" __global__ __aicore__ __vector__ void sspr_kernel(
460+ GM_ADDR x, GM_ADDR ap, const SsprTilingData tiling)
461+{
462+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
463+ 
464+ GlobalTensor<float> xGm;
465+ GlobalTensor<float> apGm;
466+ const uint64_t absIncx = static_cast<uint64_t>(tiling.incx >= 0 ? tiling.incx : -tiling.incx);
467+ xGm.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(x),
468+ static_cast<uint64_t>(tiling.n - 1U) * absIncx + 1ULL);
469+ apGm.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(ap),
470+ static_cast<uint64_t>(tiling.n) * (tiling.n + 1ULL) / 2ULL);
471+ 
472+ const uint32_t block = GetBlockIdx();
473+ const uint32_t blocks = GetBlockNum();
474+ uint32_t colStart;
475+ uint32_t colEnd;
476+ // The cost model is accurate for every n and uplo, so it replaces the previous
477+ // equal-column split unconditionally. The old code only enabled a weighted split for
478+ // strided large-n cases because its element-count model was wrong for short columns.
479+ colStart = WeightedColumnBoundary(block, blocks, tiling.n, tiling.uplo);
480+ colEnd = WeightedColumnBoundary(block + 1U, blocks, tiling.n, tiling.uplo);
481+ if (colStart >= colEnd) return;
482+ 
483+#if __DAV_C220_VEC__
484+ if (tiling.incx == 1) {
485+ UnifiedContiguousColumns(xGm, apGm, tiling, colStart, colEnd);
486+ return;
487+ }
488+ if (tiling.n <= kMaxCanonicalN &&
489+ static_cast<uint32_t>(tiling.incx > 0 ? tiling.incx : -tiling.incx) <= kMaxPackedStride) {
490+ CachedStridedColumns(reinterpret_cast<__gm__ float*>(x), apGm, tiling, colStart, colEnd);
491+ return;
492+ }
493+#endif
494+ ScalarColumns(xGm, apGm, tiling, colStart, colEnd);
495+}
496+ 
497+void sspr_kernel_do(
498+ GM_ADDR x, GM_ADDR ap, const SsprTilingData& tiling, uint32_t numBlocks, void* stream)
499+{
500+ sspr_kernel<<<numBlocks, nullptr, stream>>>(x, ap, tiling);
501+}
Ablas/spr/arch22/sspr_tiling_data.h+32-0
@@ -0,0 +1,32 @@
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+// UB-x path x vector cache size (32 KB = 8192 floats). 32KB leaves 216KB DCache for AP streaming.
16+// Ascend950 profiling shows 32KB optimal across all n values (48KB/64KB both underperform).
17+inline constexpr uint32_t UB_X_FLOATS = 8192;
18+ 
19+// n >= UB_THRESHOLD enables UB-x path. Ascend950 profiling (SPR_PROFILE_REPEAT=5, median):
20+// n<128: DCache natural caching suffices, GM path faster
21+// n>=128: DCache contention becomes significant, UB-x isolation wins
22+// Both UPPER and LOWER share the same threshold.
23+inline constexpr uint32_t UB_THRESHOLD = 128;
24+ 
25+struct SsprTilingData {
26+ uint32_t numThreads; // threads per block
27+ uint32_t columnsPerBlock; // columns per block
28+ uint32_t n; // matrix order
29+ uint32_t uplo; // ACLBLAS_UPPER(121) or ACLBLAS_LOWER(122)
30+ float alpha; // scalar alpha (dereferenced from pointer)
31+ int64_t incx; // x vector stride
32+};
Atest/spr/arch22/sspr_npu_wrapper.h+71-0
@@ -0,0 +1,71 @@
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 <algorithm>
14+#include <cstdint>
15+#include <cstdlib>
16+ 
17+#include "acl/acl.h"
18+#include "cann_ops_blas.h"
19+ 
20+// NPU wrapper for aclblasSspr — same signature as the API.
21+// For n <= 0 or null handle, passes through directly to hit host-side
22+// parameter validation / early-return paths.
23+// Otherwise, allocates device memory, copies H2D, invokes kernel,
24+// synchronises, copies D2H, and frees.
25+inline aclblasStatus_t aclblasSspr_npu(
26+ aclblasHandle_t handle,
27+ aclblasFillMode_t uplo,
28+ int n,
29+ const float* alpha,
30+ const float* x,
31+ int incx,
32+ float* ap)
33+{
34+ if (handle == nullptr || n <= 0) {
35+ return aclblasSspr(handle, uplo, n, alpha, x, incx, ap);
36+ }
37+ 
38+ const int allocN = std::max(1, n);
39+ const int absIncx = std::abs(incx);
40+ const size_t xBytes = static_cast<size_t>((allocN - 1) * absIncx + 1) * sizeof(float);
41+ const size_t apBytes = static_cast<size_t>(allocN) * (allocN + 1) / 2 * sizeof(float);
42+ 
43+ void* dX = nullptr;
44+ void* dAP = nullptr;
45+ aclError aclRet;
46+ 
47+ if (x != nullptr) {
48+ aclRet = aclrtMalloc(&dX, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);
49+ if (aclRet != ACL_SUCCESS) return ACLBLAS_STATUS_ALLOC_FAILED;
50+ aclRet = aclrtMemcpy(dX, xBytes, x, xBytes, ACL_MEMCPY_HOST_TO_DEVICE);
51+ if (aclRet != ACL_SUCCESS) { aclrtFree(dX); return ACLBLAS_STATUS_INTERNAL_ERROR; }
52+ }
53+ if (ap != nullptr) {
54+ aclRet = aclrtMalloc(&dAP, apBytes, ACL_MEM_MALLOC_HUGE_FIRST);
55+ if (aclRet != ACL_SUCCESS) { if (dX) aclrtFree(dX); return ACLBLAS_STATUS_ALLOC_FAILED; }
56+ aclRet = aclrtMemcpy(dAP, apBytes, ap, apBytes, ACL_MEMCPY_HOST_TO_DEVICE);
57+ if (aclRet != ACL_SUCCESS) { if (dX) aclrtFree(dX); aclrtFree(dAP); return ACLBLAS_STATUS_INTERNAL_ERROR; }
58+ }
59+ 
60+ aclblasStatus_t ret = aclblasSspr(handle, uplo, n, alpha,
61+ static_cast<const float*>(dX), incx, static_cast<float*>(dAP));
62+ 
63+ aclrtSynchronizeDevice();
64+ if (ret == ACLBLAS_STATUS_SUCCESS && ap != nullptr) {
65+ aclrtMemcpy(ap, apBytes, dAP, apBytes, ACL_MEMCPY_DEVICE_TO_HOST);
66+ }
67+ 
68+ if (dX) aclrtFree(dX);
69+ if (dAP) aclrtFree(dAP);
70+ return ret;
71+}
Atest/spr/arch22/sspr_test.cpp+83-0
@@ -0,0 +1,83 @@
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 <cmath>
12+#include <vector>
13+ 
14+#include "verify.h"
15+#include "blas_test.h"
16+#include "csv_loader.h"
17+#include "sspr_param.h"
18+#include "sspr_golden.h"
19+#include "sspr_npu_wrapper.h"
20+ 
21+class SsprArch22Test : public BlasTest<SsprParam> { };
22+ 
23+TEST_F(SsprArch22Test, NullHandle)
24+{
25+ float alphaVal = 1.0f;
26+ aclblasStatus_t ret = aclblasSspr(nullptr, ACLBLAS_UPPER, 4, &alphaVal, nullptr, 1, nullptr);
27+ EXPECT_EQ(ret, ACLBLAS_STATUS_HANDLE_IS_NULLPTR);
28+}
29+ 
30+INSTANTIATE_TEST_SUITE_P(
31+ Sspr, SsprArch22Test,
32+ ::testing::ValuesIn(GetCasesFromCsv<SsprParam>(ReplaceFileExtension2Csv(__FILE__))),
33+ PrintCaseInfoString<SsprParam>);
34+ 
35+TEST_P(SsprArch22Test, CsvDriven)
36+{
37+ const auto& p = GetParam();
38+ 
39+ int absIncx = std::abs(p.incx);
40+ size_t xLen = (p.n > 0) ? static_cast<size_t>((p.n - 1) * absIncx + 1) : 0;
atomgit-bot
atomgit-botatomgit-bot8 天前

🟡 Medium Priority

建议：删除未使用的 absIncx/xLen（第 39-40 行），或改为先判断 p.incx == INT_MIN 再计算，并将 (p.n - 1) * absIncx + 1 提升为 64 位运算避免有符号溢出。

likedislike
41+ size_t apLen = (p.n > 0) ? static_cast<size_t>(p.n) * (p.n + 1) / 2 : 0;
42+ 
43+ // Generate x vector (strided, supports negative incx)
44+ std::vector<float> xHost = makeBlasStrided(p.n, p.incx, p.x, p.randomSeed);
45+ 
46+ // Generate packed AP (triangular / random)
47+ // When ap uses UPPER/LOWER pattern in its BlasFillMode, makeBlasPacked preserves it.
48+ // Otherwise, use the uplo to select triangular layout.
49+ bool useUpper = (p.uplo == ACLBLAS_UPPER);
50+ std::vector<float> apHost;
51+ if (p.ap.pattern == BlasFillMode::P_UPPER || p.ap.pattern == BlasFillMode::P_LOWER) {
52+ apHost = makeBlasPacked(p.n, p.ap, p.randomSeed);
53+ } else {
54+ apHost = makeBlasTriangular(p.n, useUpper, p.ap, p.randomSeed);
55+ }
56+ 
57+ // Copy apHost for alpha=0 exact verification
58+ std::vector<float> apOrig = apHost;
59+ 
60+ const float* xPtr = xHost.empty() ? nullptr : xHost.data();
61+ float* apPtr = apHost.empty() ? nullptr : apHost.data();
62+ const float* alphaPtr = p.alphaNull ? nullptr : &p.alpha;
63+ 
64+ aclblasStatus_t ret = aclblasSspr_npu(
65+ SsprArch22Test::handle_, p.uplo, p.n, alphaPtr, xPtr, p.incx, apPtr);
66+ EXPECT_EQ(static_cast<int>(ret), static_cast<int>(p.expectResult));
67+ if (p.expectResult != ACLBLAS_STATUS_SUCCESS) return;
68+ if (p.n == 0) return;
69+ 
70+ // Compute CPU golden
71+ std::vector<float> golden = apOrig;
72+ aclblasSspr_cpu(
73+ SsprArch22Test::handle_, p.uplo, p.n, &p.alpha, xHost.data(), p.incx, golden.data());
74+ 
75+ VerifyConfig cfg;
76+ if (p.alpha == 0.0f && !p.alphaNull) {
77+ cfg.mode = PrecisionMode::EXACT;
78+ EXPECT_TRUE(Verifier::verifyVector(apPtr, golden.data(), apLen, 1, cfg, p.caseName));
79+ } else {
80+ applyMixedTolerance(cfg, ACL_FLOAT, golden.data(), apLen);
81+ EXPECT_TRUE(Verifier::verifyVector(apPtr, golden.data(), apLen, 1, cfg, p.caseName));
82+ }
83+}
Atest/spr/arch22/sspr_test.csv+59-0
@@ -0,0 +1,59 @@
1+case_name,description,uplo,n,alpha,x,incx,ap,expect_result,random_seed
2+TC_L0_02,n_negative,ACLBLAS_UPPER,-5,1.0,NULLPTR,1,NULLPTR,ACLBLAS_STATUS_INVALID_VALUE,42
3+TC_L0_03,n_zero_upper,ACLBLAS_UPPER,0,1.0,NULLPTR,1,NULLPTR,ACLBLAS_STATUS_SUCCESS,42
4+TC_L0_04,uplo_invalid,INVALID,4,1.0,NULLPTR,1,NULLPTR,ACLBLAS_STATUS_INVALID_VALUE,42
5+TC_L0_05,incx_zero,ACLBLAS_UPPER,4,1.0,NULLPTR,0,NULLPTR,ACLBLAS_STATUS_INVALID_VALUE,42
6+TC_L0_06,incx_intmin,ACLBLAS_UPPER,4,1.0,NULLPTR,-2147483648,NULLPTR,ACLBLAS_STATUS_INVALID_VALUE,42
7+TC_L0_07,alpha_null,ACLBLAS_UPPER,4,NULLPTR,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_INVALID_VALUE,42
8+TC_L0_08,x_null,ACLBLAS_UPPER,4,1.0,NULLPTR,1,RANDOM_2_2,ACLBLAS_STATUS_INVALID_VALUE,42
9+TC_L0_09,ap_null,ACLBLAS_UPPER,4,1.0,RANDOM_2_2,1,NULLPTR,ACLBLAS_STATUS_INVALID_VALUE,42
10+TC_L0_10,alpha_zero,ACLBLAS_UPPER,4,0.0,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
11+TC_L0_11,n1_upper,ACLBLAS_UPPER,1,1.0,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
12+TC_L0_12,n1_lower,ACLBLAS_LOWER,1,1.0,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
13+TC_L0_13,n4_upper,ACLBLAS_UPPER,4,1.0,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
14+TC_L0_14,n4_lower,ACLBLAS_LOWER,4,1.0,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
15+TC_L0_15,n4_upper_neg_alpha,ACLBLAS_UPPER,4,-1.0,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
16+TC_L0_16,n4_lower_neg_alpha,ACLBLAS_LOWER,4,-1.0,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
17+TC_L0_17,n4_upper_neg_incx,ACLBLAS_UPPER,4,1.0,RANDOM_2_2,-1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
18+TC_L0_18,n4_lower_neg_incx,ACLBLAS_LOWER,4,1.0,RANDOM_2_2,-1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
19+TC_L0_19,n4_upper_incx2,ACLBLAS_UPPER,4,1.0,RANDOM_2_2,2,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
20+TC_L0_20,n4_lower_incx2,ACLBLAS_LOWER,4,1.0,RANDOM_2_2,2,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
21+TC_L0_21,n64_upper_ubx,ACLBLAS_UPPER,64,0.5,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
22+TC_L0_22,n64_lower_ubx,ACLBLAS_LOWER,64,0.5,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
23+TC_L0_23,n10000_upper_gm_fallback,ACLBLAS_UPPER,10000,0.5,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
24+TC_L0_24,n10000_lower_gm_fallback,ACLBLAS_LOWER,10000,0.5,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
25+TC_L0_25,n10000_upper_incx2,ACLBLAS_UPPER,10000,1.0,RANDOM_2_2,2,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
26+TC_L0_26,n10000_lower_neg_incx,ACLBLAS_LOWER,10000,1.0,RANDOM_2_2,-1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,42
27+TC_A22_001,n7_upper_vec_tail,ACLBLAS_UPPER,7,0.75,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260851
28+TC_A22_002,n8_lower_block,ACLBLAS_LOWER,8,-0.5,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260852
29+TC_A22_003,n9_upper_vec_tail,ACLBLAS_UPPER,9,1.25,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260853
30+TC_A22_004,n39_lower_core_boundary,ACLBLAS_LOWER,39,0.75,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260854
31+TC_A22_005,n40_upper_core_boundary,ACLBLAS_UPPER,40,0.75,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260855
32+TC_A22_006,n41_lower_core_boundary,ACLBLAS_LOWER,41,0.75,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260856
33+TC_A22_007,n63_upper_repeat_tail,ACLBLAS_UPPER,63,0.625,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260857
34+TC_A22_008,n64_lower_repeat,ACLBLAS_LOWER,64,0.625,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260858
35+TC_A22_009,n65_upper_repeat_tail,ACLBLAS_UPPER,65,0.625,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260859
36+TC_A22_010,n257_upper_incx_neg2,ACLBLAS_UPPER,257,1.125,RANDOM_2_2,-2,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260860
37+TC_A22_011,n257_lower_incx_neg2,ACLBLAS_LOWER,257,1.125,RANDOM_2_2,-2,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260861
38+TC_A22_012,n513_upper_incx32,ACLBLAS_UPPER,513,0.875,RANDOM_2_2,32,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260862
39+TC_A22_013,n513_lower_incx_neg32,ACLBLAS_LOWER,513,0.875,RANDOM_2_2,-32,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260863
40+TC_A22_014,n1025_upper_incx128,ACLBLAS_UPPER,1025,-0.625,RANDOM_2_2,128,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260864
41+TC_A22_015,n1025_lower_incx_neg128,ACLBLAS_LOWER,1025,-0.625,RANDOM_2_2,-128,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260865
42+TC_A22_016,n16385_upper_incx16,ACLBLAS_UPPER,16385,0.75,RANDOM_2_2,16,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260866
43+TC_A22_017,n16385_lower_incx_neg16,ACLBLAS_LOWER,16385,0.75,RANDOM_2_2,-16,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260867
44+TC_A22_018,n257_upper_incx3,ACLBLAS_UPPER,257,1.125,RANDOM_2_2,3,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260868
45+TC_A22_019,n257_lower_incx3,ACLBLAS_LOWER,257,1.125,RANDOM_2_2,3,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260869
46+TC_A22_020,n513_upper_incx5,ACLBLAS_UPPER,513,0.875,RANDOM_2_2,5,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260870
47+TC_A22_021,n513_lower_incx_neg5,ACLBLAS_LOWER,513,0.875,RANDOM_2_2,-5,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260871
48+TC_A22_022,n1025_upper_incx7,ACLBLAS_UPPER,1025,-0.625,RANDOM_2_2,7,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260872
49+TC_A22_023,n1025_lower_incx_neg7,ACLBLAS_LOWER,1025,-0.625,RANDOM_2_2,-7,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260873
50+TC_A22_024,n4097_upper_incx9,ACLBLAS_UPPER,4097,0.75,RANDOM_2_2,9,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260874
51+TC_A22_025,n4097_lower_incx_neg9,ACLBLAS_LOWER,4097,0.75,RANDOM_2_2,-9,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260875
52+TC_A22_026,n8193_upper_incx33,ACLBLAS_UPPER,8193,1.375,RANDOM_2_2,33,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260876
53+TC_A22_027,n8193_lower_incx_neg33,ACLBLAS_LOWER,8193,1.375,RANDOM_2_2,-33,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260877
54+TC_UNIFIED_001,n127_upper_threshold_minus_one,ACLBLAS_UPPER,127,0.875,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260878
55+TC_UNIFIED_002,n127_lower_threshold_minus_one,ACLBLAS_LOWER,127,0.875,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260879
56+TC_UNIFIED_003,n128_upper_threshold_exact,ACLBLAS_UPPER,128,0.875,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260880
57+TC_UNIFIED_004,n128_lower_threshold_exact,ACLBLAS_LOWER,128,0.875,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260881
58+TC_UNIFIED_005,n129_upper_threshold_plus_one,ACLBLAS_UPPER,129,0.875,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260882
59+TC_UNIFIED_006,n129_lower_threshold_plus_one,ACLBLAS_LOWER,129,0.875,RANDOM_2_2,1,RANDOM_2_2,ACLBLAS_STATUS_SUCCESS,20260883