Star184
289
代码
代码
Issues78
Pull Requests107
Actions
项目讨论
Wiki
分析
Star184
289
已合并
Feat: 新增 TrsmBatched（StrsmBatched / CtrsmBatched）批量三角求解算子 #243
jingdemeng创建于 7月1日
Feat: 新增 TrsmBatched（StrsmBatched / CtrsmBatched）批量三角求解算子 #243
已合并
jingdemeng创建于 7月1日
从jingdemeng/ops-blasmaster合入到cann/ops-blasmaster
共 45 个文件变更+8961-0
Aexperimental/aclblasCtrsmBatched2/CMakeLists.txt+10-0
CANN-robotCANN-robot7月17日

【openlibing.ci】检测到当前PR中存在代码检查告警抑制 2 处，详情见下表，请Committer检视合理性。 / Detected 2 code check alert suppression(s) in this PR, see table below. Committers please review.

文件路径/File 行号/Line 代码片段/Snippet 工具/Tool
experimental/aclblasCtrsmBatched2/
run_full_test.py		 31 from run_full_test_common import Harness, write_config # noqa: E402 flake8,ruff
experimental/aclblasTrsmBatched/
run_full_test.py		 31 from run_full_test_common import Harness, write_config # noqa: E402 flake8,ruff
likedislike
@@ -0,0 +1,10 @@
1+cmake_minimum_required(VERSION 3.16)
2+find_package(ASC REQUIRED)
3+ 
4+project(ctrsm_batched LANGUAGES ASC CXX)
5+set(CMAKE_CXX_STANDARD 17)
6+ 
7+set(ACL_INCLUDE_DIR "$ENV{ASCEND_HOME_PATH}/aarch64-linux/include")
8+set(ACL_LIB_DIR "$ENV{ASCEND_HOME_PATH}/lib64")
9+ 
10+add_subdirectory(test)
Aexperimental/aclblasCtrsmBatched2/README.md+204-0
@@ -0,0 +1,204 @@
1+# aclblasCtrsmBatched
2+ 
3+批量复数三角矩阵求解算子（complex64），基于 Ascend C 实现，对标 cuBLAS `cublasCtrsmBatched`。
4+ 
5+## 功能描述
6+ 
7+求解批量复数三角线性方程组：
8+ 
9+- `op(A) * X = alpha * B`（side='L'）
10+- `X * op(A) = alpha * B`（side='R'）
11+ 
12+其中 A 为复数三角矩阵，alpha 为复数标量，支持：
13+- Left/Right 左右乘模式
14+- Upper/Lower 上下三角
15+- NoTrans/Trans/ConjTrans 转置模式
16+- Unit/NonUnit 单位三角矩阵
17+ 
18+## 接口定义
19+ 
20+```c
21+aclblasStatus_t aclblasCtrsmBatched(
22+ aclblasHandle_t handle,
23+ aclblasSideMode_t side,
24+ aclblasFillMode_t uplo,
25+ aclblasOperation_t transa,
26+ aclblasDiagType_t diag,
27+ int64_t m, int64_t n,
28+ const std::complex<float>* alpha,
29+ const std::complex<float>* const aArray[], int64_t lda,
30+ std::complex<float>* const bArray[], int64_t ldb,
31+ int64_t batchCount)
32+```
33+ 
34+## 架构设计
35+ 
36+采用 MIX_AIC_1_2 混合核架构（1 Cube核 + 2 Vector核）：
37+ 
38+```
39+AIV (Vector核): Panel 内三角求解（前代/回代）+ 数据格式转换（AoS<->SoA 转置）
40+AIC (Cube核): Trail 区域 GEMM 更新（Matmul 库调用）
41+同步机制: CrossCoreSetFlag/WaitFlag 跨核事件同步
42+```
43+ 
44+**核心算法流程：**
45+```
46+for each panel:
47+ [AIV] LoadPanelA -> SolveInner(nb*nb) -> WriteBack Xneg
48+ [AIV->AIC] CrossCoreSetFlag(TRSV)
49+ [AIC] DirectRankK(GEMM) -> SetFlag(GEMM)
50+ [AIV] WaitFlag(GEMM) -> WriteBackPanelRows -> LoadPanelA(next) -> 下一 panel
51+```
52+ 
53+**双 AIV 分列优化：**
54+ 
55+大矩阵场景（kDim>=128 且 nColsAligned>=128）自动启用双 AIV 分列模式：
56+- 两个 AIV 核协同处理同一矩阵，各处理一半列
57+- 消除 AIV 空转，提升 AIV 利用率
58+- Panel solve 和 WriteBack 按列并行
59+ 
60+**多核拆分优化：**
61+ 
62+小 batch 场景（batch <= AI Core 数/2）自动启用多核拆分：
63+- 每个 batch 的列方向拆分到多个 AI Core 并行处理
64+- minSplitNCols=64，充分利用空闲核心
65+ 
66+**转置优化：**
67+ 
68+AoS<->SoA 格式转换采用 3 块 buffer 轮转设计：
69+- 动态 tileCols（根据可用 UB 空间计算，大矩阵 240~280 列/tile）
70+- Gather 偏移表解交织（支持任意非对齐尺寸）
71+- Duplicate 预清零 + dstStride 控制行步长（绕过 rightPadding < 32B 限制）
72+- padOn=false 时跳过清零，减少 PipeBarrier 开销
73+ 
74+## 目录结构
75+ 
76+```
77+aclblasCtrsmBatched2/
78+├── CMakeLists.txt
79+├── README.md
80+├── run.sh # 单用例快捷脚本（编译+生成+运行+校验）
81+├── op_host/
82+│ ├── ctrsm_batched_host.cpp # Host 侧：参数校验、Tiling、Kernel 启动
83+│ └── ctrsm_batched_kernel_do.h # Kernel 启动包装声明
84+├── op_kernel/
85+│ ├── ctrsm_batched_kernel.cpp # Kernel 入口（MIX 模式，按 side/uplo/transa 分派四路径）
86+│ ├── ctrsm_batched_kernel_aic.h # AIC Cube核实现（GEMM trail 更新）
87+│ ├── ctrsm_batched_kernel_common.h # 公共常量和工具函数
88+│ ├── ctrsm_batched_tiling_data.h # Tiling 数据结构定义
89+│ │ # ---- AIV 向量核按职责拆分的协作类 ----
90+│ ├── ctrsm_batched_kernel_aiv.h # 编排类 CtrsmMixAivImpl<FORWARD,RIGHT>（含四路径别名）
91+│ ├── ctrsm_batched_kernel_aiv_cfg.h # 共享派生配置 + UB buffer 指针
92+│ ├── ctrsm_batched_kernel_aiv_convert.h # 复数 AoS<->SoA 转换/分块转置工具（路径无关）
93+│ ├── ctrsm_batched_kernel_aiv_canon_a.h # A 矩阵规范化（补零/转置/共轭）
94+│ ├── ctrsm_batched_kernel_aiv_canon_b.h # B 矩阵规范化与回写（左/右乘，模板 RIGHT）
95+│ ├── ctrsm_batched_kernel_aiv_canon_b_deinterleave.h # B 矩阵 AoS 解交织实现
96+│ └── ctrsm_batched_kernel_aiv_solver.h # Panel 三角求解（前代/回代，模板 FORWARD）
97+├── test/
98+│ ├── CMakeLists.txt
99+│ ├── ctrsm_batched_test.cpp # 测试主程序
100+│ └── data/
101+│ ├── gen_data.py # 测试数据生成（输入矩阵 + golden 参考结果）
102+│ └── verify_result.py # 精度验证
103+└── docs/
104+ ├── GPU_TEST # cuBLAS GPU 性能基准数据（320 case）
105+ └── perf_report_320cases_new.md # NPU 性能报告（对标 GPU）
106+```
107+ 
108+> AIV 实现按计算路径模板化为四个类型：`CtrsmLowerLeft` / `CtrsmLowerRight` /
109+> `CtrsmUpperLeft` / `CtrsmUpperRight`（`CtrsmMixAivImpl<FORWARD,RIGHT>` 的别名），
110+> 编译期消除 forward/right 分支；各职责类通过共享 cfg 与 UB buffer 指针协作。
111+ 
112+## 编译运行
113+ 
114+在本样例根目录下执行如下步骤，编译并执行算子。
115+ 
116+- 配置环境变量
117+ 
118+ 请根据当前环境上CANN开发套件包的安装方式，选择对应配置环境变量的命令。
119+ ```bash
120+ source ${ASCEND_HOME_PATH}/set_env.sh
121+ ```
122+ 
123+- 构建
124+ ```bash
125+ mkdir -p build && cd build
126+ cmake ..
127+ make -j8
128+ cd ..
129+ ```
130+ 
131+- 样例执行
132+ ```bash
133+ # Step 1: 生成测试数据（输入矩阵 + golden 参考结果）
134+ python3 test/data/gen_data.py 64 64 32 0 0 0 0 1.0 0.0
135+ 
136+ # Step 2: 运行算子
137+ # 参数: deviceId m n batch side uplo transa diag [alpha_re] [alpha_im]
138+ ./build/test/ctrsm_batched_test 0 64 64 32 0 0 0 0 1.0 0.0
139+ 
140+ # Step 3: 验证精度
141+ python3 test/data/verify_result.py 64 64 32
142+ ```
143+ 
144+ 执行结果如下，说明精度对比成功：
145+ ```
146+ [Success] Case accuracy verification passed.
147+ ```
148+ 
149+ 也可用 `run.sh` 一键完成上述三步（编译+生成+运行+校验）：
150+ ```bash
151+ # 参数: m n batch side uplo transa diag [alpha_re] [alpha_im] [--skip-build]
152+ bash run.sh 64 64 32 0 0 0 0 1.0 0.0
153+ ```
154+ 
155+## 参数说明
156+ 
157+gen_data.py / ctrsm_batched_test 参数顺序：
158+```
159+m n batch side uplo transa diag [alpha_re] [alpha_im]
160+```
161+ 
162+| 参数 | 取值 | 说明 |
163+|------|------|------|
164+| side | 0/1 | 0=Left, 1=Right |
165+| uplo | 0/1 | 0=Upper, 1=Lower |
166+| transa | 0/1/2 | 0=NoTrans, 1=Trans, 2=ConjTrans |
167+| diag | 0/1 | 0=NonUnit, 1=Unit |
168+| alpha_re | float | alpha 实部（默认 1.0） |
169+| alpha_im | float | alpha 虚部（默认 0.0） |
170+ 
171+## 关键参数
172+ 
173+| 参数 | 值 | 说明 |
174+|------|------|------|
175+| MIX_NB_SMALL | 16 | Panel 分块大小（kDim<=1024） |
176+| MIX_NB_LARGE | 32 | Panel 分块大小（kDim>1024） |
177+| LIM_GROUP | 16 | Xneg 分组大小 |
178+| FLOAT_ALIGN | 8 | 向量对齐宽度（float32，8 元素 = 32B） |
179+| TOTAL_AICORES | 20 | AI Core 总数（Ascend910B4） |
180+| minSplitNCols | 64 | 多核拆分最小列数 |
181+| SOC_VERSION | Ascend910B4 | 目标硬件 |
182+| dualAivMode | auto | kDim>=128 且 nColsAligned>=128 时自动启用 |
183+ 
184+## 性能数据
185+ 
186+320 Case 全量测试（对标 cuBLAS cublasCtrsmBatched）：
187+ 
188+- 平均 GPU/NPU: **0.800**
189+- NPU 更快（GPU/NPU >= 1.0）: 109/320 个 case（34%）
190+- GPU/NPU 最大: 1.776（大矩阵场景 NPU 最多快 78%）
191+ 
192+详见 `docs/perf_report_320cases_new.md`。
193+ 
194+## 精度标准
195+ 
196+相对误差阈值：`MERE < 2^-13 * 10 = 1.22e-3`
197+ 
198+精度验证覆盖：
199+- 16 种 side*uplo*transa*diag 模式组合
200+- 矩阵尺寸 16~8192
201+- batch 数 1~248
202+- 复数 alpha 缩放
203+- 非对齐尺寸（17, 20, 31, 33, 227 等）
204+- 320 Case 全量精度通过
Aexperimental/aclblasCtrsmBatched2/docs/GPU_TEST+334-0
@@ -0,0 +1,334 @@
1+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+ 
9+Compiling with nvcc...
10+Running cuBLAS cublasCtrsmBatched benchmark (per-case process, device=7)...
11+ 
12+ # m n bat mode algn alpha kernel(us)
13+------------------------------------------------------
14+ 1 16 16 118 L/U/N/N A 1.0 21.22
15+ 2 16 16 118 L/U/N/U A 1.0 18.24
16+ 3 16 16 118 L/U/T/N A 1.0 11.97
17+ 4 16 16 118 L/U/T/U A 1.0 9.73
18+ 5 16 16 118 L/L/N/N A 1.0 12.26
19+ 6 16 16 118 L/L/N/U A 1.0 9.95
20+ 7 16 16 118 L/L/T/N A 1.0 21.47
21+ 8 16 16 118 L/L/T/U A 1.0 18.46
22+ 9 16 16 118 R/U/N/N A 1.0 11.65
23+ 10 16 16 118 R/U/N/U A 1.0 9.41
24+ 11 16 16 118 R/U/T/N A 1.0 11.52
25+ 12 16 16 118 R/U/T/U A 1.0 9.28
26+ 13 16 16 118 R/L/N/N A 1.0 11.52
27+ 14 16 16 118 R/L/N/U A 1.0 9.28
28+ 15 16 16 118 R/L/T/N A 1.0 11.78
29+ 16 16 16 118 R/L/T/U A 1.0 9.47
30+ 17 224 64 87 L/U/N/N A 1.0 285.09
31+ 18 224 64 87 L/U/N/U A 1.0 252.83
32+ 19 224 64 87 L/U/T/N A 1.0 299.39
33+ 20 224 64 87 L/U/T/U A 1.0 267.74
34+ 21 224 64 87 L/L/N/N A 1.0 310.46
35+ 22 224 64 87 L/L/N/U A 1.0 278.85
36+ 23 224 64 87 L/L/T/N A 1.0 284.80
37+ 24 224 64 87 L/L/T/U A 1.0 251.42
38+ 25 224 64 87 R/U/N/N A 1.0 97.41
39+ 26 224 64 87 R/U/N/U A 1.0 86.08
40+ 27 224 64 87 R/U/T/N A 1.0 97.28
41+ 28 224 64 87 R/U/T/U A 1.0 86.08
42+ 29 224 64 87 R/L/N/N A 1.0 97.18
43+ 30 224 64 87 R/L/N/U A 1.0 86.05
44+ 31 224 64 87 R/L/T/N A 1.0 97.57
45+ 32 224 64 87 R/L/T/U A 1.0 86.30
46+ 33 224 112 71 L/U/N/N A 1.0 342.27
47+ 34 224 112 71 L/U/N/U A 1.0 305.38
48+ 35 224 112 71 L/U/T/N A 1.0 364.80
49+ 36 224 112 71 L/U/T/U A 1.0 326.91
50+ 37 224 112 71 L/L/N/N A 1.0 375.71
51+ 38 224 112 71 L/L/N/U A 1.0 336.80
52+ 39 224 112 71 L/L/T/N A 1.0 346.62
53+ 40 224 112 71 L/L/T/U A 1.0 308.19
54+ 41 224 112 71 R/U/N/N A 1.0 188.83
55+ 42 224 112 71 R/U/N/U A 1.0 169.31
56+ 43 224 112 71 R/U/T/N A 1.0 178.14
57+ 44 224 112 71 R/U/T/U A 1.0 160.22
58+ 45 224 112 71 R/L/N/N A 1.0 178.88
59+ 46 224 112 71 R/L/N/U A 1.0 160.83
60+ 47 224 112 71 R/L/T/N A 1.0 188.45
61+ 48 224 112 71 R/L/T/U A 1.0 169.41
62+ 49 227 210 248 L/U/N/N N 1.0 1542.30
63+ 50 227 210 248 L/U/N/U N 1.0 1446.53
64+ 51 227 210 248 L/U/T/N N 1.0 1673.98
65+ 52 227 210 248 L/U/T/U N 1.0 1573.18
66+ 53 227 210 248 L/L/N/N N 1.0 1721.12
67+ 54 227 210 248 L/L/N/U N 1.0 1621.02
68+ 55 227 210 248 L/L/T/N N 1.0 1631.62
69+ 56 227 210 248 L/L/T/U N 1.0 1509.60
70+ 57 227 210 248 R/U/N/N N 1.0 1636.16
71+ 58 227 210 248 R/U/N/U N 1.0 1559.14
72+ 59 227 210 248 R/U/T/N N 1.0 1492.74
73+ 60 227 210 248 R/U/T/U N 1.0 1417.34
74+ 61 227 210 248 R/L/N/N N 1.0 1508.67
75+ 62 227 210 248 R/L/N/U N 1.0 1434.37
76+ 63 227 210 248 R/L/T/N N 1.0 1628.16
77+ 64 227 210 248 R/L/T/U N 1.0 1551.46
78+ 65 184 212 10 L/U/N/N N 1.0 165.95
79+ 66 184 212 10 L/U/N/U N 1.0 138.91
80+ 67 184 212 10 L/U/T/N N 1.0 177.25
81+ 68 184 212 10 L/U/T/U N 1.0 152.26
82+ 69 184 212 10 L/L/N/N N 1.0 187.55
83+ 70 184 212 10 L/L/N/U N 1.0 161.54
84+ 71 184 212 10 L/L/T/N N 1.0 164.13
85+ 72 184 212 10 L/L/T/U N 1.0 139.68
86+ 73 184 212 10 R/U/N/N N 1.0 215.39
87+ 74 184 212 10 R/U/N/U N 1.0 187.01
88+ 75 184 212 10 R/U/T/N N 1.0 203.87
89+ 76 184 212 10 R/U/T/U N 1.0 174.50
90+ 77 184 212 10 R/L/N/N N 1.0 207.94
91+ 78 184 212 10 R/L/N/U N 1.0 177.02
92+ 79 184 212 10 R/L/T/N N 1.0 216.64
93+ 80 184 212 10 R/L/T/U N 1.0 186.11
94+ 81 896 512 1 L/U/N/N A 1.0 865.22
95+ 82 896 512 1 L/U/N/U A 1.0 743.07
96+ 83 896 512 1 L/U/T/N A 1.0 895.01
97+ 84 896 512 1 L/U/T/U A 1.0 772.00
98+ 85 896 512 1 L/L/N/N A 1.0 940.64
99+ 86 896 512 1 L/L/N/U A 1.0 819.97
100+ 87 896 512 1 L/L/T/N A 1.0 857.38
101+ 88 896 512 1 L/L/T/U A 1.0 734.14
102+ 89 896 512 1 R/U/N/N A 1.0 471.39
103+ 90 896 512 1 R/U/N/U A 1.0 403.78
104+ 91 896 512 1 R/U/T/N A 1.0 471.62
105+ 92 896 512 1 R/U/T/U A 1.0 403.10
106+ 93 896 512 1 R/L/N/N A 1.0 473.34
107+ 94 896 512 1 R/L/N/U A 1.0 399.94
108+ 95 896 512 1 R/L/T/N A 1.0 477.44
109+ 96 896 512 1 R/L/T/U A 1.0 404.61
110+ 97 464 896 124 L/U/N/N A 1.0 8589.06
111+ 98 464 896 124 L/U/N/U A 1.0 8243.23
112+ 99 464 896 124 L/U/T/N A 1.0 9506.37
113+100 464 896 124 L/U/T/U A 1.0 9214.62
114+101 464 896 124 L/L/N/N A 1.0 9542.43
115+102 464 896 124 L/L/N/U A 1.0 9234.88
116+103 464 896 124 L/L/T/N A 1.0 8825.70
117+104 464 896 124 L/L/T/U A 1.0 8404.32
118+105 464 896 124 R/U/N/N A 1.0 14619.07
119+106 464 896 124 R/U/N/U A 1.0 14401.86
120+107 464 896 124 R/U/T/N A 1.0 14479.42
121+108 464 896 124 R/U/T/U A 1.0 14122.98
122+109 464 896 124 R/L/N/N A 1.0 14513.22
123+110 464 896 124 R/L/N/U A 1.0 14307.62
124+111 464 896 124 R/L/T/N A 1.0 14805.22
125+112 464 896 124 R/L/T/U A 1.0 14351.30
126+113 441 345 82 L/U/N/N N 1.0 2420.45
127+114 441 345 82 L/U/N/U N 1.0 2303.07
128+115 441 345 82 L/U/T/N N 1.0 2419.55
129+116 441 345 82 L/U/T/U N 1.0 2280.96
130+117 441 345 82 L/L/N/N N 1.0 2481.50
131+118 441 345 82 L/L/N/U N 1.0 2364.90
132+119 441 345 82 L/L/T/N N 1.0 2467.10
133+120 441 345 82 L/L/T/U N 1.0 2323.62
134+121 441 345 82 R/U/N/N N 1.0 2128.64
135+122 441 345 82 R/U/N/U N 1.0 2041.44
136+123 441 345 82 R/U/T/N N 1.0 1998.56
137+124 441 345 82 R/U/T/U N 1.0 1904.48
138+125 441 345 82 R/L/N/N N 1.0 2015.87
139+126 441 345 82 R/L/N/U N 1.0 1917.54
140+127 441 345 82 R/L/T/N N 1.0 2119.20
141+128 441 345 82 R/L/T/U N 1.0 2035.58
142+129 425 999 22 L/U/N/N N 1.0 1876.32
143+130 425 999 22 L/U/N/U N 1.0 1770.11
144+131 425 999 22 L/U/T/N N 1.0 1960.51
145+132 425 999 22 L/U/T/U N 1.0 1843.07
146+133 425 999 22 L/L/N/N N 1.0 1992.06
147+134 425 999 22 L/L/N/U N 1.0 1878.27
148+135 425 999 22 L/L/T/N N 1.0 1954.27
149+136 425 999 22 L/L/T/U N 1.0 1822.66
150+137 425 999 22 R/U/N/N N 1.0 3449.76
151+138 425 999 22 R/U/N/U N 1.0 3312.48
152+139 425 999 22 R/U/T/N N 1.0 3321.57
153+140 425 999 22 R/U/T/U N 1.0 3181.70
154+141 425 999 22 R/L/N/N N 1.0 3336.22
155+142 425 999 22 R/L/N/U N 1.0 3192.70
156+143 425 999 22 R/L/T/N N 1.0 3449.15
157+144 425 999 22 R/L/T/U N 1.0 3301.15
158+145 529 1005 103 L/U/N/N N 1.0 10673.82
159+146 529 1005 103 L/U/N/U N 1.0 10375.90
160+147 529 1005 103 L/U/T/N N 1.0 10623.97
161+148 529 1005 103 L/U/T/U N 1.0 10417.50
162+149 529 1005 103 L/L/N/N N 1.0 10769.31
163+150 529 1005 103 L/L/N/U N 1.0 10536.93
164+151 529 1005 103 L/L/T/N N 1.0 10774.34
165+152 529 1005 103 L/L/T/U N 1.0 10451.23
166+153 529 1005 103 R/U/N/N N 1.0 17596.87
167+154 529 1005 103 R/U/N/U N 1.0 17370.11
168+155 529 1005 103 R/U/T/N N 1.0 16902.18
169+156 529 1005 103 R/U/T/U N 1.0 16689.47
170+157 529 1005 103 R/L/N/N N 1.0 16971.46
171+158 529 1005 103 R/L/N/U N 1.0 16820.61
172+159 529 1005 103 R/L/T/N N 1.0 17355.10
173+160 529 1005 103 R/L/T/U N 1.0 17318.85
174+161 3696 2240 12 L/U/N/N A 1.0 85648.55
175+162 3696 2240 12 L/U/N/U A 1.0 84656.73
176+163 3696 2240 12 L/U/T/N A 1.0 86929.63
177+164 3696 2240 12 L/U/T/U A 1.0 85933.25
178+165 3696 2240 12 L/L/N/N A 1.0 86920.39
179+166 3696 2240 12 L/L/N/U A 1.0 85809.28
180+167 3696 2240 12 L/L/T/N A 1.0 85672.64
181+168 3696 2240 12 L/L/T/U A 1.0 84589.66
182+169 3696 2240 12 R/U/N/N A 1.0 55261.25
183+170 3696 2240 12 R/U/N/U A 1.0 55712.99
184+171 3696 2240 12 R/U/T/N A 1.0 54912.93
185+172 3696 2240 12 R/U/T/U A 1.0 55213.66
186+173 3696 2240 12 R/L/N/N A 1.0 54937.92
187+174 3696 2240 12 R/L/N/U A 1.0 55848.32
188+175 3696 2240 12 R/L/T/N A 1.0 55224.99
189+176 3696 2240 12 R/L/T/U A 1.0 55633.92
190+177 2048 3056 10 L/U/N/N A 1.0 32265.89
191+178 2048 3056 10 L/U/N/U A 1.0 31689.50
192+179 2048 3056 10 L/U/T/N A 1.0 32545.57
193+180 2048 3056 10 L/U/T/U A 1.0 32237.60
194+181 2048 3056 10 L/L/N/N A 1.0 32572.48
195+182 2048 3056 10 L/L/N/U A 1.0 31961.09
196+183 2048 3056 10 L/L/T/N A 1.0 32625.28
197+184 2048 3056 10 L/L/T/U A 1.0 31882.78
198+185 2048 3056 10 R/U/N/N A 1.0 47118.34
199+186 2048 3056 10 R/U/N/U A 1.0 47349.31
200+187 2048 3056 10 R/U/T/N A 1.0 46164.19
201+188 2048 3056 10 R/U/T/U A 1.0 46005.15
202+189 2048 3056 10 R/L/N/N A 1.0 46353.95
203+190 2048 3056 10 R/L/N/U A 1.0 46460.64
204+191 2048 3056 10 R/L/T/N A 1.0 46901.60
205+192 2048 3056 10 R/L/T/U A 1.0 47032.19
206+193 3488 3984 9 L/U/N/N A 1.0 102527.39
207+194 3488 3984 9 L/U/N/U A 1.0 102334.84
208+195 3488 3984 9 L/U/T/N A 1.0 104808.80
209+196 3488 3984 9 L/U/T/U A 1.0 104526.69
210+197 3488 3984 9 L/L/N/N A 1.0 104968.73
211+198 3488 3984 9 L/L/N/U A 1.0 104753.92
212+199 3488 3984 9 L/L/T/N A 1.0 103577.12
213+200 3488 3984 9 L/L/T/U A 1.0 104086.05
214+201 3488 3984 9 R/U/N/N A 1.0 120291.20
215+202 3488 3984 9 R/U/N/U A 1.0 122157.89
216+203 3488 3984 9 R/U/T/N A 1.0 116856.28
217+204 3488 3984 9 R/U/T/U A 1.0 118841.98
218+205 3488 3984 9 R/L/N/N A 1.0 117833.80
219+206 3488 3984 9 R/L/N/U A 1.0 119511.97
220+207 3488 3984 9 R/L/T/N A 1.0 120376.45
221+208 3488 3984 9 R/L/T/U A 1.0 121777.77
222+209 1667 1941 38 L/U/N/N N 1.0 53272.04
223+210 1667 1941 38 L/U/N/U N 1.0 53484.42
224+211 1667 1941 38 L/U/T/N N 1.0 56210.95
225+212 1667 1941 38 L/U/T/U N 1.0 56814.66
226+213 1667 1941 38 L/L/N/N N 1.0 56410.79
227+214 1667 1941 38 L/L/N/U N 1.0 56884.64
228+215 1667 1941 38 L/L/T/N N 1.0 53830.72
229+216 1667 1941 38 L/L/T/U N 1.0 54005.79
230+217 1667 1941 38 R/U/N/N N 1.0 64610.88
231+218 1667 1941 38 R/U/N/U N 1.0 65375.65
232+219 1667 1941 38 R/U/T/N N 1.0 61704.00
233+220 1667 1941 38 R/U/T/U N 1.0 62640.89
234+221 1667 1941 38 R/L/N/N N 1.0 62579.90
235+222 1667 1941 38 R/L/N/U N 1.0 63028.45
236+223 1667 1941 38 R/L/T/N N 1.0 63981.18
237+224 1667 1941 38 R/L/T/U N 1.0 65156.29
238+225 3516 1213 16 L/U/N/N N 1.0 56914.18
239+226 3516 1213 16 L/U/N/U N 1.0 56911.55
240+227 3516 1213 16 L/U/T/N N 1.0 57184.26
241+228 3516 1213 16 L/U/T/U N 1.0 57743.90
242+229 3516 1213 16 L/L/N/N N 1.0 56918.59
243+230 3516 1213 16 L/L/N/U N 1.0 56856.70
244+231 3516 1213 16 L/L/T/N N 1.0 58438.56
245+232 3516 1213 16 L/L/T/U N 1.0 58054.75
246+233 3516 1213 16 R/U/N/N N 1.0 23256.61
247+234 3516 1213 16 R/U/N/U N 1.0 23204.54
248+235 3516 1213 16 R/U/T/N N 1.0 22865.31
249+236 3516 1213 16 R/U/T/U N 1.0 23085.50
250+237 3516 1213 16 R/L/N/N N 1.0 22929.92
251+238 3516 1213 16 R/L/N/U N 1.0 23164.13
252+239 3516 1213 16 R/L/T/N N 1.0 22955.01
253+240 3516 1213 16 R/L/T/U N 1.0 23150.85
254+241 4848 5456 4 L/U/N/N A 1.0 118500.10
255+242 4848 5456 4 L/U/N/U A 1.0 118494.50
256+243 4848 5456 4 L/U/T/N A 1.0 120393.02
257+244 4848 5456 4 L/U/T/U A 1.0 121084.89
258+245 4848 5456 4 L/L/N/N A 1.0 120566.66
259+246 4848 5456 4 L/L/N/U A 1.0 121182.56
260+247 4848 5456 4 L/L/T/N A 1.0 119625.53
261+248 4848 5456 4 L/L/T/U A 1.0 120280.86
262+249 4848 5456 4 R/U/N/N A 1.0 135500.12
263+250 4848 5456 4 R/U/N/U A 1.0 138115.73
264+251 4848 5456 4 R/U/T/N A 1.0 132898.78
265+252 4848 5456 4 R/U/T/U A 1.0 135684.19
266+253 4848 5456 4 R/L/N/N A 1.0 134287.00
267+254 4848 5456 4 R/L/N/U A 1.0 136519.11
268+255 4848 5456 4 R/L/T/N A 1.0 135868.95
269+256 4848 5456 4 R/L/T/U A 1.0 138065.50
270+257 4512 4416 6 L/U/N/N A 1.0 125041.05
271+258 4512 4416 6 L/U/N/U A 1.0 126066.66
272+259 4512 4416 6 L/U/T/N A 1.0 127903.39
273+260 4512 4416 6 L/U/T/U A 1.0 129417.86
274+261 4512 4416 6 L/L/N/N A 1.0 127579.77
275+262 4512 4416 6 L/L/N/U A 1.0 128757.02
276+263 4512 4416 6 L/L/T/N A 1.0 125818.27
277+264 4512 4416 6 L/L/T/U A 1.0 127075.94
278+265 4512 4416 6 R/U/N/N A 1.0 126647.13
279+266 4512 4416 6 R/U/N/U A 1.0 129172.93
280+267 4512 4416 6 R/U/T/N A 1.0 125424.12
281+268 4512 4416 6 R/U/T/U A 1.0 127995.01
282+269 4512 4416 6 R/L/N/N A 1.0 126038.56
283+270 4512 4416 6 R/L/N/U A 1.0 128781.16
284+271 4512 4416 6 R/L/T/N A 1.0 125690.37
285+272 4512 4416 6 R/L/T/U A 1.0 128672.77
286+273 5266 5889 4 L/U/N/N N 1.0 152843.03
287+274 5266 5889 4 L/U/N/U N 1.0 154440.73
288+275 5266 5889 4 L/U/T/N N 1.0 156797.83
289+276 5266 5889 4 L/U/T/U N 1.0 159061.45
290+277 5266 5889 4 L/L/N/N N 1.0 156095.94
291+278 5266 5889 4 L/L/N/U N 1.0 157649.39
292+279 5266 5889 4 L/L/T/N N 1.0 155538.84
293+280 5266 5889 4 L/L/T/U N 1.0 157420.17
294+281 5266 5889 4 R/U/N/N N 1.0 178707.81
295+282 5266 5889 4 R/U/N/U N 1.0 182535.81
296+283 5266 5889 4 R/U/T/N N 1.0 174600.80
297+284 5266 5889 4 R/U/T/U N 1.0 177594.75
298+285 5266 5889 4 R/L/N/N N 1.0 175695.17
299+286 5266 5889 4 R/L/N/U N 1.0 179542.17
300+287 5266 5889 4 R/L/T/N N 1.0 178823.88
301+288 5266 5889 4 R/L/T/U N 1.0 182232.75
302+289 5133 5564 4 L/U/N/N N 1.0 138911.86
303+290 5133 5564 4 L/U/N/U N 1.0 140006.05
304+291 5133 5564 4 L/U/T/N N 1.0 141711.58
305+292 5133 5564 4 L/U/T/U N 1.0 143658.11
306+293 5133 5564 4 L/L/N/N N 1.0 141402.88
307+294 5133 5564 4 L/L/N/U N 1.0 142816.95
308+295 5133 5564 4 L/L/T/N N 1.0 141330.42
309+296 5133 5564 4 L/L/T/U N 1.0 142665.55
310+297 5133 5564 4 R/U/N/N N 1.0 151550.28
311+298 5133 5564 4 R/U/N/U N 1.0 155652.09
312+299 5133 5564 4 R/U/T/N N 1.0 150478.97
313+300 5133 5564 4 R/U/T/U N 1.0 154308.95
314+301 5133 5564 4 R/L/N/N N 1.0 155072.42
315+302 5133 5564 4 R/L/N/U N 1.0 158661.28
316+303 5133 5564 4 R/L/T/N N 1.0 156274.28
317+304 5133 5564 4 R/L/T/U N 1.0 157852.00
318+305 6182 6412 3 L/U/N/N N 1.0 180365.95
319+306 6182 6412 3 L/U/N/U N 1.0 185667.42
320+307 6182 6412 3 L/U/T/N N 1.0 185798.17
321+308 6182 6412 3 L/U/T/U N 1.0 193007.11
322+309 6182 6412 3 L/L/N/N N 1.0 185935.45
323+310 6182 6412 3 L/L/N/U N 1.0 190123.81
324+311 6182 6412 3 L/L/T/N N 1.0 192106.59
325+312 6182 6412 3 L/L/T/U N 1.0 189200.45
326+313 6182 6412 3 R/U/N/N N 1.0 201003.58
327+314 6182 6412 3 R/U/N/U N 1.0 214975.00
328+315 6182 6412 3 R/U/T/N N 1.0 198764.48
329+316 6182 6412 3 R/U/T/U N 1.0 204046.08
330+317 6182 6412 3 R/L/N/N N 1.0 201021.09
331+318 6182 6412 3 R/L/N/U N 1.0 207048.50
332+319 6182 6412 3 R/L/T/N N 1.0 202174.27
333+320 6182 6412 3 R/L/T/U N 1.0 209951.39
334+------------------------------------------------------
Aexperimental/aclblasCtrsmBatched2/docs/perf_report_320cases_new.md+437-0
@@ -0,0 +1,437 @@
1+# CtrsmBatched 性能报告（320 Case 全量）
2+ 
3+## 测试环境
4+ 
5+- **NPU**: Ascend910B4, 频率 1650MHz
6+- **GPU**: cuBLAS cublasCtrsmBatched (对标参考)
7+- **CANN**: 9.0.0-beta.2
8+- **Profiling**: msprof op, warm-up=3
9+ 
10+## 总览
11+ 
12+- 总用例数: 320
13+- GPU/NPU >= 1.0 (NPU更快): 109 个
14+- GPU/NPU 平均: 0.800
15+- GPU/NPU 最大: 1.776
16+- GPU/NPU 最小: 0.118
17+ 
18+## Case 1-16（m=16, n=16, batch=118）
19+ 
20+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
21+|---|------|---------|---------|--------|
22+| 1 | L/U/N/N | 64.96 | 21.22 | 0.3267 |
23+| 2 | L/U/N/U | 63.28 | 18.24 | 0.2882 |
24+| 3 | L/U/T/N | 71.52 | 11.97 | 0.1674 |
25+| 4 | L/U/T/U | 73.00 | 9.73 | 0.1333 |
26+| 5 | L/L/N/N | 64.62 | 12.26 | 0.1897 |
27+| 6 | L/L/N/U | 64.66 | 9.95 | 0.1539 |
28+| 7 | L/L/T/N | 72.02 | 21.47 | 0.2981 |
29+| 8 | L/L/T/U | 72.20 | 18.46 | 0.2557 |
30+| 9 | R/U/N/N | 71.08 | 11.65 | 0.1639 |
31+| 10 | R/U/N/U | 79.50 | 9.41 | 0.1184 |
32+| 11 | R/U/T/N | 60.06 | 11.52 | 0.1918 |
33+| 12 | R/U/T/U | 63.36 | 9.28 | 0.1465 |
34+| 13 | R/L/N/N | 70.34 | 11.52 | 0.1638 |
35+| 14 | R/L/N/U | 71.04 | 9.28 | 0.1306 |
36+| 15 | R/L/T/N | 67.04 | 11.78 | 0.1757 |
37+| 16 | R/L/T/U | 70.48 | 9.47 | 0.1344 |
38+ 
39+## Case 17-32（m=224, n=64, batch=87）
40+ 
41+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
42+|---|------|---------|---------|--------|
43+| 17 | L/U/N/N | 683.71 | 285.09 | 0.4170 |
44+| 18 | L/U/N/U | 645.29 | 252.83 | 0.3918 |
45+| 19 | L/U/T/N | 925.62 | 299.39 | 0.3234 |
46+| 20 | L/U/T/U | 878.92 | 267.74 | 0.3046 |
47+| 21 | L/L/N/N | 689.27 | 310.46 | 0.4504 |
48+| 22 | L/L/N/U | 653.27 | 278.85 | 0.4269 |
49+| 23 | L/L/T/N | 919.30 | 284.80 | 0.3098 |
50+| 24 | L/L/T/U | 880.40 | 251.42 | 0.2856 |
51+| 25 | R/U/N/N | 361.55 | 97.41 | 0.2694 |
52+| 26 | R/U/N/U | 348.69 | 86.08 | 0.2469 |
53+| 27 | R/U/T/N | 342.85 | 97.28 | 0.2837 |
54+| 28 | R/U/T/U | 324.99 | 86.08 | 0.2649 |
55+| 29 | R/L/N/N | 363.71 | 97.18 | 0.2672 |
56+| 30 | R/L/N/U | 346.55 | 86.05 | 0.2483 |
57+| 31 | R/L/T/N | 345.61 | 97.57 | 0.2823 |
58+| 32 | R/L/T/U | 332.67 | 86.30 | 0.2594 |
59+ 
60+## Case 33-48（m=224, n=112, batch=71）
61+ 
62+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
63+|---|------|---------|---------|--------|
64+| 33 | L/U/N/N | 835.62 | 342.27 | 0.4096 |
65+| 34 | L/U/N/U | 797.24 | 305.38 | 0.3830 |
66+| 35 | L/U/T/N | 977.58 | 364.80 | 0.3732 |
67+| 36 | L/U/T/U | 937.18 | 326.91 | 0.3488 |
68+| 37 | L/L/N/N | 822.94 | 375.71 | 0.4565 |
69+| 38 | L/L/N/U | 791.52 | 336.80 | 0.4255 |
70+| 39 | L/L/T/N | 956.14 | 346.62 | 0.3625 |
71+| 40 | L/L/T/U | 933.98 | 308.19 | 0.3300 |
72+| 41 | R/U/N/N | 441.23 | 188.83 | 0.4280 |
73+| 42 | R/U/N/U | 423.39 | 169.31 | 0.3999 |
74+| 43 | R/U/T/N | 402.49 | 178.14 | 0.4426 |
75+| 44 | R/U/T/U | 382.59 | 160.22 | 0.4188 |
76+| 45 | R/L/N/N | 432.55 | 178.88 | 0.4135 |
77+| 46 | R/L/N/U | 412.59 | 160.83 | 0.3898 |
78+| 47 | R/L/T/N | 411.69 | 188.45 | 0.4577 |
79+| 48 | R/L/T/U | 396.27 | 169.41 | 0.4275 |
80+ 
81+## Case 49-64（m=227, n=210, batch=248）
82+ 
83+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
84+|---|------|---------|---------|--------|
85+| 49 | L/U/N/N | 4811.02 | 1542.30 | 0.3206 |
86+| 50 | L/U/N/U | 4752.32 | 1446.53 | 0.3044 |
87+| 51 | L/U/T/N | 4775.48 | 1673.98 | 0.3505 |
88+| 52 | L/U/T/U | 4737.53 | 1573.18 | 0.3321 |
89+| 53 | L/L/N/N | 4727.49 | 1721.12 | 0.3641 |
90+| 54 | L/L/N/U | 4677.69 | 1621.02 | 0.3465 |
91+| 55 | L/L/T/N | 4875.74 | 1631.62 | 0.3346 |
92+| 56 | L/L/T/U | 4824.86 | 1509.60 | 0.3129 |
93+| 57 | R/U/N/N | 4892.30 | 1636.16 | 0.3344 |
94+| 58 | R/U/N/U | 4820.38 | 1559.14 | 0.3234 |
95+| 59 | R/U/T/N | 4925.66 | 1492.74 | 0.3031 |
96+| 60 | R/U/T/U | 4879.70 | 1417.34 | 0.2905 |
97+| 61 | R/L/N/N | 4976.40 | 1508.67 | 0.3032 |
98+| 62 | R/L/N/U | 4939.74 | 1434.37 | 0.2904 |
99+| 63 | R/L/T/N | 4837.76 | 1628.16 | 0.3366 |
100+| 64 | R/L/T/U | 4786.12 | 1551.46 | 0.3242 |
101+ 
102+## Case 65-80（m=184, n=212, batch=10）
103+ 
104+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
105+|---|------|---------|---------|--------|
106+| 65 | L/U/N/N | 213.14 | 165.95 | 0.7786 |
107+| 66 | L/U/N/U | 211.94 | 138.91 | 0.6554 |
108+| 67 | L/U/T/N | 216.24 | 177.25 | 0.8197 |
109+| 68 | L/U/T/U | 213.38 | 152.26 | 0.7135 |
110+| 69 | L/L/N/N | 212.56 | 187.55 | 0.8823 |
111+| 70 | L/L/N/U | 208.62 | 161.54 | 0.7743 |
112+| 71 | L/L/T/N | 216.62 | 164.13 | 0.7577 |
113+| 72 | L/L/T/U | 218.12 | 139.68 | 0.6404 |
114+| 73 | R/U/N/N | 259.71 | 215.39 | 0.8294 |
115+| 74 | R/U/N/U | 258.37 | 187.01 | 0.7238 |
116+| 75 | R/U/T/N | 260.59 | 203.87 | 0.7824 |
117+| 76 | R/U/T/U | 264.71 | 174.50 | 0.6592 |
118+| 77 | R/L/N/N | 261.13 | 207.94 | 0.7963 |
119+| 78 | R/L/N/U | 260.37 | 177.02 | 0.6799 |
120+| 79 | R/L/T/N | 264.05 | 216.64 | 0.8205 |
121+| 80 | R/L/T/U | 258.45 | 186.11 | 0.7201 |
122+ 
123+## Case 81-96（m=896, n=512, batch=1）
124+ 
125+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
126+|---|------|---------|---------|--------|
127+| 81 | L/U/N/N | 683.67 | 865.22 | 1.2655 |
128+| 82 | L/U/N/U | 637.43 | 743.07 | 1.1657 |
129+| 83 | L/U/T/N | 1035.02 | 895.01 | 0.8647 |
130+| 84 | L/U/T/U | 976.12 | 772.00 | 0.7909 |
131+| 85 | L/L/N/N | 692.23 | 940.64 | 1.3588 |
132+| 86 | L/L/N/U | 640.89 | 819.97 | 1.2794 |
133+| 87 | L/L/T/N | 1025.70 | 857.38 | 0.8359 |
134+| 88 | L/L/T/U | 992.24 | 734.14 | 0.7399 |
135+| 89 | R/U/N/N | 514.25 | 471.39 | 0.9167 |
136+| 90 | R/U/N/U | 493.61 | 403.78 | 0.8180 |
137+| 91 | R/U/T/N | 395.97 | 471.62 | 1.1911 |
138+| 92 | R/U/T/U | 371.97 | 403.10 | 1.0837 |
139+| 93 | R/L/N/N | 503.01 | 473.34 | 0.9410 |
140+| 94 | R/L/N/U | 484.37 | 399.94 | 0.8257 |
141+| 95 | R/L/T/N | 397.45 | 477.44 | 1.2013 |
142+| 96 | R/L/T/U | 371.27 | 404.61 | 1.0898 |
143+ 
144+## Case 97-112（m=464, n=896, batch=124）
145+ 
146+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
147+|---|------|---------|---------|--------|
148+| 97 | L/U/N/N | 8759.88 | 8589.06 | 0.9805 |
149+| 98 | L/U/N/U | 8277.81 | 8243.23 | 0.9958 |
150+| 99 | L/U/T/N | 10024.32 | 9506.37 | 0.9483 |
151+| 100 | L/U/T/U | 9635.69 | 9214.62 | 0.9563 |
152+| 101 | L/L/N/N | 8810.82 | 9542.43 | 1.0830 |
153+| 102 | L/L/N/U | 8480.23 | 9234.88 | 1.0890 |
154+| 103 | L/L/T/N | 9827.86 | 8825.70 | 0.8980 |
155+| 104 | L/L/T/U | 9486.99 | 8404.32 | 0.8859 |
156+| 105 | R/U/N/N | 23799.22 | 14619.07 | 0.6143 |
157+| 106 | R/U/N/U | 23524.83 | 14401.86 | 0.6122 |
158+| 107 | R/U/T/N | 19042.10 | 14479.42 | 0.7604 |
159+| 108 | R/U/T/U | 18927.24 | 14122.98 | 0.7462 |
160+| 109 | R/L/N/N | 23522.09 | 14513.22 | 0.6170 |
161+| 110 | R/L/N/U | 23270.43 | 14307.62 | 0.6148 |
162+| 111 | R/L/T/N | 19356.05 | 14805.22 | 0.7649 |
163+| 112 | R/L/T/U | 19231.74 | 14351.30 | 0.7462 |
164+ 
165+## Case 113-128（m=441, n=345, batch=82）
166+ 
167+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
168+|---|------|---------|---------|--------|
169+| 113 | L/U/N/N | 3613.95 | 2420.45 | 0.6698 |
170+| 114 | L/U/N/U | 3585.35 | 2303.07 | 0.6424 |
171+| 115 | L/U/T/N | 3829.46 | 2419.55 | 0.6318 |
172+| 116 | L/U/T/U | 3798.28 | 2280.96 | 0.6005 |
173+| 117 | L/L/N/N | 3640.13 | 2481.50 | 0.6817 |
174+| 118 | L/L/N/U | 3580.63 | 2364.90 | 0.6605 |
175+| 119 | L/L/T/N | 3869.54 | 2467.10 | 0.6376 |
176+| 120 | L/L/T/U | 3798.22 | 2323.62 | 0.6118 |
177+| 121 | R/U/N/N | 4335.09 | 2128.64 | 0.4910 |
178+| 122 | R/U/N/U | 4297.27 | 2041.44 | 0.4751 |
179+| 123 | R/U/T/N | 4158.42 | 1998.56 | 0.4806 |
180+| 124 | R/U/T/U | 4140.56 | 1904.48 | 0.4600 |
181+| 125 | R/L/N/N | 4353.01 | 2015.87 | 0.4631 |
182+| 126 | R/L/N/U | 4302.55 | 1917.54 | 0.4457 |
183+| 127 | R/L/T/N | 4160.42 | 2119.20 | 0.5094 |
184+| 128 | R/L/T/U | 4115.16 | 2035.58 | 0.4947 |
185+ 
186+## Case 129-144（m=425, n=999, batch=22）
187+ 
188+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
189+|---|------|---------|---------|--------|
190+| 129 | L/U/N/N | 4271.73 | 1876.32 | 0.4392 |
191+| 130 | L/U/N/U | 4173.84 | 1770.11 | 0.4241 |
192+| 131 | L/U/T/N | 4309.29 | 1960.51 | 0.4550 |
193+| 132 | L/U/T/U | 4296.35 | 1843.07 | 0.4290 |
194+| 133 | L/L/N/N | 4302.99 | 1992.06 | 0.4629 |
195+| 134 | L/L/N/U | 4206.58 | 1878.27 | 0.4465 |
196+| 135 | L/L/T/N | 4483.21 | 1954.27 | 0.4359 |
197+| 136 | L/L/T/U | 4298.03 | 1822.66 | 0.4241 |
198+| 137 | R/U/N/N | 7125.36 | 3449.76 | 0.4842 |
199+| 138 | R/U/N/U | 7141.76 | 3312.48 | 0.4638 |
200+| 139 | R/U/T/N | 7013.26 | 3321.57 | 0.4736 |
201+| 140 | R/U/T/U | 7023.56 | 3181.70 | 0.4530 |
202+| 141 | R/L/N/N | 7155.60 | 3336.22 | 0.4662 |
203+| 142 | R/L/N/U | 7116.20 | 3192.70 | 0.4487 |
204+| 143 | R/L/T/N | 7046.22 | 3449.15 | 0.4895 |
205+| 144 | R/L/T/U | 7093.70 | 3301.15 | 0.4654 |
206+ 
207+## Case 145-160（m=529, n=1005, batch=103）
208+ 
209+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
210+|---|------|---------|---------|--------|
211+| 145 | L/U/N/N | 19050.24 | 10673.82 | 0.5603 |
212+| 146 | L/U/N/U | 18436.57 | 10375.90 | 0.5628 |
213+| 147 | L/U/T/N | 19159.42 | 10623.97 | 0.5545 |
214+| 148 | L/U/T/U | 18417.89 | 10417.50 | 0.5656 |
215+| 149 | L/L/N/N | 18985.74 | 10769.31 | 0.5672 |
216+| 150 | L/L/N/U | 18710.07 | 10536.93 | 0.5632 |
217+| 151 | L/L/T/N | 19025.66 | 10774.34 | 0.5663 |
218+| 152 | L/L/T/U | 19062.54 | 10451.23 | 0.5483 |
219+| 153 | R/U/N/N | 31525.31 | 17596.87 | 0.5582 |
220+| 154 | R/U/N/U | 31695.25 | 17370.11 | 0.5480 |
221+| 155 | R/U/T/N | 30567.01 | 16902.18 | 0.5530 |
222+| 156 | R/U/T/U | 30659.99 | 16689.47 | 0.5443 |
223+| 157 | R/L/N/N | 31533.33 | 16971.46 | 0.5382 |
224+| 158 | R/L/N/U | 31140.62 | 16820.61 | 0.5402 |
225+| 159 | R/L/T/N | 30774.94 | 17355.10 | 0.5639 |
226+| 160 | R/L/T/U | 30830.52 | 17318.85 | 0.5617 |
227+ 
228+## Case 161-176（m=3696, n=2240, batch=12）
229+ 
230+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
231+|---|------|---------|---------|--------|
232+| 161 | L/U/N/N | 108705.04 | 85648.55 | 0.7879 |
233+| 162 | L/U/N/U | 108400.92 | 84656.73 | 0.7810 |
234+| 163 | L/U/T/N | 116389.91 | 86929.63 | 0.7469 |
235+| 164 | L/U/T/U | 116128.60 | 85933.25 | 0.7400 |
236+| 165 | L/L/N/N | 108674.97 | 86920.39 | 0.7998 |
237+| 166 | L/L/N/U | 108279.69 | 85809.28 | 0.7925 |
238+| 167 | L/L/T/N | 116851.12 | 85672.64 | 0.7332 |
239+| 168 | L/L/T/U | 116539.83 | 84589.66 | 0.7258 |
240+| 169 | R/U/N/N | 83776.91 | 55261.25 | 0.6596 |
241+| 170 | R/U/N/U | 83865.24 | 55712.99 | 0.6643 |
242+| 171 | R/U/T/N | 80450.39 | 54912.93 | 0.6826 |
243+| 172 | R/U/T/U | 80031.18 | 55213.66 | 0.6899 |
244+| 173 | R/L/N/N | 83469.43 | 54937.92 | 0.6582 |
245+| 174 | R/L/N/U | 82893.95 | 55848.32 | 0.6737 |
246+| 175 | R/L/T/N | 80235.48 | 55224.99 | 0.6883 |
247+| 176 | R/L/T/U | 80503.53 | 55633.92 | 0.6911 |
248+ 
249+## Case 177-192（m=2048, n=3056, batch=10）
250+ 
251+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
252+|---|------|---------|---------|--------|
253+| 177 | L/U/N/N | 29891.98 | 32265.89 | 1.0794 |
254+| 178 | L/U/N/U | 30037.66 | 31689.50 | 1.0550 |
255+| 179 | L/U/T/N | 32874.24 | 32545.57 | 0.9900 |
256+| 180 | L/U/T/U | 32515.45 | 32237.60 | 0.9915 |
257+| 181 | L/L/N/N | 30155.08 | 32572.48 | 1.0802 |
258+| 182 | L/L/N/U | 30605.03 | 31961.09 | 1.0443 |
259+| 183 | L/L/T/N | 33065.76 | 32625.28 | 0.9867 |
260+| 184 | L/L/T/U | 33206.52 | 31882.78 | 0.9601 |
261+| 185 | R/U/N/N | 40298.14 | 47118.34 | 1.1692 |
262+| 186 | R/U/N/U | 39908.06 | 47349.31 | 1.1865 |
263+| 187 | R/U/T/N | 38696.65 | 46164.19 | 1.1930 |
264+| 188 | R/U/T/U | 39369.29 | 46005.15 | 1.1686 |
265+| 189 | R/L/N/N | 39527.49 | 46353.95 | 1.1727 |
266+| 190 | R/L/N/U | 40951.90 | 46460.64 | 1.1345 |
267+| 191 | R/L/T/N | 38748.84 | 46901.60 | 1.2104 |
268+| 192 | R/L/T/U | 39428.35 | 47032.19 | 1.1929 |
269+ 
270+## Case 193-208（m=3488, n=3984, batch=9）
271+ 
272+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
273+|---|------|---------|---------|--------|
274+| 193 | L/U/N/N | 93007.46 | 102527.39 | 1.1024 |
275+| 194 | L/U/N/U | 92805.36 | 102334.84 | 1.1027 |
276+| 195 | L/U/T/N | 99214.47 | 104808.80 | 1.0564 |
277+| 196 | L/U/T/U | 97832.26 | 104526.69 | 1.0684 |
278+| 197 | L/L/N/N | 92915.52 | 104968.73 | 1.1297 |
279+| 198 | L/L/N/U | 93123.27 | 104753.92 | 1.1249 |
280+| 199 | L/L/T/N | 98907.30 | 103577.12 | 1.0472 |
281+| 200 | L/L/T/U | 97859.53 | 104086.05 | 1.0636 |
282+| 201 | R/U/N/N | 116736.07 | 120291.20 | 1.0305 |
283+| 202 | R/U/N/U | 116631.91 | 122157.89 | 1.0474 |
284+| 203 | R/U/T/N | 109762.22 | 116856.28 | 1.0646 |
285+| 204 | R/U/T/U | 109452.91 | 118841.98 | 1.0858 |
286+| 205 | R/L/N/N | 117081.55 | 117833.80 | 1.0064 |
287+| 206 | R/L/N/U | 118062.12 | 119511.97 | 1.0123 |
288+| 207 | R/L/T/N | 110903.90 | 120376.45 | 1.0854 |
289+| 208 | R/L/T/U | 110267.28 | 121777.77 | 1.1044 |
290+ 
291+## Case 209-224（m=1667, n=1941, batch=38）
292+ 
293+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
294+|---|------|---------|---------|--------|
295+| 209 | L/U/N/N | 79000.72 | 53272.04 | 0.6743 |
296+| 210 | L/U/N/U | 78737.73 | 53484.42 | 0.6793 |
297+| 211 | L/U/T/N | 79340.84 | 56210.95 | 0.7085 |
298+| 212 | L/U/T/U | 79296.59 | 56814.66 | 0.7165 |
299+| 213 | L/L/N/N | 79233.56 | 56410.79 | 0.7120 |
300+| 214 | L/L/N/U | 78596.79 | 56884.64 | 0.7238 |
301+| 215 | L/L/T/N | 79604.02 | 53830.72 | 0.6762 |
302+| 216 | L/L/T/U | 79325.41 | 54005.79 | 0.6808 |
303+| 217 | R/U/N/N | 88158.83 | 64610.88 | 0.7329 |
304+| 218 | R/U/N/U | 87163.38 | 65375.65 | 0.7500 |
305+| 219 | R/U/T/N | 87104.15 | 61704.00 | 0.7084 |
306+| 220 | R/U/T/U | 86926.27 | 62640.89 | 0.7206 |
307+| 221 | R/L/N/N | 87710.23 | 62579.90 | 0.7135 |
308+| 222 | R/L/N/U | 87472.35 | 63028.45 | 0.7206 |
309+| 223 | R/L/T/N | 86978.04 | 63981.18 | 0.7356 |
310+| 224 | R/L/T/U | 86819.44 | 65156.29 | 0.7505 |
311+ 
312+## Case 225-240（m=3516, n=1213, batch=16）
313+ 
314+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
315+|---|------|---------|---------|--------|
316+| 225 | L/U/N/N | 62293.62 | 56914.18 | 0.9136 |
317+| 226 | L/U/N/U | 62002.94 | 56911.55 | 0.9179 |
318+| 227 | L/U/T/N | 63726.60 | 57184.26 | 0.8973 |
319+| 228 | L/U/T/U | 63304.05 | 57743.90 | 0.9122 |
320+| 229 | L/L/N/N | 62338.63 | 56918.59 | 0.9131 |
321+| 230 | L/L/N/U | 61966.52 | 56856.70 | 0.9175 |
322+| 231 | L/L/T/N | 63607.03 | 58438.56 | 0.9187 |
323+| 232 | L/L/T/U | 63238.64 | 58054.75 | 0.9180 |
324+| 233 | R/U/N/N | 38153.45 | 23256.61 | 0.6096 |
325+| 234 | R/U/N/U | 37913.22 | 23204.54 | 0.6120 |
326+| 235 | R/U/T/N | 37835.94 | 22865.31 | 0.6043 |
327+| 236 | R/U/T/U | 37593.61 | 23085.50 | 0.6141 |
328+| 237 | R/L/N/N | 38182.57 | 22929.92 | 0.6005 |
329+| 238 | R/L/N/U | 37845.76 | 23164.13 | 0.6121 |
330+| 239 | R/L/T/N | 38033.00 | 22955.01 | 0.6036 |
331+| 240 | R/L/T/U | 37780.66 | 23150.85 | 0.6128 |
332+ 
333+## Case 241-256（m=4848, n=5456, batch=4）
334+ 
335+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
336+|---|------|---------|---------|--------|
337+| 241 | L/U/N/N | 83442.57 | 118500.10 | 1.4201 |
338+| 242 | L/U/N/U | 83031.56 | 118494.50 | 1.4271 |
339+| 243 | L/U/T/N | 94877.18 | 120393.02 | 1.2689 |
340+| 244 | L/U/T/U | 93521.21 | 121084.89 | 1.2947 |
341+| 245 | L/L/N/N | 83689.22 | 120566.66 | 1.4406 |
342+| 246 | L/L/N/U | 82815.76 | 121182.56 | 1.4633 |
343+| 247 | L/L/T/N | 94624.47 | 119625.53 | 1.2642 |
344+| 248 | L/L/T/U | 94035.18 | 120280.86 | 1.2791 |
345+| 249 | R/U/N/N | 125892.02 | 135500.12 | 1.0763 |
346+| 250 | R/U/N/U | 123677.90 | 138115.73 | 1.1167 |
347+| 251 | R/U/T/N | 111958.86 | 132898.78 | 1.1870 |
348+| 252 | R/U/T/U | 111020.67 | 135684.19 | 1.2222 |
349+| 253 | R/L/N/N | 123240.63 | 134287.00 | 1.0896 |
350+| 254 | R/L/N/U | 121098.24 | 136519.11 | 1.1273 |
351+| 255 | R/L/T/N | 112001.80 | 135868.95 | 1.2131 |
352+| 256 | R/L/T/U | 111814.04 | 138065.50 | 1.2348 |
353+ 
354+## Case 257-272（m=4512, n=4416, batch=6）
355+ 
356+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
357+|---|------|---------|---------|--------|
358+| 257 | L/U/N/N | 73219.48 | 125041.05 | 1.7078 |
359+| 258 | L/U/N/U | 71267.91 | 126066.66 | 1.7689 |
360+| 259 | L/U/T/N | 82676.27 | 127903.39 | 1.5470 |
361+| 260 | L/U/T/U | 80556.41 | 129417.86 | 1.6065 |
362+| 261 | L/L/N/N | 75692.72 | 127579.77 | 1.6855 |
363+| 262 | L/L/N/U | 72496.25 | 128757.02 | 1.7761 |
364+| 263 | L/L/T/N | 81572.05 | 125818.27 | 1.5424 |
365+| 264 | L/L/T/U | 80581.05 | 127075.94 | 1.5770 |
366+| 265 | R/U/N/N | 97407.85 | 126647.13 | 1.3002 |
367+| 266 | R/U/N/U | 95839.46 | 129172.93 | 1.3478 |
368+| 267 | R/U/T/N | 89266.02 | 125424.12 | 1.4051 |
369+| 268 | R/U/T/U | 88244.97 | 127995.01 | 1.4505 |
370+| 269 | R/L/N/N | 96906.80 | 126038.56 | 1.3006 |
371+| 270 | R/L/N/U | 97011.48 | 128781.16 | 1.3275 |
372+| 271 | R/L/T/N | 89448.11 | 125690.37 | 1.4052 |
373+| 272 | R/L/T/U | 87919.59 | 128672.77 | 1.4635 |
374+ 
375+## Case 273-288（m=5266, n=5889, batch=4）
376+ 
377+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
378+|---|------|---------|---------|--------|
379+| 273 | L/U/N/N | 153232.61 | 152843.03 | 0.9975 |
380+| 274 | L/U/N/U | 152707.30 | 154440.73 | 1.0114 |
381+| 275 | L/U/T/N | 159815.80 | 156797.83 | 0.9811 |
382+| 276 | L/U/T/U | 158290.69 | 159061.45 | 1.0049 |
383+| 277 | L/L/N/N | 154259.88 | 156095.94 | 1.0119 |
384+| 278 | L/L/N/U | 153038.80 | 157649.39 | 1.0301 |
385+| 279 | L/L/T/N | 158804.88 | 155538.84 | 0.9794 |
386+| 280 | L/L/T/U | 157511.81 | 157420.17 | 0.9994 |
387+| 281 | R/U/N/N | 188856.52 | 178707.81 | 0.9463 |
388+| 282 | R/U/N/U | 179054.98 | 182535.81 | 1.0194 |
389+| 283 | R/U/T/N | 179061.14 | 174600.80 | 0.9751 |
390+| 284 | R/U/T/U | 176154.28 | 177594.75 | 1.0082 |
391+| 285 | R/L/N/N | 185246.12 | 175695.17 | 0.9484 |
392+| 286 | R/L/N/U | 182008.16 | 179542.17 | 0.9865 |
393+| 287 | R/L/T/N | 179401.34 | 178823.88 | 0.9968 |
394+| 288 | R/L/T/U | 177440.86 | 182232.75 | 1.0270 |
395+ 
396+## Case 289-304（m=5133, n=5564, batch=4）
397+ 
398+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
399+|---|------|---------|---------|--------|
400+| 289 | L/U/N/N | 99430.16 | 138911.86 | 1.3971 |
401+| 290 | L/U/N/U | 99111.10 | 140006.05 | 1.4126 |
402+| 291 | L/U/T/N | 104059.66 | 141711.58 | 1.3618 |
403+| 292 | L/U/T/U | 103931.62 | 143658.11 | 1.3822 |
404+| 293 | L/L/N/N | 99614.95 | 141402.88 | 1.4195 |
405+| 294 | L/L/N/U | 99115.46 | 142816.95 | 1.4409 |
406+| 295 | L/L/T/N | 104259.80 | 141330.42 | 1.3556 |
407+| 296 | L/L/T/U | 104770.75 | 142665.55 | 1.3617 |
408+| 297 | R/U/N/N | 133338.41 | 151550.28 | 1.1366 |
409+| 298 | R/U/N/U | 131676.81 | 155652.09 | 1.1821 |
410+| 299 | R/U/T/N | 128973.58 | 150478.97 | 1.1667 |
411+| 300 | R/U/T/U | 126657.95 | 154308.95 | 1.2183 |
412+| 301 | R/L/N/N | 131687.80 | 155072.42 | 1.1776 |
413+| 302 | R/L/N/U | 132118.81 | 158661.28 | 1.2009 |
414+| 303 | R/L/T/N | 129066.41 | 156274.28 | 1.2108 |
415+| 304 | R/L/T/U | 127261.64 | 157852.00 | 1.2404 |
416+ 
417+## Case 305-320（m=6182, n=6412, batch=3）
418+ 
419+| # | mode | NPU(us) | GPU(us) | GPU/NPU |
420+|---|------|---------|---------|--------|
421+| 305 | L/U/N/N | 158133.31 | 180365.95 | 1.1406 |
422+| 306 | L/U/N/U | 156182.28 | 185667.42 | 1.1888 |
423+| 307 | L/U/T/N | 165453.86 | 185798.17 | 1.1230 |
424+| 308 | L/U/T/U | 162821.73 | 193007.11 | 1.1854 |
425+| 309 | L/L/N/N | 157943.94 | 185935.45 | 1.1772 |
426+| 310 | L/L/N/U | 155997.34 | 190123.81 | 1.2188 |
427+| 311 | L/L/T/N | 165274.48 | 192106.59 | 1.1623 |
428+| 312 | L/L/T/U | 161949.34 | 189200.45 | 1.1683 |
429+| 313 | R/U/N/N | 181439.33 | 201003.58 | 1.1078 |
430+| 314 | R/U/N/U | 176329.42 | 214975.00 | 1.2192 |
431+| 315 | R/U/T/N | 172595.03 | 198764.48 | 1.1516 |
432+| 316 | R/U/T/U | 167758.52 | 204046.08 | 1.2163 |
433+| 317 | R/L/N/N | 180619.98 | 201021.09 | 1.1130 |
434+| 318 | R/L/N/U | 176612.11 | 207048.50 | 1.1723 |
435+| 319 | R/L/T/N | 174678.91 | 202174.27 | 1.1574 |
436+| 320 | R/L/T/U | 167844.05 | 209951.39 | 1.2509 |
437+ 
Aexperimental/aclblasCtrsmBatched2/op_host/ctrsm_batched_host.cpp+324-0
@@ -0,0 +1,324 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <cstdint>
12+#include <cstdio>
13+#include <cstring>
14+#include <complex>
15+#include <vector>
16+#include <algorithm>
17+#include "acl/acl.h"
18+#include "../../aclblas_minimal.h"
19+#include "tiling/platform/platform_ascendc.h"
20+#include "tiling/tiling_api.h"
21+#include "ctrsm_batched_tiling_data.h"
22+#include "ctrsm_batched_kernel_do.h"
23+ 
24+#define CHECK_RET(cond, return_expr) \
25+ do { \
26+ if (!(cond)) { \
27+ return_expr; \
28+ } \
29+ } while (0)
30+ 
31+constexpr int32_t MIX_NB_SMALL = 16;
32+constexpr int32_t MIX_NB_LARGE = 32;
33+ 
34+ 
35+static const char* GetSocVersion()
36+{
37+ return aclrtGetSocName();
38+}
39+ 
40+ 
41+static void GenerateCubeTiling(int32_t maxMN, uint8_t* tilingBuf, uint32_t* tilingSize)
42+{
43+ auto ascendcPlatform = platform_ascendc::PlatformAscendCManager::GetInstance(GetSocVersion());
44+ matmul_tiling::MultiCoreMatmulTiling tilingApi(*ascendcPlatform);
45+ tilingApi.SetDim(1);
46+ tilingApi.SetAType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND,
47+ matmul_tiling::DataType::DT_FLOAT, false);
48+ tilingApi.SetBType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND,
49+ matmul_tiling::DataType::DT_FLOAT, false);
50+ tilingApi.SetCType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND,
51+ matmul_tiling::DataType::DT_FLOAT);
52+ tilingApi.SetOrgShape(maxMN, maxMN, maxMN);
53+ tilingApi.SetShape(maxMN, maxMN, maxMN);
54+ tilingApi.SetBias(false);
55+ tilingApi.SetBufferSpace(-1, -1, -1);
56+ optiling::TCubeTiling cubeTiling;
57+ if (tilingApi.GetTiling(cubeTiling) == -1) {
58+ *tilingSize = 0;
59+ return;
60+ }
61+ *tilingSize = cubeTiling.GetDataSize();
62+ cubeTiling.SaveToBuffer(tilingBuf, *tilingSize);
63+}
64+ 
65+static bool IsValidSide(aclblasSideMode_t s) { return s == ACLBLAS_SIDE_LEFT || s == ACLBLAS_SIDE_RIGHT; }
66+static bool IsValidUplo(aclblasFillMode_t u) { return u == ACLBLAS_UPPER || u == ACLBLAS_LOWER; }
67+static bool IsValidTrans(aclblasOperation_t t) { return t == ACLBLAS_OP_N || t == ACLBLAS_OP_T || t == ACLBLAS_OP_C; }
68+static bool IsValidDiag(aclblasDiagType_t d) { return d == ACLBLAS_NON_UNIT || d == ACLBLAS_UNIT; }
69+ 
70+// 计算基本 tiling 参数（side/uplo/transa/diag 映射、useOrigA、nb 等）
71+static void FillCtrsmTilingBasicParams(aclblasSideMode_t side, aclblasFillMode_t uplo,
72+ aclblasOperation_t transa, aclblasDiagType_t diag,
73+ int64_t m, int64_t n, int64_t lda, int64_t ldb, int64_t batchCount,
74+ float alphaReal, float alphaImag,
75+ CtrsmBatchedTilingData& td,
76+ int32_t& kDim, int32_t& nCols, int32_t& kDimAligned, int32_t& nColsAligned,
77+ bool& sideRight, bool& useDualAiv)
78+{
79+ kDim = (side == ACLBLAS_SIDE_LEFT) ? (int32_t)m : (int32_t)n;
80+ nCols = (side == ACLBLAS_SIDE_LEFT) ? (int32_t)n : (int32_t)m;
81+ kDimAligned = CEIL_ALIGN(kDim, FLOAT_ALIGN);
82+ nColsAligned = CEIL_ALIGN(nCols, FLOAT_ALIGN);
83+ 
84+ sideRight = (side == ACLBLAS_SIDE_RIGHT);
85+ bool isTransN = (transa == ACLBLAS_OP_N);
86+ bool isTransT = (transa == ACLBLAS_OP_T);
87+ bool isTransC = (transa == ACLBLAS_OP_C);
88+ bool needTA = ((isTransT || isTransC) && !sideRight) || (isTransN && sideRight);
89+ bool padOn = (kDim != kDimAligned) || (nCols != nColsAligned);
90+ bool conjA = isTransC;
91+ bool useOrigA = (!needTA && !padOn && !conjA);
92+ 
93+ td.m = (int32_t)m;
94+ td.n = (int32_t)n;
95+ td.lda = (int32_t)lda;
96+ td.ldb = (int32_t)ldb;
97+ td.batchCount = (int32_t)batchCount;
98+ td.side = sideRight ? SIDE_RIGHT : SIDE_LEFT;
99+ td.uplo = (uplo == ACLBLAS_UPPER) ? UPLO_UPPER : UPLO_LOWER;
100+ td.transa = isTransN ? TRANS_N : (isTransT ? TRANS_T : TRANS_C);
101+ td.diag = (diag == ACLBLAS_NON_UNIT) ? DIAG_NONUNIT : DIAG_UNIT;
102+ td.alphaReal = alphaReal;
103+ td.alphaImag = alphaImag;
104+ td.useOrigA = useOrigA ? 1 : 0;
105+ td.aEffStride = useOrigA ? ((int32_t)lda * 2) : (kDimAligned * 2);
106+ useDualAiv = (kDim >= 128 && nColsAligned >= 128);
107+ td.dualAivMode = useDualAiv ? 1 : 0;
108+ td.nb = (kDim > 1024) ? MIX_NB_LARGE : MIX_NB_SMALL;
109+}
110+ 
111+// 计算 workspace 大小和 numBlocks
112+static void ComputeWorkspaceAndBlocks(
113+ CtrsmBatchedTilingData& td, int32_t kDimAligned, int32_t nColsAligned,
114+ int32_t numSplits, int32_t splitNColsAligned, int32_t lastNColsAligned,
115+ bool sideRight, bool useDualAiv, int64_t batchCount,
116+ int32_t& maxMN, uint32_t& numBlocks, int32_t& effectiveBatch)
117+{
118+ int32_t wsNCols = (numSplits > 1) ?
119+ ((splitNColsAligned > lastNColsAligned) ? splitNColsAligned : lastNColsAligned)
120+ : nColsAligned;
121+ int64_t aSize = (int64_t)kDimAligned * kDimAligned * 2;
122+ int64_t bSize = (int64_t)kDimAligned * wsNCols * 2;
123+ int64_t xnegSize = (int64_t)2 * 16 * 2 * td.nb * wsNCols * 2;
124+ int64_t rightTempSize = sideRight ? ((int64_t)wsNCols * kDimAligned * 2) : 0;
125+ int64_t gemmAreaSize = xnegSize;
126+ if (rightTempSize > gemmAreaSize) gemmAreaSize = rightTempSize;
127+ int64_t splitBGemmSize = (bSize + gemmAreaSize);
128+ td.splitBGemmSize = splitBGemmSize * sizeof(float);
129+ if (numSplits > 1) {
130+ td.workspaceOffset = (aSize + splitBGemmSize * numSplits) * sizeof(float);
131+ } else {
132+ td.workspaceOffset = (aSize + splitBGemmSize) * sizeof(float);
133+ }
134+ maxMN = (kDimAligned > wsNCols) ? kDimAligned : wsNCols;
135+ int32_t maxK = 2 * kDimAligned;
136+ if (maxK > maxMN) maxMN = maxK;
137+ if (numSplits > 1) {
138+ numBlocks = (uint32_t)((int32_t)batchCount * numSplits);
139+ effectiveBatch = (int32_t)batchCount;
140+ } else if (useDualAiv) {
141+ numBlocks = (uint32_t)batchCount;
142+ effectiveBatch = (int32_t)batchCount;
143+ } else {
144+ effectiveBatch = (int32_t)batchCount;
145+ if (effectiveBatch % 2 != 0) effectiveBatch++;
146+ numBlocks = (uint32_t)(effectiveBatch / 2);
147+ }
148+}
149+ 
150+// 计算 split/numBlocks/workspace 参数
151+static void FillCtrsmTilingSplitAndBlocks(
152+ CtrsmBatchedTilingData& td, int32_t kDim, int32_t nCols,
153+ int32_t kDimAligned, int32_t nColsAligned, bool sideRight, bool useDualAiv,
154+ int64_t batchCount, int32_t& maxMN, uint32_t& numBlocks, int32_t& effectiveBatch)
155+{
156+ constexpr int32_t SPLIT_ALIGN = 16;
157+ constexpr int32_t TOTAL_AICORES = 20;
158+ int32_t numSplits = 1;
159+ if (useDualAiv && (int32_t)batchCount <= TOTAL_AICORES / 2) {
160+ int32_t maxSplits = TOTAL_AICORES / (int32_t)batchCount;
161+ int32_t minSplitNCols = (td.nb > 64) ? td.nb : 64;
162+ int32_t maxBySize = nCols / minSplitNCols;
163+ if (maxSplits > maxBySize) maxSplits = maxBySize;
164+ if ((int32_t)batchCount * maxSplits > TOTAL_AICORES)
165+ maxSplits = TOTAL_AICORES / (int32_t)batchCount;
166+ if (maxSplits > 1) numSplits = maxSplits;
167+ }
168+ int32_t splitNCols = nCols;
169+ int32_t splitNColsAligned = nColsAligned;
170+ int32_t lastNCols = nCols;
171+ int32_t lastNColsAligned = nColsAligned;
172+ if (numSplits > 1) {
173+ splitNCols = (nCols / numSplits) & ~(SPLIT_ALIGN - 1);
174+ splitNColsAligned = CEIL_ALIGN(splitNCols, FLOAT_ALIGN);
175+ lastNCols = nCols - splitNCols * (numSplits - 1);
176+ lastNColsAligned = CEIL_ALIGN(lastNCols, FLOAT_ALIGN);
177+ }
178+ td.numSplits = numSplits;
179+ td.splitNCols = splitNCols;
180+ td.splitNColsAligned = splitNColsAligned;
181+ td.lastNCols = lastNCols;
182+ td.lastNColsAligned = lastNColsAligned;
183+ ComputeWorkspaceAndBlocks(td, kDimAligned, nColsAligned, numSplits,
184+ splitNColsAligned, lastNColsAligned, sideRight, useDualAiv,
185+ batchCount, maxMN, numBlocks, effectiveBatch);
186+}
187+ 
188+// 由 side/uplo/transa/diag 及尺寸推导 tiling 参数，并计算 kernel 启动参数（maxMN/numBlocks/effectiveBatch）
189+static void FillCtrsmTiling(aclblasSideMode_t side, aclblasFillMode_t uplo,
190+ aclblasOperation_t transa, aclblasDiagType_t diag,
191+ int64_t m, int64_t n, int64_t lda, int64_t ldb, int64_t batchCount,
192+ float alphaReal, float alphaImag,
193+ CtrsmBatchedTilingData& td, int32_t& maxMN,
194+ uint32_t& numBlocks, int32_t& effectiveBatch)
195+{
196+ int32_t kDim, nCols, kDimAligned, nColsAligned;
197+ bool sideRight, useDualAiv;
198+ FillCtrsmTilingBasicParams(side, uplo, transa, diag, m, n, lda, ldb, batchCount,
199+ alphaReal, alphaImag, td, kDim, nCols, kDimAligned, nColsAligned,
200+ sideRight, useDualAiv);
201+ FillCtrsmTilingSplitAndBlocks(td, kDim, nCols, kDimAligned, nColsAligned,
202+ sideRight, useDualAiv, batchCount, maxMN, numBlocks, effectiveBatch);
203+}
204+ 
205+// 分配 device 内存、拷贝入参与 cube tiling、启动 kernel 并同步、释放资源
206+static aclblasStatus_t LaunchCtrsmKernel(const CtrsmBatchedTilingData& td,
207+ int32_t maxMN, uint32_t numBlocks, int32_t effectiveBatch,
208+ const std::complex<float>* const aArray[], int64_t lda,
209+ std::complex<float>* const bArray[], int64_t batchCount, aclrtStream stream)
210+{
211+ size_t ptrArraySize = (size_t)batchCount * sizeof(void*);
212+ uint8_t *tilingDevice = nullptr, *aArrayDevice = nullptr, *bArrayDevice = nullptr;
213+ uint8_t *cubeTilingDevice = nullptr, *gemmWsDevice = nullptr, *sysWsDevice = nullptr;
214+ 
215+ CHECK_RET(aclrtMalloc((void**)&tilingDevice, sizeof(td), ACL_MEM_MALLOC_HUGE_FIRST) == ACL_SUCCESS,
216+ return ACLBLAS_STATUS_INTERNAL_ERROR);
217+ CHECK_RET(aclrtMalloc((void**)&aArrayDevice, ptrArraySize, ACL_MEM_MALLOC_HUGE_FIRST) == ACL_SUCCESS,
218+ return ACLBLAS_STATUS_INTERNAL_ERROR);
219+ CHECK_RET(aclrtMalloc((void**)&bArrayDevice, ptrArraySize, ACL_MEM_MALLOC_HUGE_FIRST) == ACL_SUCCESS,
220+ return ACLBLAS_STATUS_INTERNAL_ERROR);
221+ aclrtMemcpy(tilingDevice, sizeof(td), &td, sizeof(td), ACL_MEMCPY_HOST_TO_DEVICE);
222+ aclrtMemcpy(aArrayDevice, ptrArraySize, aArray, ptrArraySize, ACL_MEMCPY_HOST_TO_DEVICE);
223+ aclrtMemcpy(bArrayDevice, ptrArraySize, bArray, ptrArraySize, ACL_MEMCPY_HOST_TO_DEVICE);
224+ 
225+ uint8_t cubeTilingBuf[2048];
226+ uint32_t cubeTilingSize = 0;
227+ GenerateCubeTiling(maxMN, cubeTilingBuf, &cubeTilingSize);
228+ CHECK_RET(aclrtMalloc((void**)&cubeTilingDevice, sizeof(cubeTilingBuf), ACL_MEM_MALLOC_HUGE_FIRST) == ACL_SUCCESS,
229+ return ACLBLAS_STATUS_INTERNAL_ERROR);
230+ aclrtMemcpy(cubeTilingDevice, sizeof(cubeTilingBuf), cubeTilingBuf, sizeof(cubeTilingBuf), ACL_MEMCPY_HOST_TO_DEVICE);
231+ 
232+ size_t gemmWsSize = (size_t)effectiveBatch * td.workspaceOffset;
233+ CHECK_RET(aclrtMalloc((void**)&gemmWsDevice, gemmWsSize, ACL_MEM_MALLOC_HUGE_FIRST) == ACL_SUCCESS,
234+ return ACLBLAS_STATUS_INTERNAL_ERROR);
235+ aclrtMemset(gemmWsDevice, gemmWsSize, 0, gemmWsSize);
236+ 
237+ auto platformInst = platform_ascendc::PlatformAscendCManager::GetInstance(GetSocVersion());
238+ size_t sysWsSize = static_cast<size_t>(platformInst->GetLibApiWorkSpaceSize());
239+ if (sysWsSize < 16 * 1024 * 1024) sysWsSize = 16 * 1024 * 1024;
240+ CHECK_RET(aclrtMalloc((void**)&sysWsDevice, sysWsSize, ACL_MEM_MALLOC_HUGE_FIRST) == ACL_SUCCESS,
241+ return ACLBLAS_STATUS_INTERNAL_ERROR);
242+ 
243+ ctrsm_batched_mix12_kernel_do(aArrayDevice, bArrayDevice, tilingDevice, cubeTilingDevice,
244+ gemmWsDevice, sysWsDevice, numBlocks, stream);
245+ aclError syncRet = aclrtSynchronizeStream(stream);
246+ 
247+ aclrtFree(sysWsDevice);
248+ aclrtFree(gemmWsDevice);
249+ aclrtFree(cubeTilingDevice);
250+ aclrtFree(bArrayDevice);
251+ aclrtFree(aArrayDevice);
252+ aclrtFree(tilingDevice);
253+ CHECK_RET(syncRet == ACL_SUCCESS, return ACLBLAS_STATUS_INTERNAL_ERROR);
254+ return ACLBLAS_STATUS_SUCCESS;
255+}
256+ 
257+// 参数合法性校验：句柄、枚举、维度、lda/ldb。返回 SUCCESS 表示校验通过可继续
258+static aclblasStatus_t ValidateCtrsmArgs(aclblasHandle_t handle,
259+ aclblasSideMode_t side, aclblasFillMode_t uplo, aclblasOperation_t transa, aclblasDiagType_t diag,
260+ int64_t m, int64_t n, const std::complex<float>* alpha,
261+ const std::complex<float>* const aArray[], int64_t lda,
262+ std::complex<float>* const bArray[], int64_t ldb, int64_t batchCount)
263+{
264+ if (handle == nullptr) {
265+ return ACLBLAS_STATUS_HANDLE_IS_NULLPTR;
266+ }
267+ if (!IsValidSide(side) || !IsValidUplo(uplo) || !IsValidTrans(transa) || !IsValidDiag(diag)) {
268+ return ACLBLAS_STATUS_INVALID_VALUE;
269+ }
270+ if (m < 0 || n < 0 || batchCount < 0 || alpha == nullptr || aArray == nullptr || bArray == nullptr) {
271+ return ACLBLAS_STATUS_INVALID_VALUE;
272+ }
273+ int64_t minLda = (side == ACLBLAS_SIDE_LEFT) ? std::max((int64_t)1, m) : std::max((int64_t)1, n);
274+ if (lda < minLda || ldb < std::max((int64_t)1, n)) {
275+ return ACLBLAS_STATUS_INVALID_VALUE;
276+ }
277+ return ACLBLAS_STATUS_SUCCESS;
278+}
279+ 
280+aclblasStatus_t aclblasCtrsmBatched(
281+ aclblasHandle_t handle,
282+ aclblasSideMode_t side,
283+ aclblasFillMode_t uplo,
284+ aclblasOperation_t transa,
285+ aclblasDiagType_t diag,
286+ int64_t m, int64_t n,
287+ const std::complex<float>* alpha,
288+ const std::complex<float>* const aArray[], int64_t lda,
289+ std::complex<float>* const bArray[], int64_t ldb,
290+ int64_t batchCount)
291+{
292+ aclblasStatus_t vst = ValidateCtrsmArgs(handle, side, uplo, transa, diag,
293+ m, n, alpha, aArray, lda, bArray, ldb, batchCount);
294+ if (vst != ACLBLAS_STATUS_SUCCESS) {
295+ return vst;
296+ }
297+ if (m == 0 || n == 0 || batchCount == 0) {
298+ return ACLBLAS_STATUS_SUCCESS;
299+ }
300+ 
301+ float alphaReal = alpha->real();
302+ float alphaImag = alpha->imag();
303+ if (alphaReal == 0.0f && alphaImag == 0.0f) {
304+ for (int64_t i = 0; i < batchCount; i++) {
305+ size_t bytes = (size_t)m * ldb * 2 * sizeof(float);
306+ CHECK_RET(aclrtMemset(bArray[i], bytes, 0, bytes) == ACL_SUCCESS,
307+ return ACLBLAS_STATUS_INTERNAL_ERROR);
308+ }
309+ return ACLBLAS_STATUS_SUCCESS;
310+ }
311+ 
312+ auto* h = reinterpret_cast<_aclblas_handle*>(handle);
313+ aclrtStream stream = h->stream;
314+ 
315+ CtrsmBatchedTilingData td;
316+ int32_t maxMN;
317+ uint32_t numBlocks;
318+ int32_t effectiveBatch;
319+ FillCtrsmTiling(side, uplo, transa, diag, m, n, lda, ldb, batchCount,
320+ alphaReal, alphaImag, td, maxMN, numBlocks, effectiveBatch);
321+ 
322+ return LaunchCtrsmKernel(td, maxMN, numBlocks, effectiveBatch,
323+ aArray, lda, bArray, batchCount, stream);
324+}
Aexperimental/aclblasCtrsmBatched2/op_host/ctrsm_batched_kernel_do.h+18-0
@@ -0,0 +1,18 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#pragma once
12+ 
13+#include <cstdint>
14+ 
15+void ctrsm_batched_mix12_kernel_do(
16+ uint8_t* aArrayGm, uint8_t* bArrayGm, uint8_t* trsmTilingGm,
17+ uint8_t* cubeTilingGm, uint8_t* gemmWsGm, uint8_t* sysWorkspace,
18+ uint32_t numBlocks, void* stream);
Aexperimental/aclblasCtrsmBatched2/op_kernel/ctrsm_batched_kernel.cpp+69-0
@@ -0,0 +1,69 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef CTRSM_BATCHED_KERNEL_H
12+#define CTRSM_BATCHED_KERNEL_H
13+ 
14+#define ASCENDC_CUBE_ONLY
15+#include "ctrsm_batched_kernel_aiv.h"
16+#include "ctrsm_batched_kernel_aic.h"
17+ 
18+extern "C" __global__ __aicore__ void ctrsm_batched_mix12_kernel(
19+ GM_ADDR aArrayGm, GM_ADDR bArrayGm, GM_ADDR trsmTilingGm,
20+ GM_ADDR cubeTilingGm, GM_ADDR gemmWsGm, GM_ADDR sysWorkspace)
21+{
22+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);
23+ AscendC::TPipe pipe;
24+ auto trsmTiling = (__gm__ CtrsmBatchedTilingData*)trsmTilingGm;
25+ TCubeTiling cubeTiling;
26+ CopyCubeTiling(&cubeTiling, cubeTilingGm);
27+ if ASCEND_IS_AIC {
28+ CtrsmMixAic aic;
29+ aic.Init(aArrayGm, bArrayGm, gemmWsGm, trsmTiling, cubeTiling);
30+ REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), aic.mm, &aic.cubeTiling);
31+ aic.Process12();
32+ aic.mm.End();
33+ }
34+ if ASCEND_IS_AIV {
35+ // 按 side/uplo/transa 计算计算路径，分派到对应的模板路径类型
36+ bool right = (trsmTiling->side == SIDE_RIGHT);
37+ bool needTA = NeedTransA(trsmTiling);
38+ int32_t effUplo = needTA ? (trsmTiling->uplo == UPLO_UPPER ? UPLO_LOWER : UPLO_UPPER)
39+ : trsmTiling->uplo;
40+ bool forward = (effUplo == UPLO_LOWER);
41+ if (forward && !right) {
42+ CtrsmLowerLeft aiv;
43+ aiv.Init(aArrayGm, bArrayGm, gemmWsGm, trsmTiling, &pipe);
44+ aiv.Process12();
45+ } else if (forward && right) {
46+ CtrsmLowerRight aiv;
47+ aiv.Init(aArrayGm, bArrayGm, gemmWsGm, trsmTiling, &pipe);
48+ aiv.Process12();
49+ } else if (!forward && !right) {
50+ CtrsmUpperLeft aiv;
51+ aiv.Init(aArrayGm, bArrayGm, gemmWsGm, trsmTiling, &pipe);
52+ aiv.Process12();
53+ } else {
54+ CtrsmUpperRight aiv;
55+ aiv.Init(aArrayGm, bArrayGm, gemmWsGm, trsmTiling, &pipe);
56+ aiv.Process12();
57+ }
58+ }
59+}
60+ 
61+void ctrsm_batched_mix12_kernel_do(GM_ADDR aArrayGm, GM_ADDR bArrayGm, GM_ADDR trsmTilingGm,
62+ GM_ADDR cubeTilingGm, GM_ADDR gemmWsGm, GM_ADDR sysWorkspace,
63+ uint32_t numBlocks, void* stream)
64+{
65+ ctrsm_batched_mix12_kernel<<<numBlocks, nullptr, stream>>>(
66+ aArrayGm, bArrayGm, trsmTilingGm, cubeTilingGm, gemmWsGm, sysWorkspace);
67+}
68+ 
69+#endif
Aexperimental/aclblasCtrsmBatched2/op_kernel/ctrsm_batched_kernel_aic.h+418-0
@@ -0,0 +1,418 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#pragma once
12+ 
13+#include "ctrsm_batched_kernel_common.h"
14+ 
15+typedef MatmulType<AscendC::TPosition::GM, CubeFormat::ND, float> MmA;
16+typedef MatmulType<AscendC::TPosition::GM, CubeFormat::ND, float> MmB;
17+typedef MatmulType<AscendC::TPosition::GM, CubeFormat::ND, float> MmC;
18+ 
19+class CtrsmMixAic {
20+public:
21+ Matmul<MmA, MmB, MmC> mm;
22+ TCubeTiling cubeTiling;
23+ 
24+ __aicore__ inline CtrsmMixAic() {}
25+ 
26+ __aicore__ inline void Init(GM_ADDR aArrayGm, GM_ADDR bArrayGm, GM_ADDR gemmWsGm,
27+ const __gm__ CtrsmBatchedTilingData* t, const TCubeTiling& ct)
28+ {
29+ td = t;
30+ cubeTiling = ct;
31+ blockIdx = AscendC::GetBlockIdx();
32+ right = (td->side == SIDE_RIGHT);
33+ kDim = right ? td->n : td->m;
34+ nCols = right ? td->m : td->n;
35+ nb = td->nb;
36+ int32_t kDimPad = CEIL_ALIGN(kDim, FLOAT_ALIGN);
37+ int32_t nColsPad = CEIL_ALIGN(nCols, FLOAT_ALIGN);
38+ if (kDim != kDimPad || nCols != nColsPad) { kDim = kDimPad; nCols = nColsPad; }
39+ nColsAligned = CEIL_ALIGN(nCols, FLOAT_ALIGN);
40+ numSplits = td->numSplits;
41+ if (numSplits > 1) {
42+ int32_t splitIdx = (int32_t)blockIdx % numSplits;
43+ bool isLast = (splitIdx == numSplits - 1);
44+ int32_t localNColsAligned = isLast ? td->lastNColsAligned : td->splitNColsAligned;
45+ nCols = localNColsAligned;
46+ nColsAligned = localNColsAligned;
47+ }
48+ bool needTA = NeedTransA(td);
49+ int32_t effUplo = needTA ? (td->uplo == UPLO_UPPER ? UPLO_LOWER : UPLO_UPPER)
50+ : td->uplo;
51+ forward = (effUplo == UPLO_LOWER);
52+ gmAArray.SetGlobalBuffer((__gm__ uint64_t*)aArrayGm, td->batchCount);
53+ gmBArray.SetGlobalBuffer((__gm__ uint64_t*)bArrayGm, td->batchCount);
54+ gemmWsBase = gemmWsGm;
55+ useOrigA = (td->useOrigA == 1);
56+ aEffStride = td->aEffStride;
57+ }
58+ 
59+ __aicore__ inline void Process12()
60+ {
61+ int32_t numPanels = (kDim + nb - 1) / nb;
62+ if (numSplits > 1) {
63+ int32_t slot = (int32_t)blockIdx;
64+ if (numPanels <= 2) {
65+ Process12Simple(slot, -1, numPanels);
66+ } else {
67+ Process12Grouped(slot, -1, numPanels);
68+ }
69+ } else if (td->dualAivMode) {
70+ int32_t batch = (int32_t)blockIdx;
71+ if (numPanels <= 2) {
72+ Process12Simple(batch, -1, numPanels);
73+ } else {
74+ Process12Grouped(batch, -1, numPanels);
75+ }
76+ } else {
77+ int32_t g = (int32_t)blockIdx;
78+ int32_t b0 = 2 * g, b1 = 2 * g + 1;
79+ if (numPanels <= 2) {
80+ Process12Simple(b0, b1, numPanels);
81+ } else {
82+ Process12Grouped(b0, b1, numPanels);
83+ }
84+ }
85+ }
86+ 
87+private:
88+ __aicore__ inline void Ptrs(int32_t slot, __gm__ float*& eA, int32_t& eAStride,
89+ __gm__ float*& eBR, __gm__ float*& eBI,
90+ int32_t& eLdb, __gm__ float*& gmGemm)
91+ {
92+ int32_t batch = (numSplits > 1) ? (slot / numSplits) : slot;
93+ int32_t splitIdx = (numSplits > 1) ? (slot % numSplits) : 0;
94+ __gm__ float* batchWs = (__gm__ float*)gemmWsBase
95+ + (int64_t)batch * (td->workspaceOffset / sizeof(float));
96+ int64_t aPlane = (int64_t)kDim * kDim;
97+ if (useOrigA) {
98+ eA = (__gm__ float*)gmAArray.GetValue(batch);
99+ eAStride = aEffStride;
100+ } else {
101+ eA = batchWs;
102+ eAStride = 2 * kDim;
103+ }
104+ int32_t wsNCols = (numSplits > 1) ?
105+ ((td->splitNColsAligned > td->lastNColsAligned) ? td->splitNColsAligned : td->lastNColsAligned)
106+ : nColsAligned;
107+ int64_t bPlane = (int64_t)kDim * wsNCols;
108+ int64_t splitBGemmFloats = td->splitBGemmSize / sizeof(float);
109+ __gm__ float* splitWs = batchWs + 2 * aPlane + (int64_t)splitIdx * splitBGemmFloats;
110+ eBR = splitWs;
111+ eBI = eBR + bPlane;
112+ eLdb = nColsAligned;
113+ gmGemm = eBI + bPlane;
114+ }
115+ 
116+ __aicore__ inline void SetupXnegPtrs(int32_t slot, int32_t idx, int32_t numPanels,
117+ __gm__ float*& xR, __gm__ float*& xI,
118+ __gm__ float* gmGemm)
119+ {
120+ int32_t xSlot = XnegSlot(idx, numPanels, forward);
121+ int64_t xnegPlane = (int64_t)2 * LIM_GROUP * 2 * nb * nColsAligned;
122+ xR = gmGemm + (int64_t)xSlot * 2 * nb * nColsAligned;
123+ xI = gmGemm + xnegPlane + (int64_t)xSlot * 2 * nb * nColsAligned;
124+ }
125+ 
126+ __aicore__ inline void DoComplexGemm(
127+ AscendC::GlobalTensor<float>& gA,
128+ AscendC::GlobalTensor<float>& gBR, AscendC::GlobalTensor<float>& gBI,
129+ AscendC::GlobalTensor<float>& gXReal, AscendC::GlobalTensor<float>& gXImag,
130+ int32_t eAStride, int32_t eLdb,
131+ int32_t rowStart, int32_t rowCount, int32_t colStart, int32_t actualK)
132+ {
133+ int32_t aOff = rowStart * eAStride + colStart * 2;
134+ mm.SetOrgShape(rowCount, nCols, eAStride, nColsAligned, eLdb);
135+ mm.SetSingleShape(rowCount, nCols, 2 * actualK);
136+ mm.SetTensorA(gA[aOff], false);
137+ mm.SetTensorB(gXReal[0], false);
138+ mm.IterateAll(gBR[rowStart * eLdb], 1);
139+ mm.SetOrgShape(rowCount, nCols, eAStride, nColsAligned, eLdb);
140+ mm.SetSingleShape(rowCount, nCols, 2 * actualK);
141+ mm.SetTensorA(gA[aOff], false);
142+ mm.SetTensorB(gXImag[0], false);
143+ mm.IterateAll(gBI[rowStart * eLdb], 1);
144+ }
145+ 
146+ // 组装本 slot 的 A/B/Xneg 全局张量并执行一次复数 GEMM 更新。
147+ // DirectRankK / PreUpdateRankKWithinGroup / FullTrailUpdateWithinGroup 共用此流程，
148+ // 仅更新的行区间 [rowStart, rowStart+rowCount) 不同。
149+ __aicore__ inline void SetupAndGemm(int32_t slot, int32_t idx, int32_t numPanels,
150+ int32_t rowStart, int32_t rowCount,
151+ int32_t panelStart, int32_t actualNb)
152+ {
153+ __gm__ float *eA, *eBR, *eBI, *gmGemm, *xR, *xI;
154+ int32_t eAStride, eLdb;
155+ Ptrs(slot, eA, eAStride, eBR, eBI, eLdb, gmGemm);
156+ SetupXnegPtrs(slot, idx, numPanels, xR, xI, gmGemm);
157+ AscendC::GlobalTensor<float> gA, gBR, gBI, gXReal, gXImag;
158+ gA.SetGlobalBuffer(eA, (uint32_t)((int64_t)kDim * eAStride));
159+ gBR.SetGlobalBuffer(eBR, (uint32_t)((int64_t)kDim * eLdb));
160+ gBI.SetGlobalBuffer(eBI, (uint32_t)((int64_t)kDim * eLdb));
161+ gXReal.SetGlobalBuffer(xR, (uint32_t)((int64_t)2 * nb * nColsAligned));
162+ gXImag.SetGlobalBuffer(xI, (uint32_t)((int64_t)2 * nb * nColsAligned));
163+ DoComplexGemm(gA, gBR, gBI, gXReal, gXImag,
164+ eAStride, eLdb, rowStart, rowCount, panelStart, actualNb);
165+ }
166+ 
167+ __aicore__ inline void ComputeTrail(int32_t idx, int32_t numPanels,
168+ int32_t& panelStart, int32_t& actualNb,
169+ int32_t& trailStart, int32_t& trailRows)
170+ {
171+ int32_t p = forward ? idx : (numPanels - 1 - idx);
172+ panelStart = p * nb;
173+ actualNb = (nb < kDim - panelStart) ? nb : (kDim - panelStart);
174+ if (forward) { trailStart = panelStart + actualNb; trailRows = kDim - trailStart; }
175+ else { trailStart = 0; trailRows = panelStart; }
176+ }
177+ 
178+ __aicore__ inline bool DirectRankK(int32_t slot, int32_t idx, int32_t numPanels)
179+ {
180+ int32_t panelStart, actualNb, trailStart, trailRows;
181+ ComputeTrail(idx, numPanels, panelStart, actualNb, trailStart, trailRows);
182+ if (trailRows <= 0) return false;
183+ int32_t directRows = (nb < trailRows) ? nb : trailRows;
184+ int32_t directStart = forward ? trailStart : (trailStart + trailRows - directRows);
185+ SetupAndGemm(slot, idx, numPanels, directStart, directRows, panelStart, actualNb);
186+ return (trailRows > directRows);
187+ }
188+ 
189+ __aicore__ inline void PreUpdateRankKWithinGroup(int32_t slot, int32_t idx, int32_t numPanels,
190+ int32_t groupEndRow, int32_t groupBoundaryRow)
191+ {
192+ int32_t panelStart, actualNb, trailStart, trailRows;
193+ ComputeTrail(idx, numPanels, panelStart, actualNb, trailStart, trailRows);
194+ int32_t directRows = (nb < trailRows) ? nb : trailRows;
195+ int32_t preStart, preRows;
196+ if (forward) { preStart = trailStart + directRows; preRows = groupEndRow - preStart; }
197+ else {
198+ int32_t directStart = trailStart + trailRows - directRows;
199+ preStart = groupBoundaryRow;
200+ preRows = directStart - groupBoundaryRow;
201+ }
202+ if (preRows <= 0) return;
203+ SetupAndGemm(slot, idx, numPanels, preStart, preRows, panelStart, actualNb);
204+ }
205+ 
206+ __aicore__ inline void BigGroupGemm(int32_t slot, int32_t idxStart, int32_t groupSize,
207+ int32_t numPanels)
208+ {
209+ int32_t bigM, bigStart, colStart;
210+ ComputeBigGemmBounds(idxStart, groupSize, numPanels, bigM, bigStart, colStart);
211+ if (bigM <= 0) return;
212+ __gm__ float *eA, *eBR, *eBI, *gmGemm;
213+ int32_t eAStride, eLdb;
214+ Ptrs(slot, eA, eAStride, eBR, eBI, eLdb, gmGemm);
215+ int32_t g = idxStart / LIM_GROUP;
216+ int32_t groupBase = (g % 2) * LIM_GROUP;
217+ int64_t xnegPlane = (int64_t)2 * LIM_GROUP * 2 * nb * nColsAligned;
218+ __gm__ float* xR = gmGemm + (int64_t)groupBase * 2 * nb * nColsAligned;
219+ __gm__ float* xI = gmGemm + xnegPlane + (int64_t)groupBase * 2 * nb * nColsAligned;
220+ int32_t totalK = groupSize * nb;
221+ int32_t maxK = kDim - colStart;
222+ if (totalK > maxK) totalK = maxK;
223+ AscendC::GlobalTensor<float> gA, gBR, gBI, gXReal, gXImag;
224+ gA.SetGlobalBuffer(eA, (uint32_t)((int64_t)kDim * eAStride));
225+ gBR.SetGlobalBuffer(eBR, (uint32_t)((int64_t)kDim * eLdb));
226+ gBI.SetGlobalBuffer(eBI, (uint32_t)((int64_t)kDim * eLdb));
227+ gXReal.SetGlobalBuffer(xR, (uint32_t)((int64_t)2 * totalK * nColsAligned));
228+ gXImag.SetGlobalBuffer(xI, (uint32_t)((int64_t)2 * totalK * nColsAligned));
229+ DoComplexGemm(gA, gBR, gBI, gXReal, gXImag,
230+ eAStride, eLdb, bigStart, bigM, colStart, totalK);
231+ }
232+ 
233+ // 用组内 panel 子区间 [panelOffset, panelOffset+accPanels) 的 Xneg 更新远端行块
234+ // [bigStart, bigStart+bigM)。K=accPanels*nb。用于把 BigGroupGemm 按 K 拆分、提前与 AIV 并行
235+ __aicore__ inline void FarGemm(int32_t slot, int32_t idxStart, int32_t panelOffset,
236+ int32_t accPanels, int32_t bigStart, int32_t bigM)
237+ {
238+ if (bigM <= 0 || accPanels <= 0) return;
239+ __gm__ float *eA, *eBR, *eBI, *gmGemm;
240+ int32_t eAStride, eLdb;
241+ Ptrs(slot, eA, eAStride, eBR, eBI, eLdb, gmGemm);
242+ int32_t g = idxStart / LIM_GROUP;
243+ int32_t groupBase = (g % 2) * LIM_GROUP;
244+ int64_t xnegPlane = (int64_t)2 * LIM_GROUP * 2 * nb * nColsAligned;
245+ int32_t baseSlot = groupBase + panelOffset;
246+ __gm__ float* xR = gmGemm + (int64_t)baseSlot * 2 * nb * nColsAligned;
247+ __gm__ float* xI = gmGemm + xnegPlane + (int64_t)baseSlot * 2 * nb * nColsAligned;
248+ int32_t totalK = accPanels * nb;
249+ int32_t colStart = (idxStart + panelOffset) * nb;
250+ AscendC::GlobalTensor<float> gA, gBR, gBI, gXReal, gXImag;
251+ gA.SetGlobalBuffer(eA, (uint32_t)((int64_t)kDim * eAStride));
252+ gBR.SetGlobalBuffer(eBR, (uint32_t)((int64_t)kDim * eLdb));
253+ gBI.SetGlobalBuffer(eBI, (uint32_t)((int64_t)kDim * eLdb));
254+ gXReal.SetGlobalBuffer(xR, (uint32_t)((int64_t)2 * totalK * nColsAligned));
255+ gXImag.SetGlobalBuffer(xI, (uint32_t)((int64_t)2 * totalK * nColsAligned));
256+ DoComplexGemm(gA, gBR, gBI, gXReal, gXImag,
257+ eAStride, eLdb, bigStart, bigM, colStart, totalK);
258+ }
259+ 
260+ __aicore__ inline void ComputeBigGemmBounds(int32_t idxStart, int32_t groupSize,
261+ int32_t numPanels,
262+ int32_t& bigM, int32_t& bigStart, int32_t& colStart)
263+ {
264+ if (forward) {
265+ int32_t pLast = idxStart + groupSize - 1;
266+ int32_t groupEndRow = (((pLast + 2) * nb < kDim) ? (pLast + 2) * nb : kDim);
267+ bigM = kDim - groupEndRow;
268+ bigStart = groupEndRow;
269+ colStart = idxStart * nb;
270+ } else {
271+ int32_t pLast = numPanels - 1 - (idxStart + groupSize - 1);
272+ int32_t groupBoundaryRow = (((pLast - 1) * nb > 0) ? (pLast - 1) * nb : 0);
273+ bigM = groupBoundaryRow;
274+ bigStart = 0;
275+ colStart = pLast * nb;
276+ }
277+ }
278+ 
279+ __aicore__ inline bool PanelHasTrail(int32_t idx, int32_t numPanels)
280+ {
281+ int32_t panelStart, actualNb, trailStart, trailRows;
282+ ComputeTrail(idx, numPanels, panelStart, actualNb, trailStart, trailRows);
283+ return trailRows > 0;
284+ }
285+ 
286+ __aicore__ inline void Process12Simple(int32_t b0, int32_t b1, int32_t numPanels)
287+ {
288+ bool useMerged = (td->dualAivMode == 0);
289+ for (int32_t idx = 0; idx < numPanels; idx++) {
290+ if (!PanelHasTrail(idx, numPanels)) continue;
291+ AscendC::CrossCoreWaitFlag<2, PIPE_FIX>(FLAG_TRSV);
292+ if (useMerged) {
293+ FullTrailUpdateWithinGroup(b0, idx, numPanels, kDim, 0);
294+ if (b1 >= 0 && b1 < td->batchCount)
295+ FullTrailUpdateWithinGroup(b1, idx, numPanels, kDim, 0);
296+ AscendC::PipeBarrier<PIPE_FIX>();
297+ AscendC::CrossCoreSetFlag<2, PIPE_FIX>(FLAG_GEMM);
298+ } else {
299+ bool hasMore0 = DirectRankK(b0, idx, numPanels);
300+ bool hasMore1 = (b1 >= 0 && b1 < td->batchCount) ? DirectRankK(b1, idx, numPanels) : false;
301+ AscendC::PipeBarrier<PIPE_FIX>();
302+ AscendC::CrossCoreSetFlag<2, PIPE_FIX>(FLAG_GEMM);
303+ if (hasMore0) PreUpdateRankKWithinGroup(b0, idx, numPanels, kDim, 0);
304+ if (hasMore1) PreUpdateRankKWithinGroup(b1, idx, numPanels, kDim, 0);
305+ }
306+ }
307+ }
308+ 
309+ __aicore__ inline void GroupBounds(int32_t grp, int32_t numPanels,
310+ int32_t& idxStart, int32_t& idxEnd, int32_t& groupSize,
311+ int32_t& groupEndRow, int32_t& groupBoundaryRow)
312+ {
313+ ::GroupBounds(grp, numPanels, nb, kDim, forward,
314+ idxStart, idxEnd, groupSize, groupEndRow, groupBoundaryRow);
315+ }
316+ 
317+ __aicore__ inline void Process12Grouped(int32_t b0, int32_t b1, int32_t numPanels)
318+ {
319+ int32_t numGroups = (numPanels + LIM_GROUP - 1) / LIM_GROUP;
320+ for (int32_t grp = 0; grp < numGroups; grp++) {
321+ int32_t idxStart, idxEnd, groupSize, groupEndRow, groupBoundaryRow;
322+ GroupBounds(grp, numPanels, idxStart, idxEnd, groupSize, groupEndRow, groupBoundaryRow);
323+ ProcessOneGroup(b0, b1, numPanels, idxStart, idxEnd, groupSize,
324+ groupEndRow, groupBoundaryRow);
325+ }
326+ }
327+ 
328+ // 单个 panel 的 trailing 更新（merged 全量 / direct+preupdate 两条路径）
329+ __aicore__ inline void UpdateOnePanel(int32_t b0, int32_t b1, int32_t idx, int32_t numPanels,
330+ int32_t groupEndRow, int32_t groupBoundaryRow, bool useMerged)
331+ {
332+ bool hasB1 = (b1 >= 0 && b1 < td->batchCount);
333+ if (useMerged) {
334+ FullTrailUpdateWithinGroup(b0, idx, numPanels, groupEndRow, groupBoundaryRow);
335+ if (hasB1) FullTrailUpdateWithinGroup(b1, idx, numPanels, groupEndRow, groupBoundaryRow);
336+ AscendC::PipeBarrier<PIPE_FIX>();
337+ AscendC::CrossCoreSetFlag<2, PIPE_FIX>(FLAG_GEMM);
338+ return;
339+ }
340+ bool hasMore0 = DirectRankK(b0, idx, numPanels);
341+ bool hasMore1 = hasB1 ? DirectRankK(b1, idx, numPanels) : false;
342+ AscendC::PipeBarrier<PIPE_FIX>();
343+ AscendC::CrossCoreSetFlag<2, PIPE_FIX>(FLAG_GEMM);
344+ if (hasMore0) PreUpdateRankKWithinGroup(b0, idx, numPanels, groupEndRow, groupBoundaryRow);
345+ if (hasMore1) PreUpdateRankKWithinGroup(b1, idx, numPanels, groupEndRow, groupBoundaryRow);
346+ }
347+ 
348+ // 组末远端更新收尾：doSplit 补齐剩余段，否则做整块 BigGroupGemm
349+ __aicore__ inline void FinishGroupFar(int32_t b0, int32_t b1, int32_t idxStart, int32_t groupSize,
350+ int32_t numPanels, int32_t groupEndRow, int32_t groupBoundaryRow,
351+ bool doSplit, int32_t doneP, int32_t bigStartF, int32_t bigMF)
352+ {
353+ if (doSplit) {
354+ int32_t rem = groupSize - doneP;
355+ if (rem > 0) FarGemm(b0, idxStart, doneP, rem, bigStartF, bigMF);
356+ return;
357+ }
358+ int32_t bigM = forward ? (kDim - groupEndRow) : groupBoundaryRow;
359+ if (bigM <= 0) return;
360+ BigGroupGemm(b0, idxStart, groupSize, numPanels);
361+ if (b1 >= 0 && b1 < td->batchCount) BigGroupGemm(b1, idxStart, groupSize, numPanels);
362+ }
363+ 
364+ __aicore__ inline void ProcessOneGroup(int32_t b0, int32_t b1, int32_t numPanels,
365+ int32_t idxStart, int32_t idxEnd, int32_t groupSize,
366+ int32_t groupEndRow, int32_t groupBoundaryRow)
367+ {
368+ bool useMerged = (td->dualAivMode == 0);
369+ constexpr int32_t FAR_STEP = 2;
370+ // forward: 把远端 BigGroupGemm 按 K 分成多段（每 FAR_STEP 个 panel 一段），
371+ // 每段解完后立即用该段 X（K=FAR_STEP*nb）更新远端大块，与 AIV 解后续 panel 并行。
372+ // 各段 K 区间不相交 → 累加不重复；远端行 AIV 不碰 → 无 race。
373+ int32_t bigStartF = forward ? groupEndRow : 0;
374+ int32_t bigMF = forward ? (kDim - groupEndRow) : 0;
375+ bool doSplit = false;
376+ int32_t doneP = 0; // 已用 FarGemm 处理的 panel 数
377+ 
378+ for (int32_t idx = idxStart; idx <= idxEnd; idx++) {
379+ if (!PanelHasTrail(idx, numPanels)) continue;
380+ AscendC::CrossCoreWaitFlag<2, PIPE_FIX>(FLAG_TRSV);
381+ UpdateOnePanel(b0, b1, idx, numPanels, groupEndRow, groupBoundaryRow, useMerged);
382+ // 每积累满 FAR_STEP 个 panel，立即对远端大块做一段更新（与 AIV 并行）
383+ if (doSplit && (idx - idxStart + 1 - doneP) >= FAR_STEP) {
384+ FarGemm(b0, idxStart, doneP, FAR_STEP, bigStartF, bigMF);
385+ doneP += FAR_STEP;
386+ }
387+ }
388+ FinishGroupFar(b0, b1, idxStart, groupSize, numPanels, groupEndRow, groupBoundaryRow,
389+ doSplit, doneP, bigStartF, bigMF);
390+ }
391+ 
392+ __aicore__ inline void FullTrailUpdateWithinGroup(int32_t slot, int32_t idx, int32_t numPanels,
393+ int32_t groupEndRow, int32_t groupBoundaryRow)
394+ {
395+ int32_t panelStart, actualNb, trailStart, trailRows;
396+ ComputeTrail(idx, numPanels, panelStart, actualNb, trailStart, trailRows);
397+ if (trailRows <= 0) return;
398+ int32_t updateStart, updateRows;
399+ if (forward) {
400+ updateStart = trailStart;
401+ updateRows = groupEndRow - trailStart;
402+ } else {
403+ updateStart = groupBoundaryRow;
404+ updateRows = trailStart + trailRows - groupBoundaryRow;
405+ }
406+ if (updateRows <= 0) return;
407+ SetupAndGemm(slot, idx, numPanels, updateStart, updateRows, panelStart, actualNb);
408+ }
409+ 
410+ const __gm__ CtrsmBatchedTilingData* td;
411+ AscendC::GlobalTensor<uint64_t> gmAArray, gmBArray;
412+ __gm__ uint8_t* gemmWsBase;
413+ uint32_t blockIdx;
414+ int32_t kDim, nCols, nb, nColsAligned;
415+ int32_t aEffStride;
416+ int32_t numSplits;
417+ bool forward, right, useOrigA;
418+};
Aexperimental/aclblasCtrsmBatched2/op_kernel/ctrsm_batched_kernel_aiv.h+532-0
@@ -0,0 +1,532 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ctrsm_batched_kernel_aiv.h
13+ * \brief 复数三角求解 AIV 编排类，按计算路径模板化。原单一大类已按职责拆成协作的多个类：
14+ * - CtrsmAivCfg 共享派生配置 + UB buffer 指针
15+ * - CtrsmConvert 复数 AoS<->SoA 转换/分块转置工具（路径无关）
16+ * - CtrsmCanonA A 规范化（补零/转置/共轭）
17+ * - CtrsmCanonB<RIGHT> B 规范化与回写（左/右乘）
18+ * - CtrsmPanelSolver<FORWARD> Panel 三角求解（前代/回代）
19+ * - CtrsmMixAivImpl<FORWARD,RIGHT> 编排类：持有 pipe/UB buffer/cfg，装配并驱动子类
20+ * UB buffer 由编排类统一分配（保留原有别名布局以满足 192KB 预算），子类仅持指针引用。
21+ *
22+ * 四条计算路径由模板参数在编译期固定（消除运行时分支），对应文件末尾四个类型别名；
23+ * 内核入口按 side/uplo/transa 分派到对应路径类型。
24+ */
25+ 
26+#pragma once
27+ 
28+#include "ctrsm_batched_kernel_common.h"
29+#include "ctrsm_batched_kernel_aiv_cfg.h"
30+#include "ctrsm_batched_kernel_aiv_convert.h"
31+#include "ctrsm_batched_kernel_aiv_canon_a.h"
32+#include "ctrsm_batched_kernel_aiv_canon_b.h"
33+#include "ctrsm_batched_kernel_aiv_solver.h"
34+ 
35+template <bool FORWARD, bool RIGHT>
36+class CtrsmMixAivImpl {
37+public:
38+ __aicore__ inline CtrsmMixAivImpl() {}
39+ 
40+ // 初始化：解析 tiling、设置全局内存指针、分配 UB 缓冲区、装配子对象
41+ __aicore__ inline void Init(GM_ADDR aArrayGm, GM_ADDR bArrayGm, GM_ADDR gemmWsGm,
42+ const __gm__ CtrsmBatchedTilingData* t, AscendC::TPipe* pipeIn)
43+ {
44+ cfg.td = t;
45+ pipe = pipeIn;
46+ blockIdx = AscendC::GetBlockIdx();
47+ cfg.kDim = RIGHT ? t->n : t->m;
48+ cfg.nCols = RIGHT ? t->m : t->n;
49+ cfg.nb = t->nb;
50+ cfg.kDimOrig = cfg.kDim;
51+ cfg.nColsOrig = cfg.nCols;
52+ int32_t kDimPad = CEIL_ALIGN(cfg.kDim, FLOAT_ALIGN);
53+ int32_t nColsPad = CEIL_ALIGN(cfg.nCols, FLOAT_ALIGN);
54+ cfg.padOn = (cfg.kDim != kDimPad) || (cfg.nCols != nColsPad);
55+ if (cfg.padOn) { cfg.kDim = kDimPad; cfg.nCols = nColsPad; }
56+ cfg.nColsAligned = CEIL_ALIGN(cfg.nCols, FLOAT_ALIGN);
57+ cfg.nbAligned = CEIL_ALIGN(cfg.nb, FLOAT_ALIGN);
58+ cfg.needTA = NeedTransA(t);
59+ cfg.conjA = IsConjTransA(t);
60+ gmAArray.SetGlobalBuffer((__gm__ uint64_t*)aArrayGm, t->batchCount);
61+ gmBArray.SetGlobalBuffer((__gm__ uint64_t*)bArrayGm, t->batchCount);
62+ gemmWsBase = gemmWsGm;
63+ cfg.useOrigA = (t->useOrigA == 1);
64+ cfg.aEffStride = t->aEffStride;
65+ cfg.localNCols = cfg.nColsOrig;
66+ cfg.localNColsAligned = cfg.nColsAligned;
67+ cfg.nColsOffset = 0;
68+ cfg.colStart = 0;
69+ cfg.colEnd = cfg.nColsAligned;
70+ InitBuffers();
71+ WireSubObjects();
72+ convert.BuildDeinterleaveOffsets(); // gather 偏移表构建一次，后续 LoadPanelA 复用
73+ }
74+ 
75+ // 初始化UB缓冲区分配：根据矩阵维度计算各缓冲区大小并分配
76+ __aicore__ inline void InitBuffers()
77+ {
78+ int32_t kDimAligned = CEIL_ALIGN(cfg.kDim, FLOAT_ALIGN);
79+ int32_t rowMax = kDimAligned;
80+ if (cfg.nColsAligned > rowMax) rowMax = cfg.nColsAligned;
81+ if (rowMax < FLOAT_ALIGN) rowMax = FLOAT_ALIGN;
82+ rowMax = CEIL_ALIGN(rowMax, 32);
83+ pipe->InitBuffer(bufRow, rowMax * sizeof(float));
84+ pipe->InitBuffer(bufPanelA_real, cfg.nbAligned * cfg.nbAligned * sizeof(float));
85+ pipe->InitBuffer(bufPanelA_imag, cfg.nbAligned * cfg.nbAligned * sizeof(float));
86+ 
87+ int32_t fixedFloats = 2 * cfg.nbAligned * cfg.nbAligned + cfg.nbAligned + rowMax;
88+ int32_t perColFloats = 5 * cfg.nb + 1;
89+ constexpr int32_t UB_FLOATS = 192 * 1024 / (int32_t)sizeof(float);
90+ int32_t colTileMax = (UB_FLOATS - fixedFloats) / perColFloats;
91+ colTileMax = (colTileMax / FLOAT_ALIGN) * FLOAT_ALIGN;
92+ cfg.colTile = (cfg.nColsAligned <= colTileMax) ? cfg.nColsAligned : colTileMax;
93+ if (cfg.colTile < FLOAT_ALIGN) cfg.colTile = FLOAT_ALIGN;
94+ int32_t baseBuf = cfg.nb * cfg.colTile;
95+ int32_t interBuf = 2 * baseBuf; // 交错 B 存储 [实|虚] 需 nb*2*colTile
96+ int32_t panelBRealSize = (CtrsmAivCfg::TS * CtrsmAivCfg::TS > interBuf) ? CtrsmAivCfg::TS * CtrsmAivCfg::TS : interBuf;
97+ int32_t negINSize = (CtrsmAivCfg::TS * 16 > baseBuf) ? CtrsmAivCfg::TS * 16 : baseBuf;
98+ int32_t totalUsed = fixedFloats + panelBRealSize + 4 * baseBuf + negINSize + cfg.colTile;
99+ while (totalUsed > UB_FLOATS && cfg.colTile > FLOAT_ALIGN) {
100+ cfg.colTile -= FLOAT_ALIGN;
101+ baseBuf = cfg.nb * cfg.colTile;
102+ interBuf = 2 * baseBuf;
103+ panelBRealSize = (CtrsmAivCfg::TS * CtrsmAivCfg::TS > interBuf) ? CtrsmAivCfg::TS * CtrsmAivCfg::TS : interBuf;
104+ negINSize = (CtrsmAivCfg::TS * 16 > baseBuf) ? CtrsmAivCfg::TS * 16 : baseBuf;
105+ totalUsed = fixedFloats + panelBRealSize + 4 * baseBuf + negINSize + cfg.colTile;
106+ }
107+ baseBuf = cfg.nb * cfg.colTile;
108+ interBuf = 2 * baseBuf;
109+ panelBRealSize = (CtrsmAivCfg::TS * CtrsmAivCfg::TS > interBuf) ? CtrsmAivCfg::TS * CtrsmAivCfg::TS : interBuf;
110+ negINSize = (CtrsmAivCfg::TS * 16 > baseBuf) ? CtrsmAivCfg::TS * 16 : baseBuf;
111+ pipe->InitBuffer(bufPanelB_real, panelBRealSize * sizeof(float));
112+ cfg.srcBufFloats = panelBRealSize;
113+ pipe->InitBuffer(bufPanelB_imag, baseBuf * sizeof(float));
114+ pipe->InitBuffer(bufRankK, cfg.colTile * sizeof(float));
115+ pipe->InitBuffer(bufNeg_real, baseBuf * sizeof(float));
116+ pipe->InitBuffer(bufNeg_imag, baseBuf * sizeof(float));
117+ pipe->InitBuffer(bufNeg_imag_neg, negINSize * sizeof(float));
118+ int32_t gc = cfg.nb * cfg.colTile;
119+ cfg.gatherChunk = (gc < 256) ? gc : 256;
120+ cfg.gatherChunk = (cfg.gatherChunk / FLOAT_ALIGN) * FLOAT_ALIGN;
121+ if (cfg.gatherChunk < FLOAT_ALIGN) cfg.gatherChunk = FLOAT_ALIGN;
122+ // 专用 gather 偏移表 buffer（gatherChunk 个 int，固定内容）
123+ pipe->InitBuffer(bufGatherEven, cfg.gatherChunk * sizeof(int32_t));
124+ pipe->InitBuffer(bufGatherOdd, cfg.gatherChunk * sizeof(int32_t));
125+ }
126+ 
127+ // 向量核主入口：根据 numSplits/dualAivMode 选择多核拆分/分列/原始独立模式
128+ __aicore__ inline void Process12()
129+ {
130+ if (cfg.td->numSplits > 1) {
131+ int32_t aivBlock = (int32_t)blockIdx / 2;
132+ int32_t half = (int32_t)blockIdx % 2;
133+ int32_t batch = aivBlock / cfg.td->numSplits;
134+ int32_t splitIdx = aivBlock % cfg.td->numSplits;
135+ bool isLast = (splitIdx == cfg.td->numSplits - 1);
136+ cfg.localNCols = isLast ? cfg.td->lastNCols : cfg.td->splitNCols;
137+ cfg.localNColsAligned = isLast ? cfg.td->lastNColsAligned : cfg.td->splitNColsAligned;
138+ cfg.nColsOffset = splitIdx * cfg.td->splitNCols;
139+ int32_t localNColsAligned = cfg.localNColsAligned;
140+ cfg.nColsAligned = localNColsAligned;
141+ cfg.nCols = localNColsAligned;
142+ int32_t colMid = CEIL_ALIGN(localNColsAligned / 2, FLOAT_ALIGN);
143+ if (colMid > localNColsAligned) colMid = localNColsAligned;
144+ if (half == 0) { cfg.colStart = 0; cfg.colEnd = colMid; }
145+ else { cfg.colStart = colMid; cfg.colEnd = localNColsAligned; }
146+ ProcessOneBatchSplit(batch, half);
147+ } else if (cfg.td->dualAivMode) {
148+ int32_t batch = (int32_t)blockIdx / 2;
149+ int32_t half = (int32_t)blockIdx % 2;
150+ int32_t colMid = CEIL_ALIGN(cfg.nColsAligned / 2, FLOAT_ALIGN);
151+ if (colMid > cfg.nColsAligned) colMid = cfg.nColsAligned;
152+ if (half == 0) { cfg.colStart = 0; cfg.colEnd = colMid; }
153+ else { cfg.colStart = colMid; cfg.colEnd = cfg.nColsAligned; }
154+ ProcessOneBatchPadded(batch);
155+ } else {
156+ int32_t batch = (int32_t)blockIdx;
157+ cfg.colStart = 0;
158+ cfg.colEnd = cfg.nColsAligned;
159+ if (batch < cfg.td->batchCount) {
160+ cfg.buildRowStart = 0;
161+ cfg.buildRowEnd = cfg.kDim;
162+ ProcessOneBatchOriginal(batch);
163+ } else {
164+ DummyPanelSync();
165+ }
166+ }
167+ }
168+ 
169+private:
170+ __aicore__ inline void WireSubObjects()
171+ {
172+ cfg.bufPanelA_real = &bufPanelA_real; cfg.bufPanelA_imag = &bufPanelA_imag;
173+ cfg.bufPanelB_real = &bufPanelB_real; cfg.bufPanelB_imag = &bufPanelB_imag;
174+ cfg.bufNeg_real = &bufNeg_real; cfg.bufNeg_imag = &bufNeg_imag; cfg.bufNeg_imag_neg = &bufNeg_imag_neg;
175+ cfg.bufRankK = &bufRankK; cfg.bufRow = &bufRow;
176+ cfg.bufGatherEven = &bufGatherEven; cfg.bufGatherOdd = &bufGatherOdd;
177+ convert.Bind(&cfg);
178+ canonA.Bind(&cfg, &convert);
179+ canonB.Bind(&cfg, &convert);
180+ solver.Bind(&cfg, &convert);
181+ }
182+ 
183+ // 判断当前Panel是否有Trail区域需要GEMM更新
184+ __aicore__ inline bool PanelHasTrailAiv(int32_t idx, int32_t numPanels)
185+ {
186+ int32_t p = FORWARD ? idx : (numPanels - 1 - idx);
187+ int32_t panelStart = p * cfg.nb;
188+ int32_t actualNb = (cfg.nb < cfg.kDim - panelStart) ? cfg.nb : (cfg.kDim - panelStart);
189+ int32_t trailRows = FORWARD ? (cfg.kDim - (panelStart + actualNb)) : panelStart;
190+ return trailRows > 0;
191+ }
192+ 
193+ // 空batch时执行虚拟同步，配合Cube核保持跨核flag对齐
194+ __aicore__ inline void DummyPanelSync()
195+ {
196+ int32_t numPanels = (cfg.kDim + cfg.nb - 1) / cfg.nb;
197+ for (int32_t idx = 0; idx < numPanels; idx++) {
198+ if (!PanelHasTrailAiv(idx, numPanels)) continue;
199+ AscendC::PipeBarrier<PIPE_MTE3>();
200+ AscendC::CrossCoreSetFlag<2, PIPE_MTE3>(FLAG_TRSV);
201+ AscendC::CrossCoreWaitFlag<2, PIPE_MTE2>(FLAG_GEMM);
202+ }
203+ }
204+ 
205+ // 准备A矩阵工作区（转置/padding），返回有效A指针和stride
206+ __aicore__ inline void PrepareAMatrix(int32_t batch, __gm__ float* coreWs,
207+ __gm__ float* gmBc_real, __gm__ float* gmBc_imag, __gm__ float* gmGemm,
208+ __gm__ float*& effA, int32_t& effAStride)
209+ {
210+ __gm__ float* gmA = (__gm__ float*)gmAArray.GetValue(batch);
211+ __gm__ float* gmAc = coreWs;
212+ if (cfg.useOrigA) {
213+ effA = gmA;
214+ effAStride = cfg.aEffStride;
215+ } else if (cfg.needTA) {
216+ int32_t wsStride = 2 * cfg.kDim;
217+ float signIm = cfg.conjA ? -1.0f : 1.0f;
218+ convert.TransposeADirectToAoS(gmA, cfg.kDimOrig, cfg.td->lda,
219+ gmAc, wsStride, signIm, 0, cfg.kDim, cfg.padOn, cfg.kDim);
220+ effA = gmAc;
221+ effAStride = wsStride;
222+ } else {
223+ canonA.BuildPaddedA(gmA, gmAc, gmBc_real, gmBc_imag, gmGemm);
224+ effA = gmAc;
225+ effAStride = 2 * cfg.kDim;
226+ }
227+ }
228+ 
229+ // 准备B矩阵、执行求解循环、回写结果
230+ __aicore__ inline void PrepareBAndSolve(int32_t batch, __gm__ float* effA, int32_t effAStride,
231+ __gm__ float* gmBc_real, __gm__ float* gmBc_imag, __gm__ float* gmGemm)
232+ {
233+ __gm__ float* gmB = (__gm__ float*)gmBArray.GetValue(batch);
234+ if (RIGHT) {
235+ int32_t rTotal = (cfg.td->m < cfg.nColsOrig) ? cfg.td->m : cfg.nColsOrig;
236+ int32_t deintCols = (cfg.td->n < cfg.kDimOrig) ? cfg.td->n : cfg.kDimOrig;
237+ if (cfg.padOn) {
238+ convert.ZeroGmRows(gmBc_real, cfg.kDim, cfg.nColsAligned, cfg.nColsAligned);
239+ convert.ZeroGmRows(gmBc_imag, cfg.kDim, cfg.nColsAligned, cfg.nColsAligned);
240+ }
241+ convert.TransposeAoSDirectToSoA(gmB, rTotal, deintCols, cfg.td->ldb,
242+ gmBc_real, gmBc_imag, cfg.nColsAligned, 0, cfg.kDimOrig);
243+ float aRe = cfg.td->alphaReal;
244+ float aIm = cfg.td->alphaImag;
245+ if (aRe != 1.0f || aIm != 0.0f) {
246+ cfg.buildRowStart = 0;
247+ cfg.buildRowEnd = cfg.kDim;
248+ canonB.BuildBRightPhase2Alpha(gmBc_real, gmBc_imag, cfg.nColsAligned, aRe, aIm);
249+ }
250+ } else {
251+ cfg.buildRowStart = 0;
252+ cfg.buildRowEnd = cfg.kDim;
253+ canonB.BuildPaddedB(gmB, gmBc_real, gmBc_imag, gmGemm);
254+ }
255+ RunPanelLoop(effA, effAStride, gmBc_real, gmBc_imag, gmGemm, gmB);
256+ AscendC::PipeBarrier<PIPE_ALL>();
257+ if (RIGHT) {
258+ canonB.WriteBackPaddedB(gmBc_real, gmBc_imag, gmB, gmGemm, 0, cfg.td->m);
259+ }
260+ }
261+ 
262+ // 原始单AIV处理路径（无分列，workspace按blockIdx索引）
263+ __aicore__ inline void ProcessOneBatchOriginal(int32_t batch)
264+ {
265+ __gm__ float* coreWs = (__gm__ float*)gemmWsBase
266+ + (int64_t)blockIdx * (cfg.td->workspaceOffset / sizeof(float));
267+ int64_t aPlaneSize = (int64_t)cfg.kDim * cfg.kDim;
268+ __gm__ float* gmAc = coreWs;
269+ int64_t bPlaneSize = (int64_t)cfg.kDim * cfg.nColsAligned;
270+ __gm__ float* gmBc_real = coreWs + 2 * aPlaneSize;
271+ __gm__ float* gmBc_imag = gmBc_real + bPlaneSize;
272+ __gm__ float* gmGemm = gmBc_imag + bPlaneSize;
273+ __gm__ float* effA;
274+ int32_t effAStride;
275+ PrepareAMatrix(batch, coreWs, gmBc_real, gmBc_imag, gmGemm, effA, effAStride);
276+ PrepareBAndSolve(batch, effA, effAStride, gmBc_real, gmBc_imag, gmGemm);
277+ }
278+ // 处理单个batch（dual分列路径）：构建A/B工作区、执行Panel循环、回写结果
279+ __aicore__ inline void ProcessOneBatchPadded(int32_t batch)
280+ {
281+ __gm__ float* gmA = (__gm__ float*)gmAArray.GetValue(batch);
282+ __gm__ float* gmB = (__gm__ float*)gmBArray.GetValue(batch);
283+ __gm__ float* coreWs = (__gm__ float*)gemmWsBase
284+ + (int64_t)batch * (cfg.td->workspaceOffset / sizeof(float));
285+ int64_t aPlaneSize = (int64_t)cfg.kDim * cfg.kDim;
286+ __gm__ float* gmAc = coreWs;
287+ int64_t bPlaneSize = (int64_t)cfg.kDim * cfg.nColsAligned;
288+ __gm__ float* gmBc_real = coreWs + 2 * aPlaneSize;
289+ __gm__ float* gmBc_imag = gmBc_real + bPlaneSize;
290+ __gm__ float* gmGemm = gmBc_imag + bPlaneSize;
291+ int32_t half = (int32_t)blockIdx % 2;
292+ __gm__ float* effA;
293+ int32_t effAStride;
294+ 
295+ PrepareAWithSync(gmA, gmAc, gmBc_real, gmBc_imag, gmGemm, half, effA, effAStride);
296+ PrepareBWithSync(gmB, gmBc_real, gmBc_imag, gmGemm, half);
297+ 
298+ RunPanelLoop(effA, effAStride, gmBc_real, gmBc_imag, gmGemm, gmB);
299+ AscendC::PipeBarrier<PIPE_ALL>();
300+ if (RIGHT) {
301+ DualAivBarrier();
302+ // dual-AIV 写回按输出行切分：两核各写一半，消除整矩阵重复转置/回写
303+ int32_t mOrig = cfg.td->m;
304+ int32_t rowMid = ((mOrig / 2 + CtrsmAivCfg::TS - 1) / CtrsmAivCfg::TS) * CtrsmAivCfg::TS;
305+ if (rowMid > mOrig) rowMid = mOrig;
306+ int32_t wbBeg = (half == 0) ? 0 : rowMid;
307+ int32_t wbEnd = (half == 0) ? rowMid : mOrig;
308+ canonB.WriteBackPaddedB(gmBc_real, gmBc_imag, gmB, gmGemm, wbBeg, wbEnd);
309+ }
310+ }
311+ 
312+ // 多核拆分路径：每个 block 的 2 AIV 协同处理一个 split（复用 dual-AIV 同步协议）
313+ __aicore__ inline void ProcessOneBatchSplit(int32_t batch, int32_t half)
314+ {
315+ __gm__ float* gmA = (__gm__ float*)gmAArray.GetValue(batch);
316+ __gm__ float* gmBOrig = (__gm__ float*)gmBArray.GetValue(batch);
317+ __gm__ float* gmB;
318+ if (RIGHT) {
319+ gmB = gmBOrig + (int64_t)cfg.nColsOffset * cfg.td->ldb * 2;
320+ } else {
321+ gmB = gmBOrig;
322+ }
323+ int32_t savedNColsOrig = cfg.nColsOrig;
324+ cfg.nColsOrig = cfg.localNCols;
325+ int32_t aivBlock = (int32_t)blockIdx / 2;
326+ int32_t splitIdx = aivBlock % cfg.td->numSplits;
327+ __gm__ float* batchWs = (__gm__ float*)gemmWsBase
328+ + (int64_t)batch * (cfg.td->workspaceOffset / sizeof(float));
329+ int64_t aPlaneSize = (int64_t)cfg.kDim * cfg.kDim;
330+ __gm__ float* gmAc = batchWs;
331+ int32_t wsNCols = (cfg.td->splitNColsAligned > cfg.td->lastNColsAligned)
332+ ? cfg.td->splitNColsAligned : cfg.td->lastNColsAligned;
333+ int64_t bPlaneSize = (int64_t)cfg.kDim * wsNCols;
334+ int64_t splitBGemmFloats = cfg.td->splitBGemmSize / sizeof(float);
335+ __gm__ float* splitWs = batchWs + 2 * aPlaneSize + (int64_t)splitIdx * splitBGemmFloats;
336+ __gm__ float* gmBc_real = splitWs;
337+ __gm__ float* gmBc_imag = gmBc_real + bPlaneSize;
338+ __gm__ float* gmGemm = gmBc_imag + bPlaneSize;
339+ __gm__ float* effA;
340+ int32_t effAStride;
341+ 
342+ PrepareAWithSync(gmA, gmAc, gmBc_real, gmBc_imag, gmGemm, half, effA, effAStride);
343+ PrepareBSplit(gmB, gmBc_real, gmBc_imag, gmGemm, half);
344+ 
345+ RunPanelLoop(effA, effAStride, gmBc_real, gmBc_imag, gmGemm, gmB);
346+ AscendC::PipeBarrier<PIPE_ALL>();
347+ if (RIGHT) {
348+ DualAivBarrier();
349+ int32_t localM = cfg.localNCols;
350+ int32_t rowMid = ((localM / 2 + CtrsmAivCfg::TS - 1) / CtrsmAivCfg::TS) * CtrsmAivCfg::TS;
351+ if (rowMid > localM) rowMid = localM;
352+ int32_t wbBeg = (half == 0) ? 0 : rowMid;
353+ int32_t wbEnd = (half == 0) ? rowMid : localM;
354+ canonB.WriteBackPaddedB(gmBc_real, gmBc_imag, gmB, gmGemm, wbBeg, wbEnd);
355+ }
356+ cfg.nColsOrig = savedNColsOrig;
357+ }
358+ 
359+ // 共用：Left 模式下按 half 分行构建 B 工作区
360+ __aicore__ inline void PrepareBLeftHalf(__gm__ float* gmB,
361+ __gm__ float* gmBc_real, __gm__ float* gmBc_imag, __gm__ float* gmGemm,
362+ int32_t half)
363+ {
364+ int32_t totalRows = cfg.kDim;
365+ int32_t rowMid = CEIL_ALIGN(totalRows / 2, FLOAT_ALIGN);
366+ if (rowMid > totalRows) rowMid = totalRows;
367+ cfg.buildRowStart = (half == 0) ? 0 : rowMid;
368+ cfg.buildRowEnd = (half == 0) ? rowMid : totalRows;
369+ canonB.BuildPaddedB(gmB, gmBc_real, gmBc_imag, gmGemm);
370+ DualAivBarrier();
371+ }
372+ 
373+ // 共用：Right 模式下按 half 分列解交织 B 并可选 alpha 缩放
374+ __aicore__ inline void PrepareBRightHalf(__gm__ float* gmB,
375+ __gm__ float* gmBc_real, __gm__ float* gmBc_imag,
376+ int32_t half, int32_t rTotal)
377+ {
378+ int32_t nColsAligned = cfg.nColsAligned;
379+ int32_t kDimOrig = cfg.kDimOrig;
380+ int32_t kDim = cfg.kDim;
381+ int32_t cbMid = CEIL_ALIGN(kDim / 2, FLOAT_ALIGN);
382+ if (cbMid > kDim) cbMid = kDim;
383+ int32_t cbBeg = (half == 0) ? 0 : cbMid;
384+ int32_t cbEnd = (half == 0) ? cbMid : kDim;
385+ cfg.buildRowStart = cbBeg;
386+ cfg.buildRowEnd = (cbEnd < kDimOrig) ? cbEnd : kDimOrig;
387+ int32_t nOrig = cfg.td->n;
388+ int32_t deintCols = (nOrig < kDimOrig) ? nOrig : kDimOrig;
389+ convert.TransposeAoSDirectToSoA(gmB, rTotal, deintCols, cfg.td->ldb,
390+ gmBc_real, gmBc_imag, nColsAligned, cbBeg, cbEnd);
391+ float aRe = cfg.td->alphaReal;
392+ float aIm = cfg.td->alphaImag;
393+ if (aRe != 1.0f || aIm != 0.0f) {
394+ canonB.BuildBRightPhase2Alpha(gmBc_real, gmBc_imag, nColsAligned, aRe, aIm);
395+ }
396+ DualAivBarrier();
397+ }
398+ 
399+ // B 工作区准备（split 模式）：按 half 分行清零/解交织，复用 dual-AIV 协议
400+ __aicore__ inline void PrepareBSplit(__gm__ float* gmB,
401+ __gm__ float* gmBc_real, __gm__ float* gmBc_imag, __gm__ float* gmGemm,
402+ int32_t half)
403+ {
404+ if (!RIGHT) {
405+ PrepareBLeftHalf(gmB, gmBc_real, gmBc_imag, gmGemm, half);
406+ return;
407+ }
408+ int32_t nColsOrig = cfg.nColsOrig;
409+ int32_t kDimOrig = cfg.kDimOrig;
410+ int32_t rTotal = (nColsOrig < kDimOrig) ? nColsOrig : kDimOrig;
411+ PrepareBRightHalf(gmB, gmBc_real, gmBc_imag, half, rTotal);
412+ }
413+ 
414+ // A工作区准备：needTA 时单 pass 转置（两核并行各做一半输出列），否则单核直接复制
415+ __aicore__ inline void PrepareAWithSync(__gm__ float* gmA, __gm__ float* gmAc,
416+ __gm__ float* gmBc_real, __gm__ float* gmBc_imag, __gm__ float* gmGemm,
417+ int32_t half, __gm__ float*& effA, int32_t& effAStride)
418+ {
419+ if (!cfg.useOrigA) {
420+ if (cfg.needTA) {
421+ int32_t kDimOrig = cfg.kDimOrig;
422+ int32_t kDim = cfg.kDim;
423+ int32_t wsStride = 2 * kDim;
424+ float signIm = cfg.conjA ? -1.0f : 1.0f;
425+ int32_t colMid = CEIL_ALIGN(kDim / 2, FLOAT_ALIGN);
426+ if (colMid > kDim) colMid = kDim;
427+ int32_t cBeg = (half == 0) ? 0 : colMid;
428+ int32_t cEnd = (half == 0) ? colMid : kDim;
429+ convert.TransposeADirectToAoS(gmA, kDimOrig, cfg.td->lda,
430+ gmAc, wsStride, signIm, cBeg, cEnd, cfg.padOn, kDim);
431+ DualAivBarrier();
432+ } else {
433+ if (half == 0) {
434+ canonA.BuildPaddedA(gmA, gmAc, gmBc_real, gmBc_imag, gmGemm);
435+ }
436+ DualAivBarrier();
437+ }
438+ effA = gmAc;
439+ effAStride = 2 * cfg.kDim;
440+ } else {
441+ effA = gmA;
442+ effAStride = cfg.aEffStride;
443+ }
444+ }
445+ 
446+ // 双 AIV 屏障：AI Core 内两 AIV 各完成本核 workspace 构建后经 mode1 核间同步互等
447+ // （替代原 GM flag 忙等轮询，消除轮询循环与 PIPE_ALL；PIPE_MTE3 保证 GM 写对对方可见）
448+ __aicore__ inline void DualAivBarrier()
449+ {
450+ AscendC::CrossCoreSetFlag<1, PIPE_MTE3>(FLAG_DUAL_AIV);
451+ AscendC::CrossCoreWaitFlag<1, PIPE_MTE3>(FLAG_DUAL_AIV);
452+ }
453+ 
454+ // B工作区准备：设置分行范围、构建补零B，并按 left/right 与 half 完成跨AIV同步
455+ __aicore__ inline void PrepareBWithSync(__gm__ float* gmB,
456+ __gm__ float* gmBc_real, __gm__ float* gmBc_imag, __gm__ float* gmGemm,
457+ int32_t half)
458+ {
459+ if (!RIGHT) {
460+ PrepareBLeftHalf(gmB, gmBc_real, gmBc_imag, gmGemm, half);
461+ return;
462+ }
463+ // Right dual：单 pass 构建（两核各做一半输出行，含 pad）
464+ int32_t mOrig = cfg.td->m;
465+ int32_t nColsOrig = cfg.nColsOrig;
466+ int32_t rTotal = (mOrig < nColsOrig) ? mOrig : nColsOrig;
467+ PrepareBRightHalf(gmB, gmBc_real, gmBc_imag, half, rTotal);
468+ }
469+ 
470+ // 逐Panel执行三角求解，并与AIC核同步完成Trail区域GEMM更新
471+ // 优化：在等待AIC GEMM期间预加载下一Panel的A对角块（与GEMM并行）
472+ __aicore__ inline void RunPanelLoop(__gm__ float* effA, int32_t effAStride,
473+ __gm__ float* gmBcR, __gm__ float* gmBcI,
474+ __gm__ float* gmGemm, __gm__ float* gmB)
475+ {
476+ int32_t nb = cfg.nb, kDim = cfg.kDim, nColsAligned = cfg.nColsAligned;
477+ int32_t numPanels = (kDim + nb - 1) / nb;
478+ AscendC::LocalTensor<float> ubAR = bufPanelA_real.Get<float>();
479+ AscendC::LocalTensor<float> ubAI = bufPanelA_imag.Get<float>();
480+ for (int32_t idx = 0; idx < numPanels; idx++) {
481+ int32_t p = FORWARD ? idx : (numPanels - 1 - idx);
482+ int32_t panelStart = p * nb;
483+ int32_t actualNb = (nb < kDim - panelStart) ? nb : (kDim - panelStart);
484+ int32_t slot = XnegSlot(idx, numPanels, FORWARD);
485+ int64_t xnegPlaneSize = (int64_t)2 * LIM_GROUP * 2 * nb * nColsAligned;
486+ __gm__ float* xR = gmGemm + (int64_t)slot * 2 * nb * nColsAligned;
487+ __gm__ float* xI = gmGemm + xnegPlaneSize + (int64_t)slot * 2 * nb * nColsAligned;
488+ if (idx == 0) {
489+ solver.LoadPanelA(effA, effAStride, ubAR, ubAI, panelStart, actualNb);
490+ }
491+ solver.PanelTrsv(effA, effAStride, gmBcR, gmBcI, nColsAligned,
492+ xR, xI, panelStart, actualNb, PanelHasTrailAiv(idx, numPanels));
493+ if (PanelHasTrailAiv(idx, numPanels)) {
494+ AscendC::PipeBarrier<PIPE_MTE3>();
495+ AscendC::CrossCoreSetFlag<2, PIPE_MTE3>(FLAG_TRSV);
496+ if (!RIGHT) {
497+ canonB.WriteBackPanelRows(gmBcR, gmBcI, gmB, panelStart, actualNb);
498+ }
499+ int32_t nextIdx = idx + 1;
500+ if (nextIdx < numPanels) {
501+ int32_t np = FORWARD ? nextIdx : (numPanels - 1 - nextIdx);
502+ int32_t nextStart = np * nb;
503+ int32_t nextNb = (nb < kDim - nextStart) ? nb : (kDim - nextStart);
504+ solver.LoadPanelA(effA, effAStride, ubAR, ubAI, nextStart, nextNb);
505+ }
506+ AscendC::CrossCoreWaitFlag<2, PIPE_MTE2>(FLAG_GEMM);
507+ } else if (!RIGHT) {
508+ canonB.WriteBackPanelRows(gmBcR, gmBcI, gmB, panelStart, actualNb);
509+ }
510+ }
511+ }
512+ 
513+ // ---- owned resources ----
514+ AscendC::TPipe* pipe;
515+ CtrsmAivCfg cfg;
516+ CtrsmConvert convert;
517+ CtrsmCanonA canonA;
518+ CtrsmCanonB<RIGHT> canonB;
519+ CtrsmPanelSolver<FORWARD> solver;
520+ AscendC::GlobalTensor<uint64_t> gmAArray, gmBArray;
521+ __gm__ uint8_t* gemmWsBase;
522+ uint32_t blockIdx;
523+ BufVecCalc bufPanelA_real, bufPanelA_imag, bufPanelB_real, bufPanelB_imag;
524+ BufVecCalc bufNeg_real, bufNeg_imag, bufNeg_imag_neg, bufRankK, bufRow;
525+ BufVecCalc bufGatherEven, bufGatherOdd;
526+};
527+ 
528+// 四条计算路径的具体类型别名
529+using CtrsmLowerLeft = CtrsmMixAivImpl<true, false>; // 下三角前代 + 左乘
530+using CtrsmLowerRight = CtrsmMixAivImpl<true, true>; // 下三角前代 + 右乘
531+using CtrsmUpperLeft = CtrsmMixAivImpl<false, false>; // 上三角回代 + 左乘
532+using CtrsmUpperRight = CtrsmMixAivImpl<false, true>; // 上三角回代 + 右乘
Aexperimental/aclblasCtrsmBatched2/op_kernel/ctrsm_batched_kernel_aiv_canon_a.h+149-0
@@ -0,0 +1,149 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ctrsm_batched_kernel_aiv_canon_a.h
13+ * \brief A 矩阵规范化器：把用户 A 转成补零对齐的 AoS 工作区（清零→直拷/转置(+共轭)→补单位对角）。
14+ * 转置分快速路径与分块路径，均委托 CtrsmConvert 做解交织/交织/分块转置。
15+ * 与计算路径(forward/right)无关，依赖 needTA/conjA 运行时标志。
16+ */
17+ 
18+#pragma once
19+ 
20+#include "ctrsm_batched_kernel_aiv_cfg.h"
21+#include "ctrsm_batched_kernel_aiv_convert.h"
22+ 
23+class CtrsmCanonA {
24+public:
25+ __aicore__ inline CtrsmCanonA() {}
26+ 
27+ __aicore__ inline void Bind(CtrsmAivCfg* cfg, CtrsmConvert* cvt) { cfg_ = cfg; cvt_ = cvt; }
28+ 
29+ // 构建补零后的A矩阵AoS工作区：清零→直接复制（仅 NoTrans 调用）
30+ __aicore__ inline void BuildPaddedA(__gm__ float* gmA, __gm__ float* gmAc,
31+ __gm__ float* gmTmp1, __gm__ float* gmTmp2,
32+ __gm__ float* gmExtra)
33+ {
34+ int32_t wsStride = 2 * cfg_->kDim;
35+ if (cfg_->padOn) {
36+ cvt_->ZeroGmRows(gmAc, cfg_->kDim, wsStride, wsStride);
37+ }
38+ float signIm = cfg_->conjA ? -1.0f : 1.0f;
39+ CopyADirectAoS(gmA, gmAc, wsStride, signIm);
40+ if (cfg_->padOn) {
41+ PadIdentityRowsAoS(gmAc, wsStride);
42+ }
43+ }
44+ 
45+ // 将AoS工作空间逐行清零已提取到 CtrsmConvert::ZeroGmRows
46+ 
47+ // 直接复制A矩阵AoS到workspace（不转置），可选共轭翻转虚部
48+ __aicore__ inline void CopyADirectAoS(__gm__ float* gmA, __gm__ float* gmAc,
49+ int32_t wsStride, float signIm)
50+ {
51+ int32_t kDimOrig = cfg_->kDimOrig, kDim = cfg_->kDim;
52+ AscendC::GlobalTensor<float> gA, gAc;
53+ gA.SetGlobalBuffer(gmA, (uint32_t)((int64_t)kDimOrig * cfg_->td->lda * 2));
54+ gAc.SetGlobalBuffer(gmAc, (uint32_t)((int64_t)kDim * wsStride));
55+ int32_t srcRowFloats = kDimOrig * 2;
56+ int32_t srcRowAligned = CEIL_ALIGN(srcRowFloats, FLOAT_ALIGN); // = slice UB 步长
57+ int32_t batchRows = cfg_->srcBufFloats / srcRowAligned;
58+ if (batchRows > kDimOrig) batchRows = kDimOrig;
59+ if (batchRows < 1) batchRows = 1;
60+ uint32_t srcStrideBytes = (uint32_t)((cfg_->td->lda * 2 - srcRowFloats) * (int32_t)sizeof(float));
61+ if (signIm == 1.0f) {
62+ CopyDirectPlain(gA, gAc, wsStride, srcRowFloats, srcRowAligned, batchRows, srcStrideBytes);
63+ } else {
64+ CopyDirectConj(gA, gAc, wsStride, srcRowFloats, srcRowAligned, batchRows, srcStrideBytes);
65+ }
66+ }
67+ 
68+ // 纯拷贝路径：批量载入 → 批量写回
69+ __aicore__ inline void CopyDirectPlain(AscendC::GlobalTensor<float>& gA,
70+ AscendC::GlobalTensor<float>& gAc, int32_t wsStride, int32_t srcRowFloats,
71+ int32_t srcRowAligned, int32_t batchRows, uint32_t srcStrideBytes)
72+ {
73+ int32_t kDimOrig = cfg_->kDimOrig;
74+ AscendC::LocalTensor<float> ubTmp = cfg_->bufPanelB_real->Get<float>();
75+ uint32_t ubGap = (uint32_t)(((srcRowAligned - srcRowFloats) * (int32_t)sizeof(float)) / 32);
76+ uint32_t gmGap = (uint32_t)((wsStride - srcRowFloats) * (int32_t)sizeof(float));
77+ for (int32_t bStart = 0; bStart < kDimOrig; bStart += batchRows) {
78+ int32_t bCount = (batchRows < kDimOrig - bStart) ? batchRows : (kDimOrig - bStart);
79+ BatchLoadRows(gA, ubTmp, bStart, bCount, srcRowFloats, srcStrideBytes);
80+ AscendC::PipeBarrier<PIPE_ALL>();
81+ AscendC::DataCopyExtParams wbP((uint16_t)bCount,
82+ (uint32_t)(srcRowFloats * sizeof(float)), ubGap, gmGap, 0);
83+ AscendC::DataCopyPad(gAc[(int64_t)bStart * wsStride], ubTmp, wbP);
84+ AscendC::PipeBarrier<PIPE_ALL>();
85+ }
86+ }
87+ 
88+ // 共轭路径：构建 conjMask → 批量载入 → 各行独立 Mul → 各行独立写回
89+ __aicore__ inline void CopyDirectConj(AscendC::GlobalTensor<float>& gA,
90+ AscendC::GlobalTensor<float>& gAc, int32_t wsStride, int32_t srcRowFloats,
91+ int32_t srcRowAligned, int32_t batchRows, uint32_t srcStrideBytes)
92+ {
93+ int32_t kDimOrig = cfg_->kDimOrig;
94+ AscendC::LocalTensor<float> ubSrc = cfg_->bufPanelB_real->Get<float>();
95+ AscendC::LocalTensor<float> conjMask = cfg_->bufNeg_imag_neg->Get<float>();
96+ for (int32_t j = 0; j < FLOAT_ALIGN / 2; j++) {
97+ conjMask.SetValue(2 * j, 1.0f);
98+ conjMask.SetValue(2 * j + 1, -1.0f);
99+ }
100+ AscendC::PipeBarrier<PIPE_V>();
101+ for (int32_t c = FLOAT_ALIGN; c < srcRowAligned; c += FLOAT_ALIGN) {
102+ AscendC::DataCopy(conjMask[c], conjMask, FLOAT_ALIGN);
103+ AscendC::PipeBarrier<PIPE_V>();
104+ }
105+ for (int32_t bStart = 0; bStart < kDimOrig; bStart += batchRows) {
106+ int32_t bCount = (batchRows < kDimOrig - bStart) ? batchRows : (kDimOrig - bStart);
107+ BatchLoadRows(gA, ubSrc, bStart, bCount, srcRowFloats, srcStrideBytes);
108+ AscendC::PipeBarrier<PIPE_ALL>();
109+ for (int32_t j = 0; j < bCount; j++) {
110+ AscendC::Mul(ubSrc[j * srcRowAligned], ubSrc[j * srcRowAligned], conjMask, srcRowAligned);
111+ }
112+ AscendC::PipeBarrier<PIPE_ALL>();
113+ uint32_t ubGap = (uint32_t)(((srcRowAligned - srcRowFloats) * (int32_t)sizeof(float)) / 32);
114+ uint32_t gmGap = (uint32_t)((wsStride - srcRowFloats) * (int32_t)sizeof(float));
115+ AscendC::DataCopyExtParams wbP((uint16_t)bCount,
116+ (uint32_t)(srcRowFloats * sizeof(float)), ubGap, gmGap, 0);
117+ AscendC::DataCopyPad(gAc[(int64_t)bStart * wsStride], ubSrc, wbP);
118+ AscendC::PipeBarrier<PIPE_ALL>();
119+ }
120+ }
121+ 
122+ // 多行批量载入 A 的 AoS 行到 UB（dstStride=0 自动 32B 对齐打包，slice 步长 = srcRowAligned）
123+ __aicore__ inline void BatchLoadRows(AscendC::GlobalTensor<float>& gA,
124+ AscendC::LocalTensor<float>& ubDst, int32_t bStart, int32_t bCount,
125+ int32_t srcRowFloats, uint32_t srcStrideBytes)
126+ {
127+ if (bCount > 1) {
128+ AscendC::DataCopyExtParams rBatch((uint16_t)bCount,
129+ (uint32_t)(srcRowFloats * sizeof(float)), srcStrideBytes, 0, 0);
130+ AscendC::DataCopyPad(ubDst, gA[(int64_t)bStart * cfg_->td->lda * 2], rBatch, {false, 0, 0, 0});
131+ } else {
132+ AscendC::DataCopyExtParams rRow(1, (uint32_t)(srcRowFloats * sizeof(float)), 0, 0, 0);
133+ AscendC::DataCopyPad(ubDst, gA[(int64_t)bStart * cfg_->td->lda * 2], rRow, {false, 0, 0, 0});
134+ }
135+ }
136+ 
137+ // 对补零行填充单位对角线（real=1.0, imag=0.0）
138+ __aicore__ inline void PadIdentityRowsAoS(__gm__ float* gmAc, int32_t wsStride)
139+ {
140+ AscendC::LocalTensor<float> ub = cfg_->bufRow->Get<float>();
141+ AscendC::GlobalTensor<float> gAc;
142+ gAc.SetGlobalBuffer(gmAc, (uint32_t)((int64_t)cfg_->kDim * wsStride));
143+ cvt_->PadDiagonalIdentity(gAc, ub, cfg_->kDimOrig, cfg_->kDim, wsStride);
144+ }
145+ 
146+private:
147+ CtrsmAivCfg* cfg_;
148+ CtrsmConvert* cvt_;
149+};
Aexperimental/aclblasCtrsmBatched2/op_kernel/ctrsm_batched_kernel_aiv_canon_b.h+318-0
@@ -0,0 +1,318 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ctrsm_batched_kernel_aiv_canon_b.h
13+ * \brief B 矩阵规范化与回写器（按左右乘模板化）：AoS->SoA 构建、alpha 缩放，
14+ * 以及求解结果回写（Left 逐面板行回写 / Right 转置后整体交织回写）。
15+ */
16+ 
17+#pragma once
18+ 
19+#include "ctrsm_batched_kernel_aiv_cfg.h"
20+#include "ctrsm_batched_kernel_aiv_convert.h"
21+ 
22+template <bool RIGHT>
23+class CtrsmCanonB {
24+public:
25+ __aicore__ inline CtrsmCanonB() {}
26+ 
27+ __aicore__ inline void Bind(CtrsmAivCfg* cfg, CtrsmConvert* cvt) { cfg_ = cfg; cvt_ = cvt; }
28+ 
29+ // 构建补零后的B矩阵SoA工作区：清零→解交织→alpha缩放（按buildRowStart/End分行）
30+ __aicore__ inline void BuildPaddedB(__gm__ float* gmB,
31+ __gm__ float* gmBc_real, __gm__ float* gmBc_imag,
32+ __gm__ float* gmTemp)
33+ {
34+ int32_t nColsAligned = cfg_->nColsAligned;
35+ if (cfg_->padOn) {
36+ cvt_->ZeroGmRows(gmBc_real + (int64_t)cfg_->buildRowStart * nColsAligned,
37+ cfg_->buildRowEnd - cfg_->buildRowStart, nColsAligned, nColsAligned);
38+ cvt_->ZeroGmRows(gmBc_imag + (int64_t)cfg_->buildRowStart * nColsAligned,
39+ cfg_->buildRowEnd - cfg_->buildRowStart, nColsAligned, nColsAligned);
40+ }
41+ if (!RIGHT) {
42+ DeinterleaveBLeft(gmB, gmBc_real, gmBc_imag);
43+ } else {
44+ DeinterleaveBRight(gmB, gmBc_real, gmBc_imag, gmTemp);
45+ }
46+ float aRe = cfg_->td->alphaReal;
47+ float aIm = cfg_->td->alphaImag;
48+ if (aRe != 1.0f || aIm != 0.0f) {
49+ AscendC::PipeBarrier<PIPE_ALL>();
50+ AlphaScale(gmBc_real, gmBc_imag, nColsAligned, aRe, aIm);
51+ }
52+ }
53+ 
54+ // 将SoA结果写回原始AoS格式的B矩阵（仅Right模式调用）
55+ // rowBeg/rowEnd 限定输出行范围（dual-AIV 两核各写一半）
56+ __aicore__ inline void WriteBackPaddedB(__gm__ float* gmBcR, __gm__ float* gmBcI,
57+ __gm__ float* gmB, __gm__ float* gmTemp,
58+ int32_t rowBeg, int32_t rowEnd)
59+ {
60+ InterleaveBRight(gmBcR, gmBcI, gmB, gmTemp, rowBeg, rowEnd);
61+ }
62+ 
63+ // Right dual 前端阶段1：清零本核 gmBc 行段(buildRow) + 解交织 [tBeg,tEnd) 行到 temp
64+ __aicore__ inline void BuildBRightPhase1(__gm__ float* gmB,
65+ __gm__ float* gmBc_real, __gm__ float* gmBc_imag, __gm__ float* gmTemp,
66+ int32_t tBeg, int32_t tEnd)
67+ {
68+ int32_t nColsAligned = cfg_->nColsAligned;
69+ if (cfg_->padOn) {
70+ cvt_->ZeroGmRows(gmBc_real + (int64_t)cfg_->buildRowStart * nColsAligned,
71+ cfg_->buildRowEnd - cfg_->buildRowStart, nColsAligned, nColsAligned);
72+ cvt_->ZeroGmRows(gmBc_imag + (int64_t)cfg_->buildRowStart * nColsAligned,
73+ cfg_->buildRowEnd - cfg_->buildRowStart, nColsAligned, nColsAligned);
74+ }
75+ DeinterleaveBRightToTemp(gmB, gmTemp, tBeg, tEnd);
76+ }
77+ 
78+ // Right dual 前端阶段2：转置 temp→gmBc 本核输出行段 [cbBeg,cbEnd) + alpha 缩放本核行段
79+ __aicore__ inline void BuildBRightPhase2(__gm__ float* gmBc_real, __gm__ float* gmBc_imag,
80+ __gm__ float* gmTemp, int32_t cbBeg, int32_t cbEnd)
81+ {
82+ TransposeBTempToBc(gmBc_real, gmBc_imag, gmTemp, cbBeg, cbEnd);
83+ float aRe = cfg_->td->alphaReal;
84+ float aIm = cfg_->td->alphaImag;
85+ if (aRe != 1.0f || aIm != 0.0f) {
86+ AlphaScale(gmBc_real, gmBc_imag, cfg_->nColsAligned, aRe, aIm);
87+ }
88+ }
89+ 
90+ // 仅 alpha 缩放（buildRowStart/End 已在外部设置）
91+ __aicore__ inline void BuildBRightPhase2Alpha(__gm__ float* gmBc_real, __gm__ float* gmBc_imag,
92+ int32_t effLdb, float aRe, float aIm)
93+ {
94+ AlphaScale(gmBc_real, gmBc_imag, effLdb, aRe, aIm);
95+ }
96+ 
97+ // 将已求解的nb行从SoA数据交织回写到用户B矩阵（Left模式逐面板回写）
98+ // 单列分块时直接从 UB 读取（省 MTE2），多列分块退回 GM 搬入
99+ __aicore__ inline void WriteBackPanelRows(__gm__ float* gmBcR, __gm__ float* gmBcI,
100+ __gm__ float* gmB,
101+ int32_t panelStart, int32_t actualNb)
102+ {
103+ int32_t nColsAligned = cfg_->nColsAligned;
104+ int32_t mOrig = cfg_->td->m, nOrig = cfg_->td->n;
105+ int32_t endRow = panelStart + actualNb;
106+ if (endRow > mOrig) endRow = mOrig;
107+ int32_t localNCols = cfg_->localNCols;
108+ int32_t clipEnd = (localNCols < nOrig) ? localNCols : nOrig;
109+ int32_t wbColStart = (cfg_->colStart < clipEnd) ? cfg_->colStart : clipEnd;
110+ int32_t wbColEnd = (cfg_->colEnd < clipEnd) ? cfg_->colEnd : clipEnd;
111+ int32_t wbCount = wbColEnd - wbColStart;
112+ if (wbCount <= 0) return;
113+ int32_t gmBColStart = wbColStart + cfg_->nColsOffset;
114+ if (cfg_->colTile >= nColsAligned) {
115+ WriteBackPanelRowsFromUB(gmB, panelStart, endRow, gmBColStart, wbCount);
116+ } else {
117+ WriteBackPanelRowsFromGM(gmBcR, gmBcI, gmB, panelStart, endRow, wbColStart, wbCount, gmBColStart);
118+ }
119+ }
120+ 
121+ // 单列分块：直接从 UB 已有的 ubBR/ubBI 读取，批量交织后 1 条 strided DMA 写出
122+ __aicore__ inline void WriteBackPanelRowsFromUB(__gm__ float* gmB,
123+ int32_t panelStart, int32_t endRow, int32_t wbColStart, int32_t wbCount)
124+ {
125+ int32_t mOrig = cfg_->td->m;
126+ int32_t wbCountA = CEIL_ALIGN(wbCount, FLOAT_ALIGN);
127+ AscendC::LocalTensor<float> ubBR = cfg_->bufPanelB_imag->Get<float>();
128+ AscendC::LocalTensor<float> ubBI = cfg_->bufNeg_imag->Get<float>();
129+ AscendC::LocalTensor<float> ubR = cfg_->bufPanelB_real->Get<float>();
130+ AscendC::LocalTensor<float> ub = cfg_->bufRow->Get<float>();
131+ int32_t rowCount = endRow - panelStart;
132+ AscendC::DataCopy(ubR, ubBI, rowCount * wbCountA);
133+ AscendC::PipeBarrier<PIPE_V>();
134+ cvt_->BuildInterleaveOffsets(wbCount);
135+ AscendC::LocalTensor<int32_t> offInter = cfg_->bufNeg_imag->Get<int32_t>();
136+ // 批量交织：16 行全部 InterleaveRow 到 bufNeg_real（作为连续大 buffer 使用）
137+ // bufNeg_real 大小 = nb*colTile = 1024f，需要 rowCount*wbCount*2 = 16*128 = 2048f
138+ // 不够！改用 bufPanelB_real 的后半段（4096f 总量，前 rowCount*wbCountA 已被 ubR 占）
139+ // ubR 占 rowCount*wbCountA = 16*64 = 1024f，bufPanelB_real 剩余 4096-1024 = 3072f >= 2048f ✓
140+ int32_t outRowFloats = wbCount * 2;
141+ int32_t outRowAligned = CEIL_ALIGN(outRowFloats, FLOAT_ALIGN);
142+ AscendC::LocalTensor<float> ubOutAll = cfg_->bufPanelB_real->Get<float>();
143+ int32_t outBase = rowCount * wbCountA; // ubR 占用的尾部之后
144+ AscendC::LocalTensor<float> rowReal = cfg_->bufNeg_imag_neg->Get<float>();
145+ for (int32_t i = 0; i < rowCount; i++) {
146+ int32_t rowUbOff = i * wbCountA;
147+ AscendC::DataCopy(rowReal, ubBR[rowUbOff], wbCountA);
148+ AscendC::DataCopy(ub, ubR[rowUbOff], wbCountA);
149+ AscendC::LocalTensor<float> ubOutRow = ubOutAll[outBase + i * outRowAligned];
150+ cvt_->InterleaveRow(rowReal, ub, ubOutRow, offInter, wbCount);
151+ }
152+ // 统一 strided DMA 写出
153+ AscendC::GlobalTensor<float> gB;
154+ gB.SetGlobalBuffer(gmB, (uint32_t)((int64_t)mOrig * cfg_->td->ldb * 2));
155+ int64_t gbOff = (int64_t)panelStart * cfg_->td->ldb * 2 + wbColStart * 2;
156+ uint32_t ubOutGap32B = (uint32_t)(((outRowAligned - outRowFloats) * sizeof(float)) / 32);
157+ uint32_t gmGapBytes = (uint32_t)((cfg_->td->ldb * 2 - outRowFloats) * sizeof(float));
158+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(0);
159+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(0);
160+ if (((outRowAligned - outRowFloats) * (int32_t)sizeof(float)) % 32 == 0) {
161+ AscendC::DataCopyExtParams wbOut((uint16_t)rowCount,
162+ (uint32_t)(outRowFloats * sizeof(float)), ubOutGap32B, gmGapBytes, 0);
163+ AscendC::DataCopyPad(gB[gbOff], ubOutAll[outBase], wbOut);
164+ } else {
165+ AscendC::DataCopyExtParams wbOut(1, (uint32_t)(outRowFloats * sizeof(float)), 0, 0, 0);
166+ for (int32_t i = 0; i < rowCount; i++) {
167+ AscendC::DataCopyPad(gB[gbOff], ubOutAll[outBase + i * outRowAligned], wbOut);
168+ gbOff += cfg_->td->ldb * 2;
169+ }
170+ }
171+ AscendC::PipeBarrier<PIPE_ALL>();
172+ }
173+ 
174+ // 多列分块：从 GM (gmBc) 搬入 SoA 数据再交织写出
175+ // 当 wbCount 超出 UB 容量时按列分 chunk 处理
176+ __aicore__ inline void WriteBackPanelRowsFromGM(__gm__ float* gmBcR, __gm__ float* gmBcI,
177+ __gm__ float* gmB, int32_t panelStart, int32_t endRow,
178+ int32_t wbColStart, int32_t wbCount, int32_t gmBColStart)
179+ {
180+ int32_t kDim = cfg_->kDim, nColsAligned = cfg_->nColsAligned;
181+ int32_t mOrig = cfg_->td->m;
182+ AscendC::LocalTensor<float> ub = cfg_->bufRow->Get<float>();
183+ AscendC::LocalTensor<float> ubR = cfg_->bufPanelB_real->Get<float>();
184+ AscendC::LocalTensor<float> ubOut = cfg_->bufNeg_real->Get<float>();
185+ AscendC::GlobalTensor<float> gBcR, gBcI, gB;
186+ gBcR.SetGlobalBuffer(gmBcR, (uint32_t)((int64_t)kDim * nColsAligned));
187+ gBcI.SetGlobalBuffer(gmBcI, (uint32_t)((int64_t)kDim * nColsAligned));
188+ gB.SetGlobalBuffer(gmB, (uint32_t)((int64_t)mOrig * cfg_->td->ldb * 2));
189+ constexpr int32_t EVT_ID = 0;
190+ int32_t ldbStride = cfg_->td->ldb * 2;
191+ 
192+ // InterleaveRow 需要: ubR 容量 >= 2*countA, ubOut 容量 >= 2*countA
193+ // countA = CEIL_ALIGN(chunkCols, FLOAT_ALIGN)
194+ // ubR = bufPanelB_real (srcBufFloats), ubOut = bufNeg_real (nb*colTile)
195+ int32_t maxColsR = cfg_->srcBufFloats / 2;
196+ int32_t maxColsOut = cfg_->nb * cfg_->colTile / 2;
197+ int32_t maxChunkCols = (maxColsR < maxColsOut) ? maxColsR : maxColsOut;
198+ maxChunkCols = (maxChunkCols / FLOAT_ALIGN) * FLOAT_ALIGN;
199+ if (maxChunkCols < FLOAT_ALIGN) maxChunkCols = FLOAT_ALIGN;
200+ 
201+ for (int32_t cs = 0; cs < wbCount; cs += maxChunkCols) {
202+ int32_t chunkCols = (maxChunkCols < wbCount - cs) ? maxChunkCols : (wbCount - cs);
203+ cvt_->BuildInterleaveOffsets(chunkCols);
204+ AscendC::LocalTensor<int32_t> offInter = cfg_->bufNeg_imag->Get<int32_t>();
205+ int64_t gbOff = (int64_t)panelStart * ldbStride + (gmBColStart + cs) * 2;
206+ int64_t bcOff = (int64_t)panelStart * nColsAligned + wbColStart + cs;
207+ for (int32_t i = panelStart; i < endRow; i++) {
208+ AscendC::DataCopyPad(ubR, gBcR[bcOff],
209+ {1, (uint32_t)(chunkCols * sizeof(float)), 0, 0, 0}, {false, 0, 0, 0});
210+ AscendC::DataCopyPad(ub, gBcI[bcOff],
211+ {1, (uint32_t)(chunkCols * sizeof(float)), 0, 0, 0}, {false, 0, 0, 0});
212+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVT_ID);
213+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVT_ID);
214+ if (i > panelStart) {
215+ AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(EVT_ID);
216+ }
217+ cvt_->InterleaveRow(ubR, ub, ubOut, offInter, chunkCols);
218+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVT_ID);
219+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVT_ID);
220+ AscendC::DataCopyPad(gB[gbOff], ubOut,
221+ {1, (uint32_t)(chunkCols * 2 * sizeof(float)), 0, 0, 0});
222+ AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(EVT_ID);
223+ gbOff += ldbStride;
224+ bcOff += nColsAligned;
225+ }
226+ AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(EVT_ID);
227+ }
228+ }
229+ 
230+private:
231+ // 将SoA工作空间逐行清零已提取到 CtrsmConvert::ZeroGmRows
232+ 
233+ // Deinterleave helper methods split to separate file
234+ #include "ctrsm_batched_kernel_aiv_canon_b_deinterleave.h"
235+ 
236+ // 阶段2：将temp转置到gmBc工作区，仅产出[cbBeg,cbEnd)输出行（=temp列，dual两核各转一段）
237+ __aicore__ inline void TransposeBTempToBc(__gm__ float* gmBcR, __gm__ float* gmBcI,
238+ __gm__ float* gmTemp, int32_t cbBeg, int32_t cbEnd)
239+ {
240+ int32_t kDim = cfg_->kDim, kDimOrig = cfg_->kDimOrig;
241+ int32_t nColsAligned = cfg_->nColsAligned, nColsOrig = cfg_->nColsOrig;
242+ if (cbEnd > kDimOrig) cbEnd = kDimOrig;
243+ if (cbBeg < 0) cbBeg = 0;
244+ if (cbBeg >= cbEnd) return;
245+ __gm__ float* tempR = gmTemp;
246+ __gm__ float* tempI = gmTemp + (int64_t)nColsOrig * kDim;
247+ cvt_->TiledTransposePair(tempR, tempI, nColsOrig, kDimOrig, kDim, gmBcR, gmBcI, nColsAligned,
248+ cbBeg, cbEnd);
249+ }
250+ 
251+ // 对B矩阵乘以复数alpha系数（按buildRow范围分行处理）
252+ __aicore__ inline void AlphaScale(__gm__ float* gmBcR, __gm__ float* gmBcI,
253+ int32_t effLdb, float aRe, float aIm)
254+ {
255+ int32_t kDim = cfg_->kDim, nColsAligned = cfg_->nColsAligned, colTile = cfg_->colTile;
256+ AscendC::LocalTensor<float> ubR = cfg_->bufPanelB_real->Get<float>();
257+ AscendC::LocalTensor<float> ubI = cfg_->bufPanelB_imag->Get<float>();
258+ AscendC::LocalTensor<float> ubTmp = cfg_->bufRankK->Get<float>();
259+ AscendC::GlobalTensor<float> gR, gI;
260+ gR.SetGlobalBuffer(gmBcR, (uint32_t)((int64_t)kDim * effLdb));
261+ gI.SetGlobalBuffer(gmBcI, (uint32_t)((int64_t)kDim * effLdb));
262+ int32_t ct = (colTile < nColsAligned) ? colTile : nColsAligned;
263+ for (int32_t i = cfg_->buildRowStart; i < cfg_->buildRowEnd; i++) {
264+ for (int32_t cs = 0; cs < nColsAligned; cs += ct) {
265+ AlphaScaleChunk(gR, gI, ubR, ubI, ubTmp, effLdb, i, cs, ct, aRe, aIm);
266+ }
267+ }
268+ }
269+ 
270+ // 对一个分块执行复数alpha缩放：newR = aRe*R - aIm*I, newI = aRe*I + aIm*R
271+ __aicore__ inline void AlphaScaleChunk(
272+ AscendC::GlobalTensor<float>& gR, AscendC::GlobalTensor<float>& gI,
273+ AscendC::LocalTensor<float>& ubR, AscendC::LocalTensor<float>& ubI,
274+ AscendC::LocalTensor<float>& ubTmp,
275+ int32_t effLdb, int32_t row, int32_t cs, int32_t ct, float aRe, float aIm)
276+ {
277+ int32_t nColsAligned = cfg_->nColsAligned;
278+ int32_t cw = (cs + ct <= nColsAligned) ? ct : (nColsAligned - cs);
279+ int32_t cwA = CEIL_ALIGN(cw, FLOAT_ALIGN);
280+ AscendC::DataCopyExtParams rp(1, (uint32_t)(cw * sizeof(float)), 0, 0, 0);
281+ AscendC::DataCopyPadExtParams<float> pad{false, 0, 0, 0};
282+ AscendC::DataCopyPad(ubR, gR[(int64_t)row * effLdb + cs], rp, pad);
283+ AscendC::DataCopyPad(ubI, gI[(int64_t)row * effLdb + cs], rp, pad);
284+ AscendC::PipeBarrier<PIPE_ALL>();
285+ AscendC::Muls(ubTmp, ubR, aRe, cwA);
286+ AscendC::PipeBarrier<PIPE_V>();
287+ AscendC::Axpy(ubTmp, ubI, -aIm, cwA);
288+ AscendC::PipeBarrier<PIPE_V>();
289+ AscendC::Muls(ubI, ubI, aRe, cwA);
290+ AscendC::PipeBarrier<PIPE_V>();
291+ AscendC::Axpy(ubI, ubR, aIm, cwA);
292+ AscendC::PipeBarrier<PIPE_V>();
293+ AscendC::DataCopy(ubR, ubTmp, cwA);
294+ AscendC::PipeBarrier<PIPE_V>();
295+ AscendC::DataCopyPad(gR[(int64_t)row * effLdb + cs], ubR, rp);
296+ AscendC::DataCopyPad(gI[(int64_t)row * effLdb + cs], ubI, rp);
297+ AscendC::PipeBarrier<PIPE_ALL>();
298+ }
299+ 
300+ // 右侧模式回写：单 pass SoA→AoS 转置直接写回 B（替代 TiledTransposePair + 逐行 Interleave）
301+ __aicore__ inline void InterleaveBRight(__gm__ float* gmBcR, __gm__ float* gmBcI,
302+ __gm__ float* gmB, __gm__ float* gmTemp,
303+ int32_t rowBeg, int32_t rowEnd)
304+ {
305+ int32_t kDimOrig = cfg_->kDimOrig;
306+ int32_t nColsAligned = cfg_->nColsAligned, nColsOrig = cfg_->nColsOrig;
307+ int32_t mOrig = cfg_->td->m;
308+ if (rowBeg < 0) rowBeg = 0;
309+ if (rowEnd > mOrig) rowEnd = mOrig;
310+ if (rowBeg >= rowEnd) return;
311+ int32_t ldb2 = cfg_->td->ldb * 2;
312+ cvt_->TransposeSoADirectToAoS(gmBcR, gmBcI, kDimOrig, nColsOrig, nColsAligned,
313+ gmB, ldb2, rowBeg, rowEnd);
314+ }
315+ 
316+ CtrsmAivCfg* cfg_;
317+ CtrsmConvert* cvt_;
318+};