已合并
update readme #395
yangyang016创建于 13 天前
update readme #395
已合并
共 2 个文件变更+7-5
| @@ -55,8 +55,10 @@ | |||
| 55 | 55 | ||
| 56 | ## 🔥最新动态 | 56 | ## 🔥最新动态 |
| 57 | 57 | ||
| 58 | -- [2026/05] cann-samples在matmul_story新增[matmul_a16w16](Samples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a16w16/README.md) streamK等高性能example，并在既有quant_matmul_mxfp4、quant_matmul_mxfp8 MX量化矩阵乘中支持Weight NZ。 | 58 | +- [2026/08] 新增 [Kimi Delta Attention Lite](Samples/2_Performance/kimi_delta_attn_lite_story/README.md)、[Softmax RegBase 递进优化](Samples/2_Performance/softmax_regbase_story/README.md)和 [Vector Function Add](Samples/0_Introduction/vector_function_add/README.md) 样例，并上线环境自检脚本，进一步完善入门与性能调优体验。 |
| 59 | -- [2026/05] cann-samples在grouped_matmul_story新增[weight_quant_grouped_matmul_mxfp8fp4](Samples/2_Performance/grouped_matmul_story/grouped_matmul_recipes/examples/weight_quant_grouped_matmul_mxfp8fp4/README.md) MXA8W4量化example，并在既有quant_grouped_matmul_mxfp4、quant_grouped_matmul_mxfp8分组矩阵乘中支持Weight NZ。 | 59 | +- [2026/07] 新增 [Flash Attention Lite](Samples/2_Performance/flash_attn_lite_story/README.md)、[GELU + Element-wise RegBase](Samples/2_Performance/gelu_eltwise_regbase_story/README.md)、[SIMT Histogram](Samples/2_Performance/simt_histogram_story/README.md)和 [SIMT Scatter](Samples/2_Performance/simt_scatter_story/README.md) 调优实践，同时补充 [NPU 访存带宽测试](Samples/1_Features/hardware_features/mem_bandwidth/README.md)与 [CV 数据通路优化](Samples/1_Features/hardware_features/cv_datapath/README.md)样例。 |
| 60 | +- [2026/06] 建设 [SIMD VF 常见范式](Samples/2_Performance/simd_vf_story/README.md)，覆盖 Broadcast、Elementwise 与 Reduce；新增 [Reg 数据搬运场景选型指南](Samples/1_Features/memory_optimization/reg_data_movement/Reg数据搬运场景选型指南.md)及 10 个可运行示例，并补充 [KvRmsNormRopeCache RegBase](Samples/2_Performance/kv_rms_norm_rope_cache_story/README.md)优化实践。 | ||
| 61 | +- [2026/05] 矩阵乘与分组矩阵乘样例持续增强：新增 [matmul_a16w16](Samples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a16w16/README.md) StreamK、[weight_quant_grouped_matmul_mxfp8fp4](Samples/2_Performance/grouped_matmul_story/grouped_matmul_recipes/examples/weight_quant_grouped_matmul_mxfp8fp4/README.md) MXA8W4 量化等高性能样例，并为 MXFP4/MXFP8 量化场景补充 Weight NZ 支持。 | ||
| 60 | 62 | ||
| 61 | --- | 63 | --- |
| 62 | 64 | ||
MSamples/2_Performance/matmul_story/matmul_recipes/examples/quant_matmul_mxfp8/quant_matmul_mxfp8_k_tail_stepwise_copyout.asc+3-3
| @@ -31,12 +31,12 @@ | |||
| 31 | #include "host_utils/io_utils.h" | 31 | #include "host_utils/io_utils.h" |
| 32 | #include "kernel/quant_matmul_mx_kernel_swat.h" | 32 | #include "kernel/quant_matmul_mx_kernel_swat.h" |
| 33 | #include "tiling/quant_matmul_mx_tiling_swat.h" | 33 | #include "tiling/quant_matmul_mx_tiling_swat.h" |
| 34 | -#include "tiling/quant_matmul_tiling_data.h" | 34 | +#include "tiling/quant_matmul_mx_tiling_data.h" |
| 35 | 35 | ||
| 36 | template <bool TransA, bool TransB> | 36 | template <bool TransA, bool TransB> |
| 37 | __global__ __aicore__ __cube__ void QuantMatmulMxfp8KTailStepwiseCopyoutKernel( | 37 | __global__ __aicore__ __cube__ void QuantMatmulMxfp8KTailStepwiseCopyoutKernel( |
| 38 | GM_ADDR dA, GM_ADDR dB, GM_ADDR dScaleA, GM_ADDR dScaleB, GM_ADDR dC, | 38 | GM_ADDR dA, GM_ADDR dB, GM_ADDR dScaleA, GM_ADDR dScaleB, GM_ADDR dC, |
| 39 | - const QuantMatmulTilingData quantMatmulTilingData) | 39 | + const QuantMatmulMxTilingData quantMatmulTilingData) |
| 40 | { | 40 | { |
| 41 | // Keep the sample explicit about the datatype/layout combination that this | 41 | // Keep the sample explicit about the datatype/layout combination that this |
| 42 | // executable demonstrates so host tiling and kernel traits stay in sync. | 42 | // executable demonstrates so host tiling and kernel traits stay in sync. |
| @@ -116,7 +116,7 @@ int main(int argc, char* argv[]) | |||
| 116 | constexpr int32_t deviceId = 0; | 116 | constexpr int32_t deviceId = 0; |
| 117 | 117 | ||
| 118 | try { | 118 | try { |
| 119 | - QuantMatmulTilingData tilingData; | 119 | + QuantMatmulMxTilingData tilingData; |
| 120 | QuantMatmulTilingSwat<mm::DataType::DT_FLOAT8_E4M3FN, mm::DataType::DT_FLOAT8_E4M3FN> tilingEngine; | 120 | QuantMatmulTilingSwat<mm::DataType::DT_FLOAT8_E4M3FN, mm::DataType::DT_FLOAT8_E4M3FN> tilingEngine; |
| 121 | tilingEngine.GetTilingData(m, n, k, transA, transB, tilingData); | 121 | tilingEngine.GetTilingData(m, n, k, transA, transB, tilingData); |
| 122 | 122 | ||