Star181
284
代码
代码
Issues73
Pull Requests102
Actions
项目讨论
Wiki
分析
Star181
284
已开启
Feat: 新增面向arch22的aclblasCaxpy任意incx/incy支持 #334
guodong54_创建于 8月13日
Feat: 新增面向arch22的aclblasCaxpy任意incx/incy支持 #334
已开启
guodong54_创建于 8月13日
从guodong54_/ops-blasarch22-caxpy-release合入到cann/ops-blasmaster
共 16 个文件变更+2234-394
Ablas/axpy/arch22/caxpy_dispatch_policy.h+159-0
@@ -0,0 +1,159 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OPS_BLAS_CAXPY_DISPATCH_POLICY_H
12+#define OPS_BLAS_CAXPY_DISPATCH_POLICY_H
13+ 
14+#include <cstdint>
15+ 
16+enum class CaxpyStrideClass : uint32_t {
17+ DENSE = 0,
18+ X_ONLY_POSITIVE = 1,
19+ Y_ONLY_POSITIVE = 2,
20+ DUAL_SYMMETRIC_POSITIVE = 3,
21+ DUAL_ASYMMETRIC_POSITIVE = 4,
22+ NEGATIVE_MIXED = 5,
23+};
24+ 
25+enum class CaxpyKernelVariant : uint32_t {
26+ DENSE_PIPELINED_ATOMIC = 0,
27+ STRIDED_SCALAR = 1,
28+ STRIDED_SHARED_PIPELINED_ATOMIC = 2,
29+ STRIDED_DISJOINT_PIPELINED_ATOMIC = 3,
30+ STRIDED_SHARED_PIPELINED_SCALAR_TAIL = 4,
31+ STRIDED_DISJOINT_PIPELINED_SCALAR_TAIL = 5,
32+};
33+ 
34+constexpr uint64_t CAXPY_STRIDED_SCALAR_MAX_PER_BLOCK = 16;
35+constexpr uint64_t CAXPY_STRIDED_MAX_PER_BLOCK_PER_WAVE = 8192;
36+// Fixed tile capacity used by the disjoint pipeline and as the shared
37+// pipeline's UB buffer allocation size (see CAXPY_STRIDED_TILE_COUNT in
38+// caxpy_kernel.cpp, which must stay equal to this value).
39+constexpr uint64_t CAXPY_STRIDED_TILE_COUNT = 512;
40+// The shared pipeline's lowScratch UB scratch buffer used for the raw span
41+// read is 38 KiB (CAXPY_DENSE_TILE_BYTES in caxpy_kernel.cpp). A span read of
42+// `tile` complex elements at increment `incx` covers
43+// ((tile-1)*|incx|+1) complex elements == ((tile-1)*|incx|+1)*8 bytes of
44+// physical address range, which must fit in that buffer. Solving for the
45+// largest tile <= CAXPY_STRIDED_TILE_COUNT that satisfies this bound gives
46+// the runtime tile capacity the shared pipeline should use for a given incx;
47+// this generalizes the old fixed CAXPY_SHARED_X_SPAN_MAX_INCREMENT=9 cutoff
48+// (which was really just "does tile=512 fit"), letting larger |incx| still
49+// use the faster span-read shared pipeline instead of falling back to
50+// disjoint sparse reads, as long as the shrunk tile stays usable.
51+constexpr uint64_t CAXPY_SHARED_LOW_SCRATCH_BYTES = 38 * 1024;
52+constexpr uint64_t CAXPY_COMPLEX_BYTES = 8;
53+ 
54+// Largest tile (in complex elements, capped at CAXPY_STRIDED_TILE_COUNT) whose
55+// span read at the given |incx| fits in the shared pipeline's lowScratch
56+// buffer. absIncx must be >= 1 (callers only invoke this for direct-strided,
57+// non-dense cases where incx != 0).
58+constexpr uint64_t CaxpySharedTileCapacity(uint64_t absIncx)
59+{
60+ const uint64_t maxSpanComplexElements = CAXPY_SHARED_LOW_SCRATCH_BYTES / CAXPY_COMPLEX_BYTES;
61+ // tile such that (tile-1)*absIncx + 1 <= maxSpanComplexElements
62+ // => tile <= (maxSpanComplexElements - 1) / absIncx + 1
63+ const uint64_t maxTileForIncrement = (maxSpanComplexElements - 1) / absIncx + 1;
64+ return maxTileForIncrement < CAXPY_STRIDED_TILE_COUNT ? maxTileForIncrement : CAXPY_STRIDED_TILE_COUNT;
65+}
66+ 
67+// The shared pipeline's tail-splitting logic (shared_tile_bounds in
68+// caxpy_kernel.cpp) rebalances the last two tiles when the natural final tile
69+// would be shorter than CAXPY_STRIDED_SCALAR_MAX_PER_BLOCK+1
70+// (CAXPY_INTERLEAVE_COUNT+1 in caxpy_kernel.cpp, same constant) elements, by
71+// splitting their combined length roughly in half. For that split to never
72+// itself produce a half shorter than the interleave width, the worst case
73+// (combined length == tileCapacity+1, the smallest total that still triggers
74+// the split branch) must divide into two halves each >=
75+// CAXPY_STRIDED_SCALAR_MAX_PER_BLOCK+1: ceil((tileCapacity+1)/2) and
76+// floor((tileCapacity+1)/2) both >= CAXPY_STRIDED_SCALAR_MAX_PER_BLOCK+1,
77+// which holds starting at tileCapacity == 2*(CAXPY_STRIDED_SCALAR_MAX_PER_BLOCK+1)-1.
78+// A smaller usable-capacity threshold (originally CAXPY_STRIDED_SCALAR_MAX_PER_BLOCK+1
79+// itself) let the split produce a too-short half, caught by the
80+// dynamic_shared_boundary_usable regression test (n=8000, incx=303, tile=17)
81+// only in a Release build where the resulting silent wrong compute went
82+// uncaught -- a shrunk-tile-capacity variant of the same class of bug as the
83+// earlier fixed-tile computeOffset drift.
84+constexpr uint64_t CAXPY_SHARED_MIN_USABLE_TILE_CAPACITY = 2 * (CAXPY_STRIDED_SCALAR_MAX_PER_BLOCK + 1) - 1;
85+ 
86+// CAXPY_SHARED_MIN_USABLE_TILE_CAPACITY only bounds *correctness*: below it
87+// the tail-split can produce too-short tiles. It says nothing about
88+// performance. As |incx| grows, the shrunk tile's span read
89+// ((tile-1)*|incx|+1 complex elements) carries a shrinking fraction of useful
90+// data -- at incx=16 the span is ~6% useful, at incx=100 it is ~1% -- so GM
91+// bandwidth spent reading skipped-over elements eventually exceeds what the
92+// disjoint pipeline's per-element sparse reads cost instead. Measured on
93+// Ascend 910B3 / CANN 9.0 A3, n=1048576 (xlarge, enough tiles per core to
94+// reach steady state and not be dominated by per-tile dispatch overhead):
95+// shared span-read is faster than disjoint through incx<=60, comparable
96+// around incx=60-65, and progressively worse beyond that (roughly +7% at
97+// incx=70, +97% at incx=150 in one measured run). This cutoff is therefore a
98+// measured performance crossover, not a UB-buffer-fit or correctness limit;
99+// re-measure if the UB budget, tile capacity formula, or disjoint pipeline's
100+// per-element cost changes.
101+constexpr uint64_t CAXPY_SHARED_SPAN_PERFORMANCE_MAX_INCREMENT = 60;
102+ 
103+constexpr uint64_t CaxpyFinalWaveCount(uint64_t n, uint32_t numBlocks)
104+{
105+ const uint64_t waveCapacity = static_cast<uint64_t>(numBlocks) * CAXPY_STRIDED_MAX_PER_BLOCK_PER_WAVE;
106+ if (n <= waveCapacity)
107+ return n;
108+ const uint64_t remainder = n % waveCapacity;
109+ return remainder == 0 ? waveCapacity : remainder;
110+}
111+ 
112+// Preconditions are enforced by aclblasCaxpy argument validation and execution
113+// setup: incx and incy are non-zero and are not INT64_MIN; numBlocks is non-zero.
114+constexpr CaxpyStrideClass ClassifyCaxpyStride(int64_t incx, int64_t incy)
115+{
116+ if (incx < 0 || incy < 0)
117+ return CaxpyStrideClass::NEGATIVE_MIXED;
118+ if (incx == 1 && incy == 1)
119+ return CaxpyStrideClass::DENSE;
120+ if (incy == 1)
121+ return CaxpyStrideClass::X_ONLY_POSITIVE;
122+ if (incx == 1)
123+ return CaxpyStrideClass::Y_ONLY_POSITIVE;
124+ if (incx == incy)
125+ return CaxpyStrideClass::DUAL_SYMMETRIC_POSITIVE;
126+ return CaxpyStrideClass::DUAL_ASYMMETRIC_POSITIVE;
127+}
128+ 
129+constexpr uint64_t CaxpyAbsIncrement(int64_t increment)
130+{
131+ return increment > 0 ? static_cast<uint64_t>(increment) : static_cast<uint64_t>(-(increment + 1)) + 1;
132+}
133+ 
134+constexpr bool UseCaxpySharedXSpanPipeline(int64_t incx)
135+{
136+ const uint64_t absIncx = CaxpyAbsIncrement(incx);
137+ return absIncx <= CAXPY_SHARED_SPAN_PERFORMANCE_MAX_INCREMENT &&
138+ CaxpySharedTileCapacity(absIncx) >= CAXPY_SHARED_MIN_USABLE_TILE_CAPACITY;
139+}
140+ 
141+constexpr CaxpyKernelVariant SelectCaxpyKernelVariant(
142+ uint64_t n, int64_t incx, int64_t incy, bool directStrided, uint32_t numBlocks)
143+{
144+ const CaxpyStrideClass strideClass = ClassifyCaxpyStride(incx, incy);
145+ if (strideClass == CaxpyStrideClass::DENSE || !directStrided)
146+ return CaxpyKernelVariant::DENSE_PIPELINED_ATOMIC;
147+ const bool sharedXSpan = UseCaxpySharedXSpanPipeline(incx);
148+ const uint64_t finalWaveCount = CaxpyFinalWaveCount(n, numBlocks);
149+ const uint64_t finalWaveBlocks = finalWaveCount < numBlocks ? finalWaveCount : numBlocks;
150+ const uint64_t minBlockCount = finalWaveCount / finalWaveBlocks;
151+ if (minBlockCount <= CAXPY_STRIDED_SCALAR_MAX_PER_BLOCK)
152+ return n == finalWaveCount ? CaxpyKernelVariant::STRIDED_SCALAR :
153+ (sharedXSpan ? CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_SCALAR_TAIL :
154+ CaxpyKernelVariant::STRIDED_DISJOINT_PIPELINED_SCALAR_TAIL);
155+ return sharedXSpan ? CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC :
156+ CaxpyKernelVariant::STRIDED_DISJOINT_PIPELINED_ATOMIC;
157+}
158+ 
159+#endif // OPS_BLAS_CAXPY_DISPATCH_POLICY_H
Mblas/axpy/arch22/caxpy_host.cpp+510-106
@@ -14,140 +14,544 @@
14 */14 */
15 15 
16#include <cstdint>16#include <cstdint>
17+#include <cstddef>
18+#include <limits>
17#include <iostream>19#include <iostream>
18#include <vector>20#include <vector>
19#include <algorithm>21#include <algorithm>
20#include <iterator>22#include <iterator>
21#include "acl/acl.h"23#include "acl/acl.h"
24+#include "acl/acl_rt_api.h"
22#include "cann_ops_blas.h"25#include "cann_ops_blas.h"
23#include "common/helper/aclblas_handle_internal.h"26#include "common/helper/aclblas_handle_internal.h"
MM
MMaskHunter200817 小时前

[Major] Missing README update

问题

新增算子实现但未发现 README.md 变更。README 应包含：函数原型、参数表、约束说明、产品支持表、调用示例。

依据

  • interface-spec.md §5.2 要求：新算子必须更新 README

修改建议

在算子目录下的 README.md 中添加算子文档，包括函数原型、参数说明表、约束条件、产品支持表和调用示例代码。

likedislike
MMaskHunter200817 小时前

[Major] Missing public header declaration

问题

新增算子但公共头文件未修改。算子必须在 extern "C" 块中声明。

依据

  • interface-spec.md §5.1 要求：新算子必须在公共头文件中声明

修改建议

在公共头文件的 extern "C" 块中添加算子声明，与同类算子分组排列。

likedislike
MMaskHunter200816 小时前

[Major] Missing README update

问题

新增算子实现但未发现 README.md 变更。README 应包含：函数原型、参数表、约束说明、产品支持表、调用示例。

依据

  • interface-spec.md §5.2 要求：新算子必须更新 README

修改建议

在算子目录下的 README.md 中添加算子文档，包括函数原型、参数说明表、约束条件、产品支持表和调用示例代码。

likedislike
MMaskHunter200816 小时前

[Major] Missing public header declaration

问题

新增算子但公共头文件未修改。算子必须在 extern "C" 块中声明。

依据

  • interface-spec.md §5.1 要求：新算子必须在公共头文件中声明

修改建议

在公共头文件的 extern "C" 块中添加算子声明，与同类算子分组排列。

