Star759
228
代码
代码
Issues24
Pull Requests37
Actions
项目讨论
Wiki
分析
Star759
228
已合并
feat(simd_vf_story): 新增 elemwise、reduce 范式样例与 SIMD VF 性能调优 README #314
TangPC创建于 6月17日
feat(simd_vf_story): 新增 elemwise、reduce 范式样例与 SIMD VF 性能调优 README #314
已合并
TangPC创建于 6月17日
从pingchuantang/cann-samples-tpcmaster合入到cann/cann-samplesmaster
共 25 个文件变更+3888-17
MSamples/2_Performance/simd_vf_story/CMakeLists.txt+16-5
@@ -16,8 +16,10 @@ if(NOT "${NPU_ARCH}" IN_LIST SUPPORTED_NPU_ARCHS)
16 return()16 return()
17endif()17endif()
18 18 
19+# Build one executable per <case>/src/*.asc file, plus a simd_vf_story_<case> aggregate target.
20+# All targets are defined here; no per-case CMakeLists is needed.
19function(add_simd_vf_story_case CASE_NAME)21function(add_simd_vf_story_case CASE_NAME)
20- file(GLOB ASC_FILES CONFIGURE_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/src/*.asc)22+ file(GLOB ASC_FILES CONFIGURE_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/${CASE_NAME}/src/*.asc)
21 23 
22 set(CASE_TARGETS "")24 set(CASE_TARGETS "")
23 foreach(ASC_FILE ${ASC_FILES})25 foreach(ASC_FILE ${ASC_FILES})
@@ -25,6 +27,10 @@ function(add_simd_vf_story_case CASE_NAME)
25 set(TARGET_NAME "simd_vf_story_${CASE_NAME}_${TARGET_SUFFIX}")27 set(TARGET_NAME "simd_vf_story_${CASE_NAME}_${TARGET_SUFFIX}")
26 28 
27 add_executable(${TARGET_NAME} ${ASC_FILE})29 add_executable(${TARGET_NAME} ${ASC_FILE})
30+ # Keep the per-case output layout: build/.../simd_vf_story/<case>/<target>
31+ set_target_properties(${TARGET_NAME} PROPERTIES
32+ RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/${CASE_NAME}
33+ )
28 target_compile_options(${TARGET_NAME} PRIVATE34 target_compile_options(${TARGET_NAME} PRIVATE
29 "$<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${NPU_ARCH}>"35 "$<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${NPU_ARCH}>"
30 "$<$<COMPILE_LANGUAGE:ASC>:-O3>"36 "$<$<COMPILE_LANGUAGE:ASC>:-O3>"
@@ -54,10 +60,15 @@ function(add_simd_vf_story_case CASE_NAME)
54endfunction()60endfunction()
55 61 
56set_property(GLOBAL PROPERTY SIMD_VF_STORY_CASE_TARGETS "")62set_property(GLOBAL PROPERTY SIMD_VF_STORY_CASE_TARGETS "")
57-file(GLOB CHILD_CMAKES CONFIGURE_DEPENDS */CMakeLists.txt)63+ 
58-foreach(CHILD_CMAKE ${CHILD_CMAKES})64+# Each subdirectory that owns a src/ folder is a case (broadcast / elemwise / reduce).
59- get_filename_component(CHILD_DIR ${CHILD_CMAKE} DIRECTORY)65+file(GLOB CASE_SRC_DIRS CONFIGURE_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/*/src)
60- add_subdirectory(${CHILD_DIR})66+foreach(CASE_SRC_DIR ${CASE_SRC_DIRS})
67+ if(IS_DIRECTORY ${CASE_SRC_DIR})
68+ get_filename_component(CASE_DIR ${CASE_SRC_DIR} DIRECTORY)
69+ get_filename_component(CASE_NAME ${CASE_DIR} NAME)
70+ add_simd_vf_story_case(${CASE_NAME})
71+ endif()
61endforeach()72endforeach()
62 73 
63get_property(SIMD_VF_STORY_CASE_TARGETS GLOBAL PROPERTY SIMD_VF_STORY_CASE_TARGETS)74get_property(SIMD_VF_STORY_CASE_TARGETS GLOBAL PROPERTY SIMD_VF_STORY_CASE_TARGETS)
DSamples/2_Performance/simd_vf_story/broadcast/CMakeLists.txt+0-12
@@ -1,12 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under 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, INCLUDING
8-# 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-add_simd_vf_story_case(broadcast)
ASamples/2_Performance/simd_vf_story/elemwise/images/elemwise.svg+1-0
@@ -0,0 +1 @@
1+<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 820 330" font-family="Noto Sans CJK SC,Microsoft YaHei,PingFang SC,Segoe UI,Helvetica,Arial,sans-serif"><defs><marker id="a" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#1f2933"/></marker><marker id="ar" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#d64545"/></marker><marker id="ab" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#2b6cb0"/></marker><marker id="ag" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#2f855a"/></marker></defs><rect x="0" y="0" width="820" height="330" fill="white"/><text x="20" y="30" font-size="17" fill="#1f2933" text-anchor="start" font-weight="700">elemwise AXPY：z = a·x + y，lane 间无依赖（throughput-bound）</text><text x="20" y="50" font-size="12" fill="#52606d" text-anchor="start">每个 lane 独立计算，没有循环携带依赖 → 没有可缩短的串行依赖。</text><rect x="70" y="90" width="140" height="30" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="140.0" y="109.2" text-anchor="middle" font-size="12" fill="#1f2933">LoadAlign x,y</text><rect x="220" y="90" width="140" height="30" rx="6" fill="#eef2f7" stroke="#2b6cb0" stroke-width="1.5"/><text x="290.0" y="109.2" text-anchor="middle" font-size="12" fill="#1f2933">Muls a·x</text><line x1="210" y1="105" x2="220" y2="105" stroke="#1f2933" stroke-width="1.6" marker-end="url(#a)"/><rect x="370" y="90" width="140" height="30" rx="6" fill="#eef2f7" stroke="#2b6cb0" stroke-width="1.5"/><text x="440.0" y="109.2" text-anchor="middle" font-size="12" fill="#1f2933">Add +y</text><line x1="360" y1="105" x2="370" y2="105" stroke="#1f2933" stroke-width="1.6" marker-end="url(#a)"/><rect x="520" y="90" width="140" height="30" rx="6" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><text x="590.0" y="109.2" text-anchor="middle" font-size="12" fill="#1f2933">StoreAlign z</text><line x1="510" y1="105" x2="520" y2="105" stroke="#1f2933" stroke-width="1.6" marker-end="url(#a)"/><text x="70" y="142" font-size="11" fill="#52606d" text-anchor="start">中间结果全程留在 RegTensor（计算链路压缩，不回写 UB）</text><text x="20" y="180" font-size="13" fill="#1f2933" text-anchor="start" font-weight="700">三种写法 = 同一计算的不同「指令级并行」表达：</text><rect x="60" y="198" width="90" height="26" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="105.0" y="215.2" text-anchor="middle" font-size="12" fill="#1f2933">baseline</text><text x="160" y="215" font-size="12" fill="#52606d" text-anchor="start">单组寄存器，一轮 1 个 VL 块</text><rect x="60" y="234" width="90" height="26" rx="6" fill="#eef2f7" stroke="#2b6cb0" stroke-width="1.5"/><text x="105.0" y="251.2" text-anchor="middle" font-size="12" fill="#2b6cb0">unroll</text><text x="160" y="251" font-size="12" fill="#52606d" text-anchor="start">4 组独立命名寄存器，一轮 4 个 VL 块（手写暴露多条独立指令流）</text><rect x="60" y="270" width="90" height="26" rx="6" fill="#eef2f7" stroke="#2b6cb0" stroke-width="1.5"/><text x="105.0" y="287.2" text-anchor="middle" font-size="12" fill="#2b6cb0">pragma</text><text x="160" y="287" font-size="12" fill="#52606d" text-anchor="start">循环体同 baseline，#pragma unroll 4 由编译器展开</text><text x="60" y="310" font-size="12" fill="#d64545" text-anchor="start" font-weight="700">实测 PUSHQ VF：baseline 625 / unroll 646 / pragma 640，展开不增益（已吞吐打满，反增标量开销 −2~3.4%）</text></svg>
ASamples/2_Performance/simd_vf_story/elemwise/src/elemwise_axpy_baseline.asc+159-0
@@ -0,0 +1,159 @@
1+/**
2+ * 单 VF 拆分版：Elementwise AXPY: z = a * x + y（基线 VF）
3+ *
4+ * 本文件从 elemwise_axpy.asc 拆出「基线」一个 VF（ElemwiseAxpyVf），单独成可编译运行的算子，
5+ * 便于对单个 VF 做 cannsim 打点（每个 trace 干净对应一个 VF）。其余两版（手写展开 / #pragma
6+ * 展开）见同目录 elemwise_axpy_unroll.asc / elemwise_axpy_pragma.asc。
7+ *
8+ * 统一 case 规格（本目录 16 个 VF 文件全部一致，便于横向比较 trace）：
9+ * 形状 [D0, D1] = [78, 250]，总元素 N = D0*D1 = 19500。elemwise 把整批 N 个元素当一维向量算。
10+ * 该规格按「最受限算子」elemwise（需 x/y/z 三块 UB buffer）定尺寸，三块约 234KB < 256KB UB；
11+ * reduce 类只需一块输入 buffer，余量充足。D1>VL 且 %8≠0、D0%4=2，保留分块/尾块/展开尾行路径。
12+ *
13+ * 基线 VF：一轮 Repeat 处理 VL(64 lane) 个元素；LoadAlign 搬入 x、y → Muls 得 a*x → Add 叠加
14+ * y → StoreAlign 搬出。搬入不带 mask（尾部多读由后续 mask 屏蔽），Muls/Add/搬出带 mask。
15+ *
16+ * 编译 & 运行：
17+ * cmake -S . -B build -DNPU_ARCH=dav-3510 # 仓库根目录
18+ * cmake --build build --target elemwise_axpy_baseline
19+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/elemwise_axpy_baseline -s Ascend950
20+ */
21+ 
22+#include <iostream>
23+#include <vector>
24+#include <cmath>
25+#include "acl/acl.h"
26+#include "kernel_operator.h"
27+#include "../../include/vf_common.h"
28+ 
29+static constexpr float A = 2.0f; // AXPY 标量系数
30+ 
31+// ============ VF 层（基线）：一轮循环算 1 个 VL 块 ============
32+__simd_vf__ inline void ElemwiseAxpyVf(__ubuf__ float* xAddr, __ubuf__ float* yAddr,
33+ __ubuf__ float* zAddr, float a, uint32_t n, uint16_t loopNum)
34+{
35+ AscendC::Reg::RegTensor<float> vx, vy, vt;
36+ AscendC::Reg::MaskReg mask;
37+ uint32_t count = n;
38+ 
39+ for (uint16_t i = 0; i < loopNum; i++) {
40+ mask = AscendC::Reg::UpdateMask<float>(count); // 开前 count 个 lane，count 原地递减
41+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vx, xAddr + i * VL_B32); // 搬入，无 mask
42+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vy, yAddr + i * VL_B32);
43+ AscendC::Reg::Muls(vt, vx, a, mask); // vt = a * x
44+ AscendC::Reg::Add(vt, vt, vy, mask); // vt = a * x + y
45+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(zAddr + i * VL_B32, vt, mask); // 搬出，带 mask
46+ }
47+}
48+ 
49+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
50+class KernelElemwiseAxpy {
51+public:
52+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t n)
53+ {
54+ n_ = n;
55+ loopNum_ = (n_ + VL_B32 - 1) / VL_B32; // VL 块数（向上取整）
56+ uint32_t bufBytes = loopNum_ * VL_B32 * sizeof(float); // UB 按 VL 块数分配，尾轮多读的 lane 由 mask 屏蔽
57+ xGm_.SetGlobalBuffer((__gm__ float*)x, n_);
58+ yGm_.SetGlobalBuffer((__gm__ float*)y, n_);
59+ zGm_.SetGlobalBuffer((__gm__ float*)z, n_);
60+ pipe_.InitBuffer(inQueueX_, 1, bufBytes); // depth 1、单块 buffer，不开 DoubleBuffer
61+ pipe_.InitBuffer(inQueueY_, 1, bufBytes);
62+ pipe_.InitBuffer(outQueueZ_, 1, bufBytes);
63+ }
64+ 
65+ __aicore__ inline void Process()
66+ {
67+ CopyIn();
68+ Compute();
69+ CopyOut();
70+ }
71+ 
72+private:
73+ __aicore__ inline void CopyIn()
74+ {
75+ AscendC::LocalTensor<float> xUb = inQueueX_.AllocTensor<float>();
76+ AscendC::LocalTensor<float> yUb = inQueueY_.AllocTensor<float>();
77+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(n_ * sizeof(float)), 0, 0, 0}; // blockLen 单位字节
78+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
79+ AscendC::DataCopyPad(xUb, xGm_, params, padParams); // 精确搬入 n 个元素
80+ AscendC::DataCopyPad(yUb, yGm_, params, padParams);
81+ inQueueX_.EnQue(xUb);
82+ inQueueY_.EnQue(yUb);
83+ }
84+ 
85+ __aicore__ inline void Compute()
86+ {
87+ AscendC::LocalTensor<float> xUb = inQueueX_.DeQue<float>();
88+ AscendC::LocalTensor<float> yUb = inQueueY_.DeQue<float>();
89+ AscendC::LocalTensor<float> zUb = outQueueZ_.AllocTensor<float>();
90+ 
91+ auto* x = (__ubuf__ float*)xUb.GetPhyAddr();
92+ auto* y = (__ubuf__ float*)yUb.GetPhyAddr();
93+ auto* z = (__ubuf__ float*)zUb.GetPhyAddr();
94+ for (int r = 0; r < VF_REPEAT; r++) {
95+ asc_vf_call<ElemwiseAxpyVf>(x, y, z, A, n_, static_cast<uint16_t>(loopNum_));
96+ }
97+ 
98+ outQueueZ_.EnQue(zUb);
99+ inQueueX_.FreeTensor(xUb);
100+ inQueueY_.FreeTensor(yUb);
101+ }
102+ 
103+ __aicore__ inline void CopyOut()
104+ {
105+ AscendC::LocalTensor<float> zUb = outQueueZ_.DeQue<float>();
106+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(n_ * sizeof(float)), 0, 0, 0};
107+ AscendC::DataCopyPad(zGm_, zUb, params); // 精确搬出 n 个元素
108+ outQueueZ_.FreeTensor(zUb);
109+ }
110+ 
111+ AscendC::TPipe pipe_;
112+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueX_, inQueueY_;
113+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueZ_;
114+ AscendC::GlobalTensor<float> xGm_, yGm_, zGm_;
115+ uint32_t n_, loopNum_;
116+};
117+ 
118+__global__ __aicore__ __vector__ void ElemwiseAxpyKernel(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t n)
119+{
120+ KernelElemwiseAxpy op;
121+ op.Init(x, y, z, n);
122+ op.Process();
123+}
124+ 
125+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
126+int main()
127+{
128+ CHECK_ACL(aclInit(nullptr));
129+ CHECK_ACL(aclrtSetDevice(0));
130+ aclrtStream stream;
131+ CHECK_ACL(aclrtCreateStream(&stream));
132+ 
133+ size_t bytes = N * sizeof(float);
134+ std::vector<float> hX = vf::GenInput(N); // x、y 共用同一份随机输入
135+ std::vector<float> ref(N);
136+ for (uint32_t i = 0; i < N; i++) ref[i] = A * hX[i] + hX[i]; // 标杆：z = a*x + y（x=y）
137+ std::vector<float> hZ(N, 0.f);
138+ 
139+ uint8_t *dX, *dY, *dZ;
140+ CHECK_ACL(aclrtMalloc((void**)&dX, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
141+ CHECK_ACL(aclrtMalloc((void**)&dY, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
142+ CHECK_ACL(aclrtMalloc((void**)&dZ, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
143+ CHECK_ACL(aclrtMemcpy(dX, bytes, hX.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE));
144+ CHECK_ACL(aclrtMemcpy(dY, bytes, hX.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE));
145+ 
146+ ElemwiseAxpyKernel<<<1, nullptr, stream>>>(dX, dY, dZ, N);
147+ CHECK_ACL(aclrtSynchronizeStream(stream));
148+ CHECK_ACL(aclrtMemcpy(hZ.data(), bytes, dZ, bytes, ACL_MEMCPY_DEVICE_TO_HOST));
149+ 
150+ bool ok = vf::VerifyAbs(hZ, ref, 1e-4f);
151+ std::cout << "[elemwise_axpy baseline] N=" << N << " z[0..2]=" << hZ[0] << " " << hZ[1] << " " << hZ[2]
152+ << " | z[" << (N - 1) << "]=" << hZ[N - 1] << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
153+ 
154+ aclrtFree(dX); aclrtFree(dY); aclrtFree(dZ);
155+ CHECK_ACL(aclrtDestroyStream(stream));
156+ CHECK_ACL(aclrtResetDevice(0));
157+ CHECK_ACL(aclFinalize());
158+ return ok ? 0 : 1;
159+}
ASamples/2_Performance/simd_vf_story/elemwise/src/elemwise_axpy_pragma.asc+152-0
@@ -0,0 +1,152 @@
1+/**
2+ * 单 VF 拆分版：Elementwise AXPY: z = a * x + y（#pragma 展开 VF，官方惯用写法）
3+ *
4+ * 从 elemwise_axpy.asc 拆出「#pragma 展开」一个 VF（ElemwiseAxpyVfPragma）：循环体与基线一样
5+ * （单组寄存器），只在 for 前加 #pragma unroll 4 让编译器展开成多条独立指令流。代码最短、改 N
6+ * 即调展开因子（官方 gelu_eltwise 高性能样例同款）。基线 / 手写展开版见同目录另两个文件。
7+ *
8+ * 统一 case 规格（本目录 16 个 VF 文件全部一致）：[D0,D1]=[78,250]，N=D0*D1=19500。
9+ *
10+ * 编译 & 运行：
11+ * cmake --build build --target elemwise_axpy_pragma
12+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/elemwise_axpy_pragma -s Ascend950
13+ */
14+ 
15+#include <iostream>
16+#include <vector>
17+#include <cmath>
18+#include "acl/acl.h"
19+#include "kernel_operator.h"
20+#include "../../include/vf_common.h"
21+ 
22+static constexpr float A = 2.0f;
23+ 
24+// ============ VF 层（#pragma 展开）：循环体同基线，#pragma unroll 4 由编译器展开 ============
25+__simd_vf__ inline void ElemwiseAxpyVfPragma(__ubuf__ float* xAddr, __ubuf__ float* yAddr,
26+ __ubuf__ float* zAddr, float a, uint32_t n, uint16_t loopNum)
27+{
28+ AscendC::Reg::RegTensor<float> vx, vy, vt;
29+ AscendC::Reg::MaskReg mask;
30+ uint32_t count = n;
31+#pragma unroll 4
32+ for (uint16_t i = 0; i < loopNum; i++) {
33+ mask = AscendC::Reg::UpdateMask<float>(count); // 每轮自动扣 VL；展开后各副本依次扣减
34+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vx, xAddr + i * VL_B32);
35+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vy, yAddr + i * VL_B32);
36+ AscendC::Reg::Muls(vt, vx, a, mask); // vt = a * x
37+ AscendC::Reg::Add(vt, vt, vy, mask); // vt = a * x + y
38+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(zAddr + i * VL_B32, vt, mask);
39+ }
40+}
41+ 
42+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
43+class KernelElemwiseAxpy {
44+public:
45+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t n)
46+ {
47+ n_ = n;
48+ loopNum_ = (n_ + VL_B32 - 1) / VL_B32;
49+ uint32_t bufBytes = loopNum_ * VL_B32 * sizeof(float);
50+ xGm_.SetGlobalBuffer((__gm__ float*)x, n_);
51+ yGm_.SetGlobalBuffer((__gm__ float*)y, n_);
52+ zGm_.SetGlobalBuffer((__gm__ float*)z, n_);
53+ pipe_.InitBuffer(inQueueX_, 1, bufBytes);
54+ pipe_.InitBuffer(inQueueY_, 1, bufBytes);
55+ pipe_.InitBuffer(outQueueZ_, 1, bufBytes);
56+ }
57+ 
58+ __aicore__ inline void Process()
59+ {
60+ CopyIn();
61+ Compute();
62+ CopyOut();
63+ }
64+ 
65+private:
66+ __aicore__ inline void CopyIn()
67+ {
68+ AscendC::LocalTensor<float> xUb = inQueueX_.AllocTensor<float>();
69+ AscendC::LocalTensor<float> yUb = inQueueY_.AllocTensor<float>();
70+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(n_ * sizeof(float)), 0, 0, 0};
71+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
72+ AscendC::DataCopyPad(xUb, xGm_, params, padParams);
73+ AscendC::DataCopyPad(yUb, yGm_, params, padParams);
74+ inQueueX_.EnQue(xUb);
75+ inQueueY_.EnQue(yUb);
76+ }
77+ 
78+ __aicore__ inline void Compute()
79+ {
80+ AscendC::LocalTensor<float> xUb = inQueueX_.DeQue<float>();
81+ AscendC::LocalTensor<float> yUb = inQueueY_.DeQue<float>();
82+ AscendC::LocalTensor<float> zUb = outQueueZ_.AllocTensor<float>();
83+ 
84+ auto* x = (__ubuf__ float*)xUb.GetPhyAddr();
85+ auto* y = (__ubuf__ float*)yUb.GetPhyAddr();
86+ auto* z = (__ubuf__ float*)zUb.GetPhyAddr();
87+ for (int r = 0; r < VF_REPEAT; r++) {
88+ asc_vf_call<ElemwiseAxpyVfPragma>(x, y, z, A, n_, static_cast<uint16_t>(loopNum_));
89+ }
90+ 
91+ outQueueZ_.EnQue(zUb);
92+ inQueueX_.FreeTensor(xUb);
93+ inQueueY_.FreeTensor(yUb);
94+ }
95+ 
96+ __aicore__ inline void CopyOut()
97+ {
98+ AscendC::LocalTensor<float> zUb = outQueueZ_.DeQue<float>();
99+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(n_ * sizeof(float)), 0, 0, 0};
100+ AscendC::DataCopyPad(zGm_, zUb, params);
101+ outQueueZ_.FreeTensor(zUb);
102+ }
103+ 
104+ AscendC::TPipe pipe_;
105+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueX_, inQueueY_;
106+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueZ_;
107+ AscendC::GlobalTensor<float> xGm_, yGm_, zGm_;
108+ uint32_t n_, loopNum_;
109+};
110+ 
111+__global__ __aicore__ __vector__ void ElemwiseAxpyKernel(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t n)
112+{
113+ KernelElemwiseAxpy op;
114+ op.Init(x, y, z, n);
115+ op.Process();
116+}
117+ 
118+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
119+int main()
120+{
121+ CHECK_ACL(aclInit(nullptr));
122+ CHECK_ACL(aclrtSetDevice(0));
123+ aclrtStream stream;
124+ CHECK_ACL(aclrtCreateStream(&stream));
125+ 
126+ size_t bytes = N * sizeof(float);
127+ std::vector<float> hX = vf::GenInput(N); // x、y 共用同一份随机输入
128+ std::vector<float> ref(N);
129+ for (uint32_t i = 0; i < N; i++) ref[i] = A * hX[i] + hX[i]; // 标杆：z = a*x + y（x=y）
130+ std::vector<float> hZ(N, 0.f);
131+ 
132+ uint8_t *dX, *dY, *dZ;
133+ CHECK_ACL(aclrtMalloc((void**)&dX, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
134+ CHECK_ACL(aclrtMalloc((void**)&dY, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
135+ CHECK_ACL(aclrtMalloc((void**)&dZ, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
136+ CHECK_ACL(aclrtMemcpy(dX, bytes, hX.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE));
137+ CHECK_ACL(aclrtMemcpy(dY, bytes, hX.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE));
138+ 
139+ ElemwiseAxpyKernel<<<1, nullptr, stream>>>(dX, dY, dZ, N);
140+ CHECK_ACL(aclrtSynchronizeStream(stream));
141+ CHECK_ACL(aclrtMemcpy(hZ.data(), bytes, dZ, bytes, ACL_MEMCPY_DEVICE_TO_HOST));
142+ 
143+ bool ok = vf::VerifyAbs(hZ, ref, 1e-4f);
144+ std::cout << "[elemwise_axpy pragma] N=" << N << " z[0..2]=" << hZ[0] << " " << hZ[1] << " " << hZ[2]
145+ << " | z[" << (N - 1) << "]=" << hZ[N - 1] << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
146+ 
147+ aclrtFree(dX); aclrtFree(dY); aclrtFree(dZ);
148+ CHECK_ACL(aclrtDestroyStream(stream));
149+ CHECK_ACL(aclrtResetDevice(0));
150+ CHECK_ACL(aclFinalize());
151+ return ok ? 0 : 1;
152+}
ASamples/2_Performance/simd_vf_story/elemwise/src/elemwise_axpy_unroll.asc+173-0
@@ -0,0 +1,173 @@
1+/**
2+ * 单 VF 拆分版：Elementwise AXPY: z = a * x + y（手写展开 VF）
3+ *
4+ * 从 elemwise_axpy.asc 拆出「手写展开」一个 VF（ElemwiseAxpyVfUnroll）：一轮用 UNROLL=4 组独立
5+ * 命名寄存器算 UNROLL 个 VL 块，显式暴露多条独立指令流（双发 + 延迟隐藏）。基线 / #pragma 版见
6+ * 同目录 elemwise_axpy_baseline.asc / elemwise_axpy_pragma.asc。
7+ *
8+ * 统一 case 规格（本目录 16 个 VF 文件全部一致）：[D0,D1]=[78,250]，N=D0*D1=19500。
9+ *
10+ * 编译 & 运行：
11+ * cmake --build build --target elemwise_axpy_unroll
12+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/elemwise_axpy_unroll -s Ascend950
13+ */
14+ 
15+#include <iostream>
16+#include <vector>
17+#include <cmath>
18+#include "acl/acl.h"
19+#include "kernel_operator.h"
20+#include "../../include/vf_common.h"
21+ 
22+static constexpr uint16_t UNROLL = 4; // 手动展开因子：一轮算 4 个 VL 块
23+static constexpr float A = 2.0f;
24+ 
25+// ============ VF 层（手写展开）：一轮循环算 UNROLL 个 VL 块 ============
26+// 用 UNROLL 组互相独立的命名寄存器（VF 不支持 RegTensor 数组下标）。尾部凑不满 UNROLL 的块：
27+// count 扣到 0 后 UpdateMask 返回空 mask，Load 多读落在已分配 UB、Store 不写，无副作用。
28+__simd_vf__ inline void ElemwiseAxpyVfUnroll(__ubuf__ float* xAddr, __ubuf__ float* yAddr,
29+ __ubuf__ float* zAddr, float a, uint32_t n, uint16_t loopNum)
30+{
31+ static_assert(UNROLL == 4, "下面是 4 组命名寄存器手动展开，改 UNROLL 需同步增减寄存器组");
32+ AscendC::Reg::RegTensor<float> vx0, vx1, vx2, vx3, vy0, vy1, vy2, vy3, vt0, vt1, vt2, vt3;
33+ AscendC::Reg::MaskReg m0, m1, m2, m3;
34+ uint32_t count = n;
35+ const uint16_t outer = (loopNum + UNROLL - 1) / UNROLL; // 外层轮数 = 块数 / 展开因子（向上取整）
36+ 
37+ for (uint16_t i = 0; i < outer; i++) {
38+ const uint16_t base = i * UNROLL;
39+ m0 = AscendC::Reg::UpdateMask<float>(count);
40+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vx0, xAddr + (base + 0) * VL_B32);
41+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vy0, yAddr + (base + 0) * VL_B32);
42+ m1 = AscendC::Reg::UpdateMask<float>(count);
43+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vx1, xAddr + (base + 1) * VL_B32);
44+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vy1, yAddr + (base + 1) * VL_B32);
45+ m2 = AscendC::Reg::UpdateMask<float>(count);
46+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vx2, xAddr + (base + 2) * VL_B32);
47+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vy2, yAddr + (base + 2) * VL_B32);
48+ m3 = AscendC::Reg::UpdateMask<float>(count);
49+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vx3, xAddr + (base + 3) * VL_B32);
50+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(vy3, yAddr + (base + 3) * VL_B32);
51+ AscendC::Reg::Muls(vt0, vx0, a, m0); AscendC::Reg::Add(vt0, vt0, vy0, m0);
52+ AscendC::Reg::Muls(vt1, vx1, a, m1); AscendC::Reg::Add(vt1, vt1, vy1, m1);
53+ AscendC::Reg::Muls(vt2, vx2, a, m2); AscendC::Reg::Add(vt2, vt2, vy2, m2);
54+ AscendC::Reg::Muls(vt3, vx3, a, m3); AscendC::Reg::Add(vt3, vt3, vy3, m3);
55+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(zAddr + (base + 0) * VL_B32, vt0, m0);
56+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(zAddr + (base + 1) * VL_B32, vt1, m1);
57+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(zAddr + (base + 2) * VL_B32, vt2, m2);
58+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(zAddr + (base + 3) * VL_B32, vt3, m3);
59+ }
60+}
61+ 
62+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
63+class KernelElemwiseAxpy {
64+public:
65+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t n)
66+ {
67+ n_ = n;
68+ loopNum_ = (n_ + VL_B32 - 1) / VL_B32;
69+ uint32_t chunksPad = (loopNum_ + UNROLL - 1) / UNROLL * UNROLL; // 凑齐 UNROLL，覆盖手写展开的多读
70+ uint32_t bufBytes = chunksPad * VL_B32 * sizeof(float);
71+ xGm_.SetGlobalBuffer((__gm__ float*)x, n_);
72+ yGm_.SetGlobalBuffer((__gm__ float*)y, n_);
73+ zGm_.SetGlobalBuffer((__gm__ float*)z, n_);
74+ pipe_.InitBuffer(inQueueX_, 1, bufBytes);
75+ pipe_.InitBuffer(inQueueY_, 1, bufBytes);
76+ pipe_.InitBuffer(outQueueZ_, 1, bufBytes);
77+ }
78+ 
79+ __aicore__ inline void Process()
80+ {
81+ CopyIn();
82+ Compute();
83+ CopyOut();
84+ }
85+ 
86+private:
87+ __aicore__ inline void CopyIn()
88+ {
89+ AscendC::LocalTensor<float> xUb = inQueueX_.AllocTensor<float>();
90+ AscendC::LocalTensor<float> yUb = inQueueY_.AllocTensor<float>();
91+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(n_ * sizeof(float)), 0, 0, 0};
92+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
93+ AscendC::DataCopyPad(xUb, xGm_, params, padParams);
94+ AscendC::DataCopyPad(yUb, yGm_, params, padParams);
95+ inQueueX_.EnQue(xUb);
96+ inQueueY_.EnQue(yUb);
97+ }
98+ 
99+ __aicore__ inline void Compute()
100+ {
101+ AscendC::LocalTensor<float> xUb = inQueueX_.DeQue<float>();
102+ AscendC::LocalTensor<float> yUb = inQueueY_.DeQue<float>();
103+ AscendC::LocalTensor<float> zUb = outQueueZ_.AllocTensor<float>();
104+ 
105+ auto* x = (__ubuf__ float*)xUb.GetPhyAddr();
106+ auto* y = (__ubuf__ float*)yUb.GetPhyAddr();
107+ auto* z = (__ubuf__ float*)zUb.GetPhyAddr();
108+ for (int r = 0; r < VF_REPEAT; r++) {
109+ asc_vf_call<ElemwiseAxpyVfUnroll>(x, y, z, A, n_, static_cast<uint16_t>(loopNum_));
110+ }
111+ 
112+ outQueueZ_.EnQue(zUb);
113+ inQueueX_.FreeTensor(xUb);
114+ inQueueY_.FreeTensor(yUb);
115+ }
116+ 
117+ __aicore__ inline void CopyOut()
118+ {
119+ AscendC::LocalTensor<float> zUb = outQueueZ_.DeQue<float>();
120+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(n_ * sizeof(float)), 0, 0, 0};
121+ AscendC::DataCopyPad(zGm_, zUb, params);
122+ outQueueZ_.FreeTensor(zUb);
123+ }
124+ 
125+ AscendC::TPipe pipe_;
126+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueX_, inQueueY_;
127+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueZ_;
128+ AscendC::GlobalTensor<float> xGm_, yGm_, zGm_;
129+ uint32_t n_, loopNum_;
130+};
131+ 
132+__global__ __aicore__ __vector__ void ElemwiseAxpyKernel(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t n)
133+{
134+ KernelElemwiseAxpy op;
135+ op.Init(x, y, z, n);
136+ op.Process();
137+}
138+ 
139+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
140+int main()
141+{
142+ CHECK_ACL(aclInit(nullptr));
143+ CHECK_ACL(aclrtSetDevice(0));
144+ aclrtStream stream;
145+ CHECK_ACL(aclrtCreateStream(&stream));
146+ 
147+ size_t bytes = N * sizeof(float);
148+ std::vector<float> hX = vf::GenInput(N); // x、y 共用同一份随机输入
149+ std::vector<float> ref(N);
150+ for (uint32_t i = 0; i < N; i++) ref[i] = A * hX[i] + hX[i]; // 标杆：z = a*x + y（x=y）
151+ std::vector<float> hZ(N, 0.f);
152+ 
153+ uint8_t *dX, *dY, *dZ;
154+ CHECK_ACL(aclrtMalloc((void**)&dX, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
155+ CHECK_ACL(aclrtMalloc((void**)&dY, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
156+ CHECK_ACL(aclrtMalloc((void**)&dZ, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
157+ CHECK_ACL(aclrtMemcpy(dX, bytes, hX.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE));
158+ CHECK_ACL(aclrtMemcpy(dY, bytes, hX.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE));
159+ 
160+ ElemwiseAxpyKernel<<<1, nullptr, stream>>>(dX, dY, dZ, N);
161+ CHECK_ACL(aclrtSynchronizeStream(stream));
162+ CHECK_ACL(aclrtMemcpy(hZ.data(), bytes, dZ, bytes, ACL_MEMCPY_DEVICE_TO_HOST));
163+ 
164+ bool ok = vf::VerifyAbs(hZ, ref, 1e-4f);
165+ std::cout << "[elemwise_axpy unroll] N=" << N << " z[0..2]=" << hZ[0] << " " << hZ[1] << " " << hZ[2]
166+ << " | z[" << (N - 1) << "]=" << hZ[N - 1] << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
167+ 
168+ aclrtFree(dX); aclrtFree(dY); aclrtFree(dZ);
169+ CHECK_ACL(aclrtDestroyStream(stream));
170+ CHECK_ACL(aclrtResetDevice(0));
171+ CHECK_ACL(aclFinalize());
172+ return ok ? 0 : 1;
173+}
ASamples/2_Performance/simd_vf_story/include/vf_common.h+71-0
@@ -0,0 +1,71 @@
1+/**
2+ * 单 VF 拆分版的共享头：统一 case 规格 + ACL 错误检查宏 + host 侧输入生成/比对（elemwise / reduce 两个用例的 16 个 .asc 共用）。
3+ *
4+ * - 统一规格 [D0,D1]=[78,250]，N=D0*D1=19500。按最受限算子 elemwise（x/y/z 三块 UB buffer，约
5+ * 234KB<256KB UB）反推尺寸；reduce 类只需一块输入 buffer。D1>VL 且 %8≠0、D0%4=2，保留分块/
6+ * 尾块/展开尾行路径。改规格即改此处常量。
7+ * - 输入生成（GenInput）统一放这里：固定种子、16 个 VF 共用同一份随机输入。各算子的「标杆 golden」
8+ * 因实现而异（max/sum、ar/ra），留在各自 .asc 的 main 里就地计算。
9+ * - 算子专属算法常量（AXPY 的 A、展开因子 UNROLL、Max 归约初值 NEG_INF）不属于「规格」，各自留在对应文件。
10+ */
11+ 
12+#ifndef VF_COMMON_H
13+#define VF_COMMON_H
14+ 
15+#include <cstdint>
16+#include <cmath>
17+#include <iostream>
18+#include <random>
19+#include <vector>
20+#include "acl/acl.h"
21+ 
22+// ACL 调用错误检查（在 main 中使用，失败返回 1）
23+#define CHECK_ACL(call) \
24+ do { \
25+ aclError err = (call); \
26+ if (err != ACL_SUCCESS) { \
27+ std::cerr << "ACL error " << err << " at " << __LINE__ << "\n"; \
28+ return 1; \
29+ } \
30+ } while (0)
31+ 
32+// ===== 统一 case 规格 =====
33+static constexpr uint32_t VL_B32 = 256 / sizeof(float); // 64 lane
34+static constexpr uint32_t UB_ALIGN = 32 / sizeof(float); // 8：UB 行按 32B 对齐（reduce 类用）
35+static constexpr uint32_t D0 = 78; // 统一形状首维（%4=2）
36+static constexpr uint32_t D1 = 250; // 统一形状尾维（>VL 且 %8≠0）
37+static constexpr uint32_t N = D0 * D1; // 19500：elemwise 按一维向量处理
38+static constexpr int VF_REPEAT = 5; // 每个 VF 在 kernel 内连续跑的次数（profiling 取稳态）
39+ 
40+namespace vf {
41+ 
42+// 确定性随机输入（固定种子，16 个 VF 共用）；范围 [-1,1) 使 sum 良态
43+inline std::vector<float> GenInput(uint32_t n)
44+{
45+ std::vector<float> x(n);
46+ std::mt19937 gen(0u);
47+ std::uniform_real_distribution<float> dist(-1.0f, 1.0f);
48+ for (uint32_t i = 0; i < n; i++) x[i] = dist(gen);
49+ return x;
50+}
51+ 
52+// 绝对容差比对（max / elemwise：精确）
53+inline bool VerifyAbs(const std::vector<float>& got, const std::vector<float>& ref, float tol)
54+{
55+ for (size_t k = 0; k < ref.size(); k++)
56+ if (std::fabs(got[k] - ref[k]) > tol) return false;
57+ return true;
58+}
59+ 
60+// 相对容差比对（sum：浮点重排有低位差异）
61+inline bool VerifyRel(const std::vector<float>& got, const std::vector<float>& ref, double rtol)
62+{
63+ for (size_t k = 0; k < ref.size(); k++)
64+ if (std::fabs((double)got[k] - (double)ref[k]) > rtol * (std::fabs((double)ref[k]) + 1.0))
65+ return false;
66+ return true;
67+}
68+ 
69+} // namespace vf
70+ 
71+#endif // VF_COMMON_H
ASamples/2_Performance/simd_vf_story/reduce/images/ar_vs_ra.svg+1-0
@@ -0,0 +1 @@
1+<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 880 320" font-family="Noto Sans CJK SC,Microsoft YaHei,PingFang SC,Segoe UI,Helvetica,Arial,sans-serif"><defs><marker id="ab" markerWidth="10" markerHeight="10" refX="8" refY="3.2" orient="auto"><path d="M0,0 L8,3.2 L0,6.4 Z" fill="#2b6cb0"/></marker><marker id="ag" markerWidth="10" markerHeight="10" refX="8" refY="3.2" orient="auto"><path d="M0,0 L8,3.2 L0,6.4 Z" fill="#2f855a"/></marker></defs><rect width="880" height="320" fill="white"/><text x="24" y="30" font-size="17" fill="#1f2933" text-anchor="start" font-weight="700">两个归约轴：ar（axis=1，行内）vs ra（axis=0，跨行）</text><text x="24" y="52" font-size="12" fill="#52606d" text-anchor="start">输入矩阵 [D0, D1] = [78 行, 250 列]，下用 3×5 网格示意；紫色块 = 输出向量。</text><text x="40" y="86" font-size="13" fill="#2b6cb0" text-anchor="start" font-weight="700">ar：沿 D1（行内）归约 · 每行 → 1 标量 · 输出 [D0]（=78）</text><text x="171.0" y="101" font-size="10" fill="#9aa5b1" text-anchor="middle">D1（250 列）→</text><text x="84" y="156.0" font-size="10" fill="#9aa5b1" text-anchor="middle" transform="rotate(-90 84 156.0)">D0</text><rect x="96" y="108" width="30" height="30" fill="#dbeafe" stroke="#cdd5df" stroke-width="1"/><rect x="126" y="108" width="30" height="30" fill="#dbeafe" stroke="#cdd5df" stroke-width="1"/><rect x="156" y="108" width="30" height="30" fill="#dbeafe" stroke="#cdd5df" stroke-width="1"/><rect x="186" y="108" width="30" height="30" fill="#dbeafe" stroke="#cdd5df" stroke-width="1"/><rect x="216" y="108" width="30" height="30" fill="#dbeafe" stroke="#cdd5df" stroke-width="1"/><rect x="96" y="138" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="126" y="138" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="156" y="138" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="186" y="138" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="216" y="138" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="96" y="168" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="126" y="168" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="156" y="168" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="186" y="168" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="216" y="168" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><line x1="250" y1="123.0" x2="293" y2="123.0" stroke="#2b6cb0" stroke-width="2" marker-end="url(#ab)"/><rect x="296" y="109" width="30" height="28" rx="5" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><line x1="250" y1="153.0" x2="293" y2="153.0" stroke="#2b6cb0" stroke-width="2" marker-end="url(#ab)"/><rect x="296" y="139" width="30" height="28" rx="5" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><line x1="250" y1="183.0" x2="293" y2="183.0" stroke="#2b6cb0" stroke-width="2" marker-end="url(#ab)"/><rect x="296" y="169" width="30" height="28" rx="5" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><text x="311.0" y="216" font-size="12" fill="#7e57c2" text-anchor="middle" font-weight="700">[D0]</text><text x="40" y="242" font-size="11" fill="#52606d" text-anchor="start">· 每行向量内 ReduceMax → 标量；StoreUnAlign 拼标量流</text><text x="40" y="260" font-size="11" fill="#52606d" text-anchor="start">· 写法：baseline / unroll / pragma（binary 仅 reduce_sum）</text><text x="456" y="86" font-size="13" fill="#2f855a" text-anchor="start" font-weight="700">ra：沿 D0（跨行）归约 · 每列 → 1 值 · 输出 [D1]（=250）</text><text x="587.0" y="101" font-size="10" fill="#9aa5b1" text-anchor="middle">D1（250 列）→</text><text x="500" y="156.0" font-size="10" fill="#9aa5b1" text-anchor="middle" transform="rotate(-90 500 156.0)">D0</text><rect x="512" y="108" width="30" height="30" fill="#d1fae5" stroke="#cdd5df" stroke-width="1"/><rect x="542" y="108" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="572" y="108" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="602" y="108" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="632" y="108" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="512" y="138" width="30" height="30" fill="#d1fae5" stroke="#cdd5df" stroke-width="1"/><rect x="542" y="138" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="572" y="138" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="602" y="138" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="632" y="138" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="512" y="168" width="30" height="30" fill="#d1fae5" stroke="#cdd5df" stroke-width="1"/><rect x="542" y="168" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="572" y="168" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="602" y="168" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><rect x="632" y="168" width="30" height="30" fill="#f7f9fc" stroke="#cdd5df" stroke-width="1"/><line x1="527.0" y1="202" x2="527.0" y2="233" stroke="#2f855a" stroke-width="2" marker-end="url(#ag)"/><rect x="513" y="236" width="28" height="28" rx="5" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><line x1="557.0" y1="202" x2="557.0" y2="233" stroke="#2f855a" stroke-width="2" marker-end="url(#ag)"/><rect x="543" y="236" width="28" height="28" rx="5" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><line x1="587.0" y1="202" x2="587.0" y2="233" stroke="#2f855a" stroke-width="2" marker-end="url(#ag)"/><rect x="573" y="236" width="28" height="28" rx="5" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><line x1="617.0" y1="202" x2="617.0" y2="233" stroke="#2f855a" stroke-width="2" marker-end="url(#ag)"/><rect x="603" y="236" width="28" height="28" rx="5" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><line x1="647.0" y1="202" x2="647.0" y2="233" stroke="#2f855a" stroke-width="2" marker-end="url(#ag)"/><rect x="633" y="236" width="28" height="28" rx="5" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><text x="676" y="251.0" font-size="12" fill="#7e57c2" text-anchor="start" font-weight="700">[D1]</text><text x="456" y="288" font-size="11" fill="#52606d" text-anchor="start">· 跨行逐元素 Max（默认 ZEROING）；无向量内归约 → StoreAlign 整段</text><text x="456" y="306" font-size="11" fill="#52606d" text-anchor="start">· 写法：baseline / unroll / pragma（无 binary）</text><line x1="430" y1="74" x2="430" y2="310" stroke="#e3e8ee" stroke-width="1"/></svg>
ASamples/2_Performance/simd_vf_story/reduce/images/binary_fold.svg+1-0
@@ -0,0 +1 @@
1+<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 900 414" font-family="Noto Sans CJK SC,Microsoft YaHei,PingFang SC,Segoe UI,Helvetica,Arial,sans-serif"><defs><marker id="mg" viewBox="0 0 10 10" refX="9" refY="5" markerWidth="8.5" markerHeight="8.5" orient="auto" markerUnits="userSpaceOnUse"><path d="M0,0 L10,5 L0,10 z" fill="#2f855a"/></marker><marker id="mb" viewBox="0 0 10 10" refX="9" refY="5" markerWidth="8.5" markerHeight="8.5" orient="auto" markerUnits="userSpaceOnUse"><path d="M0,0 L10,5 L0,10 z" fill="#2b6cb0"/></marker><marker id="mr" viewBox="0 0 10 10" refX="9" refY="5" markerWidth="8.5" markerHeight="8.5" orient="auto" markerUnits="userSpaceOnUse"><path d="M0,0 L10,5 L0,10 z" fill="#d64545"/></marker></defs><rect width="900" height="414" fill="white"/><text x="24" y="32" font-size="18" fill="#1f2933" text-anchor="start" font-weight="700">reduce_sum_ar_binary：一行 250 元素的二分（成对）折叠</text><text x="24" y="52" font-size="12" fill="#52606d" text-anchor="start">每行 = 4 个 VL(64) 块，foldPoint=128（2 块）。配对相距 128 的块 → 串行步数 O(块)→O(log 块)，且成对求和降舍入误差。</text><rect x="40" y="86" width="176" height="30" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="128" y="106" font-size="13" fill="#1f2933" text-anchor="middle">c0 [0:64)</text><rect x="228" y="86" width="176" height="30" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="316" y="106" font-size="13" fill="#1f2933" text-anchor="middle">c1 [64:128)</text><rect x="416" y="86" width="176" height="30" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="504" y="106" font-size="13" fill="#1f2933" text-anchor="middle">c2 [128:192)</text><rect x="604" y="86" width="176" height="30" rx="6" fill="#fdeaea" stroke="#d64545" stroke-width="1.5"/><text x="692" y="106" font-size="13" fill="#1f2933" text-anchor="middle">c3 [192:250)</text><text x="780" y="78" font-size="11" fill="#d64545" text-anchor="end">尾 [250:256) padding → mask 屏蔽</text><text x="24" y="172" font-size="13" fill="#2f855a" text-anchor="start" font-weight="700">第 1 层（foldPoint=128，两条 Add 互相独立）</text><rect x="228" y="182" width="176" height="30" rx="6" fill="#eef2f7" stroke="#2f855a" stroke-width="1.5"/><text x="316" y="202" font-size="13" fill="#2f855a" text-anchor="middle">c0 + c2</text><rect x="416" y="182" width="176" height="30" rx="6" fill="#eef2f7" stroke="#2f855a" stroke-width="1.5"/><text x="504" y="201" font-size="11" fill="#2f855a" text-anchor="middle">c1 + c3 (MERGING,58)</text><path d="M 128 116 C 128 150, 292 148, 292 182" fill="none" stroke="#2f855a" stroke-width="2" marker-end="url(#mg)"/><path d="M 504 116 C 504 150, 340 148, 340 182" fill="none" stroke="#2f855a" stroke-width="2" marker-end="url(#mg)"/><path d="M 316 116 C 316 150, 468 148, 468 182" fill="none" stroke="#2b6cb0" stroke-width="2" marker-end="url(#mb)"/><path d="M 692 116 C 692 150, 540 148, 540 182" fill="none" stroke="#2b6cb0" stroke-width="2" marker-end="url(#mb)"/><text x="24" y="264" font-size="13" fill="#2f855a" text-anchor="start" font-weight="700">第 2 层：每 lane 含 4 块之和（64 lane）</text><rect x="320" y="278" width="180" height="30" rx="6" fill="#eef2f7" stroke="#2f855a" stroke-width="1.5"/><text x="410" y="298" font-size="12" fill="#2f855a" text-anchor="middle">(c0+c2) + (c1+c3)</text><path d="M 316 212 C 316 245, 388 245, 388 278" fill="none" stroke="#d64545" stroke-width="2.2" marker-end="url(#mr)"/><path d="M 504 212 C 504 245, 432 245, 432 278" fill="none" stroke="#d64545" stroke-width="2.2" marker-end="url(#mr)"/><rect x="326" y="348" width="168" height="30" rx="6" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><text x="410" y="367" font-size="10" fill="#1f2933" text-anchor="middle">向量内 ReduceSum → 行标量</text><path d="M 410 308 L 410 348" fill="none" stroke="#d64545" stroke-width="2.2" marker-end="url(#mr)"/><rect x="600" y="318" width="278" height="86" rx="8" fill="#f7faf7" stroke="#2f855a" stroke-width="1.2"/><text x="616" y="342" font-size="13" fill="#1f2933" text-anchor="start" font-weight="700">对照 baseline（线性逐块折叠）</text><text x="616" y="364" font-size="11" fill="#52606d" text-anchor="start">acc += c0→c1→c2→c3（串行步数 4），再 ReduceSum</text><text x="616" y="388" font-size="12" fill="#d64545" text-anchor="start" font-weight="700">PUSHQ VF：792 → 384 cycle（≈2×）</text></svg>
ASamples/2_Performance/simd_vf_story/reduce/images/dep_chain.svg+1-0
@@ -0,0 +1 @@
1+<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 900 372" font-family="Noto Sans CJK SC,Microsoft YaHei,PingFang SC,Segoe UI,Helvetica,Arial,sans-serif"><defs><marker id="a" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#1f2933"/></marker><marker id="ar" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#d64545"/></marker><marker id="ag" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#2f855a"/></marker><marker id="ab" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#2b6cb0"/></marker></defs><rect width="900" height="372" fill="white"/><text x="24" y="32" font-size="18" fill="#1f2933" text-anchor="start" font-weight="700">归约的两种组织方式（baseline 单路串行 vs unroll 多路并行，以 N=8 为例）</text><text x="24" y="52" font-size="12" fill="#52606d" text-anchor="start">红 = 最长串行路径（决定时延）。优化 = 缩短它。</text><text x="24" y="84" font-size="14" fill="#1f2933" text-anchor="start" font-weight="700">① baseline 单路串行顺序折叠　串行步数 = N = 8</text><rect x="64" y="96" width="50" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="89.0" y="112.2" text-anchor="middle" font-size="12" fill="#1f2933">c0</text><rect x="162" y="96" width="50" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="187.0" y="112.2" text-anchor="middle" font-size="12" fill="#1f2933">c1</text><rect x="260" y="96" width="50" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="285.0" y="112.2" text-anchor="middle" font-size="12" fill="#1f2933">c2</text><rect x="358" y="96" width="50" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="383.0" y="112.2" text-anchor="middle" font-size="12" fill="#1f2933">c3</text><rect x="456" y="96" width="50" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="481.0" y="112.2" text-anchor="middle" font-size="12" fill="#1f2933">c4</text><rect x="554" y="96" width="50" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="579.0" y="112.2" text-anchor="middle" font-size="12" fill="#1f2933">c5</text><rect x="652" y="96" width="50" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="677.0" y="112.2" text-anchor="middle" font-size="12" fill="#1f2933">c6</text><rect x="750" y="96" width="50" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="775.0" y="112.2" text-anchor="middle" font-size="12" fill="#1f2933">c7</text><rect x="64" y="152" width="50" height="24" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="89.0" y="167.85" text-anchor="middle" font-size="11" fill="#1f2933">acc</text><line x1="89.0" y1="120" x2="89.0" y2="152" stroke="#9aa5b1" stroke-width="1.3" marker-end="url(#a)"/><rect x="162" y="152" width="50" height="24" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="187.0" y="167.85" text-anchor="middle" font-size="11" fill="#1f2933">acc</text><line x1="187.0" y1="120" x2="187.0" y2="152" stroke="#9aa5b1" stroke-width="1.3" marker-end="url(#a)"/><line x1="114" y1="164.0" x2="162" y2="164.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/><rect x="260" y="152" width="50" height="24" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="285.0" y="167.85" text-anchor="middle" font-size="11" fill="#1f2933">acc</text><line x1="285.0" y1="120" x2="285.0" y2="152" stroke="#9aa5b1" stroke-width="1.3" marker-end="url(#a)"/><line x1="212" y1="164.0" x2="260" y2="164.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/><rect x="358" y="152" width="50" height="24" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="383.0" y="167.85" text-anchor="middle" font-size="11" fill="#1f2933">acc</text><line x1="383.0" y1="120" x2="383.0" y2="152" stroke="#9aa5b1" stroke-width="1.3" marker-end="url(#a)"/><line x1="310" y1="164.0" x2="358" y2="164.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/><rect x="456" y="152" width="50" height="24" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="481.0" y="167.85" text-anchor="middle" font-size="11" fill="#1f2933">acc</text><line x1="481.0" y1="120" x2="481.0" y2="152" stroke="#9aa5b1" stroke-width="1.3" marker-end="url(#a)"/><line x1="408" y1="164.0" x2="456" y2="164.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/><rect x="554" y="152" width="50" height="24" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="579.0" y="167.85" text-anchor="middle" font-size="11" fill="#1f2933">acc</text><line x1="579.0" y1="120" x2="579.0" y2="152" stroke="#9aa5b1" stroke-width="1.3" marker-end="url(#a)"/><line x1="506" y1="164.0" x2="554" y2="164.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/><rect x="652" y="152" width="50" height="24" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="677.0" y="167.85" text-anchor="middle" font-size="11" fill="#1f2933">acc</text><line x1="677.0" y1="120" x2="677.0" y2="152" stroke="#9aa5b1" stroke-width="1.3" marker-end="url(#a)"/><line x1="604" y1="164.0" x2="652" y2="164.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/><rect x="750" y="152" width="50" height="24" rx="6" fill="#eef2f7" stroke="#7b8794" stroke-width="1.5"/><text x="775.0" y="167.85" text-anchor="middle" font-size="11" fill="#1f2933">acc</text><line x1="775.0" y1="120" x2="775.0" y2="152" stroke="#9aa5b1" stroke-width="1.3" marker-end="url(#a)"/><line x1="702" y1="164.0" x2="750" y2="164.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/><rect x="830" y="152" width="56" height="24" rx="6" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><text x="858.0" y="167.85" text-anchor="middle" font-size="11" fill="#1f2933">ReduceMax</text><line x1="800" y1="164.0" x2="830" y2="164.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/><text x="24" y="212" font-size="14" fill="#1f2933" text-anchor="start" font-weight="700">② unroll 多路并行（×4）　4 路独立、最后合并　串行步数 ≈ N/4 + log2(4) = 4</text><rect x="60" y="234" width="52" height="24" rx="6" fill="#eef2f7" stroke="#2b6cb0" stroke-width="1.5"/><text x="86.0" y="249.85" text-anchor="middle" font-size="11" fill="#2b6cb0">acc0</text><rect x="150" y="234" width="54" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="177.0" y="249.85" text-anchor="middle" font-size="11" fill="#1f2933">max c0</text><rect x="250" y="234" width="54" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="277.0" y="249.85" text-anchor="middle" font-size="11" fill="#1f2933">max c4</text><line x1="112" y1="246" x2="150" y2="246" stroke="#d64545" stroke-width="2" marker-end="url(#ar)"/><line x1="204" y1="246" x2="250" y2="246" stroke="#d64545" stroke-width="2" marker-end="url(#ar)"/><rect x="60" y="268" width="52" height="24" rx="6" fill="#eef2f7" stroke="#2b6cb0" stroke-width="1.5"/><text x="86.0" y="283.85" text-anchor="middle" font-size="11" fill="#2b6cb0">acc1</text><rect x="150" y="268" width="54" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="177.0" y="283.85" text-anchor="middle" font-size="11" fill="#1f2933">max c1</text><rect x="250" y="268" width="54" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="277.0" y="283.85" text-anchor="middle" font-size="11" fill="#1f2933">max c5</text><line x1="112" y1="280" x2="150" y2="280" stroke="#d64545" stroke-width="2" marker-end="url(#ar)"/><line x1="204" y1="280" x2="250" y2="280" stroke="#d64545" stroke-width="2" marker-end="url(#ar)"/><rect x="60" y="302" width="52" height="24" rx="6" fill="#eef2f7" stroke="#2b6cb0" stroke-width="1.5"/><text x="86.0" y="317.85" text-anchor="middle" font-size="11" fill="#2b6cb0">acc2</text><rect x="150" y="302" width="54" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="177.0" y="317.85" text-anchor="middle" font-size="11" fill="#1f2933">max c2</text><rect x="250" y="302" width="54" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="277.0" y="317.85" text-anchor="middle" font-size="11" fill="#1f2933">max c6</text><line x1="112" y1="314" x2="150" y2="314" stroke="#d64545" stroke-width="2" marker-end="url(#ar)"/><line x1="204" y1="314" x2="250" y2="314" stroke="#d64545" stroke-width="2" marker-end="url(#ar)"/><rect x="60" y="336" width="52" height="24" rx="6" fill="#eef2f7" stroke="#2b6cb0" stroke-width="1.5"/><text x="86.0" y="351.85" text-anchor="middle" font-size="11" fill="#2b6cb0">acc3</text><rect x="150" y="336" width="54" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="177.0" y="351.85" text-anchor="middle" font-size="11" fill="#1f2933">max c3</text><rect x="250" y="336" width="54" height="24" rx="6" fill="#fbeec1" stroke="#caa83a" stroke-width="1.5"/><text x="277.0" y="351.85" text-anchor="middle" font-size="11" fill="#1f2933">max c7</text><line x1="112" y1="348" x2="150" y2="348" stroke="#d64545" stroke-width="2" marker-end="url(#ar)"/><line x1="204" y1="348" x2="250" y2="348" stroke="#d64545" stroke-width="2" marker-end="url(#ar)"/><line x1="304" y1="246" x2="336" y2="246" stroke="#2b6cb0" stroke-width="1.6"/><line x1="304" y1="280" x2="336" y2="280" stroke="#2b6cb0" stroke-width="1.6"/><line x1="304" y1="314" x2="336" y2="314" stroke="#2b6cb0" stroke-width="1.6"/><line x1="304" y1="348" x2="336" y2="348" stroke="#2b6cb0" stroke-width="1.6"/><line x1="336" y1="246" x2="336" y2="348" stroke="#2b6cb0" stroke-width="1.6"/><rect x="400" y="284.0" width="70" height="26" rx="6" fill="#eef2f7" stroke="#2b6cb0" stroke-width="1.5"/><text x="435.0" y="300.85" text-anchor="middle" font-size="11" fill="#2b6cb0">两两合并</text><line x1="336" y1="297.0" x2="400" y2="297.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/><rect x="495" y="284.0" width="56" height="26" rx="6" fill="#e9defa" stroke="#7e57c2" stroke-width="1.5"/><text x="523.0" y="300.85" text-anchor="middle" font-size="11" fill="#1f2933">ReduceMax</text><line x1="470" y1="297.0" x2="495" y2="297.0" stroke="#d64545" stroke-width="2.2" marker-end="url(#ar)"/></svg>
ASamples/2_Performance/simd_vf_story/reduce/images/latency_hiding.svg+1-0
@@ -0,0 +1 @@
1+<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 900 470" font-family="Noto Sans CJK SC,Microsoft YaHei,PingFang SC,Segoe UI,Helvetica,Arial,sans-serif"><defs><marker id="a" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#1f2933"/></marker><marker id="ar" markerWidth="9" markerHeight="9" refX="7.5" refY="3" orient="auto"><path d="M0,0 L7.5,3 L0,6 Z" fill="#d64545"/></marker></defs><rect width="900" height="470" fill="white"/><text x="24" y="32" font-size="18" fill="#1f2933" text-anchor="start" font-weight="700">为什么多路独立并行能提速：延迟隐藏（latency hiding）</text><text x="24" y="52" font-size="12" fill="#52606d" text-anchor="start">矢量指令「延迟 L 拍才出结果」，但矢量管线「每拍能发射 1 条」。下图每格 = 1 个发射拍（cycle），取 L=4。</text><text x="24" y="88" font-size="13" fill="#1f2933" text-anchor="start" font-weight="700">① baseline 单路串行：下一条 Max 必须等上一条出结果</text><text x="120" y="112" font-size="10" fill="#52606d" text-anchor="middle">0</text><text x="296" y="112" font-size="10" fill="#52606d" text-anchor="middle">4</text><text x="472" y="112" font-size="10" fill="#52606d" text-anchor="middle">8</text><text x="648" y="112" font-size="10" fill="#52606d" text-anchor="middle">12</text><text x="824" y="112" font-size="10" fill="#52606d" text-anchor="middle">16</text><text x="848" y="112" font-size="10" fill="#52606d" text-anchor="middle">cycle</text><rect x="120" y="120" width="44" height="30" rx="3" fill="#2b6cb0" stroke="#7b8794" stroke-width="1"/><text x="142.0" y="139.2" text-anchor="middle" font-size="12" fill="white">Max0</text><rect x="164" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="186.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="208" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="230.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="252" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="274.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="296" y="120" width="44" height="30" rx="3" fill="#2b6cb0" stroke="#7b8794" stroke-width="1"/><text x="318.0" y="139.2" text-anchor="middle" font-size="12" fill="white">Max1</text><rect x="340" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="362.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="384" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="406.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="428" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="450.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="472" y="120" width="44" height="30" rx="3" fill="#2b6cb0" stroke="#7b8794" stroke-width="1"/><text x="494.0" y="139.2" text-anchor="middle" font-size="12" fill="white">Max2</text><rect x="516" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="538.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="560" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="582.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="604" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="626.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="648" y="120" width="44" height="30" rx="3" fill="#2b6cb0" stroke="#7b8794" stroke-width="1"/><text x="670.0" y="139.2" text-anchor="middle" font-size="12" fill="white">Max3</text><rect x="692" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="714.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="736" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="758.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><rect x="780" y="120" width="44" height="30" rx="3" fill="#eceff3" stroke="#7b8794" stroke-width="1" stroke-dasharray="3 2"/><text x="802.0" y="138.15" text-anchor="middle" font-size="9" fill="#52606d">空泡</text><text x="892" y="139.0" font-size="12" fill="#d64545" text-anchor="end" font-weight="700">利用率 1/L</text><line x1="142.0" y1="156" x2="142.0" y2="170" stroke="#d64545" stroke-width="1.4"/><line x1="142.0" y1="170" x2="318.0" y2="170" stroke="#d64545" stroke-width="1.6" marker-end="url(#ar)"/><text x="328.0" y="174" font-size="11" fill="#d64545" text-anchor="start">等 Max0 结果（延迟 L=4 拍）→ 单路串行无事可发 → 空泡</text><text x="24" y="240" font-size="13" fill="#1f2933" text-anchor="start" font-weight="700">② unroll 多路并行（×4）：4 路独立的 Max 轮流填满空泡</text><text x="120" y="264" font-size="10" fill="#52606d" text-anchor="middle">0</text><text x="296" y="264" font-size="10" fill="#52606d" text-anchor="middle">4</text><text x="472" y="264" font-size="10" fill="#52606d" text-anchor="middle">8</text><text x="648" y="264" font-size="10" fill="#52606d" text-anchor="middle">12</text><text x="824" y="264" font-size="10" fill="#52606d" text-anchor="middle">16</text><text x="848" y="264" font-size="10" fill="#52606d" text-anchor="middle">cycle</text><rect x="120" y="272" width="44" height="30" rx="3" fill="#2b6cb0" stroke="#7b8794" stroke-width="1"/><text x="142.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M0·0</text><rect x="164" y="272" width="44" height="30" rx="3" fill="#2f855a" stroke="#7b8794" stroke-width="1"/><text x="186.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M1·0</text><rect x="208" y="272" width="44" height="30" rx="3" fill="#dd8b21" stroke="#7b8794" stroke-width="1"/><text x="230.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M2·0</text><rect x="252" y="272" width="44" height="30" rx="3" fill="#7e57c2" stroke="#7b8794" stroke-width="1"/><text x="274.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M3·0</text><rect x="296" y="272" width="44" height="30" rx="3" fill="#2b6cb0" stroke="#7b8794" stroke-width="1"/><text x="318.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M0·1</text><rect x="340" y="272" width="44" height="30" rx="3" fill="#2f855a" stroke="#7b8794" stroke-width="1"/><text x="362.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M1·1</text><rect x="384" y="272" width="44" height="30" rx="3" fill="#dd8b21" stroke="#7b8794" stroke-width="1"/><text x="406.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M2·1</text><rect x="428" y="272" width="44" height="30" rx="3" fill="#7e57c2" stroke="#7b8794" stroke-width="1"/><text x="450.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M3·1</text><rect x="472" y="272" width="44" height="30" rx="3" fill="#2b6cb0" stroke="#7b8794" stroke-width="1"/><text x="494.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M0·2</text><rect x="516" y="272" width="44" height="30" rx="3" fill="#2f855a" stroke="#7b8794" stroke-width="1"/><text x="538.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M1·2</text><rect x="560" y="272" width="44" height="30" rx="3" fill="#dd8b21" stroke="#7b8794" stroke-width="1"/><text x="582.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M2·2</text><rect x="604" y="272" width="44" height="30" rx="3" fill="#7e57c2" stroke="#7b8794" stroke-width="1"/><text x="626.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M3·2</text><rect x="648" y="272" width="44" height="30" rx="3" fill="#2b6cb0" stroke="#7b8794" stroke-width="1"/><text x="670.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M0·3</text><rect x="692" y="272" width="44" height="30" rx="3" fill="#2f855a" stroke="#7b8794" stroke-width="1"/><text x="714.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M1·3</text><rect x="736" y="272" width="44" height="30" rx="3" fill="#dd8b21" stroke="#7b8794" stroke-width="1"/><text x="758.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M2·3</text><rect x="780" y="272" width="44" height="30" rx="3" fill="#7e57c2" stroke="#7b8794" stroke-width="1"/><text x="802.0" y="290.5" text-anchor="middle" font-size="10" fill="white">M3·3</text><text x="892" y="291.0" font-size="12" fill="#2f855a" text-anchor="end" font-weight="700">利用率 ≈1</text><line x1="142.0" y1="308" x2="142.0" y2="322" stroke="#1f2933" stroke-width="1.2"/><line x1="142.0" y1="322" x2="318.0" y2="322" stroke="#1f2933" stroke-width="1.4" marker-end="url(#a)"/><text x="328.0" y="326" font-size="11" fill="#1f2933" text-anchor="start">M0·0 结果就绪，正好接 M0·1（中间 3 拍由 M1/M2/M3 填上）</text><rect x="24" y="372" width="852" height="84" rx="8" fill="#f7f9fc" stroke="#d7dde5"/><text x="38" y="394" font-size="13" fill="#1f2933" text-anchor="start" font-weight="700">三条优化路径，本质都在「缩短或填满」这段串行依赖：</text><text x="38" y="416" font-size="12" fill="#52606d" text-anchor="start">• unroll：指令数不变，用 4 路独立填满空泡（隐藏延迟）；还能配合硬件双发（每拍 2 条）再压。</text><text x="38" y="436" font-size="12" fill="#52606d" text-anchor="start">• binary（仅 reduce_sum）：把串行依赖直接变短（步数 4→2）+ 成对折叠令总 Add/ReduceSum 减半 → 既减步数又减量，故最快。</text><text x="38" y="454" font-size="12" fill="#d64545" text-anchor="start">• pragma：只复制循环体，仍串在同一 acc 上 → 串行步数不变、空泡照旧 → 对归约无效。</text></svg>
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_max_ar_baseline.asc+162-0
@@ -0,0 +1,162 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceMax，axis=1（沿尾维归约，输出 [D0]），基线 VF
3+ *
4+ * 从 reduce_max.asc 拆出 axis=1 基线一个 VF（ReduceMaxArVf）：每行按 VL 分块逐元素 Max 折进
5+ * 累加器（跨块必须 MERGING），再「向量内 ReduceMax」压成一个标量（落 lane0），StoreUnAlign 把
6+ * 逐行标量拼成连续输出流，循环末 StoreUnAlignPost 冲刷残留（二者成对）。
7+ * 同 axis 的多累加器手写展开 / #pragma 自动展开版见 reduce_max_ar_unroll.asc / reduce_max_ar_pragma.asc。
8+ *
9+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ar。
10+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
11+ *
12+ * 编译 & 运行：
13+ * cmake --build build --target reduce_max_ar_baseline
14+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_max_ar_baseline -s Ascend950
15+ */
16+ 
17+#include <iostream>
18+#include <vector>
19+#include <cmath>
20+#include <limits>
21+#include <algorithm>
22+#include "acl/acl.h"
23+#include "kernel_operator.h"
24+#include "../../include/vf_common.h"
25+ 
26+static constexpr float NEG_INF = -3.4028235e38f; // 约 -FLT_MAX，Max 归约初值
27+ 
28+// ============ VF 层（基线，对照用）：axis=1，沿尾维归约，逐行向量内 ReduceMax，输出 [d0] ============
29+// 每行按 VL 分块折叠（覆盖 d1 > VL）进累加器，再向量内归约成一个标量。
30+__simd_vf__ inline void ReduceMaxArVf(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
31+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
32+{
33+ AscendC::Reg::RegTensor<float> accReg, inReg, outReg;
34+ AscendC::Reg::UnalignRegForStore unalignAcc; // 非对齐散出累积器
35+ AscendC::Reg::MaskReg fullMask = AscendC::Reg::CreateMask<float, AscendC::Reg::MaskPattern::ALL>();
36+ const uint16_t rowChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
37+ 
38+ for (uint16_t i = 0; i < static_cast<uint16_t>(d0); i++) { // 外层：行循环（每行归约出一个标量）
39+ AscendC::Reg::Duplicate(accReg, NEG_INF);
40+ uint32_t remainCols = d1; // UpdateMask 每轮自动扣 VL，记录剩余列数
41+ for (uint16_t c = 0; c < rowChunks; c++) { // 内层：列块循环（按 VL 切分 d1 列，折叠进累加器）
42+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols); // 本块有效列数
43+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(
44+ inReg, xAddr + i * d1Pad + c * VL_B32);
45+ // 必须 MERGING：mask 外的 lane 保留累加器原值。默认 ZEROING 会把 mask 外清零，
46+ // 末块（mask 不满）会清掉前面整块已折叠进来的列，再 ReduceMax 就丢了一部分。
47+ AscendC::Reg::Max<float, AscendC::Reg::MaskMergeMode::MERGING>(accReg, accReg, inReg, mask); // 跨块逐元素折叠
48+ }
49+ AscendC::Reg::ReduceMax(outReg, accReg, fullMask); // 向量内归约 → 行最大值落 lane0
50+ AscendC::Reg::StoreUnAlign(zAddr, outReg, unalignAcc, 1); // 追加 1 个标量，zAddr 自动后移
51+ }
52+ AscendC::Reg::StoreUnAlignPost(zAddr, unalignAcc, 0); // 冲刷尾部不足一个 block 的残留
53+}
54+ 
55+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
56+class KernelReduceMax {
57+public:
58+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
59+ {
60+ d0_ = d0;
61+ d1_ = d1;
62+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN;
63+ outLen_ = d0_; // axis=1 输出 [d0]
64+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
65+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
66+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float);
67+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
68+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
69+ pipe_.InitBuffer(outQueue_, 1, outBytes);
70+ }
71+ 
72+ __aicore__ inline void Process()
73+ {
74+ CopyIn();
75+ Compute();
76+ CopyOut();
77+ }
78+ 
79+private:
80+ __aicore__ inline void CopyIn()
81+ {
82+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
83+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
84+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
85+ for (uint32_t i = 0; i < d0_; i++) {
86+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
87+ }
88+ inQueue_.EnQue(xUb);
89+ }
90+ 
91+ __aicore__ inline void Compute()
92+ {
93+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
94+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
95+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
96+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
97+ for (int r = 0; r < VF_REPEAT; r++) {
98+ asc_vf_call<ReduceMaxArVf>(xAddr, zAddr, d0_, d1_, d1Pad_);
99+ }
100+ outQueue_.EnQue(zUb);
101+ inQueue_.FreeTensor(xUb);
102+ }
103+ 
104+ __aicore__ inline void CopyOut()
105+ {
106+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
107+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
108+ AscendC::DataCopyPad(zGm_, zUb, params);
109+ outQueue_.FreeTensor(zUb);
110+ }
111+ 
112+ AscendC::TPipe pipe_;
113+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
114+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
115+ AscendC::GlobalTensor<float> xGm_, zGm_;
116+ uint32_t d0_, d1_, d1Pad_, outLen_;
117+};
118+ 
119+__global__ __aicore__ __vector__ void ReduceMaxKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
120+{
121+ KernelReduceMax op;
122+ op.Init(x, z, d0, d1);
123+ op.Process();
124+}
125+ 
126+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
127+int main()
128+{
129+ CHECK_ACL(aclInit(nullptr));
130+ CHECK_ACL(aclrtSetDevice(0));
131+ aclrtStream stream;
132+ CHECK_ACL(aclrtCreateStream(&stream));
133+ 
134+ const uint32_t outLen = D0;
135+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
136+ std::vector<float> ref(outLen, std::numeric_limits<float>::lowest());
137+ for (uint32_t i = 0; i < D0; i++)
138+ for (uint32_t j = 0; j < D1; j++) ref[i] = std::max(ref[i], hX[i * D1 + j]); // 标杆：每行取 Max
139+ std::vector<float> hZ(outLen, 0.f);
140+ 
141+ uint8_t *dX, *dZ;
142+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
143+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
144+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
145+ ACL_MEMCPY_HOST_TO_DEVICE));
146+ 
147+ ReduceMaxKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
148+ CHECK_ACL(aclrtSynchronizeStream(stream));
149+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
150+ ACL_MEMCPY_DEVICE_TO_HOST));
151+ 
152+ bool ok = vf::VerifyAbs(hZ, ref, 1e-4f);
153+ std::cout << "[reduce_max axis1 baseline] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
154+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
155+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
156+ 
157+ aclrtFree(dX); aclrtFree(dZ);
158+ CHECK_ACL(aclrtDestroyStream(stream));
159+ CHECK_ACL(aclrtResetDevice(0));
160+ CHECK_ACL(aclFinalize());
161+ return ok ? 0 : 1;
162+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_max_ar_pragma.asc+164-0
@@ -0,0 +1,164 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceMax，axis=1（沿尾维归约，输出 [D0]），#pragma 自动展开 VF
3+ *
4+ * 从 reduce_max.asc 拆出 axis=1 #pragma 自动展开一个 VF（ReduceMaxArVfPragma）：每行按 VL 分块逐元素 Max 折进
5+ * 累加器（跨块必须 MERGING），再「向量内 ReduceMax」压成一个标量（落 lane0），StoreUnAlign 把
6+ * 逐行标量拼成连续输出流，循环末 StoreUnAlignPost 冲刷残留（二者成对）。
7+ * 基线 / 多累加器手写展开版见 reduce_max_ar_baseline.asc / reduce_max_ar_unroll.asc。
8+ * 注意：#pragma 仅展开循环体、未打断累加串行依赖，对归约无实测加速（≈baseline，详见 README）。
9+ *
10+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ar。
11+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
12+ *
13+ * 编译 & 运行：
14+ * cmake --build build --target reduce_max_ar_pragma
15+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_max_ar_pragma -s Ascend950
16+ */
17+ 
18+#include <iostream>
19+#include <vector>
20+#include <cmath>
21+#include <limits>
22+#include <algorithm>
23+#include "acl/acl.h"
24+#include "kernel_operator.h"
25+#include "../../include/vf_common.h"
26+ 
27+static constexpr float NEG_INF = -3.4028235e38f; // 约 -FLT_MAX，Max 归约初值
28+ 
29+// ============ VF 层（基线，对照用）：axis=1，沿尾维归约，逐行向量内 ReduceMax，输出 [d0] ============
30+// 每行按 VL 分块折叠（覆盖 d1 > VL）进累加器，再向量内归约成一个标量。
31+__simd_vf__ inline void ReduceMaxArVfPragma(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
32+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
33+{
34+ AscendC::Reg::RegTensor<float> accReg, inReg, outReg;
35+ AscendC::Reg::UnalignRegForStore unalignAcc; // 非对齐散出累积器
36+ AscendC::Reg::MaskReg fullMask = AscendC::Reg::CreateMask<float, AscendC::Reg::MaskPattern::ALL>();
37+ const uint16_t rowChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
38+ 
39+ #pragma unroll 4 // 自动展开：编译器把循环体展成多条，但累加器单一、串行步数不变
40+ for (uint16_t i = 0; i < static_cast<uint16_t>(d0); i++) { // 外层：行循环（每行归约出一个标量）
41+ AscendC::Reg::Duplicate(accReg, NEG_INF);
42+ uint32_t remainCols = d1; // UpdateMask 每轮自动扣 VL，记录剩余列数
43+ for (uint16_t c = 0; c < rowChunks; c++) { // 内层：列块循环（按 VL 切分 d1 列，折叠进累加器）
44+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols); // 本块有效列数
45+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(
46+ inReg, xAddr + i * d1Pad + c * VL_B32);
47+ // 必须 MERGING：mask 外的 lane 保留累加器原值。默认 ZEROING 会把 mask 外清零，
48+ // 末块（mask 不满）会清掉前面整块已折叠进来的列，再 ReduceMax 就丢了一部分。
49+ AscendC::Reg::Max<float, AscendC::Reg::MaskMergeMode::MERGING>(accReg, accReg, inReg, mask); // 跨块逐元素折叠
50+ }
51+ AscendC::Reg::ReduceMax(outReg, accReg, fullMask); // 向量内归约 → 行最大值落 lane0
52+ AscendC::Reg::StoreUnAlign(zAddr, outReg, unalignAcc, 1); // 追加 1 个标量，zAddr 自动后移
53+ }
54+ AscendC::Reg::StoreUnAlignPost(zAddr, unalignAcc, 0); // 冲刷尾部不足一个 block 的残留
55+}
56+ 
57+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
58+class KernelReduceMax {
59+public:
60+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
61+ {
62+ d0_ = d0;
63+ d1_ = d1;
64+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN;
65+ outLen_ = d0_; // axis=1 输出 [d0]
66+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
67+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
68+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float);
69+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
70+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
71+ pipe_.InitBuffer(outQueue_, 1, outBytes);
72+ }
73+ 
74+ __aicore__ inline void Process()
75+ {
76+ CopyIn();
77+ Compute();
78+ CopyOut();
79+ }
80+ 
81+private:
82+ __aicore__ inline void CopyIn()
83+ {
84+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
85+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
86+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
87+ for (uint32_t i = 0; i < d0_; i++) {
88+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
89+ }
90+ inQueue_.EnQue(xUb);
91+ }
92+ 
93+ __aicore__ inline void Compute()
94+ {
95+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
96+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
97+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
98+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
99+ for (int r = 0; r < VF_REPEAT; r++) {
100+ asc_vf_call<ReduceMaxArVfPragma>(xAddr, zAddr, d0_, d1_, d1Pad_);
101+ }
102+ outQueue_.EnQue(zUb);
103+ inQueue_.FreeTensor(xUb);
104+ }
105+ 
106+ __aicore__ inline void CopyOut()
107+ {
108+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
109+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
110+ AscendC::DataCopyPad(zGm_, zUb, params);
111+ outQueue_.FreeTensor(zUb);
112+ }
113+ 
114+ AscendC::TPipe pipe_;
115+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
116+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
117+ AscendC::GlobalTensor<float> xGm_, zGm_;
118+ uint32_t d0_, d1_, d1Pad_, outLen_;
119+};
120+ 
121+__global__ __aicore__ __vector__ void ReduceMaxKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
122+{
123+ KernelReduceMax op;
124+ op.Init(x, z, d0, d1);
125+ op.Process();
126+}
127+ 
128+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
129+int main()
130+{
131+ CHECK_ACL(aclInit(nullptr));
132+ CHECK_ACL(aclrtSetDevice(0));
133+ aclrtStream stream;
134+ CHECK_ACL(aclrtCreateStream(&stream));
135+ 
136+ const uint32_t outLen = D0;
137+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
138+ std::vector<float> ref(outLen, std::numeric_limits<float>::lowest());
139+ for (uint32_t i = 0; i < D0; i++)
140+ for (uint32_t j = 0; j < D1; j++) ref[i] = std::max(ref[i], hX[i * D1 + j]); // 标杆：每行取 Max
141+ std::vector<float> hZ(outLen, 0.f);
142+ 
143+ uint8_t *dX, *dZ;
144+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
145+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
146+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
147+ ACL_MEMCPY_HOST_TO_DEVICE));
148+ 
149+ ReduceMaxKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
150+ CHECK_ACL(aclrtSynchronizeStream(stream));
151+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
152+ ACL_MEMCPY_DEVICE_TO_HOST));
153+ 
154+ bool ok = vf::VerifyAbs(hZ, ref, 1e-4f);
155+ std::cout << "[reduce_max axis1 pragma] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
156+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
157+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
158+ 
159+ aclrtFree(dX); aclrtFree(dZ);
160+ CHECK_ACL(aclrtDestroyStream(stream));
161+ CHECK_ACL(aclrtResetDevice(0));
162+ CHECK_ACL(aclFinalize());
163+ return ok ? 0 : 1;
164+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_max_ar_unroll.asc+180-0
@@ -0,0 +1,180 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceMax，axis=1（沿尾维归约，输出 [D0]），多累加器展开 VF
3+ *
4+ * 从 reduce_max.asc 拆出 axis=1 多累加器展开一个 VF（ReduceMaxArVfUnroll）：一轮用 4 个累加器
5+ * 并行折叠 4 行、再各自 ReduceMax，把单行 ReduceMax 的延迟相互隐藏。每行仍用单累加器按列块顺序
6+ * 折叠，归约顺序与基线一致 → 结果逐位等于基线。4 个 StoreUnAlign 按行序追加，尾部不足 4 的行单独处理。
7+ * 同 axis 的基线 / #pragma 自动展开版见 reduce_max_ar_baseline.asc / reduce_max_ar_pragma.asc。
8+ *
9+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ar。
10+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
11+ *
12+ * 编译 & 运行：
13+ * cmake --build build --target reduce_max_ar_unroll
14+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_max_ar_unroll -s Ascend950
15+ */
16+ 
17+#include <iostream>
18+#include <vector>
19+#include <cmath>
20+#include <limits>
21+#include <algorithm>
22+#include "acl/acl.h"
23+#include "kernel_operator.h"
24+#include "../../include/vf_common.h"
25+ 
26+static constexpr float NEG_INF = -3.4028235e38f; // 约 -FLT_MAX，Max 归约初值
27+ 
28+// ============ VF 层（多累加器展开）：axis=1，一轮 4 行并行折叠，再各自向量内 ReduceMax ============
29+__simd_vf__ inline void ReduceMaxArVfUnroll(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
30+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
31+{
32+ AscendC::Reg::RegTensor<float> acc0, acc1, acc2, acc3, in0, in1, in2, in3;
33+ AscendC::Reg::UnalignRegForStore unalignAcc;
34+ AscendC::Reg::MaskReg fullMask = AscendC::Reg::CreateMask<float, AscendC::Reg::MaskPattern::ALL>();
35+ const uint16_t rowChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
36+ const uint16_t groups = static_cast<uint16_t>(d0) / 4;
37+ 
38+ for (uint16_t g = 0; g < groups; g++) { // 主循环：每轮 4 行并行折叠
39+ const uint16_t r = static_cast<uint16_t>(g * 4);
40+ AscendC::Reg::Duplicate(acc0, NEG_INF); AscendC::Reg::Duplicate(acc1, NEG_INF);
41+ AscendC::Reg::Duplicate(acc2, NEG_INF); AscendC::Reg::Duplicate(acc3, NEG_INF);
42+ uint32_t remainCols = d1;
43+ for (uint16_t c = 0; c < rowChunks; c++) { // 4 行同一列结构 → 共享 mask
44+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols);
45+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in0, xAddr + (r + 0) * d1Pad + c * VL_B32);
46+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in1, xAddr + (r + 1) * d1Pad + c * VL_B32);
47+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in2, xAddr + (r + 2) * d1Pad + c * VL_B32);
48+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in3, xAddr + (r + 3) * d1Pad + c * VL_B32);
49+ AscendC::Reg::Max<float, AscendC::Reg::MaskMergeMode::MERGING>(acc0, acc0, in0, mask);
50+ AscendC::Reg::Max<float, AscendC::Reg::MaskMergeMode::MERGING>(acc1, acc1, in1, mask);
51+ AscendC::Reg::Max<float, AscendC::Reg::MaskMergeMode::MERGING>(acc2, acc2, in2, mask);
52+ AscendC::Reg::Max<float, AscendC::Reg::MaskMergeMode::MERGING>(acc3, acc3, in3, mask);
53+ }
54+ AscendC::Reg::ReduceMax(in0, acc0, fullMask); AscendC::Reg::StoreUnAlign(zAddr, in0, unalignAcc, 1); // 复用 in* 接收行标量
55+ AscendC::Reg::ReduceMax(in1, acc1, fullMask); AscendC::Reg::StoreUnAlign(zAddr, in1, unalignAcc, 1);
56+ AscendC::Reg::ReduceMax(in2, acc2, fullMask); AscendC::Reg::StoreUnAlign(zAddr, in2, unalignAcc, 1);
57+ AscendC::Reg::ReduceMax(in3, acc3, fullMask); AscendC::Reg::StoreUnAlign(zAddr, in3, unalignAcc, 1);
58+ }
59+ for (uint16_t r = static_cast<uint16_t>(groups * 4); r < static_cast<uint16_t>(d0); r++) { // 尾部 0~3 行
60+ AscendC::Reg::Duplicate(acc0, NEG_INF);
61+ uint32_t remainCols = d1;
62+ for (uint16_t c = 0; c < rowChunks; c++) {
63+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols);
64+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in0, xAddr + r * d1Pad + c * VL_B32);
65+ AscendC::Reg::Max<float, AscendC::Reg::MaskMergeMode::MERGING>(acc0, acc0, in0, mask);
66+ }
67+ AscendC::Reg::ReduceMax(in0, acc0, fullMask);
68+ AscendC::Reg::StoreUnAlign(zAddr, in0, unalignAcc, 1);
69+ }
70+ AscendC::Reg::StoreUnAlignPost(zAddr, unalignAcc, 0);
71+}
72+ 
73+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
74+class KernelReduceMax {
75+public:
76+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
77+ {
78+ d0_ = d0;
79+ d1_ = d1;
80+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN;
81+ outLen_ = d0_; // axis=1 输出 [d0]
82+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
83+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
84+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float);
85+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
86+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
87+ pipe_.InitBuffer(outQueue_, 1, outBytes);
88+ }
89+ 
90+ __aicore__ inline void Process()
91+ {
92+ CopyIn();
93+ Compute();
94+ CopyOut();
95+ }
96+ 
97+private:
98+ __aicore__ inline void CopyIn()
99+ {
100+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
101+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
102+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
103+ for (uint32_t i = 0; i < d0_; i++) {
104+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
105+ }
106+ inQueue_.EnQue(xUb);
107+ }
108+ 
109+ __aicore__ inline void Compute()
110+ {
111+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
112+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
113+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
114+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
115+ for (int r = 0; r < VF_REPEAT; r++) {
116+ asc_vf_call<ReduceMaxArVfUnroll>(xAddr, zAddr, d0_, d1_, d1Pad_);
117+ }
118+ outQueue_.EnQue(zUb);
119+ inQueue_.FreeTensor(xUb);
120+ }
121+ 
122+ __aicore__ inline void CopyOut()
123+ {
124+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
125+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
126+ AscendC::DataCopyPad(zGm_, zUb, params);
127+ outQueue_.FreeTensor(zUb);
128+ }
129+ 
130+ AscendC::TPipe pipe_;
131+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
132+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
133+ AscendC::GlobalTensor<float> xGm_, zGm_;
134+ uint32_t d0_, d1_, d1Pad_, outLen_;
135+};
136+ 
137+__global__ __aicore__ __vector__ void ReduceMaxKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
138+{
139+ KernelReduceMax op;
140+ op.Init(x, z, d0, d1);
141+ op.Process();
142+}
143+ 
144+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
145+int main()
146+{
147+ CHECK_ACL(aclInit(nullptr));
148+ CHECK_ACL(aclrtSetDevice(0));
149+ aclrtStream stream;
150+ CHECK_ACL(aclrtCreateStream(&stream));
151+ 
152+ const uint32_t outLen = D0;
153+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
154+ std::vector<float> ref(outLen, std::numeric_limits<float>::lowest());
155+ for (uint32_t i = 0; i < D0; i++)
156+ for (uint32_t j = 0; j < D1; j++) ref[i] = std::max(ref[i], hX[i * D1 + j]); // 标杆：每行取 Max
157+ std::vector<float> hZ(outLen, 0.f);
158+ 
159+ uint8_t *dX, *dZ;
160+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
161+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
162+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
163+ ACL_MEMCPY_HOST_TO_DEVICE));
164+ 
165+ ReduceMaxKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
166+ CHECK_ACL(aclrtSynchronizeStream(stream));
167+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
168+ ACL_MEMCPY_DEVICE_TO_HOST));
169+ 
170+ bool ok = vf::VerifyAbs(hZ, ref, 1e-4f);
171+ std::cout << "[reduce_max axis1 unroll] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
172+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
173+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
174+ 
175+ aclrtFree(dX); aclrtFree(dZ);
176+ CHECK_ACL(aclrtDestroyStream(stream));
177+ CHECK_ACL(aclrtResetDevice(0));
178+ CHECK_ACL(aclFinalize());
179+ return ok ? 0 : 1;
180+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_max_ra_baseline.asc+155-0
@@ -0,0 +1,155 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceMax，axis=0（沿首维归约，输出 [D1]），基线 VF
3+ *
4+ * 从 reduce_max.asc 拆出 axis=0 基线一个 VF（ReduceMaxRaVf）：按列分块，每个列块跨 D0 行做
5+ * 「逐元素 Max」累加（累加器初值 -inf），最后 StoreAlign 整段搬出（输出天然连续对齐）。
6+ * 同 axis 的多累加器手写展开 / #pragma 自动展开版见 reduce_max_ra_unroll.asc / reduce_max_ra_pragma.asc。
7+ *
8+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ra。
9+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
10+ *
11+ * 编译 & 运行：
12+ * cmake --build build --target reduce_max_ra_baseline
13+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_max_ra_baseline -s Ascend950
14+ */
15+ 
16+#include <iostream>
17+#include <vector>
18+#include <cmath>
19+#include <limits>
20+#include <algorithm>
21+#include "acl/acl.h"
22+#include "kernel_operator.h"
23+#include "../../include/vf_common.h"
24+ 
25+static constexpr float NEG_INF = -3.4028235e38f; // 约 -FLT_MAX，Max 归约初值
26+ 
27+// ============ VF 层（基线）：axis=0，沿首维归约，跨行逐元素 Max，输出 [d1] ============
28+__simd_vf__ inline void ReduceMaxRaVf(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
29+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
30+{
31+ AscendC::Reg::RegTensor<float> accReg, inReg;
32+ uint32_t remainCols = d1;
33+ const uint16_t colChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
34+ 
35+ for (uint16_t c = 0; c < colChunks; c++) { // 外层：列块循环
36+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols);
37+ AscendC::Reg::Duplicate(accReg, NEG_INF); // 累加器全 lane 置 -inf
38+ for (uint16_t i = 0; i < static_cast<uint16_t>(d0); i++) { // 内层：行循环
39+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(
40+ inReg, xAddr + i * d1Pad + c * VL_B32);
41+ AscendC::Reg::Max(accReg, accReg, inReg, mask);
42+ }
43+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(
44+ zAddr + c * VL_B32, accReg, mask); // 整段搬出（连续对齐）
45+ }
46+}
47+ 
48+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
49+class KernelReduceMax {
50+public:
51+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
52+ {
53+ d0_ = d0;
54+ d1_ = d1;
55+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN; // 行步长按 32B 对齐
56+ outLen_ = d1_; // axis=0 输出 [d1]
57+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
58+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
59+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float); // 额外留 VL 余量：尾行最后一块整轮读取不越界
60+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
61+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
62+ pipe_.InitBuffer(outQueue_, 1, outBytes);
63+ }
64+ 
65+ __aicore__ inline void Process()
66+ {
67+ CopyIn();
68+ Compute();
69+ CopyOut();
70+ }
71+ 
72+private:
73+ __aicore__ inline void CopyIn()
74+ {
75+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
76+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
77+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
78+ for (uint32_t i = 0; i < d0_; i++) { // 逐行搬到 32B 对齐槽位
79+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
80+ }
81+ inQueue_.EnQue(xUb);
82+ }
83+ 
84+ __aicore__ inline void Compute()
85+ {
86+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
87+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
88+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
89+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
90+ for (int r = 0; r < VF_REPEAT; r++) {
91+ asc_vf_call<ReduceMaxRaVf>(xAddr, zAddr, d0_, d1_, d1Pad_);
92+ }
93+ outQueue_.EnQue(zUb);
94+ inQueue_.FreeTensor(xUb);
95+ }
96+ 
97+ __aicore__ inline void CopyOut()
98+ {
99+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
100+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
101+ AscendC::DataCopyPad(zGm_, zUb, params);
102+ outQueue_.FreeTensor(zUb);
103+ }
104+ 
105+ AscendC::TPipe pipe_;
106+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
107+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
108+ AscendC::GlobalTensor<float> xGm_, zGm_;
109+ uint32_t d0_, d1_, d1Pad_, outLen_;
110+};
111+ 
112+__global__ __aicore__ __vector__ void ReduceMaxKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
113+{
114+ KernelReduceMax op;
115+ op.Init(x, z, d0, d1);
116+ op.Process();
117+}
118+ 
119+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
120+int main()
121+{
122+ CHECK_ACL(aclInit(nullptr));
123+ CHECK_ACL(aclrtSetDevice(0));
124+ aclrtStream stream;
125+ CHECK_ACL(aclrtCreateStream(&stream));
126+ 
127+ const uint32_t outLen = D1;
128+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
129+ std::vector<float> ref(outLen, std::numeric_limits<float>::lowest());
130+ for (uint32_t i = 0; i < D0; i++)
131+ for (uint32_t j = 0; j < D1; j++) ref[j] = std::max(ref[j], hX[i * D1 + j]); // 标杆：每列取 Max
132+ std::vector<float> hZ(outLen, 0.f);
133+ 
134+ uint8_t *dX, *dZ;
135+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
136+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
137+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
138+ ACL_MEMCPY_HOST_TO_DEVICE));
139+ 
140+ ReduceMaxKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
141+ CHECK_ACL(aclrtSynchronizeStream(stream));
142+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
143+ ACL_MEMCPY_DEVICE_TO_HOST));
144+ 
145+ bool ok = vf::VerifyAbs(hZ, ref, 1e-4f);
146+ std::cout << "[reduce_max axis0 baseline] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
147+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
148+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
149+ 
150+ aclrtFree(dX); aclrtFree(dZ);
151+ CHECK_ACL(aclrtDestroyStream(stream));
152+ CHECK_ACL(aclrtResetDevice(0));
153+ CHECK_ACL(aclFinalize());
154+ return ok ? 0 : 1;
155+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_max_ra_pragma.asc+157-0
@@ -0,0 +1,157 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceMax，axis=0（沿首维归约，输出 [D1]），#pragma 自动展开 VF
3+ *
4+ * 从 reduce_max.asc 拆出 axis=0 #pragma 自动展开一个 VF（ReduceMaxRaVfPragma）：按列分块，每个列块跨 D0 行做
5+ * 「逐元素 Max」累加（累加器初值 -inf），最后 StoreAlign 整段搬出（输出天然连续对齐）。
6+ * 基线 / 多累加器手写展开版见 reduce_max_ra_baseline.asc / reduce_max_ra_unroll.asc。
7+ * 注意：#pragma 仅展开循环体、未打断累加串行依赖，对归约无实测加速（≈baseline，详见 README）。
8+ *
9+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ra。
10+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
11+ *
12+ * 编译 & 运行：
13+ * cmake --build build --target reduce_max_ra_pragma
14+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_max_ra_pragma -s Ascend950
15+ */
16+ 
17+#include <iostream>
18+#include <vector>
19+#include <cmath>
20+#include <limits>
21+#include <algorithm>
22+#include "acl/acl.h"
23+#include "kernel_operator.h"
24+#include "../../include/vf_common.h"
25+ 
26+static constexpr float NEG_INF = -3.4028235e38f; // 约 -FLT_MAX，Max 归约初值
27+ 
28+// ============ VF 层（基线）：axis=0，沿首维归约，跨行逐元素 Max，输出 [d1] ============
29+__simd_vf__ inline void ReduceMaxRaVfPragma(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
30+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
31+{
32+ AscendC::Reg::RegTensor<float> accReg, inReg;
33+ uint32_t remainCols = d1;
34+ const uint16_t colChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
35+ 
36+ for (uint16_t c = 0; c < colChunks; c++) { // 外层：列块循环
37+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols);
38+ AscendC::Reg::Duplicate(accReg, NEG_INF); // 累加器全 lane 置 -inf
39+ #pragma unroll 4 // 自动展开：编译器把循环体展成多条，但累加器单一、串行步数不变
40+ for (uint16_t i = 0; i < static_cast<uint16_t>(d0); i++) { // 内层：行循环
41+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(
42+ inReg, xAddr + i * d1Pad + c * VL_B32);
43+ AscendC::Reg::Max(accReg, accReg, inReg, mask);
44+ }
45+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(
46+ zAddr + c * VL_B32, accReg, mask); // 整段搬出（连续对齐）
47+ }
48+}
49+ 
50+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
51+class KernelReduceMax {
52+public:
53+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
54+ {
55+ d0_ = d0;
56+ d1_ = d1;
57+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN; // 行步长按 32B 对齐
58+ outLen_ = d1_; // axis=0 输出 [d1]
59+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
60+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
61+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float); // 额外留 VL 余量：尾行最后一块整轮读取不越界
62+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
63+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
64+ pipe_.InitBuffer(outQueue_, 1, outBytes);
65+ }
66+ 
67+ __aicore__ inline void Process()
68+ {
69+ CopyIn();
70+ Compute();
71+ CopyOut();
72+ }
73+ 
74+private:
75+ __aicore__ inline void CopyIn()
76+ {
77+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
78+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
79+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
80+ for (uint32_t i = 0; i < d0_; i++) { // 逐行搬到 32B 对齐槽位
81+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
82+ }
83+ inQueue_.EnQue(xUb);
84+ }
85+ 
86+ __aicore__ inline void Compute()
87+ {
88+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
89+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
90+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
91+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
92+ for (int r = 0; r < VF_REPEAT; r++) {
93+ asc_vf_call<ReduceMaxRaVfPragma>(xAddr, zAddr, d0_, d1_, d1Pad_);
94+ }
95+ outQueue_.EnQue(zUb);
96+ inQueue_.FreeTensor(xUb);
97+ }
98+ 
99+ __aicore__ inline void CopyOut()
100+ {
101+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
102+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
103+ AscendC::DataCopyPad(zGm_, zUb, params);
104+ outQueue_.FreeTensor(zUb);
105+ }
106+ 
107+ AscendC::TPipe pipe_;
108+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
109+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
110+ AscendC::GlobalTensor<float> xGm_, zGm_;
111+ uint32_t d0_, d1_, d1Pad_, outLen_;
112+};
113+ 
114+__global__ __aicore__ __vector__ void ReduceMaxKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
115+{
116+ KernelReduceMax op;
117+ op.Init(x, z, d0, d1);
118+ op.Process();
119+}
120+ 
121+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
122+int main()
123+{
124+ CHECK_ACL(aclInit(nullptr));
125+ CHECK_ACL(aclrtSetDevice(0));
126+ aclrtStream stream;
127+ CHECK_ACL(aclrtCreateStream(&stream));
128+ 
129+ const uint32_t outLen = D1;
130+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
131+ std::vector<float> ref(outLen, std::numeric_limits<float>::lowest());
132+ for (uint32_t i = 0; i < D0; i++)
133+ for (uint32_t j = 0; j < D1; j++) ref[j] = std::max(ref[j], hX[i * D1 + j]); // 标杆：每列取 Max
134+ std::vector<float> hZ(outLen, 0.f);
135+ 
136+ uint8_t *dX, *dZ;
137+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
138+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
139+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
140+ ACL_MEMCPY_HOST_TO_DEVICE));
141+ 
142+ ReduceMaxKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
143+ CHECK_ACL(aclrtSynchronizeStream(stream));
144+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
145+ ACL_MEMCPY_DEVICE_TO_HOST));
146+ 
147+ bool ok = vf::VerifyAbs(hZ, ref, 1e-4f);
148+ std::cout << "[reduce_max axis0 pragma] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
149+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
150+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
151+ 
152+ aclrtFree(dX); aclrtFree(dZ);
153+ CHECK_ACL(aclrtDestroyStream(stream));
154+ CHECK_ACL(aclrtResetDevice(0));
155+ CHECK_ACL(aclFinalize());
156+ return ok ? 0 : 1;
157+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_max_ra_unroll.asc+169-0
@@ -0,0 +1,169 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceMax，axis=0（沿首维归约，输出 [D1]），多累加器展开 VF
3+ *
4+ * 从 reduce_max.asc 拆出 axis=0 多累加器展开一个 VF（ReduceMaxRaVfUnroll）：用 4 个独立累加器
5+ * acc0..acc3 按行号分组累加（行 r 进 acc[r%4]），最后两两合并 → 4 条归约流并行，利于双发/延迟
6+ * 隐藏。尾部不足 4 的行用第二个 for 处理。max 与顺序无关，结果逐位等于基线。
7+ *
8+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ra。
9+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
10+ *
11+ * 编译 & 运行：
12+ * cmake --build build --target reduce_max_ra_unroll
13+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_max_ra_unroll -s Ascend950
14+ */
15+ 
16+#include <iostream>
17+#include <vector>
18+#include <cmath>
19+#include <limits>
20+#include <algorithm>
21+#include "acl/acl.h"
22+#include "kernel_operator.h"
23+#include "../../include/vf_common.h"
24+ 
25+static constexpr float NEG_INF = -3.4028235e38f;
26+ 
27+// ============ VF 层（多累加器展开）：axis=0，4 个独立累加器并行跨行 Max ============
28+__simd_vf__ inline void ReduceMaxRaVfUnroll(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
29+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
30+{
31+ AscendC::Reg::RegTensor<float> acc0, acc1, acc2, acc3, in0, in1, in2, in3;
32+ uint32_t remainCols = d1;
33+ const uint16_t colChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
34+ 
35+ for (uint16_t c = 0; c < colChunks; c++) {
36+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols);
37+ AscendC::Reg::Duplicate(acc0, NEG_INF); AscendC::Reg::Duplicate(acc1, NEG_INF);
38+ AscendC::Reg::Duplicate(acc2, NEG_INF); AscendC::Reg::Duplicate(acc3, NEG_INF);
39+ const uint16_t groups = static_cast<uint16_t>(d0) / 4;
40+ for (uint16_t g = 0; g < groups; g++) { // 主循环：每轮 4 行 → 4 个独立累加器
41+ const uint16_t r = static_cast<uint16_t>(g * 4);
42+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in0, xAddr + (r + 0) * d1Pad + c * VL_B32);
43+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in1, xAddr + (r + 1) * d1Pad + c * VL_B32);
44+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in2, xAddr + (r + 2) * d1Pad + c * VL_B32);
45+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in3, xAddr + (r + 3) * d1Pad + c * VL_B32);
46+ AscendC::Reg::Max(acc0, acc0, in0, mask);
47+ AscendC::Reg::Max(acc1, acc1, in1, mask);
48+ AscendC::Reg::Max(acc2, acc2, in2, mask);
49+ AscendC::Reg::Max(acc3, acc3, in3, mask);
50+ }
51+ for (uint16_t r = static_cast<uint16_t>(groups * 4); r < static_cast<uint16_t>(d0); r++) { // 尾部 0~3 行
52+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in0, xAddr + r * d1Pad + c * VL_B32);
53+ AscendC::Reg::Max(acc0, acc0, in0, mask);
54+ }
55+ AscendC::Reg::Max(acc0, acc0, acc1, mask); // 合并 4 个累加器
56+ AscendC::Reg::Max(acc2, acc2, acc3, mask);
57+ AscendC::Reg::Max(acc0, acc0, acc2, mask);
58+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(zAddr + c * VL_B32, acc0, mask);
59+ }
60+}
61+ 
62+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
63+class KernelReduceMax {
64+public:
65+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
66+ {
67+ d0_ = d0;
68+ d1_ = d1;
69+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN;
70+ outLen_ = d1_; // axis=0 输出 [d1]
71+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
72+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
73+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float);
74+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
75+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
76+ pipe_.InitBuffer(outQueue_, 1, outBytes);
77+ }
78+ 
79+ __aicore__ inline void Process()
80+ {
81+ CopyIn();
82+ Compute();
83+ CopyOut();
84+ }
85+ 
86+private:
87+ __aicore__ inline void CopyIn()
88+ {
89+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
90+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
91+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
92+ for (uint32_t i = 0; i < d0_; i++) {
93+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
94+ }
95+ inQueue_.EnQue(xUb);
96+ }
97+ 
98+ __aicore__ inline void Compute()
99+ {
100+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
101+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
102+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
103+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
104+ for (int r = 0; r < VF_REPEAT; r++) {
105+ asc_vf_call<ReduceMaxRaVfUnroll>(xAddr, zAddr, d0_, d1_, d1Pad_);
106+ }
107+ outQueue_.EnQue(zUb);
108+ inQueue_.FreeTensor(xUb);
109+ }
110+ 
111+ __aicore__ inline void CopyOut()
112+ {
113+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
114+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
115+ AscendC::DataCopyPad(zGm_, zUb, params);
116+ outQueue_.FreeTensor(zUb);
117+ }
118+ 
119+ AscendC::TPipe pipe_;
120+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
121+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
122+ AscendC::GlobalTensor<float> xGm_, zGm_;
123+ uint32_t d0_, d1_, d1Pad_, outLen_;
124+};
125+ 
126+__global__ __aicore__ __vector__ void ReduceMaxKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
127+{
128+ KernelReduceMax op;
129+ op.Init(x, z, d0, d1);
130+ op.Process();
131+}
132+ 
133+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
134+int main()
135+{
136+ CHECK_ACL(aclInit(nullptr));
137+ CHECK_ACL(aclrtSetDevice(0));
138+ aclrtStream stream;
139+ CHECK_ACL(aclrtCreateStream(&stream));
140+ 
141+ const uint32_t outLen = D1;
142+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
143+ std::vector<float> ref(outLen, std::numeric_limits<float>::lowest());
144+ for (uint32_t i = 0; i < D0; i++)
145+ for (uint32_t j = 0; j < D1; j++) ref[j] = std::max(ref[j], hX[i * D1 + j]); // 标杆：每列取 Max
146+ std::vector<float> hZ(outLen, 0.f);
147+ 
148+ uint8_t *dX, *dZ;
149+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
150+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
151+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
152+ ACL_MEMCPY_HOST_TO_DEVICE));
153+ 
154+ ReduceMaxKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
155+ CHECK_ACL(aclrtSynchronizeStream(stream));
156+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
157+ ACL_MEMCPY_DEVICE_TO_HOST));
158+ 
159+ bool ok = vf::VerifyAbs(hZ, ref, 1e-4f);
160+ std::cout << "[reduce_max axis0 unroll] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
161+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
162+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
163+ 
164+ aclrtFree(dX); aclrtFree(dZ);
165+ CHECK_ACL(aclrtDestroyStream(stream));
166+ CHECK_ACL(aclrtResetDevice(0));
167+ CHECK_ACL(aclFinalize());
168+ return ok ? 0 : 1;
169+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_sum_ar_baseline.asc+163-0
@@ -0,0 +1,163 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceSum，axis=1（沿尾维归约，输出 [D0]），基线 VF
3+ *
4+ * 从 reduce_sum.asc 拆出 axis=1 基线一个 VF（ReduceSumArVf）：每行按 VL 分块逐元素 Add 折进
5+ * 累加器（跨块必须 MERGING），再「向量内 ReduceSum」压成一个标量（落 lane0），StoreUnAlign 把
6+ * 逐行标量拼成连续输出流，循环末 StoreUnAlignPost 冲刷残留（二者成对）。
7+ * 同 axis 的多累加器手写展开 / #pragma 自动展开 / 二分累加版见
8+ * reduce_sum_ar_unroll.asc / reduce_sum_ar_pragma.asc / reduce_sum_ar_binary.asc。
9+ *
10+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ar。
11+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
12+ *
13+ * 编译 & 运行：
14+ * cmake --build build --target reduce_sum_ar_baseline
15+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_sum_ar_baseline -s Ascend950
16+ */
17+ 
18+#include <iostream>
19+#include <vector>
20+#include <cmath>
21+#include "acl/acl.h"
22+#include "kernel_operator.h"
23+#include "../../include/vf_common.h"
24+ 
25+// ============ VF 层（基线，对照用）：axis=1，沿尾维归约，逐行向量内 ReduceSum，输出 [d0] ============
26+// 每行按 VL 分块折叠（覆盖 d1 > VL）进累加器，再向量内归约成一个标量。
27+__simd_vf__ inline void ReduceSumArVf(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
28+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
29+{
30+ AscendC::Reg::RegTensor<float> accReg, inReg, outReg;
31+ AscendC::Reg::UnalignRegForStore unalignAcc; // 非对齐散出累积器
32+ AscendC::Reg::MaskReg fullMask = AscendC::Reg::CreateMask<float, AscendC::Reg::MaskPattern::ALL>();
33+ const uint16_t rowChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
34+ 
35+ for (uint16_t i = 0; i < static_cast<uint16_t>(d0); i++) { // 外层：行循环（每行归约出一个标量）
36+ AscendC::Reg::Duplicate(accReg, 0.0f);
37+ uint32_t remainCols = d1; // UpdateMask 每轮自动扣 VL，记录剩余列数
38+ for (uint16_t c = 0; c < rowChunks; c++) { // 内层：列块循环（按 VL 切分 d1 列，折叠进累加器）
39+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols); // 本块有效列数
40+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(
41+ inReg, xAddr + i * d1Pad + c * VL_B32);
42+ // 必须 MERGING：mask 外的 lane 保留累加器原值。默认 ZEROING 会把 mask 外清零，
43+ // 末块（mask 不满）会清掉前面整块已累加进来的列，再 ReduceSum 就丢了一部分。
44+ AscendC::Reg::Add<float, AscendC::Reg::MaskMergeMode::MERGING>(accReg, accReg, inReg, mask); // 跨块逐元素折叠
45+ }
46+ AscendC::Reg::ReduceSum(outReg, accReg, fullMask); // 向量内归约 → 行和落 lane0（空洞为 0 不影响）
47+ AscendC::Reg::StoreUnAlign(zAddr, outReg, unalignAcc, 1); // 追加 1 个标量，zAddr 自动后移
48+ }
49+ AscendC::Reg::StoreUnAlignPost(zAddr, unalignAcc, 0); // 冲刷尾部不足一个 block 的残留
50+}
51+ 
52+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
53+class KernelReduceSum {
54+public:
55+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
56+ {
57+ d0_ = d0;
58+ d1_ = d1;
59+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN;
60+ outLen_ = d0_; // axis=1 输出 [d0]
61+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
62+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
63+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float);
64+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
65+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
66+ pipe_.InitBuffer(outQueue_, 1, outBytes);
67+ }
68+ 
69+ __aicore__ inline void Process()
70+ {
71+ CopyIn();
72+ Compute();
73+ CopyOut();
74+ }
75+ 
76+private:
77+ __aicore__ inline void CopyIn()
78+ {
79+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
80+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
81+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
82+ for (uint32_t i = 0; i < d0_; i++) {
83+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
84+ }
85+ inQueue_.EnQue(xUb);
86+ }
87+ 
88+ __aicore__ inline void Compute()
89+ {
90+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
91+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
92+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
93+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
94+ for (int r = 0; r < VF_REPEAT; r++) {
95+ asc_vf_call<ReduceSumArVf>(xAddr, zAddr, d0_, d1_, d1Pad_);
96+ }
97+ outQueue_.EnQue(zUb);
98+ inQueue_.FreeTensor(xUb);
99+ }
100+ 
101+ __aicore__ inline void CopyOut()
102+ {
103+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
104+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
105+ AscendC::DataCopyPad(zGm_, zUb, params);
106+ outQueue_.FreeTensor(zUb);
107+ }
108+ 
109+ AscendC::TPipe pipe_;
110+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
111+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
112+ AscendC::GlobalTensor<float> xGm_, zGm_;
113+ uint32_t d0_, d1_, d1Pad_, outLen_;
114+};
115+ 
116+__global__ __aicore__ __vector__ void ReduceSumKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
117+{
118+ KernelReduceSum op;
119+ op.Init(x, z, d0, d1);
120+ op.Process();
121+}
122+ 
123+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
124+int main()
125+{
126+ CHECK_ACL(aclInit(nullptr));
127+ CHECK_ACL(aclrtSetDevice(0));
128+ aclrtStream stream;
129+ CHECK_ACL(aclrtCreateStream(&stream));
130+ 
131+ const uint32_t outLen = D0;
132+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
133+ std::vector<float> ref(outLen, 0.f);
134+ {
135+ std::vector<double> acc(outLen, 0.0);
136+ for (uint32_t i = 0; i < D0; i++)
137+ for (uint32_t j = 0; j < D1; j++) acc[i] += (double)hX[i * D1 + j]; // 标杆：每行 double 求和
138+ for (uint32_t k = 0; k < outLen; k++) ref[k] = (float)acc[k];
139+ }
140+ std::vector<float> hZ(outLen, 0.f);
141+ 
142+ uint8_t *dX, *dZ;
143+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
144+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
145+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
146+ ACL_MEMCPY_HOST_TO_DEVICE));
147+ 
148+ ReduceSumKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
149+ CHECK_ACL(aclrtSynchronizeStream(stream));
150+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
151+ ACL_MEMCPY_DEVICE_TO_HOST));
152+ 
153+ bool ok = vf::VerifyRel(hZ, ref, 1e-3);
154+ std::cout << "[reduce_sum axis1 baseline] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
155+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
156+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
157+ 
158+ aclrtFree(dX); aclrtFree(dZ);
159+ CHECK_ACL(aclrtDestroyStream(stream));
160+ CHECK_ACL(aclrtResetDevice(0));
161+ CHECK_ACL(aclFinalize());
162+ return ok ? 0 : 1;
163+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_sum_ar_binary.asc+173-0
@@ -0,0 +1,173 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceSum，axis=1（沿尾维归约，输出 [D0]），二分累加 VF
3+ *
4+ * 从 reduce_sum.asc 拆出 axis=1 二分累加一个 VF（ReduceSumArVfBinary），写法参考
5+ * Samples/2_Performance/rms_norm_quant_story/src/6_binary_sum.asc 的 binary-add：把一行的求和
6+ * 元素「成对折叠」到 2 的幂折叠点再 ReduceSum，配对相距 foldPoint 的元素，既把累加串行依赖从
7+ * 线性 O(块数) 缩到 O(log 块数)，又用成对求和（pairwise summation）降低浮点舍入误差。
8+ * 二分累加只对 sum 有意义（max 选元素、精确、无精度收益，故 reduce_max 不提供此写法）。
9+ * 同 axis 的基线 / 多累加器手写展开 / #pragma 自动展开版见
10+ * reduce_sum_ar_baseline.asc / reduce_sum_ar_unroll.asc / reduce_sum_ar_pragma.asc。
11+ *
12+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ar。
13+ * 统一 case 规格（[D0,D1]=[78,250]）：每行 250 列 = 4 个 VL(64) 块，foldPoint=128（2 块、2 的幂）。
14+ * 第 1 层把上半 [128,250) 折到下半 [0,122)（块对 c0+c2、c1+c3，末块 c3 尾部空洞用 mask 屏蔽）；
15+ * 第 2 层 c0+c1 → 每 lane 含 4 块之和；最后一次向量内 ReduceSum 压成行标量。
16+ * 本写法按「每行恰 4 个 VL 块」展开，host 侧 static_assert 守护（改规格需同步调整）。
17+ *
18+ * 编译 & 运行：
19+ * cmake --build build --target reduce_sum_ar_binary
20+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_sum_ar_binary -s Ascend950
21+ */
22+ 
23+#include <iostream>
24+#include <vector>
25+#include <cmath>
26+#include "acl/acl.h"
27+#include "kernel_operator.h"
28+#include "../../include/vf_common.h"
29+ 
30+// ============ VF 层（二分累加）：axis=1，沿尾维归约，成对折叠后向量内 ReduceSum，输出 [d0] ============
31+// 统一 case 每行 4 个 VL 块，foldPoint=128（=2 块）。二分树：c0+c2、c1+c3（第 1 层）→ c0+c1（第 2 层）。
32+__simd_vf__ inline void ReduceSumArVfBinary(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
33+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
34+{
35+ AscendC::Reg::RegTensor<float> c0, c1, c2, c3, outReg;
36+ AscendC::Reg::UnalignRegForStore unalignAcc; // 非对齐散出累积器
37+ AscendC::Reg::MaskReg fullMask = AscendC::Reg::CreateMask<float, AscendC::Reg::MaskPattern::ALL>();
38+ 
39+ for (uint16_t i = 0; i < static_cast<uint16_t>(d0); i++) { // 行循环：每行二分折叠出一个标量
40+ const uint32_t base = i * d1Pad;
41+ uint32_t lastValid = d1 - 3u * VL_B32; // 末块有效列数（250-192=58）；UpdateMask 会原地递减
42+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(c0, xAddr + base + 0u * VL_B32); // [0:64)
43+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(c1, xAddr + base + 1u * VL_B32); // [64:128)
44+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(c2, xAddr + base + 2u * VL_B32); // [128:192)
45+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(c3, xAddr + base + 3u * VL_B32); // [192:256)
46+ AscendC::Reg::MaskReg lastMask = AscendC::Reg::UpdateMask<float>(lastValid); // 末块 58 个有效 lane
47+ // 第 1 层：配对距离 foldPoint=128（=2 块）。c2 整块有效；c3 尾部 [58,64) 是行外 padding，
48+ // 用 MERGING+lastMask 屏蔽（mask 外保留 c1 原值，不污染）。两条 Add 相互独立 → 串行步数减半。
49+ AscendC::Reg::Add(c0, c0, c2, fullMask); // [0:64)+[128:192)
50+ AscendC::Reg::Add<float, AscendC::Reg::MaskMergeMode::MERGING>(c1, c1, c3, lastMask); // [64:128)+[192:250)
51+ // 第 2 层：c0+c1 → 每 lane 含 4 块之和；再向量内 ReduceSum 压成行标量
52+ AscendC::Reg::Add(c0, c0, c1, fullMask);
53+ AscendC::Reg::ReduceSum(outReg, c0, fullMask); // 64 lane → 行和落 lane0
54+ AscendC::Reg::StoreUnAlign(zAddr, outReg, unalignAcc, 1); // 追加 1 个标量，zAddr 自动后移
55+ }
56+ AscendC::Reg::StoreUnAlignPost(zAddr, unalignAcc, 0); // 冲刷尾部不足一个 block 的残留
57+}
58+ 
59+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
60+class KernelReduceSum {
61+public:
62+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
63+ {
64+ d0_ = d0;
65+ d1_ = d1;
66+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN;
67+ outLen_ = d0_; // axis=1 输出 [d0]
68+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
69+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
70+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float);
71+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
72+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
73+ pipe_.InitBuffer(outQueue_, 1, outBytes);
74+ }
75+ 
76+ __aicore__ inline void Process()
77+ {
78+ CopyIn();
79+ Compute();
80+ CopyOut();
81+ }
82+ 
83+private:
84+ __aicore__ inline void CopyIn()
85+ {
86+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
87+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
88+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
89+ for (uint32_t i = 0; i < d0_; i++) {
90+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
91+ }
92+ inQueue_.EnQue(xUb);
93+ }
94+ 
95+ __aicore__ inline void Compute()
96+ {
97+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
98+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
99+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
100+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
101+ for (int r = 0; r < VF_REPEAT; r++) {
102+ asc_vf_call<ReduceSumArVfBinary>(xAddr, zAddr, d0_, d1_, d1Pad_);
103+ }
104+ outQueue_.EnQue(zUb);
105+ inQueue_.FreeTensor(xUb);
106+ }
107+ 
108+ __aicore__ inline void CopyOut()
109+ {
110+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
111+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
112+ AscendC::DataCopyPad(zGm_, zUb, params);
113+ outQueue_.FreeTensor(zUb);
114+ }
115+ 
116+ AscendC::TPipe pipe_;
117+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
118+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
119+ AscendC::GlobalTensor<float> xGm_, zGm_;
120+ uint32_t d0_, d1_, d1Pad_, outLen_;
121+};
122+ 
123+__global__ __aicore__ __vector__ void ReduceSumKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
124+{
125+ KernelReduceSum op;
126+ op.Init(x, z, d0, d1);
127+ op.Process();
128+}
129+ 
130+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
131+int main()
132+{
133+ // 本二分写法按「每行恰 4 个 VL 块、foldPoint=128」硬展开；改规格需同步调整 VF。
134+ static_assert((D1 + VL_B32 - 1) / VL_B32 == 4, "reduce_sum_ar_binary 假设每行 4 个 VL 块");
135+ 
136+ CHECK_ACL(aclInit(nullptr));
137+ CHECK_ACL(aclrtSetDevice(0));
138+ aclrtStream stream;
139+ CHECK_ACL(aclrtCreateStream(&stream));
140+ 
141+ const uint32_t outLen = D0;
142+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
143+ std::vector<float> ref(outLen, 0.f);
144+ {
145+ std::vector<double> acc(outLen, 0.0);
146+ for (uint32_t i = 0; i < D0; i++)
147+ for (uint32_t j = 0; j < D1; j++) acc[i] += (double)hX[i * D1 + j]; // 标杆：每行 double 求和
148+ for (uint32_t k = 0; k < outLen; k++) ref[k] = (float)acc[k];
149+ }
150+ std::vector<float> hZ(outLen, 0.f);
151+ 
152+ uint8_t *dX, *dZ;
153+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
154+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
155+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
156+ ACL_MEMCPY_HOST_TO_DEVICE));
157+ 
158+ ReduceSumKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
159+ CHECK_ACL(aclrtSynchronizeStream(stream));
160+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
161+ ACL_MEMCPY_DEVICE_TO_HOST));
162+ 
163+ bool ok = vf::VerifyRel(hZ, ref, 1e-3);
164+ std::cout << "[reduce_sum axis1 binary] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
165+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
166+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
167+ 
168+ aclrtFree(dX); aclrtFree(dZ);
169+ CHECK_ACL(aclrtDestroyStream(stream));
170+ CHECK_ACL(aclrtResetDevice(0));
171+ CHECK_ACL(aclFinalize());
172+ return ok ? 0 : 1;
173+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_sum_ar_pragma.asc+165-0
@@ -0,0 +1,165 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceSum，axis=1（沿尾维归约，输出 [D0]），#pragma 自动展开 VF
3+ *
4+ * 从 reduce_sum.asc 拆出 axis=1 #pragma 自动展开一个 VF（ReduceSumArVfPragma）：每行按 VL 分块逐元素 Add 折进
5+ * 累加器（跨块必须 MERGING），再「向量内 ReduceSum」压成一个标量（落 lane0），StoreUnAlign 把
6+ * 逐行标量拼成连续输出流，循环末 StoreUnAlignPost 冲刷残留（二者成对）。
7+ * 基线 / 多累加器手写展开 / 二分累加版见
8+ * reduce_sum_ar_baseline.asc / reduce_sum_ar_unroll.asc / reduce_sum_ar_binary.asc。
9+ * 注意：#pragma 仅展开循环体、未打断累加串行依赖，对归约无实测加速（≈baseline，详见 README）。
10+ *
11+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ar。
12+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
13+ *
14+ * 编译 & 运行：
15+ * cmake --build build --target reduce_sum_ar_pragma
16+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_sum_ar_pragma -s Ascend950
17+ */
18+ 
19+#include <iostream>
20+#include <vector>
21+#include <cmath>
22+#include "acl/acl.h"
23+#include "kernel_operator.h"
24+#include "../../include/vf_common.h"
25+ 
26+// ============ VF 层（基线，对照用）：axis=1，沿尾维归约，逐行向量内 ReduceSum，输出 [d0] ============
27+// 每行按 VL 分块折叠（覆盖 d1 > VL）进累加器，再向量内归约成一个标量。
28+__simd_vf__ inline void ReduceSumArVfPragma(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
29+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
30+{
31+ AscendC::Reg::RegTensor<float> accReg, inReg, outReg;
32+ AscendC::Reg::UnalignRegForStore unalignAcc; // 非对齐散出累积器
33+ AscendC::Reg::MaskReg fullMask = AscendC::Reg::CreateMask<float, AscendC::Reg::MaskPattern::ALL>();
34+ const uint16_t rowChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
35+ 
36+ #pragma unroll 4 // 自动展开：编译器把循环体展成多条，但累加器单一、串行步数不变
37+ for (uint16_t i = 0; i < static_cast<uint16_t>(d0); i++) { // 外层：行循环（每行归约出一个标量）
38+ AscendC::Reg::Duplicate(accReg, 0.0f);
39+ uint32_t remainCols = d1; // UpdateMask 每轮自动扣 VL，记录剩余列数
40+ for (uint16_t c = 0; c < rowChunks; c++) { // 内层：列块循环（按 VL 切分 d1 列，折叠进累加器）
41+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols); // 本块有效列数
42+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(
43+ inReg, xAddr + i * d1Pad + c * VL_B32);
44+ // 必须 MERGING：mask 外的 lane 保留累加器原值。默认 ZEROING 会把 mask 外清零，
45+ // 末块（mask 不满）会清掉前面整块已累加进来的列，再 ReduceSum 就丢了一部分。
46+ AscendC::Reg::Add<float, AscendC::Reg::MaskMergeMode::MERGING>(accReg, accReg, inReg, mask); // 跨块逐元素折叠
47+ }
48+ AscendC::Reg::ReduceSum(outReg, accReg, fullMask); // 向量内归约 → 行和落 lane0（空洞为 0 不影响）
49+ AscendC::Reg::StoreUnAlign(zAddr, outReg, unalignAcc, 1); // 追加 1 个标量，zAddr 自动后移
50+ }
51+ AscendC::Reg::StoreUnAlignPost(zAddr, unalignAcc, 0); // 冲刷尾部不足一个 block 的残留
52+}
53+ 
54+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
55+class KernelReduceSum {
56+public:
57+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
58+ {
59+ d0_ = d0;
60+ d1_ = d1;
61+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN;
62+ outLen_ = d0_; // axis=1 输出 [d0]
63+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
64+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
65+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float);
66+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
67+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
68+ pipe_.InitBuffer(outQueue_, 1, outBytes);
69+ }
70+ 
71+ __aicore__ inline void Process()
72+ {
73+ CopyIn();
74+ Compute();
75+ CopyOut();
76+ }
77+ 
78+private:
79+ __aicore__ inline void CopyIn()
80+ {
81+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
82+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
83+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
84+ for (uint32_t i = 0; i < d0_; i++) {
85+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
86+ }
87+ inQueue_.EnQue(xUb);
88+ }
89+ 
90+ __aicore__ inline void Compute()
91+ {
92+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
93+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
94+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
95+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
96+ for (int r = 0; r < VF_REPEAT; r++) {
97+ asc_vf_call<ReduceSumArVfPragma>(xAddr, zAddr, d0_, d1_, d1Pad_);
98+ }
99+ outQueue_.EnQue(zUb);
100+ inQueue_.FreeTensor(xUb);
101+ }
102+ 
103+ __aicore__ inline void CopyOut()
104+ {
105+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
106+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
107+ AscendC::DataCopyPad(zGm_, zUb, params);
108+ outQueue_.FreeTensor(zUb);
109+ }
110+ 
111+ AscendC::TPipe pipe_;
112+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
113+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
114+ AscendC::GlobalTensor<float> xGm_, zGm_;
115+ uint32_t d0_, d1_, d1Pad_, outLen_;
116+};
117+ 
118+__global__ __aicore__ __vector__ void ReduceSumKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
119+{
120+ KernelReduceSum op;
121+ op.Init(x, z, d0, d1);
122+ op.Process();
123+}
124+ 
125+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
126+int main()
127+{
128+ CHECK_ACL(aclInit(nullptr));
129+ CHECK_ACL(aclrtSetDevice(0));
130+ aclrtStream stream;
131+ CHECK_ACL(aclrtCreateStream(&stream));
132+ 
133+ const uint32_t outLen = D0;
134+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
135+ std::vector<float> ref(outLen, 0.f);
136+ {
137+ std::vector<double> acc(outLen, 0.0);
138+ for (uint32_t i = 0; i < D0; i++)
139+ for (uint32_t j = 0; j < D1; j++) acc[i] += (double)hX[i * D1 + j]; // 标杆：每行 double 求和
140+ for (uint32_t k = 0; k < outLen; k++) ref[k] = (float)acc[k];
141+ }
142+ std::vector<float> hZ(outLen, 0.f);
143+ 
144+ uint8_t *dX, *dZ;
145+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
146+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
147+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
148+ ACL_MEMCPY_HOST_TO_DEVICE));
149+ 
150+ ReduceSumKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
151+ CHECK_ACL(aclrtSynchronizeStream(stream));
152+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
153+ ACL_MEMCPY_DEVICE_TO_HOST));
154+ 
155+ bool ok = vf::VerifyRel(hZ, ref, 1e-3);
156+ std::cout << "[reduce_sum axis1 pragma] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
157+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
158+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
159+ 
160+ aclrtFree(dX); aclrtFree(dZ);
161+ CHECK_ACL(aclrtDestroyStream(stream));
162+ CHECK_ACL(aclrtResetDevice(0));
163+ CHECK_ACL(aclFinalize());
164+ return ok ? 0 : 1;
165+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_sum_ar_unroll.asc+181-0
@@ -0,0 +1,181 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceSum，axis=1（沿尾维归约，输出 [D0]），多累加器展开 VF
3+ *
4+ * 从 reduce_sum.asc 拆出 axis=1 多累加器展开一个 VF（ReduceSumArVfUnroll）：一轮用 4 个累加器
5+ * 并行折叠 4 行、再各自 ReduceSum，把单行 ReduceSum 的延迟相互隐藏。每行仍用单累加器按列块顺序
6+ * 折叠，归约顺序与基线一致 → 结果逐位等于基线。4 个 StoreUnAlign 按行序追加，尾部不足 4 的行单独处理。
7+ * 同 axis 的基线 / #pragma 自动展开 / 二分累加版见
8+ * reduce_sum_ar_baseline.asc / reduce_sum_ar_pragma.asc / reduce_sum_ar_binary.asc。
9+ *
10+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ar。
11+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
12+ *
13+ * 编译 & 运行：
14+ * cmake --build build --target reduce_sum_ar_unroll
15+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_sum_ar_unroll -s Ascend950
16+ */
17+ 
18+#include <iostream>
19+#include <vector>
20+#include <cmath>
21+#include "acl/acl.h"
22+#include "kernel_operator.h"
23+#include "../../include/vf_common.h"
24+ 
25+// ============ VF 层（多累加器展开）：axis=1，一轮 4 行并行折叠，再各自向量内 ReduceSum ============
26+__simd_vf__ inline void ReduceSumArVfUnroll(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
27+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
28+{
29+ AscendC::Reg::RegTensor<float> acc0, acc1, acc2, acc3, in0, in1, in2, in3;
30+ AscendC::Reg::UnalignRegForStore unalignAcc;
31+ AscendC::Reg::MaskReg fullMask = AscendC::Reg::CreateMask<float, AscendC::Reg::MaskPattern::ALL>();
32+ const uint16_t rowChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
33+ const uint16_t groups = static_cast<uint16_t>(d0) / 4;
34+ 
35+ for (uint16_t g = 0; g < groups; g++) { // 主循环：每轮 4 行并行折叠
36+ const uint16_t r = static_cast<uint16_t>(g * 4);
37+ AscendC::Reg::Duplicate(acc0, 0.0f); AscendC::Reg::Duplicate(acc1, 0.0f);
38+ AscendC::Reg::Duplicate(acc2, 0.0f); AscendC::Reg::Duplicate(acc3, 0.0f);
39+ uint32_t remainCols = d1;
40+ for (uint16_t c = 0; c < rowChunks; c++) { // 4 行同一列结构 → 共享 mask
41+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols);
42+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in0, xAddr + (r + 0) * d1Pad + c * VL_B32);
43+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in1, xAddr + (r + 1) * d1Pad + c * VL_B32);
44+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in2, xAddr + (r + 2) * d1Pad + c * VL_B32);
45+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in3, xAddr + (r + 3) * d1Pad + c * VL_B32);
46+ AscendC::Reg::Add<float, AscendC::Reg::MaskMergeMode::MERGING>(acc0, acc0, in0, mask);
47+ AscendC::Reg::Add<float, AscendC::Reg::MaskMergeMode::MERGING>(acc1, acc1, in1, mask);
48+ AscendC::Reg::Add<float, AscendC::Reg::MaskMergeMode::MERGING>(acc2, acc2, in2, mask);
49+ AscendC::Reg::Add<float, AscendC::Reg::MaskMergeMode::MERGING>(acc3, acc3, in3, mask);
50+ }
51+ AscendC::Reg::ReduceSum(in0, acc0, fullMask); AscendC::Reg::StoreUnAlign(zAddr, in0, unalignAcc, 1); // 复用 in* 接收行标量
52+ AscendC::Reg::ReduceSum(in1, acc1, fullMask); AscendC::Reg::StoreUnAlign(zAddr, in1, unalignAcc, 1);
53+ AscendC::Reg::ReduceSum(in2, acc2, fullMask); AscendC::Reg::StoreUnAlign(zAddr, in2, unalignAcc, 1);
54+ AscendC::Reg::ReduceSum(in3, acc3, fullMask); AscendC::Reg::StoreUnAlign(zAddr, in3, unalignAcc, 1);
55+ }
56+ for (uint16_t r = static_cast<uint16_t>(groups * 4); r < static_cast<uint16_t>(d0); r++) { // 尾部 0~3 行
57+ AscendC::Reg::Duplicate(acc0, 0.0f);
58+ uint32_t remainCols = d1;
59+ for (uint16_t c = 0; c < rowChunks; c++) {
60+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols);
61+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in0, xAddr + r * d1Pad + c * VL_B32);
62+ AscendC::Reg::Add<float, AscendC::Reg::MaskMergeMode::MERGING>(acc0, acc0, in0, mask);
63+ }
64+ AscendC::Reg::ReduceSum(in0, acc0, fullMask);
65+ AscendC::Reg::StoreUnAlign(zAddr, in0, unalignAcc, 1);
66+ }
67+ AscendC::Reg::StoreUnAlignPost(zAddr, unalignAcc, 0);
68+}
69+ 
70+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
71+class KernelReduceSum {
72+public:
73+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
74+ {
75+ d0_ = d0;
76+ d1_ = d1;
77+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN;
78+ outLen_ = d0_; // axis=1 输出 [d0]
79+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
80+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
81+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float);
82+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
83+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
84+ pipe_.InitBuffer(outQueue_, 1, outBytes);
85+ }
86+ 
87+ __aicore__ inline void Process()
88+ {
89+ CopyIn();
90+ Compute();
91+ CopyOut();
92+ }
93+ 
94+private:
95+ __aicore__ inline void CopyIn()
96+ {
97+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
98+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
99+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
100+ for (uint32_t i = 0; i < d0_; i++) {
101+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
102+ }
103+ inQueue_.EnQue(xUb);
104+ }
105+ 
106+ __aicore__ inline void Compute()
107+ {
108+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
109+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
110+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
111+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
112+ for (int r = 0; r < VF_REPEAT; r++) {
113+ asc_vf_call<ReduceSumArVfUnroll>(xAddr, zAddr, d0_, d1_, d1Pad_);
114+ }
115+ outQueue_.EnQue(zUb);
116+ inQueue_.FreeTensor(xUb);
117+ }
118+ 
119+ __aicore__ inline void CopyOut()
120+ {
121+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
122+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
123+ AscendC::DataCopyPad(zGm_, zUb, params);
124+ outQueue_.FreeTensor(zUb);
125+ }
126+ 
127+ AscendC::TPipe pipe_;
128+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
129+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
130+ AscendC::GlobalTensor<float> xGm_, zGm_;
131+ uint32_t d0_, d1_, d1Pad_, outLen_;
132+};
133+ 
134+__global__ __aicore__ __vector__ void ReduceSumKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
135+{
136+ KernelReduceSum op;
137+ op.Init(x, z, d0, d1);
138+ op.Process();
139+}
140+ 
141+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
142+int main()
143+{
144+ CHECK_ACL(aclInit(nullptr));
145+ CHECK_ACL(aclrtSetDevice(0));
146+ aclrtStream stream;
147+ CHECK_ACL(aclrtCreateStream(&stream));
148+ 
149+ const uint32_t outLen = D0;
150+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
151+ std::vector<float> ref(outLen, 0.f);
152+ {
153+ std::vector<double> acc(outLen, 0.0);
154+ for (uint32_t i = 0; i < D0; i++)
155+ for (uint32_t j = 0; j < D1; j++) acc[i] += (double)hX[i * D1 + j]; // 标杆：每行 double 求和
156+ for (uint32_t k = 0; k < outLen; k++) ref[k] = (float)acc[k];
157+ }
158+ std::vector<float> hZ(outLen, 0.f);
159+ 
160+ uint8_t *dX, *dZ;
161+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
162+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
163+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
164+ ACL_MEMCPY_HOST_TO_DEVICE));
165+ 
166+ ReduceSumKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
167+ CHECK_ACL(aclrtSynchronizeStream(stream));
168+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
169+ ACL_MEMCPY_DEVICE_TO_HOST));
170+ 
171+ bool ok = vf::VerifyRel(hZ, ref, 1e-3);
172+ std::cout << "[reduce_sum axis1 unroll] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
173+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
174+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
175+ 
176+ aclrtFree(dX); aclrtFree(dZ);
177+ CHECK_ACL(aclrtDestroyStream(stream));
178+ CHECK_ACL(aclrtResetDevice(0));
179+ CHECK_ACL(aclFinalize());
180+ return ok ? 0 : 1;
181+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_sum_ra_baseline.asc+158-0
@@ -0,0 +1,158 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceSum，axis=0（沿首维归约，输出 [D1]），基线 VF
3+ *
4+ * 从 reduce_sum.asc 拆出 axis=0 基线一个 VF（ReduceSumRaVf）：按列分块，每个列块跨 D0 行做
5+ * 「逐元素 Add」累加（累加器初值 0），最后 StoreAlign 整段搬出（输出天然连续对齐）。
6+ * 同 axis 的多累加器手写展开 / #pragma 自动展开版见 reduce_sum_ra_unroll.asc / reduce_sum_ra_pragma.asc。
7+ *
8+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ra。
9+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
10+ *
11+ * 编译 & 运行：
12+ * cmake --build build --target reduce_sum_ra_baseline
13+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_sum_ra_baseline -s Ascend950
14+ */
15+ 
16+#include <iostream>
17+#include <vector>
18+#include <cmath>
19+#include "acl/acl.h"
20+#include "kernel_operator.h"
21+#include "../../include/vf_common.h"
22+ 
23+// ============ VF 层（基线，对照用）：axis=0，沿首维归约，跨行逐元素 Add，输出 [d1] ============
24+// 按列分块（覆盖 d1 > VL）：每个列块跨 d0 行累加，再整段搬出。
25+__simd_vf__ inline void ReduceSumRaVf(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
26+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
27+{
28+ AscendC::Reg::RegTensor<float> accReg, inReg;
29+ uint32_t remainCols = d1; // UpdateMask 每轮自动扣 VL，记录剩余列数
30+ const uint16_t colChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
31+ 
32+ for (uint16_t c = 0; c < colChunks; c++) { // 外层：列块循环（按 VL 切分 d1 列，每块独立输出）
33+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols); // 本列块有效列数
34+ AscendC::Reg::Duplicate(accReg, 0.0f); // 累加器全 lane 置 0
35+ for (uint16_t i = 0; i < static_cast<uint16_t>(d0); i++) { // 内层：行循环（跨 d0 行累加同一列块）
36+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(
37+ inReg, xAddr + i * d1Pad + c * VL_B32); // 行已对齐到 d1Pad
38+ AscendC::Reg::Add(accReg, accReg, inReg, mask); // accReg += inReg
39+ }
40+ // 这里 mask 在整个行循环里恒定，且下面 StoreAlign 用同一 mask 搬出 → mask 外 lane 被丢弃，
41+ // 默认 ZEROING 无碍（与 axis=1 跨块折叠不同，那里必须 MERGING）
42+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(
43+ zAddr + c * VL_B32, accReg, mask); // 整段搬出（连续对齐）
44+ }
45+}
46+ 
47+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
48+class KernelReduceSum {
49+public:
50+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
51+ {
52+ d0_ = d0;
53+ d1_ = d1;
54+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN; // 行步长按 32B 对齐
55+ outLen_ = d1_; // axis=0 输出 [d1]
56+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
57+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
58+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float); // 额外留 VL 余量：尾行最后一块整轮读取不越界
59+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
60+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
61+ pipe_.InitBuffer(outQueue_, 1, outBytes);
62+ }
63+ 
64+ __aicore__ inline void Process()
65+ {
66+ CopyIn();
67+ Compute();
68+ CopyOut();
69+ }
70+ 
71+private:
72+ __aicore__ inline void CopyIn()
73+ {
74+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
75+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
76+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
77+ for (uint32_t i = 0; i < d0_; i++) { // 逐行搬到 32B 对齐槽位
78+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
79+ }
80+ inQueue_.EnQue(xUb);
81+ }
82+ 
83+ __aicore__ inline void Compute()
84+ {
85+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
86+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
87+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
88+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
89+ for (int r = 0; r < VF_REPEAT; r++) {
90+ asc_vf_call<ReduceSumRaVf>(xAddr, zAddr, d0_, d1_, d1Pad_);
91+ }
92+ outQueue_.EnQue(zUb);
93+ inQueue_.FreeTensor(xUb);
94+ }
95+ 
96+ __aicore__ inline void CopyOut()
97+ {
98+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
99+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
100+ AscendC::DataCopyPad(zGm_, zUb, params);
101+ outQueue_.FreeTensor(zUb);
102+ }
103+ 
104+ AscendC::TPipe pipe_;
105+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
106+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
107+ AscendC::GlobalTensor<float> xGm_, zGm_;
108+ uint32_t d0_, d1_, d1Pad_, outLen_;
109+};
110+ 
111+__global__ __aicore__ __vector__ void ReduceSumKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
112+{
113+ KernelReduceSum op;
114+ op.Init(x, z, d0, d1);
115+ op.Process();
116+}
117+ 
118+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
119+int main()
120+{
121+ CHECK_ACL(aclInit(nullptr));
122+ CHECK_ACL(aclrtSetDevice(0));
123+ aclrtStream stream;
124+ CHECK_ACL(aclrtCreateStream(&stream));
125+ 
126+ const uint32_t outLen = D1;
127+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
128+ std::vector<float> ref(outLen, 0.f);
129+ {
130+ std::vector<double> acc(outLen, 0.0);
131+ for (uint32_t i = 0; i < D0; i++)
132+ for (uint32_t j = 0; j < D1; j++) acc[j] += (double)hX[i * D1 + j]; // 标杆：每列 double 求和
133+ for (uint32_t k = 0; k < outLen; k++) ref[k] = (float)acc[k];
134+ }
135+ std::vector<float> hZ(outLen, 0.f);
136+ 
137+ uint8_t *dX, *dZ;
138+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
139+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
140+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
141+ ACL_MEMCPY_HOST_TO_DEVICE));
142+ 
143+ ReduceSumKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
144+ CHECK_ACL(aclrtSynchronizeStream(stream));
145+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
146+ ACL_MEMCPY_DEVICE_TO_HOST));
147+ 
148+ bool ok = vf::VerifyRel(hZ, ref, 1e-3);
149+ std::cout << "[reduce_sum axis0 baseline] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
150+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
151+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
152+ 
153+ aclrtFree(dX); aclrtFree(dZ);
154+ CHECK_ACL(aclrtDestroyStream(stream));
155+ CHECK_ACL(aclrtResetDevice(0));
156+ CHECK_ACL(aclFinalize());
157+ return ok ? 0 : 1;
158+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_sum_ra_pragma.asc+160-0
@@ -0,0 +1,160 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceSum，axis=0（沿首维归约，输出 [D1]），#pragma 自动展开 VF
3+ *
4+ * 从 reduce_sum.asc 拆出 axis=0 #pragma 自动展开一个 VF（ReduceSumRaVfPragma）：按列分块，每个列块跨 D0 行做
5+ * 「逐元素 Add」累加（累加器初值 0），最后 StoreAlign 整段搬出（输出天然连续对齐）。
6+ * 基线 / 多累加器手写展开版见 reduce_sum_ra_baseline.asc / reduce_sum_ra_unroll.asc。
7+ * 注意：#pragma 仅展开循环体、未打断累加串行依赖，对归约无实测加速（≈baseline，详见 README）。
8+ *
9+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ra。
10+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
11+ *
12+ * 编译 & 运行：
13+ * cmake --build build --target reduce_sum_ra_pragma
14+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_sum_ra_pragma -s Ascend950
15+ */
16+ 
17+#include <iostream>
18+#include <vector>
19+#include <cmath>
20+#include "acl/acl.h"
21+#include "kernel_operator.h"
22+#include "../../include/vf_common.h"
23+ 
24+// ============ VF 层（基线，对照用）：axis=0，沿首维归约，跨行逐元素 Add，输出 [d1] ============
25+// 按列分块（覆盖 d1 > VL）：每个列块跨 d0 行累加，再整段搬出。
26+__simd_vf__ inline void ReduceSumRaVfPragma(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
27+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
28+{
29+ AscendC::Reg::RegTensor<float> accReg, inReg;
30+ uint32_t remainCols = d1; // UpdateMask 每轮自动扣 VL，记录剩余列数
31+ const uint16_t colChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
32+ 
33+ for (uint16_t c = 0; c < colChunks; c++) { // 外层：列块循环（按 VL 切分 d1 列，每块独立输出）
34+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols); // 本列块有效列数
35+ AscendC::Reg::Duplicate(accReg, 0.0f); // 累加器全 lane 置 0
36+ #pragma unroll 4 // 自动展开：编译器把循环体展成多条，但累加器单一、串行步数不变
37+ for (uint16_t i = 0; i < static_cast<uint16_t>(d0); i++) { // 内层：行循环（跨 d0 行累加同一列块）
38+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(
39+ inReg, xAddr + i * d1Pad + c * VL_B32); // 行已对齐到 d1Pad
40+ AscendC::Reg::Add(accReg, accReg, inReg, mask); // accReg += inReg
41+ }
42+ // 这里 mask 在整个行循环里恒定，且下面 StoreAlign 用同一 mask 搬出 → mask 外 lane 被丢弃，
43+ // 默认 ZEROING 无碍（与 axis=1 跨块折叠不同，那里必须 MERGING）
44+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(
45+ zAddr + c * VL_B32, accReg, mask); // 整段搬出（连续对齐）
46+ }
47+}
48+ 
49+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
50+class KernelReduceSum {
51+public:
52+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
53+ {
54+ d0_ = d0;
55+ d1_ = d1;
56+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN; // 行步长按 32B 对齐
57+ outLen_ = d1_; // axis=0 输出 [d1]
58+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
59+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
60+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float); // 额外留 VL 余量：尾行最后一块整轮读取不越界
61+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
62+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
63+ pipe_.InitBuffer(outQueue_, 1, outBytes);
64+ }
65+ 
66+ __aicore__ inline void Process()
67+ {
68+ CopyIn();
69+ Compute();
70+ CopyOut();
71+ }
72+ 
73+private:
74+ __aicore__ inline void CopyIn()
75+ {
76+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
77+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
78+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
79+ for (uint32_t i = 0; i < d0_; i++) { // 逐行搬到 32B 对齐槽位
80+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
81+ }
82+ inQueue_.EnQue(xUb);
83+ }
84+ 
85+ __aicore__ inline void Compute()
86+ {
87+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
88+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
89+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
90+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
91+ for (int r = 0; r < VF_REPEAT; r++) {
92+ asc_vf_call<ReduceSumRaVfPragma>(xAddr, zAddr, d0_, d1_, d1Pad_);
93+ }
94+ outQueue_.EnQue(zUb);
95+ inQueue_.FreeTensor(xUb);
96+ }
97+ 
98+ __aicore__ inline void CopyOut()
99+ {
100+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
101+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
102+ AscendC::DataCopyPad(zGm_, zUb, params);
103+ outQueue_.FreeTensor(zUb);
104+ }
105+ 
106+ AscendC::TPipe pipe_;
107+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
108+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
109+ AscendC::GlobalTensor<float> xGm_, zGm_;
110+ uint32_t d0_, d1_, d1Pad_, outLen_;
111+};
112+ 
113+__global__ __aicore__ __vector__ void ReduceSumKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
114+{
115+ KernelReduceSum op;
116+ op.Init(x, z, d0, d1);
117+ op.Process();
118+}
119+ 
120+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
121+int main()
122+{
123+ CHECK_ACL(aclInit(nullptr));
124+ CHECK_ACL(aclrtSetDevice(0));
125+ aclrtStream stream;
126+ CHECK_ACL(aclrtCreateStream(&stream));
127+ 
128+ const uint32_t outLen = D1;
129+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
130+ std::vector<float> ref(outLen, 0.f);
131+ {
132+ std::vector<double> acc(outLen, 0.0);
133+ for (uint32_t i = 0; i < D0; i++)
134+ for (uint32_t j = 0; j < D1; j++) acc[j] += (double)hX[i * D1 + j]; // 标杆：每列 double 求和
135+ for (uint32_t k = 0; k < outLen; k++) ref[k] = (float)acc[k];
136+ }
137+ std::vector<float> hZ(outLen, 0.f);
138+ 
139+ uint8_t *dX, *dZ;
140+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
141+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
142+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
143+ ACL_MEMCPY_HOST_TO_DEVICE));
144+ 
145+ ReduceSumKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
146+ CHECK_ACL(aclrtSynchronizeStream(stream));
147+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
148+ ACL_MEMCPY_DEVICE_TO_HOST));
149+ 
150+ bool ok = vf::VerifyRel(hZ, ref, 1e-3);
151+ std::cout << "[reduce_sum axis0 pragma] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
152+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
153+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
154+ 
155+ aclrtFree(dX); aclrtFree(dZ);
156+ CHECK_ACL(aclrtDestroyStream(stream));
157+ CHECK_ACL(aclrtResetDevice(0));
158+ CHECK_ACL(aclFinalize());
159+ return ok ? 0 : 1;
160+}
ASamples/2_Performance/simd_vf_story/reduce/src/reduce_sum_ra_unroll.asc+170-0
@@ -0,0 +1,170 @@
1+/**
2+ * 单 VF 拆分版：二维 ReduceSum，axis=0（沿首维归约，输出 [D1]），多累加器展开 VF
3+ *
4+ * 从 reduce_sum.asc 拆出 axis=0 多累加器展开一个 VF（ReduceSumRaVfUnroll）：用 4 个独立累加器
5+ * acc0..acc3 按行号分组累加（行 r 进 acc[r%4]），最后两两合并 → 4 条归约流并行，利于双发/延迟
6+ * 隐藏。尾部不足 4 的行用第二个 for 处理。注：4 个部分和改变了浮点求和顺序，结果与基线有低位
7+ * 舍入差异（host golden 用 double+相对容差，仍 PASS）。
8+ *
9+ * 命名约定：ar/ra 拼出 [A,R] 布局、归约其中 R 轴 → ar=归约 axis=1（输出 [D0]）、ra=归约 axis=0（输出 [D1]）；本文件为 ra。
10+ * 统一 case 规格（本目录 16 个 VF 文件共用）：[D0,D1]=[78,250]，D1>VL 且 %8≠0、D0%4=2（覆盖分块 / 尾块 / 展开尾行各路径）。
11+ *
12+ * 编译 & 运行：
13+ * cmake --build build --target reduce_sum_ra_unroll
14+ * cannsim record ./build/Samples/2_Performance/simd_vf_story/reduce_sum_ra_unroll -s Ascend950
15+ */
16+ 
17+#include <iostream>
18+#include <vector>
19+#include <cmath>
20+#include "acl/acl.h"
21+#include "kernel_operator.h"
22+#include "../../include/vf_common.h"
23+ 
24+// ============ VF 层（多累加器展开）：axis=0，4 个独立累加器并行跨行 Add ============
25+__simd_vf__ inline void ReduceSumRaVfUnroll(__ubuf__ float* xAddr, __ubuf__ float* zAddr,
26+ uint32_t d0, uint32_t d1, uint32_t d1Pad)
27+{
28+ AscendC::Reg::RegTensor<float> acc0, acc1, acc2, acc3, in0, in1, in2, in3;
29+ uint32_t remainCols = d1;
30+ const uint16_t colChunks = static_cast<uint16_t>((d1 + VL_B32 - 1) / VL_B32);
31+ 
32+ for (uint16_t c = 0; c < colChunks; c++) {
33+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remainCols);
34+ AscendC::Reg::Duplicate(acc0, 0.0f); AscendC::Reg::Duplicate(acc1, 0.0f);
35+ AscendC::Reg::Duplicate(acc2, 0.0f); AscendC::Reg::Duplicate(acc3, 0.0f);
36+ const uint16_t groups = static_cast<uint16_t>(d0) / 4;
37+ for (uint16_t g = 0; g < groups; g++) { // 主循环：每轮 4 行 → 4 个独立累加器
38+ const uint16_t r = static_cast<uint16_t>(g * 4);
39+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in0, xAddr + (r + 0) * d1Pad + c * VL_B32);
40+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in1, xAddr + (r + 1) * d1Pad + c * VL_B32);
41+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in2, xAddr + (r + 2) * d1Pad + c * VL_B32);
42+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in3, xAddr + (r + 3) * d1Pad + c * VL_B32);
43+ AscendC::Reg::Add(acc0, acc0, in0, mask);
44+ AscendC::Reg::Add(acc1, acc1, in1, mask);
45+ AscendC::Reg::Add(acc2, acc2, in2, mask);
46+ AscendC::Reg::Add(acc3, acc3, in3, mask);
47+ }
48+ for (uint16_t r = static_cast<uint16_t>(groups * 4); r < static_cast<uint16_t>(d0); r++) { // 尾部 0~3 行
49+ AscendC::Reg::LoadAlign<float, AscendC::Reg::LoadDist::DIST_NORM>(in0, xAddr + r * d1Pad + c * VL_B32);
50+ AscendC::Reg::Add(acc0, acc0, in0, mask);
51+ }
52+ AscendC::Reg::Add(acc0, acc0, acc1, mask); // 合并 4 个累加器
53+ AscendC::Reg::Add(acc2, acc2, acc3, mask);
54+ AscendC::Reg::Add(acc0, acc0, acc2, mask);
55+ AscendC::Reg::StoreAlign<float, AscendC::Reg::StoreDist::DIST_NORM>(zAddr + c * VL_B32, acc0, mask);
56+ }
57+}
58+ 
59+// ============ Kernel：单核单 tile，标准 CopyIn → Compute → CopyOut 三段式 ============
60+class KernelReduceSum {
61+public:
62+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
63+ {
64+ d0_ = d0;
65+ d1_ = d1;
66+ d1Pad_ = (d1_ + UB_ALIGN - 1) / UB_ALIGN * UB_ALIGN;
67+ outLen_ = d1_; // axis=0 输出 [d1]
68+ xGm_.SetGlobalBuffer((__gm__ float*)x, d0_ * d1_);
69+ zGm_.SetGlobalBuffer((__gm__ float*)z, outLen_);
70+ uint32_t inBytes = (d0_ * d1Pad_ + VL_B32) * sizeof(float);
71+ uint32_t outBytes = (outLen_ * sizeof(float) + 31) / 32 * 32;
72+ pipe_.InitBuffer(inQueue_, 1, (inBytes + 31) / 32 * 32);
73+ pipe_.InitBuffer(outQueue_, 1, outBytes);
74+ }
75+ 
76+ __aicore__ inline void Process()
77+ {
78+ CopyIn();
79+ Compute();
80+ CopyOut();
81+ }
82+ 
83+private:
84+ __aicore__ inline void CopyIn()
85+ {
86+ AscendC::LocalTensor<float> xUb = inQueue_.AllocTensor<float>();
87+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(d1_ * sizeof(float)), 0, 0, 0};
88+ AscendC::DataCopyPadExtParams<float> padParams{false, 0, 0, 0.f};
89+ for (uint32_t i = 0; i < d0_; i++) {
90+ AscendC::DataCopyPad(xUb[i * d1Pad_], xGm_[i * d1_], params, padParams);
91+ }
92+ inQueue_.EnQue(xUb);
93+ }
94+ 
95+ __aicore__ inline void Compute()
96+ {
97+ AscendC::LocalTensor<float> xUb = inQueue_.DeQue<float>();
98+ AscendC::LocalTensor<float> zUb = outQueue_.AllocTensor<float>();
99+ auto* xAddr = (__ubuf__ float*)xUb.GetPhyAddr();
100+ auto* zAddr = (__ubuf__ float*)zUb.GetPhyAddr();
101+ for (int r = 0; r < VF_REPEAT; r++) {
102+ asc_vf_call<ReduceSumRaVfUnroll>(xAddr, zAddr, d0_, d1_, d1Pad_);
103+ }
104+ outQueue_.EnQue(zUb);
105+ inQueue_.FreeTensor(xUb);
106+ }
107+ 
108+ __aicore__ inline void CopyOut()
109+ {
110+ AscendC::LocalTensor<float> zUb = outQueue_.DeQue<float>();
111+ AscendC::DataCopyExtParams params{1, static_cast<uint32_t>(outLen_ * sizeof(float)), 0, 0, 0};
112+ AscendC::DataCopyPad(zGm_, zUb, params);
113+ outQueue_.FreeTensor(zUb);
114+ }
115+ 
116+ AscendC::TPipe pipe_;
117+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue_;
118+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue_;
119+ AscendC::GlobalTensor<float> xGm_, zGm_;
120+ uint32_t d0_, d1_, d1Pad_, outLen_;
121+};
122+ 
123+__global__ __aicore__ __vector__ void ReduceSumKernel(GM_ADDR x, GM_ADDR z, uint32_t d0, uint32_t d1)
124+{
125+ KernelReduceSum op;
126+ op.Init(x, z, d0, d1);
127+ op.Process();
128+}
129+ 
130+// ============ Host 侧：随机输入（头文件共用）+ 本地 golden 校验 ============
131+int main()
132+{
133+ CHECK_ACL(aclInit(nullptr));
134+ CHECK_ACL(aclrtSetDevice(0));
135+ aclrtStream stream;
136+ CHECK_ACL(aclrtCreateStream(&stream));
137+ 
138+ const uint32_t outLen = D1;
139+ std::vector<float> hX = vf::GenInput(D0 * D1); // 16 个 VF 共用同一份随机输入
140+ std::vector<float> ref(outLen, 0.f);
141+ {
142+ std::vector<double> acc(outLen, 0.0);
143+ for (uint32_t i = 0; i < D0; i++)
144+ for (uint32_t j = 0; j < D1; j++) acc[j] += (double)hX[i * D1 + j]; // 标杆：每列 double 求和
145+ for (uint32_t k = 0; k < outLen; k++) ref[k] = (float)acc[k];
146+ }
147+ std::vector<float> hZ(outLen, 0.f);
148+ 
149+ uint8_t *dX, *dZ;
150+ CHECK_ACL(aclrtMalloc((void**)&dX, hX.size() * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
151+ CHECK_ACL(aclrtMalloc((void**)&dZ, outLen * sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST));
152+ CHECK_ACL(aclrtMemcpy(dX, hX.size() * sizeof(float), hX.data(), hX.size() * sizeof(float),
153+ ACL_MEMCPY_HOST_TO_DEVICE));
154+ 
155+ ReduceSumKernel<<<1, nullptr, stream>>>(dX, dZ, D0, D1);
156+ CHECK_ACL(aclrtSynchronizeStream(stream));
157+ CHECK_ACL(aclrtMemcpy(hZ.data(), outLen * sizeof(float), dZ, outLen * sizeof(float),
158+ ACL_MEMCPY_DEVICE_TO_HOST));
159+ 
160+ bool ok = vf::VerifyRel(hZ, ref, 1e-3);
161+ std::cout << "[reduce_sum axis0 unroll] [" << D0 << "x" << D1 << "] outLen=" << outLen << " z[0..2]=" << hZ[0]
162+ << " " << hZ[1] << " " << hZ[2] << " | z[" << (outLen - 1) << "]=" << hZ[outLen - 1]
163+ << " -> " << (ok ? "PASSED" : "FAILED") << std::endl;
164+ 
165+ aclrtFree(dX); aclrtFree(dZ);
166+ CHECK_ACL(aclrtDestroyStream(stream));
167+ CHECK_ACL(aclrtResetDevice(0));
168+ CHECK_ACL(aclFinalize());
169+ return ok ? 0 : 1;
170+}