Star573
199
代码
代码
Issues23
Pull Requests31
Actions
项目讨论
Wiki
分析
Star573
199
已合并
[matmul_tutorial] tutorial文档补充负载均衡图片以及全载模板说明 #158
Hu1L1创建于 4月13日
[matmul_tutorial] tutorial文档补充负载均衡图片以及全载模板说明 #158
已合并
Hu1L1创建于 4月13日
从Hu1L1/ops-samplesgraph合入到cann/cann-samplesmaster
共 16 个文件变更+45-43
DSamples/2_Performance/matmul_story/docs/images/0-nopingpong-pipeline.png+0-0
DSamples/2_Performance/matmul_story/docs/images/1-pingpong-pipeline.png+0-0
ASamples/2_Performance/matmul_story/docs/images/chart_tail_load_balance.png+0-0
DSamples/2_Performance/matmul_story/docs/images/demo0_base.png+0-0
DSamples/2_Performance/matmul_story/docs/images/demo2_tail_time.png+0-0
DSamples/2_Performance/matmul_story/docs/images/demo3_fix_wait.png+0-0
DSamples/2_Performance/matmul_story/docs/images/demo3_tail_time.png+0-0
DSamples/2_Performance/matmul_story/docs/images/demo4_unitflag.png+0-0
ASamples/2_Performance/matmul_story/docs/images/perf_demo_steps.png+0-0
ASamples/2_Performance/matmul_story/docs/images/unitflag_compare.png+0-0
MSamples/2_Performance/matmul_story/docs/quant_matmul_mxfp4_tutorials.md+39-37
@@ -40,7 +40,7 @@
40数据路径可概括为：**GM →（MTE2）→ L1 →（MTE1）→ L0 → MMAD → L0C →（FIXPIPE）→ GM**。40数据路径可概括为：**GM →（MTE2）→ L1 →（MTE1）→ L0 → MMAD → L0C →（FIXPIPE）→ GM**。
41 41 
42<div align="center">42<div align="center">
43 <img src="images/image23.png" width="800" alt="MXFP4 数据搬运总览" />43 <img src="images/image23.png" width="1500" alt="MXFP4 数据搬运总览" />
44</div>44</div>
45 45 
46### 符号与约定46### 符号与约定
@@ -51,6 +51,18 @@
51- `blockNum`：Device 侧的AIC核数。51- `blockNum`：Device 侧的AIC核数。
52- `round`：各个核进行计算的轮数，并行计算情况下，一轮中32个核会完成32个tile块的计算，不满32个tile块则仅有部分核参与计算。52- `round`：各个核进行计算的轮数，并行计算情况下，一轮中32个核会完成32个tile块的计算，不满32个tile块则仅有部分核参与计算。
53 53 
54### 优化总览与注意事项
55 
56下图在四个代表性 Shape 上汇总各优化 Step 的耗时轨迹，括号内为相对**上一Step**的加速变化比例。
57 
58在完整的代码实现中，是否使用A全载模板属于 **多约束下的联合决策**，并非由单一静态Shape决定，而在`tutorial`下，为简化代码实现与方便演示，**并未完全实现Tiling侧的逻辑**，因此图中三条Case无demo7数据。`tutorial demo7` 和`recipe A_full_load`模板在特定Shape下行为的差异，以`recipe`为准。
59 
60<div align="center">
61 <img src="images/perf_demo_steps.png" width="1500" alt="total_steps" />
62</div>
63 
64特别说明，性能数据相差在1us以内可以认为是硬件执行带来的轻微波动，属于正常现象。
65 
54---66---
55 67 
56<a id="tutorial-step-0"></a>68<a id="tutorial-step-0"></a>
@@ -98,7 +110,7 @@ Step 0 实现一个基础的mxfp4的matmul调用Demo，承担 **基准** 角色
98显然L1、L0级上需要开启双Buffer。MTE1读取L1上ping块数据时，MTE2写入L1pong块缓存。MTE1读取L1上pong块数据时，MTE2写入L1ping块缓存。由此，L1 级 Ping-Pong 使 MTE2 与 MTE1 重叠。MMAD计算读取L0A/L0Bping块缓存时，MTE1写入L0A/L0Bpong块缓存，MMAD计算读取L0A/L0Bpong块缓存时，MTE1写入L0A/L0Bping块缓存。由此，L0 级 Ping-Pong 使 MTE1 与 MMAD 重叠。未开启PingPong和开启Pingpong的预期流水如下。110显然L1、L0级上需要开启双Buffer。MTE1读取L1上ping块数据时，MTE2写入L1pong块缓存。MTE1读取L1上pong块数据时，MTE2写入L1ping块缓存。由此，L1 级 Ping-Pong 使 MTE2 与 MTE1 重叠。MMAD计算读取L0A/L0Bping块缓存时，MTE1写入L0A/L0Bpong块缓存，MMAD计算读取L0A/L0Bpong块缓存时，MTE1写入L0A/L0Bping块缓存。由此，L0 级 Ping-Pong 使 MTE1 与 MMAD 重叠。未开启PingPong和开启Pingpong的预期流水如下。
99 111 
100<div align="center">112<div align="center">
101 <img src="images/double-buffer.png" width="800" alt="Ping-Pong 流水预期" />113 <img src="images/double-buffer.png" width="1500" alt="Ping-Pong 流水预期" />
102</div>114</div>
103 115 
104### 1.3 典型 Shape 116### 1.3 典型 Shape
@@ -139,19 +151,19 @@ Step 0 实现一个基础的mxfp4的matmul调用Demo，承担 **基准** 角色
139 151 
140### 2.4 示例 Case 与流水分析152### 2.4 示例 Case 与流水分析
141 153 
142**Case_1280_512_4096（多 tile、便于对比行列访问对流水的影响）** 154**Case_8192_1024_1024（多 tile、便于对比行列访问对流水的影响）**
143- Shape：`m=1280, k=512, n=4096`。 155- Shape：`m=8192, k=1024, n=1024`。
144- 在 `baseM=baseN=256` 下，`totalCnt=5×16=80`。 156- 在 `baseM=baseN=256` 下，`totalCnt=32×4=128`。
145 157 
146 158 
147<div align="center">159<div align="center">
148 <img src="images/image28.png" width="1800" alt="SWAT 效果对比" />160 <img src="images/image28.png" width="1500" alt="SWAT 效果对比" />
149</div>161</div>
150 162 
151 163 
152### 2.5 代码索引比对164### 2.5 代码索引比对
153 165 
154- **样例根目录**：[3_block_swat](../matmul_tutorials/3_block_swat/)166- **样例根目录**：[2_block_swat](../matmul_tutorials/2_block_swat/)
155 167 
156---168---
157 169 
@@ -163,10 +175,6 @@ Step 0 实现一个基础的mxfp4的matmul调用Demo，承担 **基准** 角色
163 175 
164在 **Step 3** 中，我们对 block 的调度进行了优化。但注意到，由于 **单个 tile 大小** 在调度阶段确定（`baseM×baseN`及边界 `tailM`/`tailN`），固定的切分策略下，**尾轮** 容易出现 tile 块个数小于 `blockNum` 核数的情况，在此场景下，若各 tile 仍作为单块分配至不同核，则尾轮可参与的核数会小于 `blockNum`，部分核无法参与计算，导致尾轮吞吐量下降，对于此种情况，我们需要考虑一种新的切分策略。176在 **Step 3** 中，我们对 block 的调度进行了优化。但注意到，由于 **单个 tile 大小** 在调度阶段确定（`baseM×baseN`及边界 `tailM`/`tailN`），固定的切分策略下，**尾轮** 容易出现 tile 块个数小于 `blockNum` 核数的情况，在此场景下，若各 tile 仍作为单块分配至不同核，则尾轮可参与的核数会小于 `blockNum`，部分核无法参与计算，导致尾轮吞吐量下降，对于此种情况，我们需要考虑一种新的切分策略。
165 177 
166<div align="center">
167 <img src="images/demo2_tail_time.png" width="1500" alt="demo2_tail_time" />
168</div>
169 
170### 3.2 优化思路178### 3.2 优化思路
171 179 
172在 **沿用 Step 3 式多核排布**（或同类 SWAT 调度）的前提下，对尾轮 tile 块实施 **`mTailTile×nTailTile` 二次切分**，如图所示，减小尾轮 tile 块大小，从而增加尾轮 tile 块数量，分配到更多的核上计算，**提高并行度**。180在 **沿用 Step 3 式多核排布**（或同类 SWAT 调度）的前提下，对尾轮 tile 块实施 **`mTailTile×nTailTile` 二次切分**，如图所示，减小尾轮 tile 块大小，从而增加尾轮 tile 块数量，分配到更多的核上计算，**提高并行度**。
@@ -179,23 +187,23 @@ Step 0 实现一个基础的mxfp4的matmul调用Demo，承担 **基准** 角色
179 187 
180- 尾轮 **待处理 tile 数** 显著小于 `blockNum`，例如 32 核下尾轮仅有 4 个 tile 块，通常保证二次切分后新 tile 数量可增至接近核数。 188- 尾轮 **待处理 tile 数** 显著小于 `blockNum`，例如 32 核下尾轮仅有 4 个 tile 块，通常保证二次切分后新 tile 数量可增至接近核数。
181 189 
182### 3.4 示例 Case 与流水分析190### 3.4 示例 Case
183 191 
184**Case_1280_512_4096（多 tile、便于对比行列访问对流水的影响）** 192**Case_8448_4096_4096（多 tile、便于对比行列访问对流水的影响）**
185- Shape：`m=1280, k=512, n=4096`。 193- Shape：`m=8448, k=512, n=4096`。
186- 在 `baseM=baseN=256` 下，`totalCnt=5×16=80`。 194- 在 `baseM=baseN=256` 下，`totalCnt=33*16=528`。
187- 若 `blockNum=32`，则 `round=3`：**尾轮余 16 个 tile**，每个tile**大小为（256， 256）**，如果无尾轮负载均衡，则出现16个核计算，16个核空闲的情况，而在尾轮负载均衡情况下（当前固定n方向切分成2块），则**剩余的16个tile**被切分为**32个tile**，每个**tile大小为（256, 128）**。 195- 若 `blockNum=32`，则 `round=17`：**尾轮余 16 个 tile**，每个tile**大小为（256， 256）**，如果无尾轮负载均衡，则出现16个核计算，16个核空闲的情况，而在尾轮负载均衡情况下（当前固定n方向切分成2块），则**剩余的16个tile**被切分为**32个tile**，每个**tile大小为（256, 128）**。
188- **流水关注点**：**Step 4** 引入尾块子划分后，尾轮 **并行 MMAD 实例数与耗时** 的变化。196- **流水关注点**：**Step 4** 引入尾块子划分后，尾轮 **并行 MMAD 实例数与耗时** 的变化。
189 197 
190从下图中可以看出，在对尾轮tile块进行拆分后，对比拆分前，尾轮tile块计算的耗时显著减少198从下图中可以看出，在对尾轮tile块进行拆分后，尾轮中32个核全部得到了利用，负载更加均衡，最长耗时也相应下降。
Hu1L1
Hu1L1Hu1L14月13日