likedislike
24#include "common/helper/host_utils.h"27#include "common/helper/host_utils.h"
28+#include "caxpy_dispatch_policy.h"
29+#include "caxpy_kernel.h"
30+#include "caxpy_tiling_data.h"
25 31 
26-void caxpy_kernel_do(uint8_t* x, uint8_t* maskBuf, uint8_t* y, uint8_t* workSpace, uint8_t* tilingGm,32+constexpr uint32_t FLOATS_PER_COMPLEX = 2;
27- uint32_t numBlocks, void *stream);
28- 
29-constexpr uint32_t COMPLEX_NUM = 2;
30constexpr uint32_t K_FACTOR_4 = 4;33constexpr uint32_t K_FACTOR_4 = 4;
31constexpr uint32_t DEFAULT_VECTOR_NUM = 40;34constexpr uint32_t DEFAULT_VECTOR_NUM = 40;
32constexpr uint32_t MAX_DATA_COUNT = 38 * 1024 / sizeof(float);35constexpr uint32_t MAX_DATA_COUNT = 38 * 1024 / sizeof(float);
36+constexpr uint32_t STRIDED_TILE_COMPLEX_COUNT = 512;
37+constexpr uint32_t BLOCKED_FLOATS_PER_COMPLEX = 8;
38+constexpr uint32_t GATHER_PADDING_COUNT = 64;
39+constexpr uint32_t X_NORMALIZE_OFFSET_COUNT =
40+ STRIDED_TILE_COMPLEX_COUNT * BLOCKED_FLOATS_PER_COMPLEX + GATHER_PADDING_COUNT;
41+constexpr uint32_t STORE_OFFSET_COUNT = STRIDED_TILE_COMPLEX_COUNT * FLOATS_PER_COMPLEX + GATHER_PADDING_COUNT;
42+constexpr uint32_t X_NORMALIZE_OFFSET_START = MAX_DATA_COUNT;
43+constexpr uint32_t STORE_OFFSET_START = X_NORMALIZE_OFFSET_START + X_NORMALIZE_OFFSET_COUNT;
44+constexpr uint32_t Y_GATHER_OFFSET_COUNT = STRIDED_TILE_COMPLEX_COUNT * FLOATS_PER_COMPLEX + GATHER_PADDING_COUNT;
45+constexpr uint32_t Y_GATHER_OFFSET_START = STORE_OFFSET_START + STORE_OFFSET_COUNT;
46+constexpr uint32_t DENSE_SWAP_OFFSET_START = Y_GATHER_OFFSET_START + Y_GATHER_OFFSET_COUNT;
47+constexpr uint32_t MASK_DATA_COUNT = DENSE_SWAP_OFFSET_START + MAX_DATA_COUNT;
48+static_assert(STORE_OFFSET_START + STORE_OFFSET_COUNT <= MASK_DATA_COUNT);
49+static_assert(Y_GATHER_OFFSET_START + Y_GATHER_OFFSET_COUNT <= MASK_DATA_COUNT);
50+static_assert(DENSE_SWAP_OFFSET_START + MAX_DATA_COUNT <= MASK_DATA_COUNT);
33 51 
34-struct CaxpyTilingData {52+// CaxpyTilingData is defined once in the shared header and included by both the
35- uint32_t n;53+// host and the device kernel; see caxpy_tiling_data.h.
36- float alphaReal;
37- float alphaImag;
38- uint32_t startOffset[40];
39- uint32_t calNum[40];
40-};
41 54 
42-CaxpyTilingData CalTilingData(uint32_t n, uint32_t vecCoreNum, float alphaReal, float alphaImag)55+CaxpyTilingData CalTilingData(
56+ uint32_t n, uint32_t vecCoreNum, float alphaReal, float alphaImag, int64_t incx, int64_t incy,
57+ uint32_t logicalBase = 0, uint32_t totalN = 0)
43{58{
44- CaxpyTilingData tilingData;59+ CaxpyTilingData tilingData{};
45- tilingData.n = n * COMPLEX_NUM;
46- tilingData.alphaReal = alphaReal;
47- tilingData.alphaImag = alphaImag;
48- 
49 if (vecCoreNum == 0) {60 if (vecCoreNum == 0) {
50 vecCoreNum = 1;61 vecCoreNum = 1;
51 }62 }
52 vecCoreNum = vecCoreNum > DEFAULT_VECTOR_NUM ? DEFAULT_VECTOR_NUM : vecCoreNum;63 vecCoreNum = vecCoreNum > DEFAULT_VECTOR_NUM ? DEFAULT_VECTOR_NUM : vecCoreNum;
53- 64+ tilingData.totalN = totalN == 0 ? n : totalN;
54- for (uint32_t i = 0; i < vecCoreNum; i++) {65+ tilingData.activeBlocks = vecCoreNum;
55- tilingData.startOffset[i] = 0;66+ tilingData.alphaReal = alphaReal;
56- tilingData.calNum[i] = 0;67+ tilingData.alphaImag = alphaImag;
57- }68+ tilingData.incx = incx;
58- 69+ tilingData.incy = incy;
59- uint32_t rowNumEachCore = n / vecCoreNum;70+ (void)logicalBase;
60- uint32_t remainRowNum = n % vecCoreNum;
61- 
62- if (rowNumEachCore == 0) {
63- for (uint32_t i = 0; i < remainRowNum; i++) {
64- tilingData.calNum[i] = COMPLEX_NUM;
65- tilingData.startOffset[i] = i * COMPLEX_NUM;
66- }
67- } else {
68- uint32_t currOffset = 0;
69- uint32_t currNum;
70- for (uint32_t i = 0; i < vecCoreNum; i++) {
71- if (i < remainRowNum) {
72- currNum = rowNumEachCore + 1;
73- } else {
74- currNum = rowNumEachCore;
75- }
76- tilingData.calNum[i] = currNum * COMPLEX_NUM;
77- tilingData.startOffset[i] = currOffset;
78- currOffset += currNum * COMPLEX_NUM;
79- }
80- }
81- 
82 return tilingData;71 return tilingData;
83}72}
84 73 
85-void GenMaskData(uint32_t* maskData)74+// Magnitude of a BLAS increment. Computed as -(inc+1)+1 rather than -inc so
75+// that INT64_MIN would not overflow -- though callers reject INT64_MIN upstream
76+// (ValidateCaxpyArguments), this keeps the helper correct in isolation.
77+uint64_t AbsIncrement(int64_t increment)
86{78{
87- uint32_t offsetNum = COMPLEX_NUM;79+ return increment > 0 ? static_cast<uint64_t>(increment) : static_cast<uint64_t>(-(increment + 1)) + 1;
88- uint32_t complexCount = MAX_DATA_COUNT / COMPLEX_NUM;80+}
81+ 
82+bool CheckedVectorSpan(int64_t n, int64_t increment, uint64_t& span)
83+{
84+ const uint64_t absIncrement = AbsIncrement(increment);
85+ if (n <= 1) {
86+ span = 1;
87+ return true;
88+ }
89+ const uint64_t logicalDistance = static_cast<uint64_t>(n - 1);
90+ if (logicalDistance > (std::numeric_limits<uint64_t>::max() - 1) / absIncrement) {
91+ return false;
92+ }
93+ span = 1 + logicalDistance * absIncrement;
94+ return span <= std::numeric_limits<size_t>::max() / sizeof(aclblasComplex);
95+}
96+ 
97+// Maps a logical element index (0..n-1) to its physical index in memory. For a
98+// negative increment BLAS walks the vector backwards, so logical 0 sits at the
99+// far end.
100+uint64_t PhysicalIndex(uint64_t logicalIndex, uint64_t n, int64_t increment)
101+{
102+ const uint64_t absIncrement = AbsIncrement(increment);
103+ return increment > 0 ? logicalIndex * absIncrement : (n - 1 - logicalIndex) * absIncrement;
104+}
105+ 
106+// The mask buffer holds precomputed gather/scatter offset tables so the device
107+// kernel never has to derive strided addresses on its hot path. Each Gen*Offsets
108+// routine fills one section (see the *_OFFSET_START layout constants above).
109+ 
110+// Interleaved -> planar: real parts land in the first half, imaginary in the
111+// second, so downstream vector ops see two contiguous real-valued lanes.
112+void GenComputeOffsets(uint32_t* offsets)
113+{
114+ const uint32_t complexCount = MAX_DATA_COUNT / FLOATS_PER_COMPLEX;
89 for (uint32_t i = 0; i < complexCount; i++) {115 for (uint32_t i = 0; i < complexCount; i++) {
90- maskData[offsetNum * i] = K_FACTOR_4 * i;116+ offsets[FLOATS_PER_COMPLEX * i] = K_FACTOR_4 * i;
91- maskData[offsetNum * i + 1] = K_FACTOR_4 * (i + complexCount);117+ offsets[FLOATS_PER_COMPLEX * i + 1] = K_FACTOR_4 * (i + complexCount);
92 }118 }
93}119}
94 120 
95-aclblasStatus_t aclblasCaxpy(121+// Swaps real/imaginary byte positions within each complex element in place.
96- aclblasHandle_t handle, int n, const aclblasComplex* alpha, const aclblasComplex* x, int incx,122+void GenDenseSwapOffsets(uint32_t* offsets)
97- aclblasComplex* y, int incy)
98{123{
99- if (alpha == nullptr) {124+ const uint32_t complexCount = MAX_DATA_COUNT / FLOATS_PER_COMPLEX;
125+ for (uint32_t i = 0; i < complexCount; ++i) {
126+ offsets[FLOATS_PER_COMPLEX * i] = (FLOATS_PER_COMPLEX * i + 1) * sizeof(float);
127+ offsets[FLOATS_PER_COMPLEX * i + 1] = FLOATS_PER_COMPLEX * i * sizeof(float);
128+ }
129+}
130+ 
131+// Gathers strided x into the blocked (8-float-per-complex) working layout. For
132+// increments where the shared pipeline's span read still fits its lowScratch
133+// UB buffer (UseCaxpySharedXSpanPipeline), the span read brings the data in
134+// contiguously and only the blocked layout stride remains; larger increments
135+// fall back to the disjoint pipeline, which reads each element individually
136+// and does not use this span-read addressing at all.
137+void GenXNormalizeOffsets(uint32_t* offsets, uint32_t xIncrement)
138+{
139+ for (uint32_t i = 0; i < STRIDED_TILE_COMPLEX_COUNT; ++i) {
140+ const uint32_t sourceByteOffset =
141+ UseCaxpySharedXSpanPipeline(static_cast<int64_t>(xIncrement)) ?
142+ i * xIncrement * FLOATS_PER_COMPLEX * sizeof(float) :
143+ i * BLOCKED_FLOATS_PER_COMPLEX * sizeof(float);
144+ for (uint32_t lane = 0; lane < BLOCKED_FLOATS_PER_COMPLEX; ++lane) {
145+ offsets[i * BLOCKED_FLOATS_PER_COMPLEX + lane] =
146+ lane == 1 ? sourceByteOffset + sizeof(float) : sourceByteOffset;
147+ }
148+ }
149+}
150+ 
151+// Scatters results from the blocked layout back to packed interleaved complex.
152+void GenStoreOffsets(uint32_t* offsets)
153+{
154+ for (uint32_t i = 0; i < STRIDED_TILE_COMPLEX_COUNT; ++i) {
155+ const uint32_t sourceByteOffset = i * BLOCKED_FLOATS_PER_COMPLEX * sizeof(float);
156+ offsets[i * FLOATS_PER_COMPLEX] = sourceByteOffset;
157+ offsets[i * FLOATS_PER_COMPLEX + 1] = sourceByteOffset + sizeof(float);
158+ }
159+}
160+ 
161+// Gathers strided y (real, imag) pairs for the accumulation step.
162+void GenYGatherOffsets(uint32_t* offsets, uint32_t yIncrement)
163+{
164+ for (uint32_t i = 0; i < STRIDED_TILE_COMPLEX_COUNT; ++i) {
165+ const uint32_t sourceByteOffset = i * yIncrement * FLOATS_PER_COMPLEX * sizeof(float);
166+ offsets[i * FLOATS_PER_COMPLEX] = sourceByteOffset;
167+ offsets[i * FLOATS_PER_COMPLEX + 1] = sourceByteOffset + sizeof(float);
168+ }
169+}
170+ 
171+void GenMaskData(uint32_t* maskData, uint32_t xIncrement, uint32_t yIncrement)
172+{
173+ std::fill(maskData, maskData + MASK_DATA_COUNT, 0);
174+ GenComputeOffsets(maskData);
175+ GenXNormalizeOffsets(maskData + X_NORMALIZE_OFFSET_START, xIncrement);
176+ GenStoreOffsets(maskData + STORE_OFFSET_START);
177+ GenYGatherOffsets(maskData + Y_GATHER_OFFSET_START, yIncrement);
178+ GenDenseSwapOffsets(maskData + DENSE_SWAP_OFFSET_START);
179+}
180+ 
181+aclblasStatus_t GetCaxpyMaskCache(aclblasHandle_t handle, uint64_t absIncx, uint64_t absIncy, uint8_t** maskDevice)
182+{
183+ if (handle->caxpy_mask_cache == nullptr) {
184+ void* buffer = nullptr;
185+ const aclError aclRet = aclrtMalloc(&buffer, MASK_DATA_COUNT * sizeof(uint32_t), ACL_MEM_MALLOC_HUGE_FIRST);
186+ CHECK_RET(
187+ aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet);
188+ return ACLBLAS_STATUS_ALLOC_FAILED);
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
189+ handle->caxpy_mask_cache = buffer;
190+ }
191+ if (handle->caxpy_mask_cache_incx != absIncx || handle->caxpy_mask_cache_incy != absIncy) {
192+ // Kernels from earlier calls may still be reading the cached mask; drain the
193+ // stream before overwriting it in place.
194+ const aclblasStatus_t syncStatus = SynchronizeHandleStream(handle);
195+ CHECK_RET(syncStatus == ACLBLAS_STATUS_SUCCESS, return syncStatus);
196+ std::vector<uint32_t> maskHost(MASK_DATA_COUNT);
197+ GenMaskData(maskHost.data(), static_cast<uint32_t>(absIncx), static_cast<uint32_t>(absIncy));
198+ const aclError aclRet = aclrtMemcpy(
199+ handle->caxpy_mask_cache, MASK_DATA_COUNT * sizeof(uint32_t), maskHost.data(),
200+ MASK_DATA_COUNT * sizeof(uint32_t), ACL_MEMCPY_HOST_TO_DEVICE);
201+ CHECK_RET(
202+ aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", aclRet);
203+ return ACLBLAS_STATUS_INTERNAL_ERROR);
204+ handle->caxpy_mask_cache_incx = absIncx;
205+ handle->caxpy_mask_cache_incy = absIncy;
206+ }
207+ *maskDevice = reinterpret_cast<uint8_t*>(handle->caxpy_mask_cache);
208+ return ACLBLAS_STATUS_SUCCESS;
209+}
210+ 
211+struct CaxpyExecution {
212+ int64_t n;
213+ int64_t incx;
214+ int64_t incy;
215+ uint32_t numBlocks;
216+ bool useStridedKernel;
217+ bool packX;
218+ bool packY;
219+ CaxpyKernelVariant kernelVariant;
220+ aclrtStream stream;
221+ uint8_t* tilingDevice;
222+ uint8_t* maskDevice;
223+ CaxpyTilingData tiling;
224+};
225+ 
226+struct CaxpyPackedBuffers {
227+ uint8_t* x = nullptr;
228+ uint8_t* y = nullptr;
229+};
230+ 
231+aclblasStatus_t ValidateCaxpyArguments(
232+ aclblasHandle_t handle, int64_t n, const aclblasComplex* alpha, const aclblasComplex* x, int64_t incx,
233+ const aclblasComplex* y, int64_t incy)
234+{
235+ if (handle == nullptr) {
236+ OP_LOGE("aclblasCaxpy", "handle is nullptr");
237+ return ACLBLAS_STATUS_HANDLE_IS_NULLPTR;
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
238+ }
239+ if (n < 0) {
240+ OP_LOGE("aclblasCaxpy", "n must be >= 0, got %lld", static_cast<long long>(n));
241+ return ACLBLAS_STATUS_INVALID_VALUE;
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
242+ }
243+ if (alpha == nullptr) {
244+ OP_LOGE("aclblasCaxpy", "alpha must not be nullptr");
245+ return ACLBLAS_STATUS_INVALID_VALUE;
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
246+ }
247+ if (incx == 0 || incy == 0) {
248+ OP_LOGE("aclblasCaxpy", "incx and incy must not be zero");
249+ return ACLBLAS_STATUS_INVALID_VALUE;
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
250+ }
251+ if (incx == std::numeric_limits<int64_t>::min() || incy == std::numeric_limits<int64_t>::min()) {
252+ OP_LOGE("aclblasCaxpy", "incx and incy must not be INT64_MIN");
253+ return ACLBLAS_STATUS_INVALID_VALUE;
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
254+ }
255+ if (n > 0 && (x == nullptr || y == nullptr)) {
256+ OP_LOGE("aclblasCaxpy", "x and y must not be nullptr when n > 0");
257+ return ACLBLAS_STATUS_INVALID_VALUE;
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
258+ }
259+ if (n > std::numeric_limits<uint32_t>::max() / FLOATS_PER_COMPLEX) {
260+ OP_LOGE("aclblasCaxpy", "n is too large: %lld", static_cast<long long>(n));
261+ return ACLBLAS_STATUS_INVALID_VALUE;
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
262+ }
263+ uint64_t span = 0;
264+ if (n > 0 && (!CheckedVectorSpan(n, incx, span) || !CheckedVectorSpan(n, incy, span))) {
265+ OP_LOGE("aclblasCaxpy", "vector span overflows addressable memory");
100 return ACLBLAS_STATUS_INVALID_VALUE;266 return ACLBLAS_STATUS_INVALID_VALUE;
101 }267 }
102- auto* h = handle;
103- aclrtStream useStream = h->stream;
104- 
105- uint32_t numBlocks = DEFAULT_VECTOR_NUM;
106- 
107- float alphaReal = alpha->real;
108- float alphaImag = alpha->imag;
109- 
110- CaxpyTilingData tiling = CalTilingData(static_cast<uint32_t>(n), numBlocks, alphaReal, alphaImag);
111- 
112- uint32_t maskSize = MAX_DATA_COUNT * sizeof(uint32_t) * COMPLEX_NUM;
113- 
114- uint8_t* maskDevice = nullptr;
115- uint8_t* tilingDevice = nullptr;
116- 
117- uint32_t* maskHost = new uint32_t[MAX_DATA_COUNT * COMPLEX_NUM];
118- GenMaskData(maskHost);
119- 
120- aclError aclRet = aclrtMalloc((void**)&maskDevice, maskSize, ACL_MEM_MALLOC_HUGE_FIRST);
121- CHECK_RET(
122- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet); delete[] maskHost;
123- return ACLBLAS_STATUS_ALLOC_FAILED);
124- 
125- aclRet = aclrtMalloc((void**)&tilingDevice, sizeof(CaxpyTilingData), ACL_MEM_MALLOC_HUGE_FIRST);
126- CHECK_RET(
127- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet); aclrtFree(maskDevice);
128- delete[] maskHost; return ACLBLAS_STATUS_ALLOC_FAILED);
129- 
130- aclRet = aclrtMemcpy(maskDevice, maskSize, maskHost, maskSize, ACL_MEMCPY_HOST_TO_DEVICE);
131- CHECK_RET(
132- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", aclRet); aclrtFree(tilingDevice);
133- aclrtFree(maskDevice); delete[] maskHost; return ACLBLAS_STATUS_INTERNAL_ERROR);
134- 
135- aclRet =
136- aclrtMemcpy(tilingDevice, sizeof(CaxpyTilingData), &tiling, sizeof(CaxpyTilingData), ACL_MEMCPY_HOST_TO_DEVICE);
137- CHECK_RET(
138- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", aclRet); aclrtFree(tilingDevice);
139- aclrtFree(maskDevice); delete[] maskHost; return ACLBLAS_STATUS_INTERNAL_ERROR);
140- 
141- caxpy_kernel_do(reinterpret_cast<uint8_t*>(const_cast<aclblasComplex*>(x)), maskDevice, reinterpret_cast<uint8_t*>(y), nullptr,
142- tilingDevice, numBlocks, useStream);
143- aclRet = aclrtSynchronizeStream(useStream);
144- CHECK_RET(
145- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", aclRet); aclrtFree(tilingDevice);
146- aclrtFree(maskDevice); delete[] maskHost; return ACLBLAS_STATUS_INTERNAL_ERROR);
147- 
148- aclrtFree(maskDevice);
149- aclrtFree(tilingDevice);
150- delete[] maskHost;
151- 
152 return ACLBLAS_STATUS_SUCCESS;268 return ACLBLAS_STATUS_SUCCESS;
153-}269+}
270+ 
271+aclblasStatus_t BuildCaxpyExecution(
272+ aclblasHandle_t handle, int64_t n, const aclblasComplex& alpha, int64_t incx, int64_t incy,
273+ CaxpyExecution& execution)
274+{
275+ uint32_t numBlocks = GetAivCoreCount();
276+ CHECK_RET(numBlocks != 0, LOG_PRINT("GetAivCoreCount failed.\n"); return ACLBLAS_STATUS_EXECUTION_FAILED);
277+ numBlocks = std::min(numBlocks, DEFAULT_VECTOR_NUM);
278+ numBlocks = std::min(numBlocks, static_cast<uint32_t>(n));
279+ const uint64_t absIncx = AbsIncrement(incx);
280+ const uint64_t absIncy = AbsIncrement(incy);
281+ const uint64_t maxIncrement = std::max(absIncx, absIncy);
282+ const bool incrementFits = maxIncrement <= std::numeric_limits<uint32_t>::max() / sizeof(aclblasComplex);
283+ const bool addressFits =
284+ n <= 1 || (incrementFits && static_cast<uint64_t>(n - 1) <=
285+ std::numeric_limits<uint32_t>::max() / (FLOATS_PER_COMPLEX * maxIncrement));
286+ const bool contiguous = incx == 1 && incy == 1;
287+ const bool useStridedKernel = !contiguous && incrementFits && addressFits;
288+ const CaxpyKernelVariant kernelVariant =
289+ SelectCaxpyKernelVariant(static_cast<uint64_t>(n), incx, incy, useStridedKernel, numBlocks);
290+ CHECK_RET(
291+ sizeof(CaxpyTilingData) <= GetEffectiveWorkspaceSize(handle),
292+ LOG_PRINT("workspace need %zu, available %zu\n", sizeof(CaxpyTilingData), GetEffectiveWorkspaceSize(handle));
293+ return ACLBLAS_STATUS_EXECUTION_FAILED);
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
294+ uint8_t* maskDevice = nullptr;
295+ if (kernelVariant != CaxpyKernelVariant::STRIDED_SCALAR) {
296+ const aclblasStatus_t maskStatus = GetCaxpyMaskCache(handle, absIncx, absIncy, &maskDevice);
297+ CHECK_RET(maskStatus == ACLBLAS_STATUS_SUCCESS, return maskStatus);
298+ }
299+ // When a vector is strided but we are not taking the strided kernel path,
300+ // fall back to packing it into a contiguous scratch buffer and running the
301+ // dense kernel. Unit-stride vectors need no packing.
302+ const bool packX = !contiguous && !useStridedKernel && incx != 1;
303+ const bool packY = !contiguous && !useStridedKernel && incy != 1;
304+ execution = {
305+ n,
306+ incx,
307+ incy,
308+ numBlocks,
309+ useStridedKernel,
310+ packX,
311+ packY,
312+ kernelVariant,
313+ handle->stream,
314+ reinterpret_cast<uint8_t*>(GetEffectiveWorkspace(handle)),
315+ maskDevice,
316+ CalTilingData(
317+ static_cast<uint32_t>(n), numBlocks, alpha.real, alpha.imag, useStridedKernel ? incx : 1,
318+ useStridedKernel ? incy : 1)};
319+ return ACLBLAS_STATUS_SUCCESS;
320+}
321+ 
322+void FreePackedBuffers(CaxpyPackedBuffers& buffers)
323+{
324+ if (buffers.y != nullptr) {
325+ aclrtFree(buffers.y);
326+ }
327+ if (buffers.x != nullptr) {
328+ aclrtFree(buffers.x);
329+ }
330+ buffers = {};
331+}
332+ 
333+aclblasStatus_t AllocatePackedBuffers(const CaxpyExecution& execution, CaxpyPackedBuffers& buffers)
334+{
335+ const size_t bytes = static_cast<size_t>(execution.n) * sizeof(aclblasComplex);
336+ if (execution.packX) {
337+ const aclError result = aclrtMalloc(reinterpret_cast<void**>(&buffers.x), bytes, ACL_MEM_MALLOC_HUGE_FIRST);
338+ CHECK_RET(
339+ result == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", result);
340+ return ACLBLAS_STATUS_ALLOC_FAILED);
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
341+ }
342+ if (execution.packY) {
343+ const aclError result = aclrtMalloc(reinterpret_cast<void**>(&buffers.y), bytes, ACL_MEM_MALLOC_HUGE_FIRST);
344+ CHECK_RET(
345+ result == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", result); FreePackedBuffers(buffers);
346+ return ACLBLAS_STATUS_ALLOC_FAILED);
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
347+ }
348+ return ACLBLAS_STATUS_SUCCESS;
349+}
350+ 
351+aclblasStatus_t CopyStridedVector(
352+ uint8_t* packed, uint8_t* strided, uint64_t n, int64_t increment, aclrtStream stream, bool unpack)
353+{
354+ if (n == 0) {
355+ return ACLBLAS_STATUS_SUCCESS;
356+ }
357+ const size_t elementBytes = sizeof(aclblasComplex);
358+ const uint64_t absIncrement = AbsIncrement(increment);
359+ if (increment > 0) {
360+ uint8_t* destination = unpack ? strided : packed;
361+ const uint8_t* source = unpack ? packed : strided;
362+ const size_t destinationPitch = unpack ? absIncrement * elementBytes : elementBytes;
363+ const size_t sourcePitch = unpack ? elementBytes : absIncrement * elementBytes;
364+ const aclError result = aclrtMemcpy2dAsync(
365+ destination, destinationPitch, source, sourcePitch, elementBytes, n, ACL_MEMCPY_DEVICE_TO_DEVICE, stream);
366+ CHECK_RET(
367+ result == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy2dAsync failed. ERROR: %d\n", result);
368+ return ACLBLAS_STATUS_INTERNAL_ERROR);
369+ return ACLBLAS_STATUS_SUCCESS;
370+ }
371+ 
372+ std::vector<uint8_t*> destinations(n);
373+ std::vector<uint8_t*> sources(n);
374+ std::vector<size_t> destinationSizes(n, elementBytes);
375+ std::vector<size_t> copySizes(n, elementBytes);
376+ for (uint64_t i = 0; i < n; ++i) {
377+ const size_t packedOffset = static_cast<size_t>(i) * elementBytes;
378+ const size_t stridedOffset = static_cast<size_t>((n - 1 - i) * absIncrement) * elementBytes;
379+ destinations[i] = unpack ? strided + stridedOffset : packed + packedOffset;
380+ sources[i] = unpack ? packed + packedOffset : strided + stridedOffset;
381+ }
382+ aclrtMemcpyBatchAttr attribute{};
383+ attribute.dstLoc.type = ACL_MEM_LOCATION_TYPE_DEVICE;
384+ attribute.srcLoc.type = ACL_MEM_LOCATION_TYPE_DEVICE;
385+ const aclError result = aclrtMemcpyBatchAsync(
386+ destinations.data(), destinationSizes.data(), sources.data(), copySizes.data(), n, attribute, stream);
387+ CHECK_RET(
388+ result == ACL_SUCCESS, LOG_PRINT("aclrtMemcpyBatchAsync failed. ERROR: %d\n", result);
389+ return ACLBLAS_STATUS_INTERNAL_ERROR);
390+ return ACLBLAS_STATUS_SUCCESS;
391+}
392+ 
393+aclblasStatus_t PreparePackedInputs(
394+ const CaxpyExecution& execution, const aclblasComplex* x, aclblasComplex* y, CaxpyPackedBuffers& buffers)
395+{
396+ aclblasStatus_t status = AllocatePackedBuffers(execution, buffers);
397+ CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, return status);
398+ if (execution.packX) {
399+ status = CopyStridedVector(
400+ buffers.x, const_cast<uint8_t*>(reinterpret_cast<const uint8_t*>(x)), execution.n, execution.incx,
401+ execution.stream, false);
402+ CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, FreePackedBuffers(buffers); return status);
403+ }
404+ if (execution.packY) {
405+ status = CopyStridedVector(
406+ buffers.y, reinterpret_cast<uint8_t*>(y), execution.n, execution.incy, execution.stream, false);
407+ CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, FreePackedBuffers(buffers); return status);
408+ }
409+ return ACLBLAS_STATUS_SUCCESS;
410+}
411+ 
412+aclblasStatus_t UploadTiling(const CaxpyExecution& execution, const CaxpyTilingData& tiling)
MM
MMaskHunter200817 小时前

[Minor] Missing enum validation

问题

host 实现中未对枚举参数（uplo/trans/diag/side）进行校验。

依据

  • interface-spec.md §2 校验顺序第2步：非法枚举 → INVALID_ENUM

修改建议

if ((uplo != ACLBLAS_UPPER && uplo != ACLBLAS_LOWER) ||
    (trans != ACLBLAS_OP_N && trans != ACLBLAS_OP_T && trans != ACLBLAS_OP_C) ||
    (diag != ACLBLAS_UNIT && diag != ACLBLAS_NON_UNIT)) {
    return ACLBLAS_STATUS_INVALID_ENUM;
}
likedislike
MMaskHunter200816 小时前

[Minor] Missing enum validation

问题

host 实现中未对枚举参数（uplo/trans/diag/side）进行校验。

依据

  • interface-spec.md §2 校验顺序第2步：非法枚举 → INVALID_ENUM

修改建议

if ((uplo != ACLBLAS_UPPER && uplo != ACLBLAS_LOWER) ||
    (trans != ACLBLAS_OP_N && trans != ACLBLAS_OP_T && trans != ACLBLAS_OP_C) ||
    (diag != ACLBLAS_UNIT && diag != ACLBLAS_NON_UNIT)) {
    return ACLBLAS_STATUS_INVALID_ENUM;
}
likedislike
413+{
414+ const aclError result = aclrtMemcpyAsync(
415+ execution.tilingDevice, sizeof(tiling), &tiling, sizeof(tiling), ACL_MEMCPY_HOST_TO_DEVICE, execution.stream);
416+ CHECK_RET(
417+ result == ACL_SUCCESS, LOG_PRINT("aclrtMemcpyAsync failed. ERROR: %d\n", result);
418+ return ACLBLAS_STATUS_INTERNAL_ERROR);
419+ return ACLBLAS_STATUS_SUCCESS;
420+}
421+ 
422+CaxpyTilingData BuildStridedTiling(
423+ const CaxpyExecution& execution, const aclblasComplex& alpha, uint32_t count, uint32_t blocks)
424+{
425+ CaxpyTilingData tiling = CalTilingData(
426+ count, blocks, alpha.real, alpha.imag, execution.incx, execution.incy, 0, count);
427+ 
428+ // Only the shared pipeline's run_shared_pipeline/shared_tile_bounds read
429+ // this field; the disjoint and scalar kernel entry points ignore it and
430+ // keep using the fixed CAXPY_STRIDED_TILE_COUNT capacity, so it is safe
431+ // to always populate it here for every strided variant.
432+ tiling.sharedTileCapacity = static_cast<uint32_t>(CaxpySharedTileCapacity(AbsIncrement(execution.incx)));
433+ return tiling;
434+}
435+ 
436+uint64_t CaxpySegmentStorageOffset(uint64_t totalN, uint64_t segmentStart, uint64_t segmentN, int64_t increment)
437+{
438+ const uint64_t absIncrement = AbsIncrement(increment);
439+ return increment > 0 ? segmentStart * absIncrement : (totalN - segmentStart - segmentN) * absIncrement;
440+}
441+ 
442+void LaunchCaxpyStridedPipelined(const CaxpyExecution& execution, uint8_t* x, uint8_t* y, uint32_t blocks)
443+{
444+ if (execution.kernelVariant == CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC ||
445+ execution.kernelVariant == CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_SCALAR_TAIL) {
446+ caxpy_strided_shared_kernel_do(
447+ x, execution.maskDevice, y, nullptr, execution.tilingDevice, blocks, execution.stream);
448+ return;
449+ }
450+ caxpy_strided_disjoint_kernel_do(
451+ x, execution.maskDevice, y, nullptr, execution.tilingDevice, blocks, execution.stream);
452+}
453+ 
454+aclblasStatus_t LaunchCaxpyStridedSegment(
455+ const CaxpyExecution& execution, const aclblasComplex& alpha, const aclblasComplex* x, aclblasComplex* y,
456+ uint64_t segmentStart, uint64_t segmentN, bool scalar)
457+{
458+ const uint32_t blocks = static_cast<uint32_t>(std::min<uint64_t>(execution.numBlocks, segmentN));
459+ CaxpyTilingData tiling = BuildStridedTiling(execution, alpha, static_cast<uint32_t>(segmentN), blocks);
460+ const aclblasStatus_t status = UploadTiling(execution, tiling);
461+ CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, return status);
462+ 
463+ const uint64_t totalN = static_cast<uint64_t>(execution.n);
464+ const uint64_t xOffset = CaxpySegmentStorageOffset(totalN, segmentStart, segmentN, execution.incx);
465+ const uint64_t yOffset = CaxpySegmentStorageOffset(totalN, segmentStart, segmentN, execution.incy);
466+ auto* segmentX = const_cast<aclblasComplex*>(x + xOffset);
467+ auto* segmentY = y + yOffset;
468+ if (scalar) {
469+ caxpy_strided_scalar_kernel_do(
470+ reinterpret_cast<uint8_t*>(segmentX), reinterpret_cast<uint8_t*>(segmentY), execution.tilingDevice, blocks,
471+ execution.stream);
472+ } else {
473+ LaunchCaxpyStridedPipelined(
474+ execution, reinterpret_cast<uint8_t*>(segmentX), reinterpret_cast<uint8_t*>(segmentY), blocks);
475+ }
476+ return ACLBLAS_STATUS_SUCCESS;
477+}
478+ 
479+aclblasStatus_t LaunchCaxpy(
480+ const CaxpyExecution& execution, const aclblasComplex& alpha, const aclblasComplex* x, aclblasComplex* y,
481+ const CaxpyPackedBuffers& buffers)
482+{
483+ if (execution.useStridedKernel) {
484+ if (execution.kernelVariant == CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_SCALAR_TAIL ||
485+ execution.kernelVariant == CaxpyKernelVariant::STRIDED_DISJOINT_PIPELINED_SCALAR_TAIL) {
486+ const uint64_t tailN = CaxpyFinalWaveCount(static_cast<uint64_t>(execution.n), execution.numBlocks);
487+ const uint64_t prefixN = static_cast<uint64_t>(execution.n) - tailN;
488+ aclblasStatus_t status = LaunchCaxpyStridedSegment(execution, alpha, x, y, 0, prefixN, false);
489+ CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, return status);
490+ status = LaunchCaxpyStridedSegment(execution, alpha, x, y, prefixN, tailN, true);
491+ CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, return status);
492+ return ACLBLAS_STATUS_SUCCESS;
493+ }
494+ const uint32_t blocks = static_cast<uint32_t>(std::min<uint64_t>(execution.numBlocks, execution.n));
495+ CaxpyTilingData tiling = BuildStridedTiling(execution, alpha, static_cast<uint32_t>(execution.n), blocks);
496+ const aclblasStatus_t status = UploadTiling(execution, tiling);
497+ CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, return status);
498+ if (execution.kernelVariant == CaxpyKernelVariant::STRIDED_SCALAR) {
499+ caxpy_strided_scalar_kernel_do(
500+ const_cast<uint8_t*>(reinterpret_cast<const uint8_t*>(x)), reinterpret_cast<uint8_t*>(y),
501+ execution.tilingDevice, blocks, execution.stream);
502+ } else {
503+ LaunchCaxpyStridedPipelined(
504+ execution, const_cast<uint8_t*>(reinterpret_cast<const uint8_t*>(x)), reinterpret_cast<uint8_t*>(y),
505+ blocks);
506+ }
507+ return ACLBLAS_STATUS_SUCCESS;
508+ }
509+ const aclblasStatus_t status = UploadTiling(execution, execution.tiling);
510+ CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, return status);
511+ caxpy_kernel_do(
512+ execution.packX ? buffers.x : const_cast<uint8_t*>(reinterpret_cast<const uint8_t*>(x)), execution.maskDevice,
513+ execution.packY ? buffers.y : reinterpret_cast<uint8_t*>(y), nullptr, execution.tilingDevice,
514+ execution.numBlocks, execution.stream);
515+ return ACLBLAS_STATUS_SUCCESS;
516+}
517+ 
518+aclblasStatus_t FinishCaxpy(const CaxpyExecution& execution, aclblasComplex* y, CaxpyPackedBuffers& buffers)
519+{
520+ if (execution.packY) {
521+ const aclblasStatus_t status = CopyStridedVector(
522+ buffers.y, reinterpret_cast<uint8_t*>(y), execution.n, execution.incy, execution.stream, true);
523+ CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, FreePackedBuffers(buffers); return status);
524+ }
525+ if (execution.packX || execution.packY) {
526+ const aclError result = aclrtSynchronizeStream(execution.stream);
527+ CHECK_RET(
528+ result == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", result);
529+ FreePackedBuffers(buffers); return ACLBLAS_STATUS_INTERNAL_ERROR);
530+ }
531+ FreePackedBuffers(buffers);
532+ return ACLBLAS_STATUS_SUCCESS;
533+}
534+ 
535+aclblasStatus_t aclblasCaxpy(
536+ aclblasHandle_t handle, int n, const aclblasComplex* alpha, const aclblasComplex* x, int incx, aclblasComplex* y,
537+ int incy)
538+{
539+ if (handle == nullptr) {
540+ OP_LOGE("aclblasCaxpy", "handle is nullptr");
541+ return ACLBLAS_STATUS_HANDLE_IS_NULLPTR;
MM
MMaskHunter200817 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
MMaskHunter200816 小时前

[Minor] Missing error log in return path

问题

此行错误返回路径缺少 OP_LOGE 日志，不利于问题定位。

依据

  • code-quality.md §6.3 要求：错误返回路径应包含 OP_LOGE 并附带函数名

修改建议

// 修改前:
return ACLBLAS_STATUS_INVALID_VALUE;

// 修改后:
OP_LOGE("aclblasXxx: invalid parameter");
return ACLBLAS_STATUS_INVALID_VALUE;
likedislike
542+ }
543+ if (n == 0) {
544+ return ACLBLAS_STATUS_SUCCESS;
545+ }
546+ const aclblasStatus_t argumentStatus = ValidateCaxpyArguments(handle, n, alpha, x, incx, y, incy);
547+ CHECK_RET(argumentStatus == ACLBLAS_STATUS_SUCCESS, return argumentStatus);
548+ CaxpyExecution execution{};
549+ const aclblasStatus_t buildStatus = BuildCaxpyExecution(handle, n, *alpha, incx, incy, execution);
550+ CHECK_RET(buildStatus == ACLBLAS_STATUS_SUCCESS, return buildStatus);
551+ CaxpyPackedBuffers buffers;
552+ const aclblasStatus_t prepareStatus = PreparePackedInputs(execution, x, y, buffers);
553+ CHECK_RET(prepareStatus == ACLBLAS_STATUS_SUCCESS, return prepareStatus);
554+ const aclblasStatus_t launchStatus = LaunchCaxpy(execution, *alpha, x, y, buffers);
555+ CHECK_RET(launchStatus == ACLBLAS_STATUS_SUCCESS, FreePackedBuffers(buffers); return launchStatus);
556+ return FinishCaxpy(execution, y, buffers);
557+}
Mblas/axpy/arch22/caxpy_kernel.cpp+847-148
@@ -1,3 +1,4 @@
1+ 
1/**2/**
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.3 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of4 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
@@ -18,10 +19,49 @@
18#endif19#endif
19 20 
20#include "kernel_operator.h"21#include "kernel_operator.h"
22+#include "c_api/asc_simd.h"
23+#include "caxpy_tiling_data.h"
24+#include "caxpy_dispatch_policy.h"
21 25 
22#if __DAV_C220_VEC__26#if __DAV_C220_VEC__
23-__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub_uint32(__ubuf__ uint32_t *dst,27+// Unified Buffer capacity on 910B3 (192 KiB). All *_ub_layout() functions below
24- __gm__ uint32_t *src, uint32_t len)28+// must fit their working set within this budget; the static_asserts next to
29+// each layout enforce it at compile time instead of relying on manual review.
30+constexpr uint32_t CAXPY_UB_BUDGET_BYTES = 192 * 1024;
31+ 
32+// --- Dense (contiguous incx=1, incy=1) kernel UB layout -------------------
33+// Two ping-pong slots for the load buffer (ubX) and the output buffer (ubOut),
34+// each sized to one full tile (38 KiB), followed by a full-tile offset table
35+// (ubOffset) and a small 8-float coefficient table (ubCoefficient).
36+constexpr uint32_t CAXPY_DENSE_TILE_FLOATS = 38 * 1024 / sizeof(float);
37+constexpr uint32_t CAXPY_DENSE_TILE_BYTES = 38 * 1024;
38+constexpr uint32_t CAXPY_DENSE_UB_OUT0_BYTES = 0 * CAXPY_DENSE_TILE_BYTES;
39+constexpr uint32_t CAXPY_DENSE_UB_OUT1_BYTES = 1 * CAXPY_DENSE_TILE_BYTES;
40+constexpr uint32_t CAXPY_DENSE_UB_X0_BYTES = 2 * CAXPY_DENSE_TILE_BYTES;
41+constexpr uint32_t CAXPY_DENSE_UB_X1_BYTES = 3 * CAXPY_DENSE_TILE_BYTES;
42+constexpr uint32_t CAXPY_DENSE_UB_OFFSET_BYTES = 4 * CAXPY_DENSE_TILE_BYTES;
43+constexpr uint32_t CAXPY_DENSE_UB_COEFFICIENT_BYTES = 5 * CAXPY_DENSE_TILE_BYTES;
44+constexpr uint32_t CAXPY_DENSE_UB_COEFFICIENT_COUNT = 8;
45+// Offset table region: one uint32 pair-swap offset per float in a full tile.
46+// The trailing terms mirror the strided-kernel GM mask layout below this
47+// function (CAXPY_X_OFFSET_COUNT and 2*CAXPY_SPAN_OFFSET_COUNT, both defined
48+// later in this file for the strided kernels' own UB tables): the host lays
49+// out gmAug as [compute offsets][x-normalize offsets][store offsets]
50+// [y-gather offsets][dense swap offsets], so this dense-kernel offset into
51+// gmAug must skip past the strided sections that precede it in that layout,
52+// even though the dense and strided kernels are otherwise independent code
53+// paths. See caxpy_host.cpp's DENSE_SWAP_OFFSET_START for the authoritative
54+// host-side computation this mirrors.
55+constexpr uint32_t CAXPY_DENSE_SWAP_SKIP_X_OFFSET_WORDS = 512 * 8 + 64; // == CAXPY_X_OFFSET_COUNT
56+constexpr uint32_t CAXPY_DENSE_SWAP_SKIP_SPAN_OFFSET_WORDS = 2 * (512 * 2 + 64); // == 2 * CAXPY_SPAN_OFFSET_COUNT
57+constexpr uint32_t CAXPY_DENSE_SWAP_OFFSET_START =
58+ CAXPY_DENSE_TILE_FLOATS + CAXPY_DENSE_SWAP_SKIP_X_OFFSET_WORDS + CAXPY_DENSE_SWAP_SKIP_SPAN_OFFSET_WORDS;
59+static_assert(
60+ CAXPY_DENSE_UB_COEFFICIENT_BYTES + CAXPY_DENSE_UB_COEFFICIENT_COUNT * sizeof(float) <= CAXPY_UB_BUDGET_BYTES,
61+ "CAXPY dense UB layout exceeds the 910B3 Unified Buffer budget");
62+ 
63+__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub_uint32(
64+ __ubuf__ uint32_t* dst, __gm__ uint32_t* src, uint32_t len)
25{65{
26 uint16_t nBurst = 1;66 uint16_t nBurst = 1;
27 uint32_t lenBurst = len * sizeof(uint32_t);67 uint32_t lenBurst = len * sizeof(uint32_t);
@@ -29,13 +69,11 @@ __aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub_uint32(
29 uint8_t rightPaddingNum = 0;69 uint8_t rightPaddingNum = 0;
30 uint32_t srcGap = 0;70 uint32_t srcGap = 0;
31 uint32_t dstGap = 0;71 uint32_t dstGap = 0;
32- copy_gm_to_ubuf_align_b32(dst, src,72+ copy_gm_to_ubuf_align_b32(dst, src, 0, nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
33- 0,
34- nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
35}73}
36 74 
37-__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub(__ubuf__ float *dst, __gm__ float *src,75+__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub(
38- uint32_t len)76+ __ubuf__ float* dst, __gm__ float* src, uint32_t len)
39{77{
40 uint16_t nBurst = 1;78 uint16_t nBurst = 1;
41 uint32_t lenBurst = len * sizeof(float);79 uint32_t lenBurst = len * sizeof(float);
@@ -43,13 +81,11 @@ __aicore__ __inline__ __attribute__((always_inline)) void copy_vec_gm2ub(__ubuf_
43 uint8_t rightPaddingNum = 0;81 uint8_t rightPaddingNum = 0;
44 uint32_t srcGap = 0;82 uint32_t srcGap = 0;
45 uint32_t dstGap = 0;83 uint32_t dstGap = 0;
46- copy_gm_to_ubuf_align_b32(dst, src,84+ copy_gm_to_ubuf_align_b32(dst, src, 0, nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
47- 0,
48- nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
49}85}
50 86 
51-__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_ub2gm(__gm__ float *dst, __ubuf__ float *src,87+__aicore__ __inline__ __attribute__((always_inline)) void copy_vec_ub2gm(
52- uint32_t len)88+ __gm__ float* dst, __ubuf__ float* src, uint32_t len)
53{89{
54 uint16_t nBurst = 1;90 uint16_t nBurst = 1;
55 uint32_t lenBurst = len * sizeof(float);91 uint32_t lenBurst = len * sizeof(float);
@@ -57,145 +93,677 @@ __aicore__ __inline__ __attribute__((always_inline)) void copy_vec_ub2gm(__gm__
57 uint8_t rightPaddingNum = 0;93 uint8_t rightPaddingNum = 0;
58 uint32_t srcGap = 0;94 uint32_t srcGap = 0;
59 uint32_t dstGap = 0;95 uint32_t dstGap = 0;
60- copy_ubuf_to_gm_align_b32(dst, src,96+ copy_ubuf_to_gm_align_b32(dst, src, 0, nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
61- 0,
62- nBurst, lenBurst, leftPaddingNum, rightPaddingNum, srcGap, dstGap);
63}97}
64 98 
65-__aicore__ __inline__ __attribute__((always_inline)) void caxpy_compute_aiv(99+__aicore__ __inline__ __attribute__((always_inline)) uint32_t abs_increment(int64_t increment)
66- __gm__ float *gm_in, __gm__ float *gm_out, __ubuf__ float *ub_in, __ubuf__ float *ub_out,
67- __ubuf__ uint32_t *ub_offset, float s_real, float s_imag, uint32_t copy_len, uint32_t len, uint32_t event_id)
68{100{
69- uint32_t repeatTime = (len + 63) / 64;101+ return increment > 0 ? static_cast<uint32_t>(increment) : static_cast<uint32_t>(-increment);
70- uint32_t computeRepeat = (len / 2 + 63) / 64;
71- 
72- uint32_t real_offset = 0;
73- uint32_t imag_offset = 38 * 1024 / sizeof(float) / 2;
74- 
75- auto ub_out_real = ub_out;
76- auto ub_out_imag = ub_out + imag_offset;
77- 
78- copy_vec_gm2ub(ub_in, gm_in, copy_len);
79- 
80- set_flag(PIPE_MTE2, PIPE_V, event_id);
81- wait_flag(PIPE_MTE2, PIPE_V, event_id);
82- 
83- vreducev2(reinterpret_cast<__ubuf__ uint32_t *>(ub_out_real),
84- reinterpret_cast<__ubuf__ uint32_t *>(ub_in),
85- nullptr, repeatTime,
86- 1,
87- 1,
88- 8, 8);
89- 
90- vreducev2(reinterpret_cast<__ubuf__ uint32_t *>(ub_out_imag),
91- reinterpret_cast<__ubuf__ uint32_t *>(ub_in),
92- nullptr, repeatTime,
93- 1,
94- 2,
95- 8, 8);
96- 
97- pipe_barrier(PIPE_V);
98- 
99- vmuls(ub_in, ub_out_real, s_real, computeRepeat, 1, 1, 8, 8);
100- vmuls(ub_in + imag_offset, ub_out_real, s_imag, computeRepeat, 1, 1, 8, 8);
101- 
102- vmuls(ub_out_real, ub_out_imag, s_imag, computeRepeat, 1, 1, 8, 8);
103- 
104- pipe_barrier(PIPE_V);
105- vsub(ub_in, ub_in, ub_out_real, computeRepeat, 1, 1, 1, 8, 8, 8);
106- 
107- vmuls(ub_out_imag, ub_out_imag, s_real, computeRepeat, 1, 1, 8, 8);
108- 
109- pipe_barrier(PIPE_V);
110- vadd(ub_in + imag_offset, ub_out_imag, ub_in + imag_offset, computeRepeat, 1, 1, 1, 8, 8, 8);
111- 
112- pipe_barrier(PIPE_V);
113- 
114- vgather(reinterpret_cast<__ubuf__ uint32_t *>(ub_out), reinterpret_cast<__ubuf__ uint32_t *>(ub_offset),
115- (uintptr_t)(ub_in),
116- 8,
117- repeatTime
118- );
119- 
120- set_flag(PIPE_V, PIPE_MTE3, event_id);
121- wait_flag(PIPE_V, PIPE_MTE3, event_id);
122- 
123- copy_vec_ub2gm(gm_out, ub_out, copy_len);
124}102}
125 103 
126-__aicore__ __inline__ __attribute__((always_inline)) void caxpy_aiv(__gm__ float *gm_in, __gm__ uint32_t *gm_aug,104+__aicore__ __inline__ __attribute__((always_inline)) bool copy_complex_gm2ub_signed(
127- __gm__ float *gm_out, float s_real, float s_imag,105+ __ubuf__ float* dst, __ubuf__ float* denseTemp, __gm__ float* source, uint32_t logicalOffset,
128- uint32_t offset, uint32_t calNum)106+ uint32_t totalComplexCount, uint32_t complexCount, int64_t signedIncrement)
129{107{
108+ constexpr uint32_t maxBurstCount = 4088;
109+ const uint32_t increment = abs_increment(signedIncrement);
110+ const uint32_t physicalOffset = signedIncrement > 0 ?
111+ logicalOffset * increment :
112+ (totalComplexCount - logicalOffset - complexCount) * increment;
113+ auto src = source + physicalOffset * 2;
114+ if (increment == 1) {
115+ copy_gm_to_ubuf_align_b32(denseTemp, src, 0, 1, complexCount * 2 * sizeof(float), 0, 0, 0, 0);
116+ return true;
117+ }
118+ const uint32_t gapBytes = (increment - 1) * 2 * sizeof(float);
119+ uint32_t copied = 0;
120+ while (copied < complexCount) {
121+ const uint32_t batch = (complexCount - copied) > maxBurstCount ? maxBurstCount : complexCount - copied;
122+ asc_copy_gm2ub_align(
123+ dst + copied * 8, src + copied * increment * 2, static_cast<uint16_t>(batch), 2 * sizeof(float), 0, 0,
124+ gapBytes, 0);
125+ copied += batch;
126+ }
127+ return false;
128+}
129+ 
130+__aicore__ __inline__ __attribute__((always_inline)) void normalize_loaded_complexes(
131+ __ubuf__ float* dst, __ubuf__ float* denseTemp, uint32_t complexCount, bool denseLoaded, bool reverse)
132+{
133+ if (denseLoaded) {
134+ for (uint32_t i = 0; i < complexCount; ++i) {
135+ const uint32_t sourceIndex = reverse ? complexCount - 1 - i : i;
136+ dst[i * 8] = denseTemp[sourceIndex * 2];
137+ dst[i * 8 + 1] = denseTemp[sourceIndex * 2 + 1];
138+ }
139+ return;
140+ }
141+ if (reverse) {
142+ for (uint32_t left = 0; left < complexCount / 2; ++left) {
143+ const uint32_t right = complexCount - 1 - left;
144+ const float leftReal = dst[left * 8];
145+ const float leftImag = dst[left * 8 + 1];
146+ dst[left * 8] = dst[right * 8];
147+ dst[left * 8 + 1] = dst[right * 8 + 1];
148+ dst[right * 8] = leftReal;
149+ dst[right * 8 + 1] = leftImag;
150+ }
151+ }
152+}
153+ 
154+__aicore__ __inline__ __attribute__((always_inline)) void copy_complex_ub2gm_signed(
155+ __gm__ float* destination, __ubuf__ float* src, uint32_t logicalOffset, uint32_t totalComplexCount,
156+ uint32_t complexCount, int64_t signedIncrement)
157+{
158+ constexpr uint32_t maxBurstCount = 4088;
159+ const uint32_t increment = abs_increment(signedIncrement);
160+ const uint32_t physicalOffset = signedIncrement > 0 ?
161+ logicalOffset * increment :
162+ (totalComplexCount - logicalOffset - complexCount) * increment;
163+ auto dst = destination + physicalOffset * 2;
164+ if (increment == 1) {
165+ copy_ubuf_to_gm_align_b32(dst, src, 0, 1, complexCount * 2 * sizeof(float), 0, 0, 0, 0);
166+ return;
167+ }
168+ const uint32_t gapBytes = (increment - 1) * 2 * sizeof(float);
169+ uint32_t copied = 0;
170+ while (copied < complexCount) {
171+ const uint32_t batch = (complexCount - copied) > maxBurstCount ? maxBurstCount : complexCount - copied;
172+ asc_copy_ub2gm_align(
173+ dst + copied * increment * 2, src + copied * 8, static_cast<uint16_t>(batch), 2 * sizeof(float), 0, 0, 0,
174+ gapBytes);
175+ copied += batch;
176+ }
177+}
178+ 
179+__aicore__ __inline__ __attribute__((always_inline)) void caxpy_compute_planar(
180+ __ubuf__ float* resultReal, __ubuf__ float* resultImag, __ubuf__ float* sourceReal,
181+ __ubuf__ float* sourceImag, float alphaReal, float alphaImag, uint32_t repeat)
182+{
183+ // Keep every vector address 32-byte aligned. An interleaved implementation using
184+ // source + 1 is mathematically equivalent, but faults on C220 vector instructions.
185+ vmuls(resultReal, sourceReal, alphaReal, repeat, 1, 1, 8, 8);
186+ vmuls(resultImag, sourceImag, alphaReal, repeat, 1, 1, 8, 8);
187+ pipe_barrier(PIPE_V);
188+ vaxpy(resultReal, sourceImag, -alphaImag, repeat, 1, 1, 8, 8);
189+ vaxpy(resultImag, sourceReal, alphaImag, repeat, 1, 1, 8, 8);
190+ pipe_barrier(PIPE_V);
191+}
192+ 
193+ 
194+__aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_load(
195+ __gm__ float* gmX, __ubuf__ float* ubX, uint32_t copyLength, uint32_t eventId)
196+{
197+ copy_vec_gm2ub(ubX, gmX, copyLength);
198+ set_flag(PIPE_MTE2, PIPE_V, eventId);
199+}
200+ 
201+__aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_compute(
202+ __ubuf__ float* ubX, __ubuf__ float* ubOut, __ubuf__ uint32_t* ubOffset,
203+ __ubuf__ float* ubCoefficient, float alphaReal, uint32_t computeLength, uint32_t eventId)
204+{
205+ const uint32_t repeat = (computeLength + 63) / 64;
206+ 
207+ wait_flag(PIPE_MTE2, PIPE_V, eventId);
208+ // Turn [real, imag] into [imag, real]. The alternating coefficient then
209+ // applies [-alphaImag, +alphaImag] without splitting into planar buffers.
210+ vgather(
211+ reinterpret_cast<__ubuf__ uint32_t*>(ubOut), ubOffset, (uintptr_t)ubX, 8, repeat);
212+ pipe_barrier(PIPE_V);
213+ vmuls(ubX, ubX, alphaReal, repeat, 1, 1, 8, 8);
214+ vmul(ubOut, ubOut, ubCoefficient, repeat, 1, 1, 0, 8, 8, 0);
215+ pipe_barrier(PIPE_V);
216+ vadd(ubOut, ubX, ubOut, repeat, 1, 1, 1, 8, 8, 8);
217+ set_flag(PIPE_V, PIPE_MTE3, eventId);
218+}
219+ 
220+__aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_store(
221+ __gm__ float* gmY, __ubuf__ float* ubOut, uint32_t copyLength, uint32_t eventId)
222+{
223+ wait_flag(PIPE_V, PIPE_MTE3, eventId);
224+ copy_vec_ub2gm(gmY, ubOut, copyLength);
225+ set_flag(PIPE_MTE3, PIPE_MTE2, eventId);
226+}
227+ 
228+__aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_prepare(
229+ __gm__ uint32_t* gmAug, __ubuf__ uint32_t*& ubOffset, __ubuf__ float*& ubCoefficient, float alphaImag)
230+{
231+ constexpr uint32_t tileFloats = CAXPY_DENSE_TILE_FLOATS;
232+ ubOffset = reinterpret_cast<__ubuf__ uint32_t*>((uintptr_t)CAXPY_DENSE_UB_OFFSET_BYTES);
233+ ubCoefficient = reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_DENSE_UB_COEFFICIENT_BYTES);
234+ copy_vec_gm2ub_uint32(ubOffset, gmAug + CAXPY_DENSE_SWAP_OFFSET_START, tileFloats);
235+ for (uint32_t index = 0; index < CAXPY_DENSE_UB_COEFFICIENT_COUNT; ++index)
236+ ubCoefficient[index] = (index & 1) == 0 ? -alphaImag : alphaImag;
237+ set_flag(PIPE_S, PIPE_V, EVENT_ID2);
238+ wait_flag(PIPE_S, PIPE_V, EVENT_ID2);
239+}
240+ 
241+__aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_load_next_if_any(
242+ __gm__ float* gmX, __ubuf__ float* ubX[2], uint32_t offset, uint32_t calNum, uint32_t tileFloats,
243+ uint32_t tile, uint32_t tileCount, uint32_t current)
244+{
245+ if (tile + 1 >= tileCount)
246+ return;
247+ const uint32_t next = 1 - current;
248+ const uint32_t nextEvent = next == 0 ? EVENT_ID0 : EVENT_ID1;
249+ const uint32_t nextOffset = (tile + 1) * tileFloats;
250+ const uint32_t nextCount = calNum - nextOffset > tileFloats ? tileFloats : calNum - nextOffset;
251+ wait_flag(PIPE_MTE3, PIPE_MTE2, nextEvent);
252+ caxpy_dense_load(gmX + offset + nextOffset, ubX[next], nextCount, nextEvent);
253+}
254+ 
255+__aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_run(
Z
Zzhuzhiming9 天前

【编码规范 R7】函数非空非注释行数 ≤ 50。

caxpy_dense_run 以及后面的 run_shared_pipeline / run_disjoint_pipeline 明显超标。请按 load / compute / store / pipeline 拆分子函数。

likedislike
256+ __gm__ float* gmX, __gm__ uint32_t* gmAug, __gm__ float* gmY, float alphaReal, float alphaImag,
257+ uint32_t offset, uint32_t calNum)
258+{
259+ constexpr uint32_t tileFloats = CAXPY_DENSE_TILE_FLOATS;
260+ __ubuf__ float* ubOut[2] = {
261+ reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_DENSE_UB_OUT0_BYTES),
262+ reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_DENSE_UB_OUT1_BYTES)};
263+ __ubuf__ float* ubX[2] = {
264+ reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_DENSE_UB_X0_BYTES),
265+ reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_DENSE_UB_X1_BYTES)};
266+ __ubuf__ uint32_t* ubOffset = nullptr;
267+ __ubuf__ float* ubCoefficient = nullptr;
268+ caxpy_dense_prepare(gmAug, ubOffset, ubCoefficient, alphaImag);
269+ 
130 set_atomic_add();270 set_atomic_add();
131 set_atomic_f32();271 set_atomic_f32();
272+ set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
273+ set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
132 274 
133- auto ub_out_ping = reinterpret_cast<__ubuf__ float *>((uintptr_t)0 * 1024);275+ const uint32_t tileCount = (calNum + tileFloats - 1) / tileFloats;
134- auto ub_out_pong = reinterpret_cast<__ubuf__ float *>((uintptr_t)38 * 1024);276+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
135- auto ub_in_ping = reinterpret_cast<__ubuf__ float *>((uintptr_t)76 * 1024);277+ const uint32_t firstCount = calNum > tileFloats ? tileFloats : calNum;
136- auto ub_in_pong = reinterpret_cast<__ubuf__ float *>((uintptr_t)114 * 1024);278+ caxpy_dense_load(gmX + offset, ubX[0], firstCount, EVENT_ID0);
137- auto ub_offset = reinterpret_cast<__ubuf__ uint32_t *>((uintptr_t)152 * 1024);
138 279 
139- uint32_t ping_flag = 1;280+ for (uint32_t tile = 0; tile < tileCount; ++tile) {
281+ const uint32_t current = tile & 1;
282+ const uint32_t currentEvent = current == 0 ? EVENT_ID0 : EVENT_ID1;
283+ const uint32_t currentOffset = tile * tileFloats;
284+ const uint32_t currentCount =
285+ calNum - currentOffset > tileFloats ? tileFloats : calNum - currentOffset;
140 286 
141- uint32_t maxDataCount = 38 * 1024 / sizeof(float);287+ caxpy_dense_load_next_if_any(gmX, ubX, offset, calNum, tileFloats, tile, tileCount, current);
142 288 
143- uint32_t repeatTime = calNum / maxDataCount;289+ caxpy_dense_compute(
144- uint32_t remainNum = calNum % maxDataCount;290+ ubX[current], ubOut[current], ubOffset, ubCoefficient, alphaReal, tileFloats, currentEvent);
145- 291+ caxpy_dense_store(gmY + offset + currentOffset, ubOut[current], currentCount, currentEvent);
146- copy_vec_gm2ub_uint32(ub_offset, gm_aug, maxDataCount);
147- 
148- uint32_t curr_offset = offset;
149- 
150- if (repeatTime > 0) {
151- set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
152- set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
153- for (uint32_t i = 0; i < repeatTime; i++) {
154- auto ub_in = ping_flag ? ub_in_ping : ub_in_pong;
155- auto ub_out = ping_flag ? ub_out_ping : ub_out_pong;
156- 
157- auto event_id = ping_flag ? EVENT_ID0 : EVENT_ID1;
158- 
159- wait_flag(PIPE_MTE3, PIPE_MTE2, event_id);
160- 
161- caxpy_compute_aiv(gm_in + curr_offset, gm_out + curr_offset, ub_in, ub_out, ub_offset, s_real, s_imag,
162- maxDataCount, maxDataCount, event_id);
163- 
164- set_flag(PIPE_MTE3, PIPE_MTE2, event_id);
165- 
166- curr_offset += maxDataCount;
167- ping_flag = 1 - ping_flag;
168- }
169- wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
170- wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
171- }
172- 
173- if (remainNum > 0) {
174- set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
175- set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
176- auto ub_in = ping_flag ? ub_in_ping : ub_in_pong;
177- auto ub_out = ping_flag ? ub_out_ping : ub_out_pong;
178- 
179- auto event_id = ping_flag ? EVENT_ID0 : EVENT_ID1;
180- 
181- wait_flag(PIPE_MTE3, PIPE_MTE2, event_id);
182- 
183- caxpy_compute_aiv(gm_in + curr_offset, gm_out + curr_offset, ub_in, ub_out, ub_offset, s_real, s_imag,
184- remainNum, maxDataCount, event_id);
185- 
186- set_flag(PIPE_MTE3, PIPE_MTE2, event_id);
187- 
188- wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
189- wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
190 }292 }
191 293 
294+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
295+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
192 pipe_barrier(PIPE_ALL);296 pipe_barrier(PIPE_ALL);
193 set_atomic_none();297 set_atomic_none();
194}298}
299+ 
300+ 
301+constexpr uint32_t CAXPY_MAX_DATA_COUNT = 38 * 1024 / sizeof(float);
302+constexpr uint32_t CAXPY_UB_COMPLEX_CAPACITY = CAXPY_MAX_DATA_COUNT / 2;
303+// CAXPY_STRIDED_TILE_COUNT is defined once in caxpy_dispatch_policy.h (shared
304+// with the host's shared-tile-capacity computation); this file must not
305+// redefine it, or the two could silently drift apart the way the host's
306+// GenXNormalizeOffsets 16-element compute-offset table drifted from the
307+// device's copy size before that was fixed.
308+static_assert(CAXPY_STRIDED_TILE_COUNT == 512, "CAXPY_STRIDED_TILE_COUNT definition moved or changed unexpectedly");
309+constexpr uint32_t CAXPY_INTERLEAVE_COUNT = 16;
310+constexpr uint32_t CAXPY_BLOCKED_FLOATS = 8;
311+constexpr uint32_t CAXPY_GATHER_PADDING = 64;
312+constexpr uint32_t CAXPY_X_OFFSET_COUNT = CAXPY_STRIDED_TILE_COUNT * CAXPY_BLOCKED_FLOATS + CAXPY_GATHER_PADDING;
313+constexpr uint32_t CAXPY_SPAN_OFFSET_COUNT = CAXPY_STRIDED_TILE_COUNT * 2 + CAXPY_GATHER_PADDING;
314+static_assert(
315+ CAXPY_DENSE_SWAP_SKIP_X_OFFSET_WORDS == CAXPY_X_OFFSET_COUNT,
316+ "CAXPY_DENSE_SWAP_OFFSET_START's skip term no longer matches CAXPY_X_OFFSET_COUNT");
317+static_assert(
318+ CAXPY_DENSE_SWAP_SKIP_SPAN_OFFSET_WORDS == 2 * CAXPY_SPAN_OFFSET_COUNT,
319+ "CAXPY_DENSE_SWAP_OFFSET_START's skip term no longer matches CAXPY_SPAN_OFFSET_COUNT");
320+constexpr uint32_t CAXPY_STORE_OFFSET_START = CAXPY_MAX_DATA_COUNT + CAXPY_X_OFFSET_COUNT;
321+constexpr uint32_t CAXPY_Y_GATHER_OFFSET_START = CAXPY_STORE_OFFSET_START + CAXPY_SPAN_OFFSET_COUNT;
322+ 
323+// --- Strided kernel UB byte layout -----------------------------------------
324+// Named byte offsets for the "shared" ping-pong buffer pair (strided_ub_layout,
325+// used directly by run_shared_pipeline slot 0, and by disjoint/shared variants
326+// via the overrides below). Each of x/work/y/lowScratch gets a 38 KiB region
327+// (CAXPY_DENSE_TILE_BYTES, reused here rather than duplicated: the strided
328+// blocked layout only needs CAXPY_STRIDED_TILE_COUNT*CAXPY_BLOCKED_FLOATS=16 KiB
329+// per region, but the historical layout budgets a full 38 KiB slot per region
330+// to match the dense kernel's tile size, leaving headroom after each xDense/
331+// workDense sub-offset). The offset-table region follows after 4 such slots.
332+constexpr uint32_t CAXPY_STRIDED_UB_X_BYTES = 0 * CAXPY_DENSE_TILE_BYTES;
333+constexpr uint32_t CAXPY_STRIDED_UB_WORK_BYTES = 1 * CAXPY_DENSE_TILE_BYTES;
334+constexpr uint32_t CAXPY_STRIDED_UB_Y_BYTES = 2 * CAXPY_DENSE_TILE_BYTES;
335+constexpr uint32_t CAXPY_STRIDED_UB_LOW_SCRATCH_BYTES = 3 * CAXPY_DENSE_TILE_BYTES;
336+constexpr uint32_t CAXPY_STRIDED_UB_OFFSET_TABLES_BYTES = 4 * CAXPY_DENSE_TILE_BYTES;
337+// Offset table sub-regions, laid out consecutively starting at
338+// CAXPY_STRIDED_UB_OFFSET_TABLES_BYTES: xNormalizeOffset, then computeOffset
339+// (sized for one CAXPY_INTERLEAVE_COUNT sub-chunk: convertRepeat<=2, one
340+// offset pair per element -> 2*8=16 words, rounded up to the 2*64 historical
341+// allocation for headroom), then yGatherOffset, then storeOffset.
342+constexpr uint32_t CAXPY_COMPUTE_OFFSET_COUNT = 2 * 64;
343+constexpr uint32_t CAXPY_STRIDED_UB_COMPUTE_OFFSET_BYTES =
344+ CAXPY_STRIDED_UB_OFFSET_TABLES_BYTES + CAXPY_X_OFFSET_COUNT * sizeof(uint32_t);
345+constexpr uint32_t CAXPY_STRIDED_UB_Y_GATHER_OFFSET_BYTES =
346+ CAXPY_STRIDED_UB_COMPUTE_OFFSET_BYTES + CAXPY_COMPUTE_OFFSET_COUNT * sizeof(uint32_t);
347+constexpr uint32_t CAXPY_STRIDED_UB_STORE_OFFSET_BYTES =
348+ CAXPY_STRIDED_UB_Y_GATHER_OFFSET_BYTES + CAXPY_SPAN_OFFSET_COUNT * sizeof(uint32_t);
349+constexpr uint32_t CAXPY_STRIDED_UB_Y_COMPACT_BYTES =
350+ CAXPY_STRIDED_UB_STORE_OFFSET_BYTES + CAXPY_SPAN_OFFSET_COUNT * sizeof(uint32_t);
351+static_assert(
352+ CAXPY_STRIDED_UB_OFFSET_TABLES_BYTES == 152 * 1024,
353+ "CAXPY strided offset-table region start moved; update disjoint/shared_ub_layout overrides");
354+static_assert(
355+ CAXPY_STRIDED_UB_Y_COMPACT_BYTES + CAXPY_STRIDED_TILE_COUNT * 2 * sizeof(float) <= CAXPY_UB_BUDGET_BYTES,
356+ "CAXPY strided UB layout exceeds the 910B3 Unified Buffer budget");
357+ 
358+struct CaxpyStridedUb {
359+ __ubuf__ float* x;
360+ __ubuf__ float* work;
361+ __ubuf__ float* y;
362+ __ubuf__ uint32_t* lowScratch;
363+ __ubuf__ float* xDense;
364+ __ubuf__ float* yDense;
365+ __ubuf__ float* workDense;
366+ __ubuf__ uint32_t* xNormalizeOffset;
367+ __ubuf__ uint32_t* computeOffset;
368+ __ubuf__ uint32_t* yGatherOffset;
369+ __ubuf__ uint32_t* storeOffset;
370+ __ubuf__ float* yCompact;
371+};
372+ 
373+__aicore__ __inline__ __attribute__((always_inline)) CaxpyStridedUb strided_ub_layout()
374+{
375+ CaxpyStridedUb ub;
376+ ub.x = reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_STRIDED_UB_X_BYTES);
377+ ub.work = reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_STRIDED_UB_WORK_BYTES);
378+ ub.y = reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_STRIDED_UB_Y_BYTES);
379+ ub.lowScratch = reinterpret_cast<__ubuf__ uint32_t*>((uintptr_t)CAXPY_STRIDED_UB_LOW_SCRATCH_BYTES);
380+ ub.xDense = ub.x + CAXPY_STRIDED_TILE_COUNT * CAXPY_BLOCKED_FLOATS;
381+ ub.yDense = ub.y + CAXPY_STRIDED_TILE_COUNT * CAXPY_BLOCKED_FLOATS;
382+ ub.workDense = ub.work + CAXPY_STRIDED_TILE_COUNT * CAXPY_BLOCKED_FLOATS;
383+ ub.xNormalizeOffset = reinterpret_cast<__ubuf__ uint32_t*>((uintptr_t)CAXPY_STRIDED_UB_OFFSET_TABLES_BYTES);
384+ ub.computeOffset = reinterpret_cast<__ubuf__ uint32_t*>((uintptr_t)CAXPY_STRIDED_UB_COMPUTE_OFFSET_BYTES);
385+ ub.yGatherOffset = reinterpret_cast<__ubuf__ uint32_t*>((uintptr_t)CAXPY_STRIDED_UB_Y_GATHER_OFFSET_BYTES);
386+ ub.storeOffset = reinterpret_cast<__ubuf__ uint32_t*>((uintptr_t)CAXPY_STRIDED_UB_STORE_OFFSET_BYTES);
387+ ub.yCompact = reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_STRIDED_UB_Y_COMPACT_BYTES);
388+ return ub;
389+}
390+ 
391+__aicore__ __inline__ __attribute__((always_inline)) void strided_preload_offsets(
392+ CaxpyStridedUb& ub, __gm__ uint32_t* gmAug, bool useXSpanGather, bool useYSpanGather)
393+{
394+ copy_vec_gm2ub_uint32(ub.computeOffset, gmAug, CAXPY_COMPUTE_OFFSET_COUNT);
395+ if (useXSpanGather) {
396+ copy_vec_gm2ub_uint32(ub.xNormalizeOffset, gmAug + CAXPY_MAX_DATA_COUNT, CAXPY_X_OFFSET_COUNT);
397+ }
398+ if (useYSpanGather) {
399+ copy_vec_gm2ub_uint32(ub.yGatherOffset, gmAug + CAXPY_Y_GATHER_OFFSET_START, CAXPY_SPAN_OFFSET_COUNT);
400+ copy_vec_gm2ub_uint32(ub.storeOffset, gmAug + CAXPY_STORE_OFFSET_START, CAXPY_SPAN_OFFSET_COUNT);
401+ }
402+}
403+ 
404+__aicore__ __inline__ __attribute__((always_inline)) bool strided_load_x_span(
405+ CaxpyStridedUb& ub, __gm__ float* gmX, uint32_t logicalOffset, uint32_t totalCount, uint32_t count, int64_t incx,
406+ bool useSpanGather)
407+{
408+ if (!useSpanGather) {
409+ return copy_complex_gm2ub_signed(ub.x, ub.xDense, gmX, logicalOffset, totalCount, count, incx);
410+ }
411+ const uint32_t increment = abs_increment(incx);
412+ const uint32_t physicalOffset =
413+ incx > 0 ? logicalOffset * increment : (totalCount - logicalOffset - count) * increment;
414+ copy_vec_gm2ub(
415+ reinterpret_cast<__ubuf__ float*>(ub.lowScratch), gmX + physicalOffset * 2, (count - 1) * increment * 2 + 2);
416+ return false;
417+}
418+ 
419+__aicore__ __inline__ __attribute__((always_inline)) void strided_normalize_inputs(
420+ CaxpyStridedUb& ub, uint32_t count, int64_t incx, int64_t incy, bool xDense, bool yDense, bool useXSpanGather,
421+ bool useYSpanGather)
422+{
423+ if (useXSpanGather) {
424+ vgather(
425+ reinterpret_cast<__ubuf__ uint32_t*>(ub.x), ub.xNormalizeOffset, (uintptr_t)ub.lowScratch, 8,
426+ (count + 7) / 8);
427+ }
428+ if (useYSpanGather) {
429+ vgather(
430+ reinterpret_cast<__ubuf__ uint32_t*>(ub.yCompact), ub.yGatherOffset, (uintptr_t)ub.y, 8,
431+ (count * 2 + 63) / 64);
432+ }
433+ const bool normalizeX = xDense || incx < 0;
434+ const bool normalizeY = incy != 1 && !useYSpanGather && (yDense || incy < 0);
435+ if (!normalizeX && !normalizeY)
436+ return;
437+ set_flag(PIPE_V, PIPE_S, EVENT_ID0);
438+ wait_flag(PIPE_V, PIPE_S, EVENT_ID0);
439+ if (normalizeX)
440+ normalize_loaded_complexes(ub.x, ub.xDense, count, xDense, incx < 0);
441+ if (normalizeY)
442+ normalize_loaded_complexes(ub.y, ub.yDense, count, yDense, incy < 0);
443+ set_flag(PIPE_S, PIPE_V, EVENT_ID0);
444+ wait_flag(PIPE_S, PIPE_V, EVENT_ID0);
445+}
446+ 
447+__aicore__ __inline__ __attribute__((always_inline)) void strided_compute_tile(
448+ CaxpyStridedUb& ub, float real, float imag, uint32_t count, int64_t incy, bool useYSpanGather)
449+{
450+ auto workReal = ub.workDense;
451+ auto workImag = workReal + CAXPY_UB_COMPLEX_CAPACITY;
452+ for (uint32_t converted = 0; converted < count; converted += CAXPY_INTERLEAVE_COUNT) {
453+ const uint32_t convertCount =
454+ count - converted > CAXPY_INTERLEAVE_COUNT ? CAXPY_INTERLEAVE_COUNT : count - converted;
455+ const uint32_t convertRepeat = (convertCount + 7) / 8;
456+ const uint32_t computeRepeat = (convertCount + 63) / 64;
457+ const uint32_t blockedOffset = converted * CAXPY_BLOCKED_FLOATS;
458+ auto xReal = ub.x + blockedOffset;
459+ auto xImag = xReal + CAXPY_UB_COMPLEX_CAPACITY;
460+ vreducev2(
461+ reinterpret_cast<__ubuf__ uint32_t*>(workReal), reinterpret_cast<__ubuf__ uint32_t*>(xReal), nullptr,
462+ convertRepeat, 1, 1, 8, 8);
463+ vreducev2(
464+ reinterpret_cast<__ubuf__ uint32_t*>(workImag), reinterpret_cast<__ubuf__ uint32_t*>(xReal), nullptr,
465+ convertRepeat, 1, 2, 8, 8);
466+ pipe_barrier(PIPE_V);
467+ caxpy_compute_planar(xReal, xImag, workReal, workImag, real, imag, computeRepeat);
468+ vgather(
469+ reinterpret_cast<__ubuf__ uint32_t*>(ub.work + blockedOffset), ub.computeOffset, (uintptr_t)xReal, 8,
470+ convertRepeat);
471+ pipe_barrier(PIPE_V);
472+ if (incy != 1 && !useYSpanGather) {
473+ vadd(
474+ ub.work + blockedOffset, ub.work + blockedOffset, ub.y + blockedOffset, convertRepeat, 1, 1, 1, 8, 8,
475+ 8);
476+ }
477+ }
478+}
479+ 
480+__aicore__ __inline__ __attribute__((always_inline)) CaxpyStridedUb disjoint_ub_layout(uint32_t index)
481+{
482+ CaxpyStridedUb ub = strided_ub_layout();
483+ ub.x = reinterpret_cast<__ubuf__ float*>(
484+ (uintptr_t)(index == 0 ? CAXPY_STRIDED_UB_X_BYTES : CAXPY_STRIDED_UB_WORK_BYTES));
485+ ub.work = reinterpret_cast<__ubuf__ float*>(
486+ (uintptr_t)(index == 0 ? CAXPY_STRIDED_UB_Y_BYTES : CAXPY_STRIDED_UB_LOW_SCRATCH_BYTES));
487+ ub.xDense = ub.x + CAXPY_STRIDED_TILE_COUNT * CAXPY_BLOCKED_FLOATS;
488+ ub.workDense = ub.work + CAXPY_STRIDED_TILE_COUNT * CAXPY_BLOCKED_FLOATS;
489+ return ub;
490+}
491+ 
492+__aicore__ __inline__ __attribute__((always_inline)) void disjoint_load_tile(
493+ CaxpyStridedUb& ub, __gm__ float* gmX, uint32_t logicalOffset, uint32_t totalCount, uint32_t count,
494+ int64_t incx, uint32_t eventId)
495+{
496+ copy_complex_gm2ub_signed(ub.x, ub.xDense, gmX, logicalOffset, totalCount, count, incx);
497+ set_flag(PIPE_MTE2, PIPE_V, eventId);
498+}
499+ 
500+__aicore__ __inline__ __attribute__((always_inline)) void disjoint_compute_tile(
501+ CaxpyStridedUb& ub, float alphaReal, float alphaImag, uint32_t count, int64_t incx, uint32_t eventId)
502+{
503+ wait_flag(PIPE_MTE2, PIPE_V, eventId);
504+ strided_normalize_inputs(ub, count, incx, 1, false, false, false, false);
505+ strided_compute_tile(ub, alphaReal, alphaImag, count, 1, false);
506+}
507+ 
508+__aicore__ __inline__ __attribute__((always_inline)) void strided_store_tile(
509+ CaxpyStridedUb& ub, __gm__ float* gmY, uint32_t logicalOffset, uint32_t totalCount, uint32_t count,
510+ int64_t incy, uint32_t eventId)
511+{
512+ __ubuf__ float* source = ub.work;
513+ if (incy == 1) {
514+ const uint32_t repeat = (count * 2 + 63) / 64;
515+ vgather(
516+ reinterpret_cast<__ubuf__ uint32_t*>(ub.workDense), ub.storeOffset, (uintptr_t)ub.work, 8, repeat);
517+ pipe_barrier(PIPE_V);
518+ source = ub.workDense;
519+ }
520+ set_flag(PIPE_V, PIPE_MTE3, eventId);
521+ wait_flag(PIPE_V, PIPE_MTE3, eventId);
522+ copy_complex_ub2gm_signed(gmY, source, logicalOffset, totalCount, count, incy);
523+ set_flag(PIPE_MTE3, PIPE_MTE2, eventId);
524+}
525+ 
526+__aicore__ __inline__ __attribute__((always_inline)) void disjoint_preload_offsets(
527+ CaxpyStridedUb& ping, __gm__ uint32_t* gmAug, int64_t incy)
528+{
529+ copy_vec_gm2ub_uint32(ping.computeOffset, gmAug, CAXPY_COMPUTE_OFFSET_COUNT);
530+ if (incy == 1)
531+ copy_vec_gm2ub_uint32(
532+ ping.storeOffset, gmAug + CAXPY_STORE_OFFSET_START, CAXPY_SPAN_OFFSET_COUNT);
533+}
534+ 
535+__aicore__ __inline__ __attribute__((always_inline)) void disjoint_load_next_if_any(
536+ CaxpyStridedUb& ping, CaxpyStridedUb& pong, __gm__ float* gmX, uint32_t logicalOffset,
537+ uint32_t totalComplexCount, uint32_t complexCount, uint32_t tileCapacity, int64_t incx, uint32_t tile,
538+ uint32_t tileCount, uint32_t current)
539+{
540+ if (tile + 1 >= tileCount)
541+ return;
542+ const uint32_t next = 1 - current;
543+ const uint32_t nextEvent = next == 0 ? EVENT_ID0 : EVENT_ID1;
544+ CaxpyStridedUb& nextUb = next == 0 ? ping : pong;
545+ const uint32_t nextOffset = (tile + 1) * tileCapacity;
546+ const uint32_t nextCount =
547+ complexCount - nextOffset > tileCapacity ? tileCapacity : complexCount - nextOffset;
548+ wait_flag(PIPE_MTE3, PIPE_MTE2, nextEvent);
549+ disjoint_load_tile(
550+ nextUb, gmX, logicalOffset + nextOffset, totalComplexCount, nextCount, incx, nextEvent);
551+}
552+ 
553+__aicore__ __inline__ __attribute__((always_inline)) void run_disjoint_pipeline(
554+ __gm__ float* gmX, __gm__ uint32_t* gmAug, __gm__ float* gmY, float alphaReal, float alphaImag,
555+ uint32_t logicalOffset, uint32_t totalComplexCount, uint32_t complexCount, int64_t incx, int64_t incy)
556+{
557+ CaxpyStridedUb ping = disjoint_ub_layout(0);
558+ CaxpyStridedUb pong = disjoint_ub_layout(1);
559+ disjoint_preload_offsets(ping, gmAug, incy);
560+ 
561+ set_atomic_add();
562+ set_atomic_f32();
563+ set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
564+ set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
565+ 
566+ constexpr uint32_t tileCapacity = CAXPY_STRIDED_TILE_COUNT;
567+ const uint32_t tileCount = (complexCount + tileCapacity - 1) / tileCapacity;
568+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
569+ const uint32_t firstCount = complexCount > tileCapacity ? tileCapacity : complexCount;
570+ disjoint_load_tile(ping, gmX, logicalOffset, totalComplexCount, firstCount, incx, EVENT_ID0);
571+ 
572+ for (uint32_t tile = 0; tile < tileCount; ++tile) {
573+ const uint32_t current = tile & 1;
574+ const uint32_t currentEvent = current == 0 ? EVENT_ID0 : EVENT_ID1;
575+ CaxpyStridedUb& currentUb = current == 0 ? ping : pong;
576+ const uint32_t currentOffset = tile * tileCapacity;
577+ const uint32_t currentCount =
578+ complexCount - currentOffset > tileCapacity ? tileCapacity : complexCount - currentOffset;
579+ 
580+ disjoint_load_next_if_any(
581+ ping, pong, gmX, logicalOffset, totalComplexCount, complexCount, tileCapacity, incx, tile, tileCount,
582+ current);
583+ 
584+ disjoint_compute_tile(currentUb, alphaReal, alphaImag, currentCount, incx, currentEvent);
585+ strided_store_tile(
586+ currentUb, gmY, logicalOffset + currentOffset, totalComplexCount, currentCount, incy, currentEvent);
587+ }
588+ 
589+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
590+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
591+ pipe_barrier(PIPE_ALL);
592+ set_atomic_none();
593+}
594+ 
595+__aicore__ __inline__ __attribute__((always_inline)) CaxpyStridedUb shared_ub_layout(uint32_t index)
596+{
597+ CaxpyStridedUb ub = strided_ub_layout();
598+ if (index != 0) {
599+ ub.work = reinterpret_cast<__ubuf__ float*>((uintptr_t)CAXPY_STRIDED_UB_Y_BYTES);
600+ ub.workDense = ub.work + CAXPY_STRIDED_TILE_COUNT * CAXPY_BLOCKED_FLOATS;
601+ }
602+ return ub;
603+}
604+ 
605+__aicore__ __inline__ __attribute__((always_inline)) void shared_load_tile(
606+ CaxpyStridedUb& ub, __gm__ float* gmX, uint32_t logicalOffset, uint32_t totalCount, uint32_t count,
607+ int64_t incx, uint32_t waitEvent, uint32_t readyEvent, bool waitForRaw)
608+{
609+ // MTE3 owns each work buffer after a store. Hand it back through MTE2,
610+ // which is the event direction used by the stable dense pipeline.
611+ wait_flag(PIPE_MTE3, PIPE_MTE2, readyEvent);
612+ if (waitForRaw)
613+ wait_flag(PIPE_V, PIPE_MTE2, waitEvent);
614+ strided_load_x_span(ub, gmX, logicalOffset, totalCount, count, incx, true);
615+ set_flag(PIPE_MTE2, PIPE_V, readyEvent);
616+}
617+ 
618+__aicore__ __inline__ __attribute__((always_inline)) void shared_compute_tile(
619+ CaxpyStridedUb& ub, float alphaReal, float alphaImag, uint32_t count, int64_t incx, uint32_t eventId)
620+{
621+ wait_flag(PIPE_MTE2, PIPE_V, eventId);
622+ strided_normalize_inputs(ub, count, incx, 1, false, false, true, false);
623+ // raw span is dead immediately after gather/normalization. Release it to
624+ // MTE2 before the remaining complex arithmetic finishes on PIPE_V.
625+ set_flag(PIPE_V, PIPE_MTE2, eventId);
626+ strided_compute_tile(ub, alphaReal, alphaImag, count, 1, false);
627+}
628+ 
629+__aicore__ __inline__ __attribute__((always_inline)) void shared_tile_bounds(
630+ uint32_t complexCount, uint32_t tileCount, uint32_t tile, uint32_t tileCapacity, uint32_t& tileOffset,
631+ uint32_t& tileLength)
632+{
633+ // A final tile shorter than the interleave width is mis-staged by the span-read
634+ // path, so avoid producing one: when the natural schedule would leave such a
635+ // remainder, split the last two tiles evenly instead. Coverage stays exact and
636+ // contiguous and no tile exceeds the capacity, so buffer sizing is unchanged.
637+ // tileCapacity is a runtime value (see run_shared_pipeline) that is always
638+ // >= CAXPY_SHARED_MIN_USABLE_TILE_CAPACITY (CAXPY_INTERLEAVE_COUNT + 1), so
639+ // this rebalancing always has room to produce two tiles of at least
640+ // CAXPY_INTERLEAVE_COUNT/2 elements each.
641+ const uint32_t remainder = complexCount - (tileCount - 1) * tileCapacity;
642+ if (tileCount >= 2 && remainder < CAXPY_INTERLEAVE_COUNT + 1) {
643+ const uint32_t splitBase = (tileCount - 2) * tileCapacity;
644+ if (tile < tileCount - 2) {
645+ tileOffset = tile * tileCapacity;
646+ tileLength = tileCapacity;
647+ return;
648+ }
649+ const uint32_t splitTotal = complexCount - splitBase;
650+ const uint32_t firstHalf = (splitTotal + 1) / 2;
651+ if (tile == tileCount - 2) {
652+ tileOffset = splitBase;
653+ tileLength = firstHalf;
654+ } else {
655+ tileOffset = splitBase + firstHalf;
656+ tileLength = splitTotal - firstHalf;
657+ }
658+ return;
659+ }
660+ tileOffset = tile * tileCapacity;
661+ tileLength = complexCount - tileOffset > tileCapacity ? tileCapacity : complexCount - tileOffset;
662+}
663+ 
664+__aicore__ __inline__ __attribute__((always_inline)) void shared_preload_offsets(
665+ CaxpyStridedUb& ping, __gm__ uint32_t* gmAug, int64_t incy)
666+{
667+ strided_preload_offsets(ping, gmAug, true, false);
668+ if (incy == 1)
669+ copy_vec_gm2ub_uint32(
670+ ping.storeOffset, gmAug + CAXPY_STORE_OFFSET_START, CAXPY_SPAN_OFFSET_COUNT);
671+}
672+ 
673+__aicore__ __inline__ __attribute__((always_inline)) void shared_load_next_if_any(
674+ CaxpyStridedUb& ping, CaxpyStridedUb& pong, __gm__ float* gmX, uint32_t logicalOffset,
675+ uint32_t totalComplexCount, uint32_t complexCount, uint32_t tileCapacity, int64_t incx, uint32_t tile,
676+ uint32_t tileCount, uint32_t current)
677+{
678+ if (tile + 1 >= tileCount)
679+ return;
680+ const uint32_t next = 1 - current;
681+ const uint32_t nextEvent = next == 0 ? EVENT_ID0 : EVENT_ID1;
682+ const uint32_t currentEvent = current == 0 ? EVENT_ID0 : EVENT_ID1;
683+ CaxpyStridedUb& nextUb = next == 0 ? ping : pong;
684+ uint32_t nextOffset = 0;
685+ uint32_t nextCount = 0;
686+ shared_tile_bounds(complexCount, tileCount, tile + 1, tileCapacity, nextOffset, nextCount);
687+ shared_load_tile(
688+ nextUb, gmX, logicalOffset + nextOffset, totalComplexCount, nextCount, incx, currentEvent, nextEvent, true);
689+}
690+ 
691+__aicore__ __inline__ __attribute__((always_inline)) void run_shared_pipeline(
692+ __gm__ float* gmX, __gm__ uint32_t* gmAug, __gm__ float* gmY, float alphaReal, float alphaImag,
693+ uint32_t logicalOffset, uint32_t totalComplexCount, uint32_t complexCount, int64_t incx, int64_t incy,
694+ uint32_t tileCapacity)
695+{
696+ CaxpyStridedUb ping = shared_ub_layout(0);
697+ CaxpyStridedUb pong = shared_ub_layout(1);
698+ shared_preload_offsets(ping, gmAug, incy);
699+ 
700+ set_atomic_add();
701+ set_atomic_f32();
702+ set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
703+ set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
704+ 
705+ const uint32_t tileCount = (complexCount + tileCapacity - 1) / tileCapacity;
706+ uint32_t firstOffset = 0;
707+ uint32_t firstCount = 0;
708+ shared_tile_bounds(complexCount, tileCount, 0, tileCapacity, firstOffset, firstCount);
709+ shared_load_tile(
710+ ping, gmX, logicalOffset + firstOffset, totalComplexCount, firstCount, incx, EVENT_ID0, EVENT_ID0, false);
711+ 
712+ for (uint32_t tile = 0; tile < tileCount; ++tile) {
713+ const uint32_t current = tile & 1;
714+ const uint32_t currentEvent = current == 0 ? EVENT_ID0 : EVENT_ID1;
715+ CaxpyStridedUb& currentUb = current == 0 ? ping : pong;
716+ uint32_t currentOffset = 0;
717+ uint32_t currentCount = 0;
718+ shared_tile_bounds(complexCount, tileCount, tile, tileCapacity, currentOffset, currentCount);
719+ 
720+ shared_compute_tile(currentUb, alphaReal, alphaImag, currentCount, incx, currentEvent);
721+ 
722+ shared_load_next_if_any(
723+ ping, pong, gmX, logicalOffset, totalComplexCount, complexCount, tileCapacity, incx, tile, tileCount,
724+ current);
725+ 
726+ strided_store_tile(
727+ currentUb, gmY, logicalOffset + currentOffset, totalComplexCount, currentCount, incy, currentEvent);
728+ }
729+ 
730+ const uint32_t lastRawEvent = ((tileCount - 1) & 1) == 0 ? EVENT_ID0 : EVENT_ID1;
731+ wait_flag(PIPE_V, PIPE_MTE2, lastRawEvent);
732+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
733+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID1);
734+ pipe_barrier(PIPE_ALL);
735+ set_atomic_none();
736+}
737+ 
738+__aicore__ __inline__ __attribute__((always_inline)) void caxpy_strided_shared_run(
739+ __gm__ float* gmX, __gm__ uint32_t* gmAug, __gm__ float* gmY, float alphaReal, float alphaImag,
740+ uint32_t logicalOffset, uint32_t totalComplexCount, uint32_t complexCount, int64_t incx, int64_t incy,
741+ uint32_t tileCapacity)
742+{
743+ run_shared_pipeline(
744+ gmX, gmAug, gmY, alphaReal, alphaImag, logicalOffset, totalComplexCount, complexCount, incx, incy,
745+ tileCapacity);
746+}
747+ 
748+__aicore__ __inline__ __attribute__((always_inline)) void caxpy_strided_disjoint_run(
749+ __gm__ float* gmX, __gm__ uint32_t* gmAug, __gm__ float* gmY, float alphaReal, float alphaImag,
750+ uint32_t logicalOffset, uint32_t totalComplexCount, uint32_t complexCount, int64_t incx, int64_t incy,
751+ uint32_t tileCapacity)
752+{
753+ // The disjoint pipeline never span-reads, so it has no UB-buffer-fit
754+ // constraint to shrink its tile for; it always uses the fixed
755+ // CAXPY_STRIDED_TILE_COUNT capacity and ignores the runtime tileCapacity
756+ // value threaded through only so both kernel entries share one signature.
757+ (void)tileCapacity;
758+ run_disjoint_pipeline(
759+ gmX, gmAug, gmY, alphaReal, alphaImag, logicalOffset, totalComplexCount, complexCount, incx, incy);
760+}
761+ 
195#endif762#endif
196 763 
197-extern "C" __global__ __aicore__ __vector__ void caxpy(__gm__ float *__restrict__ x, __gm__ uint32_t *__restrict__ aug,764+extern "C" __global__ __aicore__ __vector__ void caxpy(
198- __gm__ float *__restrict__ y, __gm__ float *__restrict__ tiling_gm)765+ __gm__ float* __restrict__ x, __gm__ uint32_t* __restrict__ aug, __gm__ float* __restrict__ y,
766+ __gm__ float* __restrict__ tiling_gm)
199{767{
200#if __DAV_C220_VEC__768#if __DAV_C220_VEC__
201 auto vec_idx = AscendC::GetBlockIdx();769 auto vec_idx = AscendC::GetBlockIdx();
@@ -203,25 +771,156 @@ extern "C" __global__ __aicore__ __vector__ void caxpy(__gm__ float *__restrict_
203 set_mask_norm();771 set_mask_norm();
204 set_vector_mask((uint64_t)-1, (uint64_t)-1);772 set_vector_mask((uint64_t)-1, (uint64_t)-1);
205 773 
206- auto tiling_buf = reinterpret_cast<__gm__ uint8_t *>(tiling_gm);774+ auto tiling = reinterpret_cast<__gm__ CaxpyTilingData*>(tiling_gm);
207 775 
208- uint32_t n = (*(__gm__ uint32_t *)((__gm__ uint8_t *)tiling_buf));776+ const uint32_t numBlocks = tiling->activeBlocks;
209- float alphaReal = (*(__gm__ float *)((__gm__ uint8_t *)tiling_buf + 4));777+ const uint32_t totalN = tiling->totalN;
210- float alphaImag = (*(__gm__ float *)((__gm__ uint8_t *)tiling_buf + 8));778+ if (numBlocks == 0 || vec_idx >= numBlocks) {
211- 779+ return;
212- uint32_t offset = (*(__gm__ uint32_t *)((__gm__ uint8_t *)tiling_buf + 12 + 4 * vec_idx));
213- uint32_t calNum = (*(__gm__ uint32_t *)((__gm__ uint8_t *)tiling_buf + 12 + 40 * 4 + 4 * vec_idx));
214- 
215- if (calNum > 0) {
216- caxpy_aiv(x, aug, y, alphaReal, alphaImag, offset, calNum);
217 }780 }
781+ const uint32_t perBlock = totalN / numBlocks;
782+ const uint32_t remainder = totalN % numBlocks;
783+ const uint32_t complexCount = perBlock + (vec_idx < remainder ? 1 : 0);
784+ const uint32_t complexOffset = vec_idx * perBlock + (vec_idx < remainder ? vec_idx : remainder);
785+ caxpy_dense_run(
786+ x, aug, y, tiling->alphaReal, tiling->alphaImag, complexOffset * 2, complexCount * 2);
218#endif787#endif
219}788}
220 789 
221-void caxpy_kernel_do(GM_ADDR x, GM_ADDR maskBuf, GM_ADDR y, GM_ADDR workSpace, GM_ADDR tilingGm,790+#if __DAV_C220_VEC__
222- uint32_t numBlocks, void *stream)791+#define DEFINE_CAXPY_STRIDED_KERNEL(KERNEL, PIPELINE) \
792+extern "C" __global__ __aicore__ __vector__ void KERNEL( \
793+ __gm__ float* __restrict__ x, __gm__ uint32_t* __restrict__ aug, __gm__ float* __restrict__ y, \
794+ __gm__ float* __restrict__ tilingGm) \
795+{ \
796+ const uint32_t vecIdx = AscendC::GetBlockIdx(); \
797+ set_mask_norm(); set_vector_mask((uint64_t)-1, (uint64_t)-1); \
798+ auto tiling = reinterpret_cast<__gm__ CaxpyTilingData*>(tilingGm); \
799+ const uint32_t numBlocks = tiling->activeBlocks; \
800+ const float alphaReal = tiling->alphaReal; \
801+ const float alphaImag = tiling->alphaImag; \
802+ const uint32_t totalN = tiling->totalN; \
803+ const int64_t incx = tiling->incx; \
804+ const int64_t incy = tiling->incy; \
805+ const uint32_t tileCapacity = tiling->sharedTileCapacity; \
806+ if (numBlocks == 0) { return; } \
807+ const uint32_t waveCapacity = numBlocks * CAXPY_STRIDED_MAX_PER_BLOCK_PER_WAVE; \
808+ for (uint32_t logicalBase = 0; logicalBase < totalN; logicalBase += waveCapacity) { \
809+ const uint32_t waveN = totalN - logicalBase > waveCapacity ? waveCapacity : totalN - logicalBase; \
810+ const uint32_t waveBlocks = numBlocks < waveN ? numBlocks : waveN; \
811+ if (vecIdx >= waveBlocks) { continue; } \
812+ const uint32_t rowNum = waveN / waveBlocks; const uint32_t remain = waveN % waveBlocks; \
813+ const uint32_t myCount = rowNum + (vecIdx < remain ? 1 : 0); \
814+ const uint32_t myStart = logicalBase + vecIdx * rowNum + (vecIdx < remain ? vecIdx : remain); \
815+ if (incy > 0) { PIPELINE(x, aug, y, alphaReal, alphaImag, myStart, totalN, myCount, incx, incy, tileCapacity); } \
816+ else { const uint32_t mirroredStart = totalN - myStart - myCount; PIPELINE(x, aug, y, alphaReal, alphaImag, mirroredStart, totalN, myCount, -incx, -incy, tileCapacity); } \
817+ } \
818+}
819+#else
820+#define DEFINE_CAXPY_STRIDED_KERNEL(KERNEL, PIPELINE) \
821+extern "C" __global__ __aicore__ __vector__ void KERNEL( \
822+ __gm__ float*, __gm__ uint32_t*, __gm__ float*, __gm__ float*) {}
823+#endif
824+DEFINE_CAXPY_STRIDED_KERNEL(caxpy_strided_shared, caxpy_strided_shared_run)
825+DEFINE_CAXPY_STRIDED_KERNEL(caxpy_strided_disjoint, caxpy_strided_disjoint_run)
826+#undef DEFINE_CAXPY_STRIDED_KERNEL
827+ 
828+void caxpy_kernel_do(
829+ GM_ADDR x, GM_ADDR maskBuf, GM_ADDR y, GM_ADDR workSpace, GM_ADDR tilingGm, uint32_t numBlocks, void* stream)
223{830{
224- caxpy<<<numBlocks, nullptr, stream>>>((float *)x,831+ caxpy<<<numBlocks, nullptr, stream>>>((float*)x, (uint32_t*)maskBuf, (float*)y, (float*)tilingGm);
225- (uint32_t *)maskBuf,832+}
226- (float *)y, (float *)tilingGm);833+ 
227-}834+void caxpy_strided_shared_kernel_do(
835+ GM_ADDR x, GM_ADDR maskBuf, GM_ADDR y, GM_ADDR workSpace, GM_ADDR tilingGm, uint32_t numBlocks, void* stream)
836+{
837+ caxpy_strided_shared<<<numBlocks, nullptr, stream>>>((float*)x, (uint32_t*)maskBuf, (float*)y, (float*)tilingGm);
838+}
839+ 
840+void caxpy_strided_disjoint_kernel_do(
841+ GM_ADDR x, GM_ADDR maskBuf, GM_ADDR y, GM_ADDR workSpace, GM_ADDR tilingGm, uint32_t numBlocks, void* stream)
842+{
843+ caxpy_strided_disjoint<<<numBlocks, nullptr, stream>>>(
844+ (float*)x, (uint32_t*)maskBuf, (float*)y, (float*)tilingGm);
845+}
846+ 
847+#if __DAV_C220_VEC__
848+__aicore__ __inline__ __attribute__((always_inline)) void scalar_copy_complex_in(
849+ __ubuf__ float* dst, __gm__ float* src)
850+{
851+ copy_gm_to_ubuf_align_b32(dst, src, 0, 1, 2 * sizeof(float), 0, 0, 0, 0);
852+}
853+ 
854+__aicore__ __inline__ __attribute__((always_inline)) void scalar_copy_complex_out(
855+ __gm__ float* dst, __ubuf__ float* src)
856+{
857+ copy_ubuf_to_gm_align_b32(dst, src, 0, 1, 2 * sizeof(float), 0, 0, 0, 0);
858+}
859+ 
860+__aicore__ __inline__ __attribute__((always_inline)) void scalar_compute_all(
861+ __gm__ float* gmX, __gm__ float* gmY, float alphaReal, float alphaImag, uint32_t logicalOffset,
862+ uint32_t totalComplexCount, uint32_t complexCount, int64_t incx, int64_t incy)
863+{
864+ const uint32_t absIncx = abs_increment(incx);
865+ const uint32_t absIncy = abs_increment(incy);
866+ auto ubX = reinterpret_cast<__ubuf__ float*>((uintptr_t)0);
867+ auto ubY = reinterpret_cast<__ubuf__ float*>((uintptr_t)32);
868+ 
869+ set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
870+ for (uint32_t local = 0; local < complexCount; ++local) {
871+ const uint32_t logical = logicalOffset + local;
872+ const uint32_t xPhysical =
873+ incx > 0 ? logical * absIncx : (totalComplexCount - 1 - logical) * absIncx;
874+ const uint32_t yPhysical =
875+ incy > 0 ? logical * absIncy : (totalComplexCount - 1 - logical) * absIncy;
876+ 
877+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
878+ scalar_copy_complex_in(ubX, gmX + xPhysical * 2);
879+ scalar_copy_complex_in(ubY, gmY + yPhysical * 2);
880+ set_flag(PIPE_MTE2, PIPE_S, EVENT_ID0);
881+ wait_flag(PIPE_MTE2, PIPE_S, EVENT_ID0);
882+ 
883+ const float xReal = ubX[0];
884+ const float xImag = ubX[1];
885+ const float yReal = ubY[0];
886+ const float yImag = ubY[1];
887+ ubY[0] = yReal + alphaReal * xReal - alphaImag * xImag;
888+ ubY[1] = yImag + alphaReal * xImag + alphaImag * xReal;
889+ set_flag(PIPE_S, PIPE_MTE3, EVENT_ID0);
890+ wait_flag(PIPE_S, PIPE_MTE3, EVENT_ID0);
891+ scalar_copy_complex_out(gmY + yPhysical * 2, ubY);
892+ set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
893+ }
894+ wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0);
895+}
896+#endif
897+ 
898+extern "C" __global__ __aicore__ __vector__ void caxpy_strided_scalar(
899+ __gm__ float* __restrict__ x, __gm__ float* __restrict__ y, __gm__ float* __restrict__ tilingGm)
900+{
901+#if __DAV_C220_VEC__
902+ const uint32_t vecIdx = AscendC::GetBlockIdx();
903+ auto tiling = reinterpret_cast<__gm__ CaxpyTilingData*>(tilingGm);
904+ const uint32_t numBlocks = tiling->activeBlocks;
905+ const float alphaReal = tiling->alphaReal;
906+ const float alphaImag = tiling->alphaImag;
907+ const uint32_t totalN = tiling->totalN;
908+ const int64_t incx = tiling->incx;
909+ const int64_t incy = tiling->incy;
910+ if (numBlocks == 0 || vecIdx >= numBlocks) {
911+ return;
912+ }
913+ 
914+ const uint32_t rowNum = totalN / numBlocks;
915+ const uint32_t remain = totalN % numBlocks;
916+ const uint32_t count = rowNum + (vecIdx < remain ? 1 : 0);
917+ const uint32_t start = vecIdx * rowNum + (vecIdx < remain ? vecIdx : remain);
918+ scalar_compute_all(x, y, alphaReal, alphaImag, start, totalN, count, incx, incy);
919+#endif
920+}
921+ 
922+void caxpy_strided_scalar_kernel_do(
923+ GM_ADDR x, GM_ADDR y, GM_ADDR tilingGm, uint32_t numBlocks, void* stream)
924+{
925+ caxpy_strided_scalar<<<numBlocks, nullptr, stream>>>((float*)x, (float*)y, (float*)tilingGm);
926+}
Ablas/axpy/arch22/caxpy_kernel.h+27-0
@@ -0,0 +1,27 @@
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+void caxpy_kernel_do(
16+ uint8_t* x, uint8_t* maskBuf, uint8_t* y, uint8_t* workSpace, uint8_t* tilingGm, uint32_t numBlocks,
17+ void* stream);
18+ 
19+void caxpy_strided_shared_kernel_do(
20+ uint8_t* x, uint8_t* maskBuf, uint8_t* y, uint8_t* workSpace, uint8_t* tilingGm, uint32_t numBlocks,
21+ void* stream);
22+ 
23+void caxpy_strided_disjoint_kernel_do(
24+ uint8_t* x, uint8_t* maskBuf, uint8_t* y, uint8_t* workSpace, uint8_t* tilingGm, uint32_t numBlocks,
25+ void* stream);
26+ 
27+void caxpy_strided_scalar_kernel_do(uint8_t* x, uint8_t* y, uint8_t* tilingGm, uint32_t numBlocks, void* stream);
Ablas/axpy/arch22/caxpy_tiling_data.h+36-0
@@ -0,0 +1,36 @@
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+// Single source of truth for the CAXPY host->device tiling block. BOTH the host
16+// (caxpy_host.cpp) and the device kernel (caxpy_kernel.cpp) include this header
17+// and access fields by name. Do not re-declare this layout anywhere else and do
18+// not read it via hard-coded byte offsets -- that is exactly the drift hazard
19+// this shared definition exists to remove.
20+ 
21+struct CaxpyTilingData {
22+ uint32_t totalN;
23+ uint32_t activeBlocks;
24+ float alphaReal;
25+ float alphaImag;
26+ int64_t incx;
27+ int64_t incy;
28+ // Shared-pipeline span-read tile capacity, in complex elements. Computed by
29+ // the host from incx so that the span read
30+ // ((sharedTileCapacity-1)*|incx|+1) complex elements always fits in the
31+ // 38 KiB lowScratch UB buffer -- see CaxpySharedTileCapacity in
32+ // caxpy_dispatch_policy.h for the formula. Unused (left at its default of
33+ // CAXPY_STRIDED_TILE_COUNT) by the disjoint and scalar kernel variants,
34+ // which never span-read.
35+ uint32_t sharedTileCapacity;
36+};
Mblas/common/helper/aclblas_auxiliary.cpp+7-0
@@ -64,6 +64,13 @@ aclblasStatus_t aclblasDestroy(aclblasHandle_t handle)
64 64 
65 FreeLibraryWorkspace(h);65 FreeLibraryWorkspace(h);
66 66 
67+ if (h->caxpy_mask_cache != nullptr) {
68+ aclrtFree(h->caxpy_mask_cache);
69+ h->caxpy_mask_cache = nullptr;
70+ h->caxpy_mask_cache_incx = 0;
71+ h->caxpy_mask_cache_incy = 0;
72+ }
73+ 
67 h->stream = nullptr;74 h->stream = nullptr;
68 75 
69 delete h;76 delete h;
Mblas/common/helper/aclblas_handle_internal.h+9-0
@@ -55,6 +55,15 @@ struct _aclblas_handle {
55 55 
56 /** Size of the last library-managed workspace buffer (preserved across user switches). */56 /** Size of the last library-managed workspace buffer (preserved across user switches). */
57 size_t library_workspace_size = 0;57 size_t library_workspace_size = 0;
58+ 
59+ /* ========== Cached caxpy gather-offset buffer ========== */
60+ /** Device buffer holding the caxpy gather-offset table; owned by the handle and
61+ * built once per |incx| instead of being re-allocated and re-uploaded per call. */
62+ void* caxpy_mask_cache = nullptr;
63+ /** The |incx| value the cached mask was built for; 0 means the cache content is invalid. */
64+ uint64_t caxpy_mask_cache_incx = 0;
65+ /** The |incy| value the cached mask was built for; 0 means the cache content is invalid. */
66+ uint64_t caxpy_mask_cache_incy = 0;
58};67};
59 68 
60/**69/**
Mtest/axpy/caxpy/CMakeLists.txt+11-1
@@ -8,4 +8,14 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------------------------------------9# ----------------------------------------------------------------------------------------------------------
10 10 
11-ops_blas_add_tests(${OPS_BLAS})11+ops_blas_add_gtest_tests(
12+ ${OPS_BLAS}
13+ caxpy_test)
14+ 
15+ops_blas_add_tests(
16+ ${OPS_BLAS}
17+ caxpy_dispatch_policy_test)
18+ 
19+target_include_directories(
20+ caxpy_dispatch_policy_test PRIVATE
21+ ${CMAKE_SOURCE_DIR}/blas/axpy/arch22)
Atest/axpy/caxpy/arch22/caxpy_dispatch_policy_test.cpp+121-0
@@ -0,0 +1,121 @@
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 <array>
12+#include <cstdint>
13+#include <iostream>
14+ 
15+#include "caxpy_dispatch_policy.h"
16+ 
17+namespace {
18+struct ClassificationCase {
19+ int64_t incx;
20+ int64_t incy;
21+ CaxpyStrideClass expected;
22+};
23+ 
24+struct PolicyCase {
25+ uint64_t n;
26+ int64_t incx;
27+ int64_t incy;
28+ bool directStrided;
29+ uint32_t numBlocks;
30+ CaxpyKernelVariant expected;
31+};
32+ 
33+ 
34+constexpr std::array<ClassificationCase, 15> CASES{{
35+ {1, 1, CaxpyStrideClass::DENSE},
36+ {2, 1, CaxpyStrideClass::X_ONLY_POSITIVE},
37+ {9, 1, CaxpyStrideClass::X_ONLY_POSITIVE},
38+ {1, 2, CaxpyStrideClass::Y_ONLY_POSITIVE},
39+ {1, 10, CaxpyStrideClass::Y_ONLY_POSITIVE},
40+ {2, 2, CaxpyStrideClass::DUAL_SYMMETRIC_POSITIVE},
41+ {10, 10, CaxpyStrideClass::DUAL_SYMMETRIC_POSITIVE},
42+ {2, 3, CaxpyStrideClass::DUAL_ASYMMETRIC_POSITIVE},
43+ {10, 11, CaxpyStrideClass::DUAL_ASYMMETRIC_POSITIVE},
44+ {-1, 1, CaxpyStrideClass::NEGATIVE_MIXED},
45+ {1, -1, CaxpyStrideClass::NEGATIVE_MIXED},
46+ {-2, 3, CaxpyStrideClass::NEGATIVE_MIXED},
47+ {5, -3, CaxpyStrideClass::NEGATIVE_MIXED},
48+ {-2, -3, CaxpyStrideClass::NEGATIVE_MIXED},
49+ {-10, -10, CaxpyStrideClass::NEGATIVE_MIXED},
50+}};
51+ 
52+constexpr std::array<PolicyCase, 23> POLICY_CASES{{
53+ {1, 1, 1, false, 1, CaxpyKernelVariant::DENSE_PIPELINED_ATOMIC},
54+ {8192, 1, 1, false, 40, CaxpyKernelVariant::DENSE_PIPELINED_ATOMIC},
55+ {640, 2, 1, true, 40, CaxpyKernelVariant::STRIDED_SCALAR},
56+ {640, 1, 3, true, 40, CaxpyKernelVariant::STRIDED_SCALAR},
57+ {640, 5, 5, true, 40, CaxpyKernelVariant::STRIDED_SCALAR},
58+ {640, 2, 3, true, 40, CaxpyKernelVariant::STRIDED_SCALAR},
59+ {640, -2, 3, true, 40, CaxpyKernelVariant::STRIDED_SCALAR},
60+ {641, 5, 1, true, 40, CaxpyKernelVariant::STRIDED_SCALAR},
61+ {679, 5, 1, true, 40, CaxpyKernelVariant::STRIDED_SCALAR},
62+ {17, 65537, -32769, true, 40, CaxpyKernelVariant::STRIDED_SCALAR},
63+ {680, 9, 1, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC},
64+ {680, 1, 9, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC},
65+ {680, 5, 5, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC},
66+ {680, 2, 3, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC},
67+ {680, -2, -3, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC},
68+ // incx=10 now stays on the shared span-read pipeline: it exceeds the old
69+ // fixed CAXPY_SHARED_X_SPAN_MAX_INCREMENT=9 cutoff but is well within
70+ // both the tile-capacity correctness floor and the measured performance
71+ // cap (CAXPY_SHARED_SPAN_PERFORMANCE_MAX_INCREMENT=60). This case used to
72+ // assert STRIDED_DISJOINT_PIPELINED_ATOMIC under the old fixed threshold;
73+ // that expectation is now wrong on purpose -- see caxpy_dispatch_policy.h
74+ // for the crossover measurement this depends on.
75+ {680, 10, 1, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC},
76+ // incx=61 exceeds the performance cap and correctly falls back to
77+ // disjoint even though its shrunk tile (capacity 80) would still be
78+ // usable by the correctness floor alone.
79+ {680, 61, 1, true, 40, CaxpyKernelVariant::STRIDED_DISJOINT_PIPELINED_ATOMIC},
80+ // incx=304 shrinks the tile capacity to 32, below
81+ // CAXPY_SHARED_MIN_USABLE_TILE_CAPACITY (33) -- the correctness floor,
82+ // not the performance cap, is what forces disjoint here. Kept as
83+ // defense-in-depth coverage of that floor in case the performance cap is
84+ // ever raised past this incx.
85+ {680, 304, 1, true, 40, CaxpyKernelVariant::STRIDED_DISJOINT_PIPELINED_ATOMIC},
86+ {327680, 2, 3, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC},
87+ {327681, 2, 3, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_SCALAR_TAIL},
88+ {328359, -2, 3, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_SCALAR_TAIL},
89+ {328360, 2, -3, true, 40, CaxpyKernelVariant::STRIDED_SHARED_PIPELINED_ATOMIC},
90+ {640, 5, 5, false, 40, CaxpyKernelVariant::DENSE_PIPELINED_ATOMIC},
91+}};
92+ 
93+ 
94+} // namespace
95+ 
96+int main()
97+{
98+ bool passed = true;
99+ for (const ClassificationCase& test : CASES) {
100+ const CaxpyStrideClass actual = ClassifyCaxpyStride(test.incx, test.incy);
101+ if (actual == test.expected)
102+ continue;
103+ std::cerr << "classification mismatch: incx=" << test.incx << " incy=" << test.incy
104+ << " expected=" << static_cast<uint32_t>(test.expected)
105+ << " actual=" << static_cast<uint32_t>(actual) << '\n';
106+ passed = false;
107+ }
108+ for (const PolicyCase& test : POLICY_CASES) {
109+ const CaxpyKernelVariant actual =
110+ SelectCaxpyKernelVariant(test.n, test.incx, test.incy, test.directStrided, test.numBlocks);
111+ if (actual == test.expected)
112+ continue;
113+ std::cerr << "policy mismatch: n=" << test.n << " incx=" << test.incx << " incy=" << test.incy
114+ << " direct_strided=" << test.directStrided
115+ << " expected=" << static_cast<uint32_t>(test.expected)
116+ << " actual=" << static_cast<uint32_t>(actual) << '\n';
117+ passed = false;
118+ }
119+ 
120+ return passed ? 0 : 1;
121+}
Atest/axpy/caxpy/arch22/caxpy_npu_wrapper.h+117-0
@@ -0,0 +1,117 @@
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 <limits>
15+#include <vector>
16+ 
17+#include "acl/acl.h"
18+#include "cann_ops_blas.h"
19+#include "fill.h"
20+ 
21+static inline bool CaxpyNeedPassThrough(aclblasHandle_t handle, int n)
22+{
23+ return handle == nullptr || n <= 0;
24+}
25+ 
26+static inline aclError CaxpyAllocCopyH2D(void*& dPtr, const void* hPtr, size_t bytes)
27+{
28+ dPtr = nullptr;
29+ if (hPtr == nullptr || bytes == 0) {
30+ return ACL_SUCCESS;
31+ }
32+ aclError ret = aclrtMalloc(&dPtr, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
33+ if (ret != ACL_SUCCESS) {
34+ return ret;
35+ }
36+ ret = aclrtMemcpy(dPtr, bytes, hPtr, bytes, ACL_MEMCPY_HOST_TO_DEVICE);
37+ if (ret != ACL_SUCCESS) {
38+ aclrtFree(dPtr);
39+ dPtr = nullptr;
40+ }
41+ return ret;
42+}
43+ 
44+static inline void CaxpyFreeAll(void* dX, void* dY)
45+{
46+ if (dX) {
47+ aclrtFree(dX);
48+ }
49+ if (dY) {
50+ aclrtFree(dY);
51+ }
52+}
53+ 
54+// Physical storage span (in complex elements) covered by a length-n vector
55+// with the given (possibly negative) increment.
56+static inline size_t CaxpySpanElements(int n, int increment)
57+{
58+ if (n <= 0) {
59+ return 0;
60+ }
61+ const uint64_t absIncrement =
62+ increment >= 0 ? static_cast<uint64_t>(increment)
63+ : static_cast<uint64_t>(-(static_cast<int64_t>(increment) + 1)) + 1;
64+ const uint64_t logicalDistance = static_cast<uint64_t>(n - 1);
65+ if (absIncrement != 0 && logicalDistance >
66+ (std::numeric_limits<size_t>::max() / sizeof(aclblasComplex) - 1) /
67+ absIncrement) {
68+ return 0;
69+ }
70+ return static_cast<size_t>(logicalDistance * absIncrement + 1);
71+}
72+ 
73+inline aclblasStatus_t aclblasCaxpy_npu(
74+ aclblasHandle_t handle, int n, const aclblasComplex* alpha, const aclblasComplex* x, int incx,
75+ aclblasComplex* y, int incy)
76+{
77+ if (CaxpyNeedPassThrough(handle, n)) {
78+ return aclblasCaxpy(handle, n, alpha, x, incx, y, incy);
79+ }
80+ 
81+ const size_t xElements = CaxpySpanElements(n, incx);
82+ const size_t yElements = CaxpySpanElements(n, incy);
83+ if ((n > 0 && xElements == 0) || (n > 0 && yElements == 0)) {
84+ return ACLBLAS_STATUS_INVALID_VALUE;
85+ }
86+ const size_t xBytes = xElements * sizeof(aclblasComplex);
87+ const size_t yBytes = yElements * sizeof(aclblasComplex);
88+ 
89+ void* dX = nullptr;
90+ void* dY = nullptr;
91+ 
92+ if (CaxpyAllocCopyH2D(dX, x, xBytes) != ACL_SUCCESS) {
93+ return ACLBLAS_STATUS_ALLOC_FAILED;
94+ }
95+ if (CaxpyAllocCopyH2D(dY, y, yBytes) != ACL_SUCCESS) {
96+ CaxpyFreeAll(dX, dY);
97+ return ACLBLAS_STATUS_ALLOC_FAILED;
98+ }
99+ 
100+ aclblasStatus_t ret = aclblasCaxpy(
101+ handle, n, alpha, static_cast<const aclblasComplex*>(dX), incx, static_cast<aclblasComplex*>(dY), incy);
102+ 
103+ if (aclrtSynchronizeDevice() != ACL_SUCCESS) {
104+ CaxpyFreeAll(dX, dY);
105+ return ACLBLAS_STATUS_INTERNAL_ERROR;
106+ }
107+ 
108+ if (ret == ACLBLAS_STATUS_SUCCESS && y != nullptr && dY != nullptr) {
109+ if (aclrtMemcpy(y, yBytes, dY, yBytes, ACL_MEMCPY_DEVICE_TO_HOST) != ACL_SUCCESS) {
110+ CaxpyFreeAll(dX, dY);
111+ return ACLBLAS_STATUS_INTERNAL_ERROR;
112+ }
113+ }
114+ 
115+ CaxpyFreeAll(dX, dY);
116+ return ret;
117+}
Mtest/axpy/caxpy/arch22/caxpy_test.cpp+116-139
@@ -8,158 +8,135 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-/* !
12- * \file caxpy_test.cpp
13- * \brief Test for aclblasCaxpy interface
14- */
15- 
16-#include <cstdint>
17-#include <iostream>
18#include <vector>11#include <vector>
19-#include <algorithm>12+#include <iostream>
20-#include <iterator>13+#include <iomanip>
21-#include "acl/acl.h"
22-#include "cann_ops_blas.h"
23-#include "complex.h"
24 14 
25-#define CHECK_RET(cond, return_expr) \15+#include "verify.h"
26- do { \16+#include "blas_test.h"
27- if (!(cond)) { \17+#include "csv_loader.h"
28- return_expr; \18+#include "fill.h"
29- } \19+#include "caxpy_param.h"
30- } while (0)20+#include "caxpy_golden.h"
21+#include "caxpy_npu_wrapper.h"
22+#include "npu_event_timer.h"
31 23 
32-#define LOG_PRINT(message, ...) \24+namespace {
33- do { \
34- printf(message, ##__VA_ARGS__); \
35- } while (0)
36 25 
37-uint32_t VerifyCaxpyResult(std::vector<aclblasComplex>& output, std::vector<aclblasComplex>& golden)26+constexpr aclblasComplex GUARD_VALUE{12345.25F, -54321.5F};
27+constexpr size_t GUARD_COUNT = 64;
28+ 
29+// Deterministic per-index complex value generator: independent of BlasFillMode
30+// so cross-checking against the reference cblas-style AXPY is exact, matching
31+// the layout the legacy hardcoded-array caxpy_test used before the CSV
32+// migration.
33+aclblasComplex MakeStridedValue(size_t index, int32_t salt)
38{34{
39- auto printTensor = [](std::vector<aclblasComplex>& tensor, const char* name) {35+ const float real = 0.125F * static_cast<float>(static_cast<int32_t>(index % 29) - 14 + salt);
40- constexpr size_t maxPrintSize = 10;36+ const float imag = 0.0625F * static_cast<float>(static_cast<int32_t>(index % 23) - 11 - salt);
41- std::cout << name << ": ";37+ return {real, imag};
42- for (size_t i = 0; i < std::min(tensor.size(), maxPrintSize); i++) {
43- std::cout << "(" << tensor[i].real << "," << tensor[i].imag << ") ";
44- }
45- if (tensor.size() > maxPrintSize) {
46- std::cout << "...";
47- }
48- std::cout << std::endl;
49- };
50- printTensor(output, "Output");
51- printTensor(golden, "Golden");
52- constexpr float EPSILON = 1e-3f;
53- for (size_t i = 0; i < output.size(); i++) {
54- float diff = blasComplexAbs(output[i] - golden[i]);
55- if (diff > EPSILON) {
56- std::cout << "[Failed] Caxpy Index " << i << ": output=(" << output[i].real << "," << output[i].imag
57- << ") golden=(" << golden[i].real << "," << golden[i].imag << ") diff=" << diff << std::endl;
58- return 1;
59- }
60- }
61- std::cout << "[Success] Caxpy accuracy verification passed." << std::endl;
62- return 0;
63}38}
64 39 
65-int32_t TestCaxpy(aclblasHandle handle, aclrtStream stream)40+// Build a length-n complex vector (logical, stride-independent index order)
41+// with GUARD_COUNT sentinel padding on both sides so any out-of-bounds
42+// read/write by the operator under test corrupts a detectable guard region.
43+std::vector<aclblasComplex> MakeGuardedVector(int n, int increment, int32_t salt)
66{44{
67- constexpr uint32_t totalLength = 8 * 2048;45+ const size_t span = CaxpySpanElements(n, increment);
68- constexpr aclblasComplex valueX{1.0f, 0.5f};46+ std::vector<aclblasComplex> data(GUARD_COUNT + span + GUARD_COUNT, GUARD_VALUE);
69- constexpr aclblasComplex valueY{2.0f, 1.0f};47+ for (size_t i = 0; i < span; ++i)
70- constexpr aclblasComplex alpha{2.0f, 1.0f};48+ data[GUARD_COUNT + i] = MakeStridedValue(i, salt);
71- std::vector<aclblasComplex> x(totalLength, valueX);49+ return data;
72- std::vector<aclblasComplex> y(totalLength, valueY);
73- int incx = 1;
74- int incy = 1;
75- 
76- aclblasComplex* xDevice = nullptr;
77- aclblasComplex* yDevice = nullptr;
78- size_t xByteSize = totalLength * sizeof(aclblasComplex);
79- size_t yByteSize = totalLength * sizeof(aclblasComplex);
80- aclError aclRet = aclrtMalloc((void**)&xDevice, xByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
81- CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc xDevice failed. ERROR: %d\n", aclRet); return aclRet);
82- aclRet = aclrtMalloc((void**)&yDevice, yByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
83- CHECK_RET(
84- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc yDevice failed. ERROR: %d\n", aclRet); aclrtFree(xDevice);
85- return aclRet);
86- aclRet = aclrtMemcpy(xDevice, xByteSize, x.data(), xByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
87- CHECK_RET(
88- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy xDevice failed. ERROR: %d\n", aclRet); aclrtFree(xDevice);
89- aclrtFree(yDevice); return aclRet);
90- aclRet = aclrtMemcpy(yDevice, yByteSize, y.data(), yByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
91- CHECK_RET(
92- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy yDevice failed. ERROR: %d\n", aclRet); aclrtFree(xDevice);
93- aclrtFree(yDevice); return aclRet);
94- 
95- std::cout << "========== Testing aclblasCaxpy ==========" << std::endl;
96- std::cout << "Formula: y = alpha * x + y" << std::endl;
97- std::cout << "alpha = (" << alpha.real << ", " << alpha.imag << ")" << std::endl;
98- std::cout << "x = (" << valueX.real << ", " << valueX.imag << ") * " << totalLength << std::endl;
99- std::cout << "y = (" << valueY.real << ", " << valueY.imag << ") * " << totalLength << std::endl;
100- 
101- auto ret = aclblasCaxpy(handle, totalLength, &alpha, xDevice, incx, yDevice, incy);
102- CHECK_RET(
103- ret == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCaxpy failed. ERROR: %d\n", ret); aclrtFree(xDevice);
104- aclrtFree(yDevice); return ret);
105- 
106- aclRet = aclrtSynchronizeStream(stream);
107- CHECK_RET(
108- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", aclRet); aclrtFree(xDevice);
109- aclrtFree(yDevice); return aclRet);
110- aclRet = aclrtMemcpy(y.data(), yByteSize, yDevice, yByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
111- CHECK_RET(
112- aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy y failed. ERROR: %d\n", aclRet); aclrtFree(xDevice);
113- aclrtFree(yDevice); return aclRet);
114- aclrtFree(xDevice);
115- aclrtFree(yDevice);
116- 
117- std::vector<aclblasComplex> golden(totalLength);
118- for (size_t i = 0; i < totalLength; i++) {
119- golden[i] = alpha * x[i] + valueY;
120- }
121- 
122- return VerifyCaxpyResult(y, golden);
123}50}
124 51 
125-int32_t main(int32_t argc, char* argv[])52+} // namespace
126-{
127- int32_t deviceId = 0;
128 53 
129- aclInit(nullptr);54+// ── Test fixture ──────────────────────────────────────────────────────────────
130- aclrtSetDevice(deviceId);55+class CaxpyArch22Test : public BlasTest<CaxpyParam> { };
131 56 
132- aclblasHandle_t handle = nullptr;57+// ── TEST_F: null handle / null alpha (not expressible via CSV) ───────────────
133- auto ret = aclblasCreate(&handle);58+TEST_F(CaxpyArch22Test, NullHandle) {
134- CHECK_RET(ret == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasCreate failed. ERROR: %d\n", ret); return ret);59+ aclblasComplex alpha{1.0F, 0.0F};
60+ aclblasStatus_t ret = aclblasCaxpy_npu(nullptr, 1, &alpha, nullptr, 1, nullptr, 1);
61+ EXPECT_EQ(static_cast<int>(ret), static_cast<int>(ACLBLAS_STATUS_HANDLE_IS_NULLPTR));
62+}
135 63 
136- aclrtStream stream = nullptr;64+TEST_F(CaxpyArch22Test, NullAlpha) {
137- aclrtCreateStream(&stream);65+ aclblasComplex x{1.0F, 0.0F};
138- ret = aclblasSetStream(handle, stream);66+ aclblasComplex y{1.0F, 0.0F};
139- CHECK_RET(ret == ACLBLAS_STATUS_SUCCESS, LOG_PRINT("aclblasSetStream failed. ERROR: %d\n", ret); return ret);67+ aclblasStatus_t ret = aclblasCaxpy_npu(CaxpyArch22Test::handle_, 1, nullptr, &x, 1, &y, 1);
68+ EXPECT_EQ(static_cast<int>(ret), static_cast<int>(ACLBLAS_STATUS_INVALID_VALUE));
69+}
140 70 
141- int32_t result = 0;71+// ── CSV parameterised test suite ─────────────────────────────────────────────
72+INSTANTIATE_TEST_SUITE_P(
73+ Caxpy, CaxpyArch22Test,
74+ ::testing::ValuesIn(GetCasesFromCsv<CaxpyParam>(ReplaceFileExtension2Csv(__FILE__))),
75+ PrintCaseInfoString<CaxpyParam>);
142 76 
143- result = TestCaxpy(handle, stream);77+// ── TEST_P: 5-step CSV-driven flow ───────────────────────────────────────────
144- if (result != 0) {78+TEST_P(CaxpyArch22Test, CsvDriven) {
145- std::cout << "[FAIL] Caxpy test failed" << std::endl;79+ const auto& p = GetParam();
146- aclblasDestroy(handle);80+ const aclblasComplex alpha{p.alphaReal, p.alphaImag};
147- aclrtDestroyStream(stream);81+ 
148- aclrtResetDevice(deviceId);82+ // Step 1: Generate host data (guarded x/y storage, deterministic values)
149- aclFinalize();83+ std::vector<aclblasComplex> xHost;
150- return result;84+ if (p.nullx == 0 && p.n > 0)
85+ xHost = MakeGuardedVector(p.n, p.incx, 3);
86+ std::vector<aclblasComplex> yHost;
87+ if (p.nully == 0 && p.n > 0)
88+ yHost = MakeGuardedVector(p.n, p.incy, -5);
89+ const std::vector<aclblasComplex> initialY = yHost;
90+ 
91+ const aclblasComplex* xPtr = xHost.empty() ? nullptr : xHost.data() + GUARD_COUNT;
92+ aclblasComplex* yPtr = yHost.empty() ? nullptr : yHost.data() + GUARD_COUNT;
93+ 
94+ // Step 2: Execute on NPU (wrapper handles device memory, nullptr passthrough)
95+ aclblasStatus_t ret = aclblasCaxpy_npu(CaxpyArch22Test::handle_, p.n, &alpha, xPtr, p.incx, yPtr, p.incy);
96+ 
97+ // Step 3: Verify expected return code
98+ EXPECT_EQ(static_cast<int>(ret), static_cast<int>(p.expectResult));
99+ if (p.expectResult != ACLBLAS_STATUS_SUCCESS) return;
100+ if (p.n == 0) return;
101+ 
102+ // Step 4: Compute golden on CPU over the same guarded storage, so any
103+ // stray write into the guard region is caught by the full-buffer compare.
104+ std::vector<aclblasComplex> xGolden = p.nullx == 0 ? MakeGuardedVector(p.n, p.incx, 3) : std::vector<aclblasComplex>{};
105+ std::vector<aclblasComplex> yGolden = p.nully == 0 ? MakeGuardedVector(p.n, p.incy, -5) : std::vector<aclblasComplex>{};
106+ const aclblasComplex* xGoldenPtr = xGolden.empty() ? nullptr : xGolden.data() + GUARD_COUNT;
107+ aclblasComplex* yGoldenPtr = yGolden.empty() ? nullptr : yGolden.data() + GUARD_COUNT;
108+ aclblasStatus_t cpuRet = aclblasCaxpy_cpu(CaxpyArch22Test::handle_, p.n, &alpha, xGoldenPtr, p.incx, yGoldenPtr, p.incy);
109+ EXPECT_EQ(static_cast<int>(cpuRet), static_cast<int>(ACLBLAS_STATUS_SUCCESS));
110+ 
111+ // Step 5: Precision verification, real/imag components split (MERE/MARE).
112+ VerifyConfig cfg;
113+ cfg.mode = PrecisionMode::MERE_MARE;
114+ cfg.mereThreshold = p.mereThreshold;
115+ cfg.mareMultiplier = p.mareMultiplier;
116+ 
117+ EXPECT_TRUE(Verifier::verifyMereMareComplexFloat(
118+ yHost.data(), yGolden.data(), yHost.size(), cfg.mereThreshold, cfg.mareMultiplier, 0.0, p.caseName));
119+ 
120+ if (p.measure != 0 && blas_test::MeasureRequested()) {
121+ std::vector<double> samples;
122+ const int measureRet = blas_test::MeasureNpuEvent(
123+ stream_, p.warmup, p.measurements,
124+ [&]() {
125+ yHost = initialY;
126+ return true;
127+ },
128+ [&]() {
129+ const aclblasComplex* timedX = xHost.empty() ? nullptr : xHost.data() + GUARD_COUNT;
130+ aclblasComplex* timedY = yHost.empty() ? nullptr : yHost.data() + GUARD_COUNT;
131+ return aclblasCaxpy_npu(
132+ CaxpyArch22Test::handle_, p.n, &alpha, timedX, p.incx, timedY, p.incy) ==
133+ ACLBLAS_STATUS_SUCCESS;
134+ },
135+ samples);
136+ ASSERT_EQ(measureRet, 0);
137+ std::cout << std::setprecision(9) << "PERFORMANCE_SUMMARY_JSON={\"case\":\"" << p.caseName
138+ << "\",\"n\":" << p.n << ",\"p50_us\":" << blas_test::MedianUs(samples)
139+ << ",\"avg_us\":" << blas_test::MeanUs(samples) << ",\"warmups\":" << p.warmup
140+ << ",\"measurements\":" << p.measurements << ",\"timing\":\"npu_event\"}\n";
151 }141 }
152- std::cout << "[PASS] Caxpy test passed" << std::endl;142+}
153- 
154- aclblasDestroy(handle);
155- aclrtDestroyStream(stream);
156- aclrtResetDevice(deviceId);
157- aclFinalize();
158- 
159- std::cout << "========================================" << std::endl;
160- std::cout << "Test Summary:" << std::endl;
161- std::cout << " Passed: 1 - Caxpy" << std::endl;
162- std::cout << "========================================" << std::endl;
163- 
164- return 0;
165-}
Atest/axpy/caxpy/arch22/caxpy_test.csv+79-0
@@ -0,0 +1,79 @@
1+case_name,description,n,incx,incy,alpha_real,alpha_imag,nullx,nully,expect_result,mere_threshold,mare_multiplier,random_seed,measure,warmup,measurements,classification,flops,profile
2+TC_zero,zero,0,1,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
3+TC_one,one,1,1,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
4+TC_lane_2,lane 2,2,1,1,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
5+TC_lane_3,lane 3,3,1,1,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
6+TC_lane_31,lane 31,31,1,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
7+TC_lane_32,lane 32,32,1,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
8+TC_lane_33,lane 33,33,1,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
9+TC_lane_63,lane 63,63,1,1,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
10+TC_lane_64,lane 64,64,1,1,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
11+TC_lane_65,lane 65,65,1,1,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
12+TC_core_39,core 39,39,1,1,0.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
13+TC_core_40,core 40,40,1,1,0.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
14+TC_core_41,core 41,41,1,1,0.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
15+TC_strided_x,strided x,257,2,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
16+TC_vector_reorder_tail_1,vector reorder tail 1,1,2,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
17+TC_vector_reorder_tail_7,vector reorder tail 7,7,2,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
18+TC_vector_reorder_tail_8,vector reorder tail 8,8,2,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
19+TC_vector_reorder_tail_9,vector reorder tail 9,9,2,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
20+TC_vector_reorder_tile,vector reorder tile,512,2,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
21+TC_scalar_block_15,scalar block 15,600,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
22+TC_mixed_scalar_block_15_16,mixed scalar block 15 16,601,-2,3,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
23+TC_scalar_block_16,scalar block 16,640,2,-3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
24+TC_mixed_scalar_block_16_17,mixed scalar block 16 17,641,5,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
25+TC_scalar_upper_bound_16_17,scalar upper bound 16 17,679,-5,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
26+TC_scalar_wide_increment,scalar wide increment,17,65537,-32769,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
27+TC_vector_reorder_negative_tail,vector reorder negative tail,513,-10,-11,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
28+TC_strided_y,strided y,257,1,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
29+TC_strided_xy,strided xy,257,2,3,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
30+TC_stride_5_x,stride 5 x,257,5,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
31+TC_stride_5_y,stride 5 y,257,1,5,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
32+TC_stride_5_9,stride 5 9,257,5,9,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
33+TC_stride_9_x,stride 9 x,257,9,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
34+TC_stride_10_x_fallback,stride 10 x fallback,257,10,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
35+TC_stride_5_multicore,stride 5 multicore,4097,5,5,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
36+TC_shared_raw_inc2_multitile_tail,shared raw inc2 multitile tail,40961,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
37+TC_shared_tail_1,shared tail 1,41000,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
38+TC_shared_tail_8,shared tail 8,41280,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
39+TC_shared_tail_9,shared tail 9,41320,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
40+TC_shared_tail_16,shared tail 16,41600,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
41+TC_shared_tail_17,shared tail 17,41640,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
42+TC_shared_tail_1_incy1,shared tail 1 incy1,41000,2,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
43+TC_shared_tail_1_negative,shared tail 1 negative,41000,-2,-3,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
44+TC_shared_tail_multi_4tile,shared tail multi 4tile,61441,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
45+TC_shared_raw_inc5_multitile_tail,shared raw inc5 multitile tail,40961,5,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
46+TC_shared_raw_inc9_multitile_tail,shared raw inc9 multitile tail,40961,9,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
47+TC_negative_x,negative x,257,-2,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
48+TC_negative_unit_y,negative unit y,257,2,-1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
49+TC_negative_unit_y_multitile,negative unit y multitile,4097,2,-1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
50+TC_negative_y,negative y,257,1,-3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
51+TC_negative_x_positive_y,negative x positive y,257,-2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
52+TC_positive_x_negative_y,positive x negative y,257,5,-3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
53+TC_negative_xy,negative xy,257,-2,-3,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
54+TC_negative_9_x,negative 9 x,257,-9,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
55+TC_negative_x_multicore,negative x multicore,4097,-2,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
56+TC_negative_y_multicore,negative y multicore,4097,1,-3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
57+TC_opposite_sign_multicore,opposite sign multicore,4097,-2,3,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
58+TC_negative_xy_multicore,negative xy multicore,4097,-2,-3,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
59+TC_wide_stride_multitile,wide stride multitile,40961,10,11,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
60+TC_wide_stride_negative_multitile,wide stride negative multitile,40963,-10,-11,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
61+TC_dynamic_shared_incx16,dynamic shared incx16,40961,16,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
62+TC_dynamic_shared_incx32,dynamic shared incx32,40961,32,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
63+TC_dynamic_shared_incx41,dynamic shared incx41,40961,41,3,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
64+TC_dynamic_shared_boundary_usable,dynamic shared boundary usable,8000,60,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
65+TC_dynamic_shared_boundary_disjoint,dynamic shared boundary disjoint,8000,61,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
66+TC_strided_tile_plus_1,strided tile plus 1,194561,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
67+TC_multiwave_strided,multiwave strided,400000,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
68+TC_multiwave_negative,multiwave negative,400001,-2,-3,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
69+TC_multiwave_3_waves,multiwave 3 waves,700000,1,2,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
70+TC_pipeline_full_wave,pipeline full wave,327680,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
71+TC_hybrid_tail_1,hybrid tail 1,327681,2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
72+TC_hybrid_tail_1_negative_y,hybrid tail 1 negative y,327681,2,-3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
73+TC_hybrid_tail_679_negative_x,hybrid tail 679 negative x,328359,-2,3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
74+TC_hybrid_tail_679_negative_xy,hybrid tail 679 negative xy,328359,-2,-3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
75+TC_pipeline_tail_680_negative_y,pipeline tail 680 negative y,328360,2,-3,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
76+TC_regular,regular,16384,1,1,0.75,-0.25,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,1,100,200,comparison,131072,1
77+TC_tile_minus_1,tile minus 1,194559,1,1,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
78+TC_tile,tile,194560,1,1,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,1,100,200,comparison,1556480,0
79+TC_tile_plus_1,tile plus 1,194561,1,1,2.0,1.0,0,0,ACLBLAS_STATUS_SUCCESS,0.0001220703125,10.0,42,0,30,50,smoke,0,0
Atest/axpy/caxpy/caxpy_golden.h+69-0
@@ -0,0 +1,69 @@
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 "cann_ops_blas.h"
16+ 
17+inline aclblasStatus_t CaxpyValidateParams(
18+ aclblasHandle_t handle, int n, const aclblasComplex* alpha, const aclblasComplex* x, int incx,
19+ const aclblasComplex* y, int incy)
20+{
21+ if (handle == nullptr || alpha == nullptr) {
22+ return ACLBLAS_STATUS_HANDLE_IS_NULLPTR;
23+ }
24+ if (n < 0 || incx == 0 || incy == 0) {
25+ return ACLBLAS_STATUS_INVALID_VALUE;
26+ }
27+ if (n > 0 && (x == nullptr || y == nullptr)) {
28+ return ACLBLAS_STATUS_INVALID_VALUE;
29+ }
30+ return ACLBLAS_STATUS_SUCCESS;
31+}
32+ 
33+inline aclblasComplex CaxpyComplexMac(aclblasComplex alpha, aclblasComplex x, aclblasComplex y)
34+{
35+ return {alpha.real * x.real - alpha.imag * x.imag + y.real, alpha.real * x.imag + alpha.imag * x.real + y.imag};
36+}
37+ 
38+// Logical-index -> physical storage index for a strided complex vector.
39+// Positive stride reads forward from the base; negative stride is BLAS
40+// convention: the last logical element is physically first.
41+inline int64_t CaxpyStorageIndex(int64_t logicalIndex, int64_t n, int64_t increment)
42+{
43+ const uint64_t absIncrement =
44+ increment > 0 ? static_cast<uint64_t>(increment) : static_cast<uint64_t>(-increment);
45+ return increment > 0 ? logicalIndex * static_cast<int64_t>(absIncrement)
46+ : (n - 1 - logicalIndex) * static_cast<int64_t>(absIncrement);
47+}
48+ 
49+// y = alpha * x + y, complex, arbitrary (possibly negative) incx/incy.
50+// x and y are the caller-owned base storage pointers (not offset by any guard region).
51+inline aclblasStatus_t aclblasCaxpy_cpu(
52+ aclblasHandle_t handle, int n, const aclblasComplex* alpha, const aclblasComplex* x, int incx,
53+ aclblasComplex* y, int incy)
54+{
55+ aclblasStatus_t st = CaxpyValidateParams(handle, n, alpha, x, incx, y, incy);
56+ if (st != ACLBLAS_STATUS_SUCCESS) {
57+ return st;
58+ }
59+ if (n == 0) {
60+ return ACLBLAS_STATUS_SUCCESS;
61+ }
62+ 
63+ for (int64_t i = 0; i < n; i++) {
64+ const int64_t xIdx = CaxpyStorageIndex(i, n, incx);
65+ const int64_t yIdx = CaxpyStorageIndex(i, n, incy);
66+ y[yIdx] = CaxpyComplexMac(*alpha, x[xIdx], y[yIdx]);
67+ }
68+ return ACLBLAS_STATUS_SUCCESS;
69+}
Atest/axpy/caxpy/caxpy_param.h+42-0
@@ -0,0 +1,42 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#pragma once
12+ 
13+#include <string>
14+#include "cann_ops_blas.h"
15+#include "csv_loader.h"
16+ 
17+struct CaxpyParam : public BlasTestParamBase {
18+ int n = 0;
19+ int incx = 1;
20+ int incy = 1;
21+ float alphaReal = 1.0f;
22+ float alphaImag = 0.0f;
23+ int nullx = 0;
24+ int nully = 0;
25+ int measure = 0;
26+ int warmup = 30;
27+ int measurements = 50;
28+ 
29+ CaxpyParam(const csv_map& csv) : BlasTestParamBase(csv)
30+ {
31+ n = parseInt(ReadMap(csv, "n", "0"));
32+ incx = parseInt(ReadMap(csv, "incx", "1"));
33+ incy = parseInt(ReadMap(csv, "incy", "1"));
34+ alphaReal = parseFloat(ReadMap(csv, "alpha_real", "1.0"));
35+ alphaImag = parseFloat(ReadMap(csv, "alpha_imag", "0.0"));
36+ nullx = parseInt(ReadMap(csv, "nullx", "0"));
37+ nully = parseInt(ReadMap(csv, "nully", "0"));
38+ measure = parseInt(ReadMap(csv, "measure", "0"));
39+ warmup = parseInt(ReadMap(csv, "warmup", "30"));
40+ measurements = parseInt(ReadMap(csv, "measurements", "50"));
41+ }
42+};
Atest/axpy/caxpy/manifest.yaml+15-0
@@ -0,0 +1,15 @@
1+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+ 
9+name: caxpy
10+soc: ascend910b3
11+arch: arch22
12+cases: arch22/caxpy_test.csv
13+performance_metrics:
14+ flop_model: "8 floating-point operations per complex element for one complex multiply-add"
15+ measured_peak_tflops: 1.0
Atest/frame/npu_event_timer.h+69-0
@@ -0,0 +1,69 @@
1+#pragma once
2+/**
3+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+ * CANN Open Software License Agreement Version 2.0 (the "License").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+#include <algorithm>
13+#include <cstdlib>
14+#include <functional>
15+#include <vector>
16+ 
17+#include "acl/acl.h"
18+#include "acl/acl_rt.h"
19+ 
20+namespace blas_test {
21+ 
22+inline bool MeasureRequested() { return std::getenv("BLAS_MEASURE") != nullptr; }
23+ 
24+inline int MeasureNpuEvent(
25+ aclrtStream stream, int warmup, int measurements, const std::function<bool()>& resetA,
26+ const std::function<bool()>& launch, std::vector<double>& outUs)
27+{
28+ if (stream == nullptr || warmup <= 0 || measurements <= 0) return 1;
29+ aclrtEvent evStart = nullptr;
30+ aclrtEvent evEnd = nullptr;
31+ if (aclrtCreateEvent(&evStart) != ACL_SUCCESS || aclrtCreateEvent(&evEnd) != ACL_SUCCESS) return 2;
32+ const auto fail = [&](int code) {
33+ aclrtDestroyEvent(evStart);
34+ aclrtDestroyEvent(evEnd);
35+ return code;
36+ };
37+ for (int i = 0; i < warmup; ++i) {
38+ if (!resetA() || !launch() || aclrtSynchronizeStream(stream) != ACL_SUCCESS) return fail(3);
39+ }
40+ outUs.clear();
41+ outUs.reserve(static_cast<size_t>(measurements));
42+ for (int i = 0; i < measurements; ++i) {
43+ if (!resetA()) return fail(4);
44+ if (aclrtRecordEvent(evStart, stream) != ACL_SUCCESS || !launch() ||
45+ aclrtRecordEvent(evEnd, stream) != ACL_SUCCESS || aclrtSynchronizeStream(stream) != ACL_SUCCESS)
46+ return fail(5);
47+ float ms = 0.0F;
48+ if (aclrtEventElapsedTime(&ms, evStart, evEnd) != ACL_SUCCESS) return fail(6);
49+ outUs.push_back(static_cast<double>(ms) * 1000.0);
50+ }
51+ return fail(0);
52+}
53+ 
54+inline double MedianUs(std::vector<double> values)
55+{
56+ if (values.empty()) return 0.0;
57+ std::sort(values.begin(), values.end());
58+ return values[values.size() / 2];
59+}
60+ 
61+inline double MeanUs(const std::vector<double>& values)
62+{
63+ if (values.empty()) return 0.0;
64+ double sum = 0.0;
65+ for (double value : values) sum += value;
66+ return sum / static_cast<double>(values.size());
67+}
68+ 
69+} // namespace blas_test