前后case对齐一下

likedislike
191 199 
192<div align="center">200<div align="center">
193 <img src="images/demo3_tail_time.png" width="1500" alt="demo3_tail_time" />201 <img src="images/chart_tail_load_balance.png" width="1500" alt="tail_load_balance" />
194</div>202</div>
195 203 
196### 3.5 代码索引比对204### 3.5 代码索引比对
197 205 
198- **样例根目录**：[4_last_round_tile_balance](../matmul_tutorials/4_last_round_tile_balance/)206- **样例根目录**：[3_last_round_tile_balance](../matmul_tutorials/3_last_round_tile_balance/)
199 207 
200---208---
201 209 
@@ -207,32 +215,25 @@ Step 0 实现一个基础的mxfp4的matmul调用Demo，承担 **基准** 角色
207 215 
208未开启 **UnitFlag** 时，**`FIXPIPE` 须待 `MMAD` 指令全部执行完毕后才能开始将 `L0C` 结果搬出**。因此 **`MMAD` 整段计算** 与 **`FIXPIPE` 搬出** 之间是 **整段完成依赖**：须先满足 **全量 `MMAD` 结束** 这一条件，`FIXPIPE` 方可启动。在 **无法用 L0C Double Buffer** 等机制缓冲、掩盖该依赖时，**`FIXPIPE` 在 `MMAD` 运行期间处于等待**，**两段流水难以时间重叠**，造成性能损失，如图所示。216未开启 **UnitFlag** 时，**`FIXPIPE` 须待 `MMAD` 指令全部执行完毕后才能开始将 `L0C` 结果搬出**。因此 **`MMAD` 整段计算** 与 **`FIXPIPE` 搬出** 之间是 **整段完成依赖**：须先满足 **全量 `MMAD` 结束** 这一条件，`FIXPIPE` 方可启动。在 **无法用 L0C Double Buffer** 等机制缓冲、掩盖该依赖时，**`FIXPIPE` 在 `MMAD` 运行期间处于等待**，**两段流水难以时间重叠**，造成性能损失，如图所示。
209 217 
210<div align="center">
211 <img src="images/demo3_fix_wait.png" width="1500" alt="demo3_fix_wait" />
212</div>
213 
214### 4.2 优化思路218### 4.2 优化思路
215 219 
216**UnitFlag** 为 **`MMAD` 计算指令** 与 **`FIXPIPE` 数据搬运指令** 提供基于内存访问的 **细粒度同步（512B 粒度）**。开启后，`MMAD` **每计算完 512B 数据**，`FIXPIPE` **即可搬出对应数据块**，从而在 **无法开启 L0C Double-Buffer** 的条件下 **提高计算与搬出流水的并行度**。220**UnitFlag** 为 **`MMAD` 计算指令** 与 **`FIXPIPE` 数据搬运指令** 提供基于内存访问的 **细粒度同步（512B 粒度）**。开启后，`MMAD` **每计算完 512B 数据**，`FIXPIPE` **即可搬出对应数据块**，从而在 **无法开启 L0C Double-Buffer** 的条件下 **提高计算与搬出流水的并行度**。开启**UnitFlag**前后流水对比如下示意图所示：
217 221 
218<div align="center">222<div align="center">
219 <img src="images/image34.png" width="1500" alt="UnitFlag 流水预期（摘自《性能优化指南》）" />223 <img src="images/image34.png" width="1500" alt="UnitFlag 流水预期（摘自《性能优化指南》）" />
220</div>224</div>
221 225 
222### 4.3 典型 Shape 226### 4.3 流水分析
223 227 
224本节不绑定具体 `m,k,n` 数值。228本节不绑定特定Shape，优化前后流水图对比如下所示，可以明显看到在UnitFlag优化以后，MMAD的执行不在等待FIXPIPE，很好的解决了MMAD断流的问题：
225 229 
226### 4.4 示例 Case 与流水分析
227 
228在下图中，相邻基本块的MMAD指令不再强制等待FIXPIPE，一定程度上提高了流水并行度
229<div align="center">230<div align="center">
230 <img src="images/demo4_unitflag.png" width="1500" alt="demo4_unitflag" />231 <img src="images/unitflag_compare.png" width="1500" alt="total_steps" />
231</div>232</div>
232 233 
233### 4.5 代码索引比对234### 4.4 代码索引比对
234 235 
235- **样例根目录**：[5_unit_flag](../matmul_tutorials/5_unit_flag/)236- **样例根目录**：[4_unit_flag](../matmul_tutorials/4_unit_flag/)
236 237 
237---238---
238 239 
@@ -250,7 +251,7 @@ L1上存在2个bank，前面一半是和后一半分别是2个bank。如果写
250 251 
251下图是 解决L1-BANK冲突前后L1的数据排布图。252下图是 解决L1-BANK冲突前后L1的数据排布图。
252<div align="center">253<div align="center">
253 <img src="images/L1-bankconflict-Graph.png" width="800" alt="修改bank冲突后L1数据排布图" />254 <img src="images/L1-bankconflict-Graph.png" width="1500" alt="修改bank冲突后L1数据排布图" />
254</div>255</div>
255 256 
256 257 
@@ -282,7 +283,7 @@ L1上存在2个bank，前面一半是和后一半分别是2个bank。如果写
282 283 
283现在研究访存bound的case的性能优化。比如`[128, 8192, 4096]`，表现在流水上 **MTE2 段长时间连续、与后续 MTE1/MMAD 难以充分重叠**。对于这种场景，优化mte2 的搬运速度，才是关键。这类case的简化流水图如下所示：284现在研究访存bound的case的性能优化。比如`[128, 8192, 4096]`，表现在流水上 **MTE2 段长时间连续、与后续 MTE1/MMAD 难以充分重叠**。对于这种场景，优化mte2 的搬运速度，才是关键。这类case的简化流水图如下所示：
284<div align="center">285<div align="center">
285 <img src="images/memory-bound-flowchat.png" width="800" alt="scale一次载入一块流水图" />286 <img src="images/memory-bound-flowchat.png" width="1500" alt="scale一次载入一块流水图" />
286</div>287</div>
287 288 
288本例中一次的搬运量如下，可以看出scaleA和scaleB的搬运量小于20K。搬运指令，若总数据量小于20K的话，对芯片搬运不友好，会导致速度降低。289本例中一次的搬运量如下，可以看出scaleA和scaleB的搬运量小于20K。搬运指令，若总数据量小于20K的话，对芯片搬运不友好，会导致速度降低。
@@ -351,7 +352,7 @@ $$
351 352 
352在本例中，采用**Scale访问合并** 策略后，scaleA和scaleB只在首次会加载到L1，后续只有A和B矩阵加载到L1。353在本例中，采用**Scale访问合并** 策略后，scaleA和scaleB只在首次会加载到L1，后续只有A和B矩阵加载到L1。
353<div align="center">354<div align="center">
354 <img src="images/memory-bound-flowchat-co.png" width="800" alt="scale一次载入多块流水图" />355 <img src="images/memory-bound-flowchat-co.png" width="1500" alt="scale一次载入多块流水图" />
355</div>356</div>
356 357 
357 358 
@@ -384,7 +385,7 @@ $$
384 385 
385上面优化的各个case，A矩阵和B矩阵需要在K方向进行切分，一块一块地搬入到L1/L0,完成K方向累加，最终将单个Tile块输入到HBM上。这种策略我们叫做**非全载策略**。但是如果A矩阵比较小而B矩阵特别大，若仍然采用上述策略，那么肯定会导致A矩阵重复搬运，很有可能会导致性能劣化。这回拿`[128, 4096, 81920]`来举例。使用非全载策略，流水示意图如下。矩阵A、scaleA、矩阵B和scaleB矩阵每次载入一个切块，到L1/L0,完成K方向累加。386上面优化的各个case，A矩阵和B矩阵需要在K方向进行切分，一块一块地搬入到L1/L0,完成K方向累加，最终将单个Tile块输入到HBM上。这种策略我们叫做**非全载策略**。但是如果A矩阵比较小而B矩阵特别大，若仍然采用上述策略，那么肯定会导致A矩阵重复搬运，很有可能会导致性能劣化。这回拿`[128, 4096, 81920]`来举例。使用非全载策略，流水示意图如下。矩阵A、scaleA、矩阵B和scaleB矩阵每次载入一个切块，到L1/L0,完成K方向累加。
386<div align="center">387<div align="center">
387 <img src="images/memory-bound-flowchat1.png" width="800" alt="非全载策略的流水图" />388 <img src="images/memory-bound-flowchat1.png" width="1500" alt="非全载策略的流水图" />
388</div>389</div>
389 390 
390让我们计算下左矩阵所占的总空间大小。通过计算发现左矩阵只占L1的一半空间。那么左矩阵是不是可以全载？391让我们计算下左矩阵所占的总空间大小。通过计算发现左矩阵只占L1的一半空间。那么左矩阵是不是可以全载？
@@ -407,7 +408,7 @@ $$
407 408 
408因此，引入**A 矩阵全载（A-Full-Load）** 策略将左矩阵 A 常驻在L1空间上。这里采用的方式是按照切块大小去填满A矩阵在L1的空间，不是一次性加载完成。scaleA矩阵是一次性加载到L1的。右矩阵B和scaleB矩阵是按照切块大小一块一块地载入到L1。每计算输出一个切块，都可以复用驻留在L1上数据，这样可以最大限度地减少数据的重复搬运。采用该策略的流水示意图如下。409因此，引入**A 矩阵全载（A-Full-Load）** 策略将左矩阵 A 常驻在L1空间上。这里采用的方式是按照切块大小去填满A矩阵在L1的空间，不是一次性加载完成。scaleA矩阵是一次性加载到L1的。右矩阵B和scaleB矩阵是按照切块大小一块一块地载入到L1。每计算输出一个切块，都可以复用驻留在L1上数据，这样可以最大限度地减少数据的重复搬运。采用该策略的流水示意图如下。
409<div align="center">410<div align="center">
410 <img src="images/a-fullload-flowchart.png" width="800" alt="全载策略的流水图" />411 <img src="images/a-fullload-flowchart.png" width="1500" alt="全载策略的流水图" />
411</div>412</div>
412 413 
413### 7.3 典型 Shape414### 7.3 典型 Shape
@@ -552,6 +553,7 @@ $$
552 553 
553非全载与全载的 **搬运速度** 同一量级；全载通过去掉 A、scaleA 的重复搬运，在 **总搬运量更小** 的前提下得到 **更短 MTE2 时间** 与 **更优端到端性能**。二者对比说明：在访存 bound 场景下，**减少左矩阵重复搬运** 能有效提升实际表现。554非全载与全载的 **搬运速度** 同一量级；全载通过去掉 A、scaleA 的重复搬运，在 **总搬运量更小** 的前提下得到 **更短 MTE2 时间** 与 **更优端到端性能**。二者对比说明：在访存 bound 场景下，**减少左矩阵重复搬运** 能有效提升实际表现。
554 555 
556 
555### 7.5 代码目录557### 7.5 代码目录
556 558 
557- **样例根目录**：[7_fullload](../matmul_tutorials/7_fullload/)559- **样例根目录**：[7_fullload](../matmul_tutorials/7_fullload/)
MSamples/2_Performance/matmul_story/matmul_tutorials/1_pingpong/include/kernel/quant_matmul_mx_kernel_impl_base.h+1-1
@@ -195,7 +195,7 @@ __global__ __aicore__ __cube__ void QuantMatmulMxfp4BaseKernel(uint64_t m, uint6
195 constexpr uint32_t BASE_N = 256;195 constexpr uint32_t BASE_N = 256;
196 constexpr uint32_t BASE_K = 256;196 constexpr uint32_t BASE_K = 256;
197 constexpr uint32_t PINGPONG_NUM = 2;197 constexpr uint32_t PINGPONG_NUM = 2;
198 constexpr uint32_t L1_BUFFER_NUM = 3;198 constexpr uint32_t L1_BUFFER_NUM = 2;
199 Params params;199 Params params;
200 params.problemShape.m = static_cast<int64_t>(m);200 params.problemShape.m = static_cast<int64_t>(m);
201 params.problemShape.n = static_cast<int64_t>(n);201 params.problemShape.n = static_cast<int64_t>(n);
MSamples/2_Performance/matmul_story/matmul_tutorials/2_block_swat/include/kernel/quant_matmul_mx_kernel_swat_impl.h+1-1
@@ -181,7 +181,7 @@ __global__ __aicore__ __cube__ void QuantMatmulMxfp4SwatKernel(uint64_t m, uint6
181 constexpr uint32_t BASE_N = 256;181 constexpr uint32_t BASE_N = 256;
182 constexpr uint32_t BASE_K = 256; // 128 / sizeof(fp4x2_e2m1_t)182 constexpr uint32_t BASE_K = 256; // 128 / sizeof(fp4x2_e2m1_t)
183 constexpr uint32_t PINGPONG_NUM = 2;183 constexpr uint32_t PINGPONG_NUM = 2;
184 constexpr uint32_t L1_BUFFER_NUM = 3;184 constexpr uint32_t L1_BUFFER_NUM = 2;
185 185 
186 Params params;186 Params params;
187 params.problemShape.m = static_cast<int64_t>(m);187 params.problemShape.m = static_cast<int64_t>(m);
MSamples/2_Performance/matmul_story/matmul_tutorials/3_last_round_tile_balance/include/kernel/quant_matmul_mx_kernel_last_round_tile_balance_impl.h+1-1
@@ -188,7 +188,7 @@ __global__ __aicore__ __cube__ void QuantMatmulMxfp4LastRoundTileBalanceKernel(u
188 constexpr uint32_t PINGPONG_NUM = 2;188 constexpr uint32_t PINGPONG_NUM = 2;
189 constexpr uint32_t M_TAIL_TILE = 1;189 constexpr uint32_t M_TAIL_TILE = 1;
190 constexpr uint32_t N_TAIL_TILE = 2;190 constexpr uint32_t N_TAIL_TILE = 2;
191 constexpr uint32_t L1_BUFFER_NUM = 3;191 constexpr uint32_t L1_BUFFER_NUM = 2;
192 192 
193 Params params;193 Params params;
194 params.problemShape.m = static_cast<int64_t>(m);194 params.problemShape.m = static_cast<int64_t>(m);
MSamples/2_Performance/matmul_story/matmul_tutorials/4_unit_flag/include/kernel/quant_matmul_mx_kernel_swat.h+1-1
@@ -185,7 +185,7 @@ __global__ __aicore__ __cube__ void QuantMatmulMxfp4UnitFlagKernel(uint64_t m, u
185 constexpr uint32_t PINGPONG_NUM = 2;185 constexpr uint32_t PINGPONG_NUM = 2;
186 constexpr uint32_t M_TAIL_TILE = 1;186 constexpr uint32_t M_TAIL_TILE = 1;
187 constexpr uint32_t N_TAIL_TILE = 2;187 constexpr uint32_t N_TAIL_TILE = 2;
188 constexpr uint32_t L1_BUFFER_NUM = 3;188 constexpr uint32_t L1_BUFFER_NUM = 2;
189 189 
190 Params params;190 Params params;
191 params.problemShape.m = static_cast<int64_t>(m);191 params.problemShape.m = static_cast<int64_t>(m);
MSamples/2_Performance/matmul_story/matmul_tutorials/5_halfl1_ping_halfl1_pong/include/kernel/quant_matmul_mx_kernel_impl_base.h+2-2
@@ -201,8 +201,8 @@ __global__ __aicore__ __cube__ void QuantMatmulMxfp4BaseKernel(uint64_t m, uint6
201 constexpr uint32_t BASE_K = 256;201 constexpr uint32_t BASE_K = 256;
202 constexpr uint32_t PINGPONG_NUM = 2;202 constexpr uint32_t PINGPONG_NUM = 2;
203 constexpr uint32_t M_TAIL_TILE = 1;203 constexpr uint32_t M_TAIL_TILE = 1;
204 constexpr uint32_t N_TAIL_TILE = 1;204 constexpr uint32_t N_TAIL_TILE = 2;
205 constexpr uint32_t L1_BUFFER_NUM = 3;205 constexpr uint32_t L1_BUFFER_NUM = 2;
206 206 
207 Params params;207 Params params;
208 params.problemShape.m = static_cast<int64_t>(m);208 params.problemShape.m = static_cast<int64_t>(m);