Star596
205
代码
代码
Issues11
Pull Requests40
Actions
项目讨论
Wiki
分析
Star596
205
已合并
新增 matmul_a32w32 swat 与 streamk 模板 #295
wuyufei创建于 6月9日
新增 matmul_a32w32 swat 与 streamk 模板 #295
已合并
wuyufei创建于 6月9日
从wuyufei/cann-samplesmaster合入到cann/cann-samplesmaster
共 27 个文件变更+1725-185
MSamples/2_Performance/matmul_story/docs/matmul_performance.md+1-1
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本文档系统阐述非量化矩阵乘算子的实现原理、性能建模方法及优化实践，覆盖Float16、BFloat16数据类型场景。通过系统性的优化策略，帮助开发者快速掌握算子性能调优的核心技术，提升算子在昇腾平台上的执行效率。5+本文档系统阐述非量化矩阵乘算子的实现原理、性能建模方法及优化实践，以A16W16为例进行说明，覆盖Float16、BFloat16数据类型场景。通过系统性的优化策略，帮助开发者快速掌握算子性能调优的核心技术，提升算子在昇腾平台上的执行效率。
6 6 
7**矩阵乘法**是神经网络和大模型的"底层计算引擎"，从特征传递到注意力机制实现，再到大规模参数运算，所有核心流程都依赖其完成。没有矩阵乘法就无法支撑模型的高效运行与规模突破。7**矩阵乘法**是神经网络和大模型的"底层计算引擎"，从特征传递到注意力机制实现，再到大规模参数运算，所有核心流程都依赖其完成。没有矩阵乘法就无法支撑模型的高效运行与规模突破。
8 8 
MSamples/2_Performance/matmul_story/matmul_recipes/CMakeLists.txt+6-1
@@ -41,6 +41,11 @@ add_matmul_recipe(matmul_a16w16_swat
41add_matmul_recipe(matmul_a16w16_streamk41add_matmul_recipe(matmul_a16w16_streamk
42 matmul_a16w16/matmul_a16w16_streamk.asc matmul_a16w16)42 matmul_a16w16/matmul_a16w16_streamk.asc matmul_a16w16)
43 43 
44+add_matmul_recipe(matmul_a32w32_swat
45+ matmul_a32w32/matmul_a32w32_swat.asc matmul_a32w32)
46+add_matmul_recipe(matmul_a32w32_streamk
47+ matmul_a32w32/matmul_a32w32_streamk.asc matmul_a32w32)
48+ 
44add_matmul_recipe(quant_matmul_mxfp8_swat49add_matmul_recipe(quant_matmul_mxfp8_swat
45 quant_matmul_mxfp8/quant_matmul_mxfp8_swat.asc quant_matmul_mxfp8)50 quant_matmul_mxfp8/quant_matmul_mxfp8_swat.asc quant_matmul_mxfp8)
46add_matmul_recipe(quant_matmul_mxfp8_swat_4_buffer51add_matmul_recipe(quant_matmul_mxfp8_swat_4_buffer
@@ -63,7 +68,7 @@ add_matmul_recipe(weight_quant_matmul_mxfp8fp4_swat_4_buffer
63# Install helper scripts into the same directory as the executables. Host code68# Install helper scripts into the same directory as the executables. Host code
64# uses the executable directory as baseDir and runs `python3 verify_result.py`69# uses the executable directory as baseDir and runs `python3 verify_result.py`
65# there; gen_data.py is also documented as run from that cwd.70# there; gen_data.py is also documented as run from that cwd.
66-set(RECIPE_VARIANTS quant_matmul_mxfp4 quant_matmul_mxfp8 matmul_a16w16 quant_matmul_hifp8 weight_quant_matmul_mxfp8fp4)71+set(RECIPE_VARIANTS quant_matmul_mxfp4 quant_matmul_mxfp8 matmul_a16w16 matmul_a32w32 quant_matmul_hifp8 weight_quant_matmul_mxfp8fp4)
67foreach(variant IN LISTS RECIPE_VARIANTS)72foreach(variant IN LISTS RECIPE_VARIANTS)
68 install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/examples/${variant}/scripts/73 install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/examples/${variant}/scripts/
69 DESTINATION ${MATMUL_INSTALL_PREFIX}/matmul_recipes/${variant}74 DESTINATION ${MATMUL_INSTALL_PREFIX}/matmul_recipes/${variant}
MSamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a16w16/matmul_a16w16_streamk.asc+29-28
@@ -29,14 +29,14 @@
29#include "block/block_scheduler_policy.h"29#include "block/block_scheduler_policy.h"
30#include "host_utils/common_utils.h"30#include "host_utils/common_utils.h"
31#include "host_utils/io_utils.h"31#include "host_utils/io_utils.h"
32-#include "kernel/matmul_a16w16_kernel_streamk.h"32+#include "kernel/matmul_kernel_streamk.h"
33-#include "tiling/matmul_a16w16_tiling_streamk.h"33+#include "tiling/matmul_tiling_streamk.h"
34-#include "tiling/matmul_a16w16_tiling_data.h"34+#include "tiling/matmul_tiling_data.h"
35-#include "epilogue/matmul_a16w16_block_epilogue_streamk.h"35+#include "epilogue/matmul_block_epilogue_streamk.h"
36 36 
37template <class LAYOUT_A, class LAYOUT_B, class LAYOUT_C>37template <class LAYOUT_A, class LAYOUT_B, class LAYOUT_C>
38-__global__ __aicore__ __mix__(1, 2) void MatMulA16W16StreamKKernel(38+__global__ __aicore__ __mix__(1, 2) void MatMulStreamKKernel(
39- GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, GM_ADDR dWorkSpace, MatmulA16W16TilingData matmulA16W16TilingData)39+ GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, GM_ADDR dWorkSpace, MatmulTilingData matmulTilingData)
40{40{
41 using TypeA = bfloat16_t;41 using TypeA = bfloat16_t;
42 using TypeB = bfloat16_t;42 using TypeB = bfloat16_t;
@@ -45,22 +45,22 @@ __global__ __aicore__ __mix__(1, 2) void MatMulA16W16StreamKKernel(
45 using LayoutB = LAYOUT_B;45 using LayoutB = LAYOUT_B;
46 using LayoutC = LAYOUT_C;46 using LayoutC = LAYOUT_C;
47 47 
48- using BlockScheduler = MatmulA16W16StreamKScheduler;48+ using BlockScheduler = MatmulStreamKScheduler;
49- using DispatchPolicy = MatmulA16W16MultiBlockWithStreamK;49+ using DispatchPolicy = MatmulMultiBlockWithStreamK;
50 using BlockMmad = Block::BlockMmad<DispatchPolicy, TypeA, LayoutA, TypeB, LayoutB, TypeC, LayoutC>;50 using BlockMmad = Block::BlockMmad<DispatchPolicy, TypeA, LayoutA, TypeB, LayoutB, TypeC, LayoutC>;
51 using ProblemShape = MatmulShape;51 using ProblemShape = MatmulShape;
52 using BlockEpilogue = Block::BlockEpilogueStreamK<float, TypeC>;52 using BlockEpilogue = Block::BlockEpilogueStreamK<float, TypeC>;
53- using MatmulKernelImpl = Kernel::MatmulA16W16KernelStreamK<ProblemShape, BlockMmad, BlockScheduler, BlockEpilogue>;53+ using MatmulKernelImpl = Kernel::MatmulKernelStreamK<ProblemShape, BlockMmad, BlockScheduler, BlockEpilogue>;
54 using Params = typename MatmulKernelImpl::Params;54 using Params = typename MatmulKernelImpl::Params;
55 Params params = {55 Params params = {
56- {matmulA16W16TilingData.m, matmulA16W16TilingData.n, matmulA16W16TilingData.k, 1UL},56+ {matmulTilingData.m, matmulTilingData.n, matmulTilingData.k, 1UL},
57 {dA, dB, dC, dWorkSpace},57 {dA, dB, dC, dWorkSpace},
58- {matmulA16W16TilingData.usedCoreNum, matmulA16W16TilingData.baseM, matmulA16W16TilingData.baseN,58+ {matmulTilingData.usedCoreNum, matmulTilingData.baseM, matmulTilingData.baseN,
59- matmulA16W16TilingData.baseK, matmulA16W16TilingData.kL1, matmulA16W16TilingData.skSingleCoreK},59+ matmulTilingData.baseK, matmulTilingData.kL1, matmulTilingData.skSingleCoreK},
60 {dC, dWorkSpace},60 {dC, dWorkSpace},
61- {matmulA16W16TilingData.mL1, matmulA16W16TilingData.nL1, matmulA16W16TilingData.kL1,61+ {matmulTilingData.mL1, matmulTilingData.nL1, matmulTilingData.kL1,
62- matmulA16W16TilingData.baseM, matmulA16W16TilingData.baseN, matmulA16W16TilingData.baseK,62+ matmulTilingData.baseM, matmulTilingData.baseN, matmulTilingData.baseK,
63- matmulA16W16TilingData.usedCoreNum}};63+ matmulTilingData.usedCoreNum, matmulTilingData.isHf32}};
64 MatmulKernelImpl matmulKernelImpl;64 MatmulKernelImpl matmulKernelImpl;
65 matmulKernelImpl(params);65 matmulKernelImpl(params);
66}66}
@@ -68,7 +68,7 @@ __global__ __aicore__ __mix__(1, 2) void MatMulA16W16StreamKKernel(
68namespace {68namespace {
69template <bool TransA, bool TransB>69template <bool TransA, bool TransB>
70void LaunchKernel(70void LaunchKernel(
71- uint8_t* dA, uint8_t* dB, uint8_t* dC, uint8_t* dWorkSpace, const MatmulA16W16TilingData& tilingData,71+ uint8_t* dA, uint8_t* dB, uint8_t* dC, uint8_t* dWorkSpace, const MatmulTilingData& tilingData,
72 aclrtStream stream)72 aclrtStream stream)
73{73{
74 using LAYOUT_A = std::conditional_t<74 using LAYOUT_A = std::conditional_t<
@@ -77,7 +77,7 @@ void LaunchKernel(
77 using LAYOUT_B = std::conditional_t<77 using LAYOUT_B = std::conditional_t<
78 TransB, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,78 TransB, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,
79 AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>;79 AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>;
80- MatMulA16W16StreamKKernel<LAYOUT_A, LAYOUT_B, AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>80+ MatMulStreamKKernel<LAYOUT_A, LAYOUT_B, AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>
81 <<<tilingData.usedCoreNum, nullptr, stream>>>(dA, dB, dC, dWorkSpace, tilingData);81 <<<tilingData.usedCoreNum, nullptr, stream>>>(dA, dB, dC, dWorkSpace, tilingData);
82}82}
83} // namespace83} // namespace
@@ -115,20 +115,21 @@ int main(int argc, char* argv[])
115 }115 }
116 };116 };
117 try {117 try {
118- MatmulA16W16TilingData tilingData;118+ MatmulTilingData tilingData;
119- std::unique_ptr<MatmulA16W16TilingBase> tilingEngine = std::make_unique<MatmulA16W16TilingStreamK>();119+ std::unique_ptr<MatmulTilingBase> tilingEngine = std::make_unique<MatmulTilingStreamK>();
120- tilingEngine->GetTilingData(m, n, k, transA, transB, tilingData);120+ uint64_t dataTypeSize = sizeof(half);
121+ tilingEngine->GetTilingData(m, n, k, dataTypeSize, transA, transB, tilingData);
122+ CHECK_COND(aclInit(nullptr) == ACL_SUCCESS, "Failed to initialize ACL runtime.");
123+ aclInitialized = true;
121 uint32_t deviceCount = 0;124 uint32_t deviceCount = 0;
122 CHECK_COND(aclrtGetDeviceCount(&deviceCount) == ACL_SUCCESS, "Failed to query ACL device count.");125 CHECK_COND(aclrtGetDeviceCount(&deviceCount) == ACL_SUCCESS, "Failed to query ACL device count.");
123 CHECK_COND(deviceCount > 0U, "No ACL devices are available.");126 CHECK_COND(deviceCount > 0U, "No ACL devices are available.");
124- CHECK_COND(aclInit(nullptr) == ACL_SUCCESS, "Failed to initialize ACL runtime.");
125- aclInitialized = true;
126 CHECK_COND(aclrtSetDevice(deviceId) == ACL_SUCCESS, "Failed to set the ACL device.");127 CHECK_COND(aclrtSetDevice(deviceId) == ACL_SUCCESS, "Failed to set the ACL device.");
127 deviceSet = true;128 deviceSet = true;
128 CHECK_COND(aclrtCreateStream(&stream) == ACL_SUCCESS, "Failed to create the ACL stream.");129 CHECK_COND(aclrtCreateStream(&stream) == ACL_SUCCESS, "Failed to create the ACL stream.");
129- size_t sizeA = m * k * sizeof(half);130+ size_t sizeA = m * k * dataTypeSize;
130- size_t sizeB = k * n * sizeof(half);131+ size_t sizeB = k * n * dataTypeSize;
131- size_t sizeC = m * n * sizeof(half);132+ size_t sizeC = m * n * dataTypeSize;
132 size_t sizeWorkSpace = tilingEngine->GetWorkSpace();133 size_t sizeWorkSpace = tilingEngine->GetWorkSpace();
133 char exePath[PATH_MAX];134 char exePath[PATH_MAX];
134 ssize_t len = readlink("/proc/self/exe", exePath, sizeof(exePath) - 1);135 ssize_t len = readlink("/proc/self/exe", exePath, sizeof(exePath) - 1);
@@ -211,10 +212,10 @@ int main(int argc, char* argv[])
211 CHECK_COND(212 CHECK_COND(
212 aclrtMemcpyAsync(hC, sizeC, dC, sizeC, ACL_MEMCPY_DEVICE_TO_HOST, stream) == ACL_SUCCESS,213 aclrtMemcpyAsync(hC, sizeC, dC, sizeC, ACL_MEMCPY_DEVICE_TO_HOST, stream) == ACL_SUCCESS,
213 "Failed to copy output C from device to host.");214 "Failed to copy output C from device to host.");
215+ CHECK_COND(
216+ aclrtSynchronizeStream(stream) == ACL_SUCCESS,
217+ "Failed to synchronize the ACL stream after kernel execution.");
214 }218 }
215- CHECK_COND(
216- aclrtSynchronizeStream(stream) == ACL_SUCCESS,
217- "Failed to synchronize the ACL stream after kernel execution.");
218 WriteFile(outputDir + "/npu_out.bin", hC, sizeC);219 WriteFile(outputDir + "/npu_out.bin", hC, sizeC);
219 std::string cmd =220 std::string cmd =
220 "cd \"" + baseDir + "\" && python3 verify_result.py " + std::to_string(m) + " " + std::to_string(n);221 "cd \"" + baseDir + "\" && python3 verify_result.py " + std::to_string(m) + " " + std::to_string(n);
MSamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a16w16/matmul_a16w16_swat.asc+31-30
@@ -29,13 +29,13 @@
29#include "block/block_scheduler_policy.h"29#include "block/block_scheduler_policy.h"
30#include "host_utils/common_utils.h"30#include "host_utils/common_utils.h"
31#include "host_utils/io_utils.h"31#include "host_utils/io_utils.h"
32-#include "kernel/matmul_a16w16_kernel_swat.h"32+#include "kernel/matmul_kernel_swat.h"
33-#include "tiling/matmul_a16w16_tiling_swat.h"33+#include "tiling/matmul_tiling_swat.h"
34-#include "tiling/matmul_a16w16_tiling_data.h"34+#include "tiling/matmul_tiling_data.h"
35 35 
36template <class LAYOUT_A, class LAYOUT_B, class LAYOUT_C>36template <class LAYOUT_A, class LAYOUT_B, class LAYOUT_C>
37-__global__ __aicore__ __cube__ void MatMulA16W16SwatKernel(37+__global__ __aicore__ __cube__ void MatMulSwatKernel(
38- GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, MatmulA16W16TilingData matmulA16W16TilingData)38+ GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, MatmulTilingData matmulTilingData)
39{39{
40 // Data type of the matrix (half or bfloat16_t)40 // Data type of the matrix (half or bfloat16_t)
41 using TypeA = bfloat16_t;41 using TypeA = bfloat16_t;
@@ -46,32 +46,32 @@ __global__ __aicore__ __cube__ void MatMulA16W16SwatKernel(
46 using LayoutB = LAYOUT_B;46 using LayoutB = LAYOUT_B;
47 using LayoutC = LAYOUT_C;47 using LayoutC = LAYOUT_C;
48 48 
49- using BlockScheduler = MatmulA16W16SwatScheduler<NO_FULL_LOAD_MODE>;49+ using BlockScheduler = MatmulSwatScheduler<NO_FULL_LOAD_MODE>;
50- using DispatchPolicy = MatmulA16W16MultiBlockWithSwat<NO_FULL_LOAD_MODE>;50+ using DispatchPolicy = MatmulMultiBlockWithSwat<NO_FULL_LOAD_MODE>;
51 using BlockMmad = Block::BlockMmad<DispatchPolicy, TypeA, LayoutA, TypeB, LayoutB, TypeC, LayoutC>;51 using BlockMmad = Block::BlockMmad<DispatchPolicy, TypeA, LayoutA, TypeB, LayoutB, TypeC, LayoutC>;
52 using ProblemShape = MatmulShape;52 using ProblemShape = MatmulShape;
53- using MatmulKernelImpl = Kernel::MatmulA16W16KernelSwat<ProblemShape, BlockMmad, BlockScheduler>;53+ using MatmulKernelImpl = Kernel::MatmulKernelSwat<ProblemShape, BlockMmad, BlockScheduler>;
54 using Params = typename MatmulKernelImpl::Params;54 using Params = typename MatmulKernelImpl::Params;
55 // Translate the serialized host tiling data into the strongly typed kernel55 // Translate the serialized host tiling data into the strongly typed kernel
56 // parameter bundle expected by the device-side implementation.56 // parameter bundle expected by the device-side implementation.
57 Params params = {57 Params params = {
58- {matmulA16W16TilingData.m, matmulA16W16TilingData.n, matmulA16W16TilingData.k, 1UL},58+ {matmulTilingData.m, matmulTilingData.n, matmulTilingData.k, 1UL},
59 {dA, dB, dC},59 {dA, dB, dC},
60- {matmulA16W16TilingData.mL1, matmulA16W16TilingData.nL1, matmulA16W16TilingData.kL1,60+ {matmulTilingData.mL1, matmulTilingData.nL1, matmulTilingData.kL1,
61- matmulA16W16TilingData.baseM, matmulA16W16TilingData.baseN, matmulA16W16TilingData.baseK,61+ matmulTilingData.baseM, matmulTilingData.baseN, matmulTilingData.baseK,
62- matmulA16W16TilingData.mBaseTailSplitCnt, matmulA16W16TilingData.nBaseTailSplitCnt,62+ matmulTilingData.mBaseTailSplitCnt, matmulTilingData.nBaseTailSplitCnt,
63- matmulA16W16TilingData.mTailMain, matmulA16W16TilingData.nTailMain, matmulA16W16TilingData.mTailCnt,63+ matmulTilingData.mTailMain, matmulTilingData.nTailMain, matmulTilingData.mTailCnt,
64- matmulA16W16TilingData.nTailCnt},64+ matmulTilingData.nTailCnt},
65- {matmulA16W16TilingData.mL1, matmulA16W16TilingData.nL1, matmulA16W16TilingData.kL1,65+ {matmulTilingData.mL1, matmulTilingData.nL1, matmulTilingData.kL1,
66- matmulA16W16TilingData.baseM, matmulA16W16TilingData.baseN, matmulA16W16TilingData.baseK,66+ matmulTilingData.baseM, matmulTilingData.baseN, matmulTilingData.baseK,
67- matmulA16W16TilingData.l0cDB}};67+ matmulTilingData.isHf32, matmulTilingData.l0cDB}};
68 MatmulKernelImpl matmulKernelImpl;68 MatmulKernelImpl matmulKernelImpl;
69 matmulKernelImpl(params);69 matmulKernelImpl(params);
70}70}
71 71 
72namespace {72namespace {
73template <bool TransA, bool TransB>73template <bool TransA, bool TransB>
74-void LaunchKernel(uint8_t* dA, uint8_t* dB, uint8_t* dC, const MatmulA16W16TilingData& tilingData, aclrtStream stream)74+void LaunchKernel(uint8_t* dA, uint8_t* dB, uint8_t* dC, const MatmulTilingData& tilingData, aclrtStream stream)
75{75{
76 using LAYOUT_A = std::conditional_t<76 using LAYOUT_A = std::conditional_t<
77 TransA, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,77 TransA, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,
@@ -79,7 +79,7 @@ void LaunchKernel(uint8_t* dA, uint8_t* dB, uint8_t* dC, const MatmulA16W16Tilin
79 using LAYOUT_B = std::conditional_t<79 using LAYOUT_B = std::conditional_t<
80 TransB, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,80 TransB, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,
81 AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>;81 AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>;
82- MatMulA16W16SwatKernel<LAYOUT_A, LAYOUT_B, AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>82+ MatMulSwatKernel<LAYOUT_A, LAYOUT_B, AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>
83 <<<tilingData.usedCoreNum, nullptr, stream>>>(dA, dB, dC, tilingData);83 <<<tilingData.usedCoreNum, nullptr, stream>>>(dA, dB, dC, tilingData);
84}84}
85} // namespace85} // namespace
@@ -117,22 +117,23 @@ int main(int argc, char* argv[])
117 }117 }
118 };118 };
119 try {119 try {
120- MatmulA16W16TilingData tilingData;120+ MatmulTilingData tilingData;
121 // Host tiling picks the block shape, tail strategy, and buffering plan121 // Host tiling picks the block shape, tail strategy, and buffering plan
122 // // that will later be consumed by the device kernel.122 // // that will later be consumed by the device kernel.
123- std::unique_ptr<MatmulA16W16TilingBase> tilingEngine = std::make_unique<MatmulA16W16TilingSwat>();123+ std::unique_ptr<MatmulTilingBase> tilingEngine = std::make_unique<MatmulTilingSwat>();
124- tilingEngine->GetTilingData(m, n, k, transA, transB, tilingData);124+ uint64_t dataTypeSize = sizeof(half);
125+ tilingEngine->GetTilingData(m, n, k, dataTypeSize, transA, transB, tilingData);
126+ CHECK_COND(aclInit(nullptr) == ACL_SUCCESS, "Failed to initialize ACL runtime.");
127+ aclInitialized = true;
125 uint32_t deviceCount = 0;128 uint32_t deviceCount = 0;
126 CHECK_COND(aclrtGetDeviceCount(&deviceCount) == ACL_SUCCESS, "Failed to query ACL device count.");129 CHECK_COND(aclrtGetDeviceCount(&deviceCount) == ACL_SUCCESS, "Failed to query ACL device count.");
127 CHECK_COND(deviceCount > 0U, "No ACL devices are available.");130 CHECK_COND(deviceCount > 0U, "No ACL devices are available.");
128- CHECK_COND(aclInit(nullptr) == ACL_SUCCESS, "Failed to initialize ACL runtime.");
129- aclInitialized = true;
130 CHECK_COND(aclrtSetDevice(deviceId) == ACL_SUCCESS, "Failed to set the ACL device.");131 CHECK_COND(aclrtSetDevice(deviceId) == ACL_SUCCESS, "Failed to set the ACL device.");
131 deviceSet = true;132 deviceSet = true;
132 CHECK_COND(aclrtCreateStream(&stream) == ACL_SUCCESS, "Failed to create the ACL stream.");133 CHECK_COND(aclrtCreateStream(&stream) == ACL_SUCCESS, "Failed to create the ACL stream.");
133- size_t sizeA = m * k * sizeof(half);134+ size_t sizeA = m * k * dataTypeSize;
134- size_t sizeB = k * n * sizeof(half);135+ size_t sizeB = k * n * dataTypeSize;
135- size_t sizeC = m * n * sizeof(half);136+ size_t sizeC = m * n * dataTypeSize;
136 char exePath[PATH_MAX];137 char exePath[PATH_MAX];
137 ssize_t len = readlink("/proc/self/exe", exePath, sizeof(exePath) - 1);138 ssize_t len = readlink("/proc/self/exe", exePath, sizeof(exePath) - 1);
138 std::string baseDir = ".";139 std::string baseDir = ".";
@@ -206,10 +207,10 @@ int main(int argc, char* argv[])
206 CHECK_COND(207 CHECK_COND(
207 aclrtMemcpyAsync(hC, sizeC, dC, sizeC, ACL_MEMCPY_DEVICE_TO_HOST, stream) == ACL_SUCCESS,208 aclrtMemcpyAsync(hC, sizeC, dC, sizeC, ACL_MEMCPY_DEVICE_TO_HOST, stream) == ACL_SUCCESS,
208 "Failed to copy output C from device to host.");209 "Failed to copy output C from device to host.");
210+ CHECK_COND(
211+ aclrtSynchronizeStream(stream) == ACL_SUCCESS,
212+ "Failed to synchronize the ACL stream after kernel execution.");
209 }213 }
210- CHECK_COND(
211- aclrtSynchronizeStream(stream) == ACL_SUCCESS,
212- "Failed to synchronize the ACL stream after kernel execution.");
213 WriteFile(outputDir + "/npu_out.bin", hC, sizeC);214 WriteFile(outputDir + "/npu_out.bin", hC, sizeC);
214 std::string cmd =215 std::string cmd =
215 "cd \"" + baseDir + "\" && python3 verify_result.py " + std::to_string(m) + " " + std::to_string(n);216 "cd \"" + baseDir + "\" && python3 verify_result.py " + std::to_string(m) + " " + std::to_string(n);
MSamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a16w16/scripts/matmul_a16w16_algorithm_recommend.py+1-2
@@ -195,8 +195,7 @@ def resolve_gen_data_script(script_dir: Path) -> Path:
195def generate_input(script_dir: Path, m: int, k: int, n: int, trans_a: bool, trans_b: bool) -> None:195def generate_input(script_dir: Path, m: int, k: int, n: int, trans_a: bool, trans_b: bool) -> None:
196 script_path = resolve_gen_data_script(script_dir)196 script_path = resolve_gen_data_script(script_dir)
197 args = [sys.executable, str(script_path), str(m), str(k), str(n)]197 args = [sys.executable, str(script_path), str(m), str(k), str(n)]
198- if trans_a or trans_b:198+ args.extend([str(trans_a).lower(), str(trans_b).lower()])
199- args.extend([str(trans_a).lower(), str(trans_b).lower()])
200 with tempfile.TemporaryFile(mode="w+t", encoding="utf-8") as log_file:199 with tempfile.TemporaryFile(mode="w+t", encoding="utf-8") as log_file:
201 result = subprocess.run(200 result = subprocess.run(
202 args,201 args,
ASamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a32w32/README.md+191-0
@@ -0,0 +1,191 @@
1+# A32W32非量化矩阵乘算子
2+ 
3+## 概述
4+ 
5+本示例展示了A32W32非量化矩阵乘算子在昇腾AI处理器上的完整实现，包含基于SWAT模板和StreamK模板的高性能优化方案。A32W32支持HF32、Float32数据类型，是深度学习中最基础且最重要的计算操作，广泛应用于各种神经网络层，包括全连接层、注意力机制等。
6+ 
7+当前目录提供以下能力：
8+ 
9+- `matmul_a32w32_swat`：基于SWAT模板的实现。
10+- `matmul_a32w32_streamk`：基于StreamK模板的实现。
11+- `gen_data.py`：生成输入数据和CPU golden结果。
12+- `verify_result.py`：校验NPU输出与CPU golden是否一致。
13+- `matmul_a32w32_algorithm_recommend.py`：对当前目录下可执行算法进行兼容性筛选和耗时排序。
14+ 
15+## 使用约束
16+ 
17+当前样例支持以下场景：
18+ 
19+- A的形状为`[M, K]`(非转置)或`[K, M]`(转置)
20+- B的形状为`[K, N]`(非转置)或`[N, K]`(转置)。
21+- 输出矩阵C的形状为`[M, N]`。
22+- 支持数据类型：Float32，矩阵A,B,C数据类型需保持一致。
23+ 
24+### HF32 模式说明
25+ 
26+默认启用HF32（Hybrid Float32）模式以获得更好的性能。HF32是一种兼顾精度与性能的浮点格式，在矩阵乘法场景下相比传统Float32可获得更高的计算吞吐。
27+ 
28+如需切换为传统Float32模式，需修改以下两个源文件：
29+ 
30+1. **`matmul_a32w32_streamk.asc`**（第121行附近）
31+2. **`matmul_a32w32_swat.asc`**（第119行附近）
32+ 
33+在各自的 `tilingEngine->GetTilingData(m, n, k, dataTypeSize, transA, transB, tilingData);` 语句后面新增一行：
34+ 
35+```
36+tilingData.isHf32 = false;
37+```
38+ 
39+### StreamK 输入范围限制
40+ 
41+`matmul_a32w32_streamk` 对输入参数有特定的范围要求，需满足以下条件之一：
42+ 
43+1. **SK 模式**：
44+ - K 维度需足够大（建议 K ≥ 8192）
45+ - M 和 N 的切分块数乘积不超过 AIC 核数的一半
46+ 
47+2. **DPSK 模式**：
48+ - M 和 N 需为 256 的倍数
49+ - K 维度需足够大（建议 K ≥ 8192）
50+ - M 和 N 的块数乘积满足特定的负载均衡条件
51+ 
52+如果输入参数不满足上述条件，直接运行 `matmul_a32w32_streamk` 将会报错退出。建议使用算法推荐脚本自动选择适合当前形状的算法。
53+ 
54+## 支持架构
55+ 
56+NPU ARCH 3510
57+ 
58+## 性能优化指南
59+ 
60+关于算子涉及的模板实现及优化策略，请参考[非量化矩阵乘算子性能优化指南](../../../docs/matmul_performance.md)
61+ 
62+## API参考
63+ 
64+[Ascend C API文档](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/850/API/ascendcopapi/atlasascendc_api_07_0003.html)
65+ 
66+## 参数说明
67+ 
68+两个可执行文件的命令行参数格式一致：
69+ 
70+```text
71+<program> m k n [transA transB]
72+```
73+ 
74+- `m`：矩阵A的行数。
75+- `k`：矩阵A的列数，同时也是矩阵B的归约维。
76+- `n`：矩阵B的列数，对应输出矩阵的列数。
77+- `transA`：矩阵A的转置信息（可选，默认为false）。
78+- `transB`：矩阵B的转置信息（可选，默认为true）。
79+ 
80+在当前布局下，如果不传入转置信息，默认A矩阵非转置，B矩阵转置：
81+ 
82+- A按`[M, K]`组织。
83+- B按`[N, K]`组织。
84+- 输出矩阵C的形状为`[M, N]`。
85+ 
86+## 数据与校验
87+ 
88+`gen_data.py`会在当前目录下生成以下文件：
89+ 
90+- `input/input_a.bin`
91+- `input/input_b.bin`
92+- `output/cpu_output.bin`
93+ 
94+样例执行完成后会额外生成：
95+ 
96+- `output/npu_out.bin`
97+ 
98+两个可执行文件在运行结束后都会自动调用`verify_result.py`，将NPU输出与CPU golden进行一致性校验。
99+ 
100+## 一键运行（推荐）
101+ 
102+仓库提供 `run.sh`（位于 `matmul_recipes/examples/matmul_a32w32/scripts/`），可一键串联 **构建 → 数据生成 → 算子执行 → 结果校验** 全流程。
103+推荐先进入样例目录再执行，命令更短：
104+ 
105+```bash
106+cd Samples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a32w32
107+ 
108+# 自动构建 + 自动推荐最优算法 + 运行
109+bash scripts/run.sh 128 16384 128
110+ 
111+# 带转置参数
112+bash scripts/run.sh 128 16384 128 false true
113+ 
114+# 指定目标可执行文件，跳过重新构建
115+bash scripts/run.sh \
116+ --target matmul_a32w32_swat --skip-build 128 16384 128
117+ 
118+# 查看完整帮助
119+bash scripts/run.sh --help
120+```
121+ 
122+### run.sh 参数说明
123+ 
124+| 参数 | 说明 |
125+|------|------|
126+| `m k n` | 矩阵维度（必填）。 |
127+| `transA transB` | 转置参数（可选）。默认 `false true`，即 A 不转置、B 转置。 |
128+| `--target <name>` | 指定要运行的可执行文件名。省略时自动调用推荐脚本选择最优目标。 |
129+| `--skip-build` | 跳过构建/安装阶段，复用已有 `build_out`。 |
130+| `-h, --help` | 显示帮助信息。 |
131+ 
132+如需查看完整算法推荐排名（含耗时表格），请在安装目录下直接运行 `matmul_a32w32_algorithm_recommend.py`（见下文「手动构建与运行」）。
133+ 
134+## 手动构建与运行
135+ 
136+如需手动控制各步骤，可在仓库根目录下完成编译和安装后，进入当前样例目录：
137+ 
138+```bash
139+cmake -S . -B build -DNPU_ARCH=dav-3510
140+cmake --build build --parallel
141+cmake --install build --prefix ./build_out
142+cd build_out/2_Performance/matmul_story/matmul_recipes/matmul_a32w32
143+```
144+ 
145+### 1. 生成测试数据
146+ 
147+```bash
148+python3 gen_data.py 128 16384 128 false true
149+```
150+ 
151+### 2. 运行单个算法样例
152+ 
153+```bash
154+./matmul_a32w32_swat 128 16384 128
155+```
156+ 
157+或：
158+ 
159+```bash
160+./matmul_a32w32_streamk 128 16384 128
161+```
162+ 
163+带转置参数：
164+ 
165+```bash
166+./matmul_a32w32_swat 128 16384 128 false true
167+```
168+ 
169+### 3. 运行算法推荐脚本
170+ 
171+```bash
172+python3 matmul_a32w32_algorithm_recommend.py 128 16384 128
173+```
174+ 
175+或带转置参数：
176+ 
177+```bash
178+python3 matmul_a32w32_algorithm_recommend.py 128 16384 128 false true
179+```
180+ 
181+下图为推荐脚本输出的**结构示意**（数值为虚构，仅说明版式）：
182+ 
183+```text
184+[Profile Breakdown]
185++---------------------------+----------+---------+----------+---------+---------+------------+--------------+
186+| algorithm |kernel(us)| mac(us) |scalar(us)| mte1(us)| mte2(us)|fixpipe(us) |icache_miss(%)|
187++===========================+==========+=========+==========+=========+=========+============+==============+
188+| matmul_a32w32_swat | 12.345| 1.234 | 0.567| 0.123 | 0.456 | 0.789 | 0.100 |
189+| matmul_a32w32_streamk | 15.678| 2.100 | 0.800| 0.200 | 0.300 | 0.500 | 0.250 |
190++---------------------------+----------+---------+----------+---------+---------+------------+--------------+
191+```
ASamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a32w32/matmul_a32w32_streamk.asc+233-0
@@ -0,0 +1,233 @@
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 matmul_a32w32_streamk.cpp
13+ * \brief Sample launcher for A32W32 MATMUL StreamK streaming example.
14+ */
15+#include <cstdint>
16+#include <cstdlib>
17+#include <exception>
18+#include <iomanip>
19+#include <iostream>
20+#include <limits.h>
21+#include <memory>
22+#include <stdexcept>
23+#include <string>
24+#include <sys/stat.h>
25+#include <unistd.h>
26+#include "acl/acl.h"
27+#include "kernel_operator.h"
28+#include "block/block_mmad.h"
29+#include "block/block_scheduler_policy.h"
30+#include "host_utils/common_utils.h"
31+#include "host_utils/io_utils.h"
32+#include "kernel/matmul_kernel_streamk.h"
33+#include "tiling/matmul_tiling_streamk.h"
34+#include "tiling/matmul_tiling_data.h"
35+#include "epilogue/matmul_block_epilogue_streamk.h"
36+ 
37+template <class LAYOUT_A, class LAYOUT_B, class LAYOUT_C>
38+__global__ __aicore__ __mix__(1, 2) void MatMulA32W32StreamKKernel(
39+ GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, GM_ADDR dWorkSpace, MatmulTilingData matmulTilingData)
40+{
41+ using TypeA = float;
42+ using TypeB = float;
43+ using TypeC = float;
44+ using LayoutA = LAYOUT_A;
45+ using LayoutB = LAYOUT_B;
46+ using LayoutC = LAYOUT_C;
47+ 
48+ using BlockScheduler = MatmulStreamKScheduler;
49+ using DispatchPolicy = MatmulMultiBlockWithStreamK;
50+ using BlockMmad = Block::BlockMmad<DispatchPolicy, TypeA, LayoutA, TypeB, LayoutB, TypeC, LayoutC>;
51+ using ProblemShape = MatmulShape;
52+ using BlockEpilogue = Block::BlockEpilogueStreamK<float, TypeC>;
53+ using MatmulKernelImpl = Kernel::MatmulKernelStreamK<ProblemShape, BlockMmad, BlockScheduler, BlockEpilogue>;
54+ using Params = typename MatmulKernelImpl::Params;
55+ Params params = {
56+ {matmulTilingData.m, matmulTilingData.n, matmulTilingData.k, 1UL},
57+ {dA, dB, dC, dWorkSpace},
58+ {matmulTilingData.usedCoreNum, matmulTilingData.baseM, matmulTilingData.baseN,
59+ matmulTilingData.baseK, matmulTilingData.kL1, matmulTilingData.skSingleCoreK},
60+ {dC, dWorkSpace},
61+ {matmulTilingData.mL1, matmulTilingData.nL1, matmulTilingData.kL1,
62+ matmulTilingData.baseM, matmulTilingData.baseN, matmulTilingData.baseK,
63+ matmulTilingData.usedCoreNum, matmulTilingData.isHf32}};
64+ MatmulKernelImpl matmulKernelImpl;
65+ matmulKernelImpl(params);
66+}
67+ 
68+namespace {
69+template <bool TransA, bool TransB>
70+void LaunchKernel(
71+ uint8_t* dA, uint8_t* dB, uint8_t* dC, uint8_t* dWorkSpace, const MatmulTilingData& tilingData,
72+ aclrtStream stream)
73+{
74+ using LAYOUT_A = std::conditional_t<
75+ TransA, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,
76+ AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>;
77+ using LAYOUT_B = std::conditional_t<
78+ TransB, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,
79+ AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>;
80+ MatMulA32W32StreamKKernel<LAYOUT_A, LAYOUT_B, AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>
81+ <<<tilingData.usedCoreNum, nullptr, stream>>>(dA, dB, dC, dWorkSpace, tilingData);
82+}
83+} // namespace
84+ 
85+int main(int argc, char* argv[])
86+{
87+ uint64_t m = 0;
88+ uint64_t k = 0;
89+ uint64_t n = 0;
90+ bool transA = false;
91+ bool transB = false;
92+ try {
93+ ParseArguments(argc, argv, m, k, n, transA, transB);
94+ } catch (const std::exception& e) {
95+ std::cerr << e.what() << std::endl;
96+ PrintUsage(argv[0]);
97+ return 1;
98+ }
99+ constexpr int32_t deviceId = 0;
100+ aclrtStream stream = nullptr;
101+ bool aclInitialized = false;
102+ bool deviceSet = false;
103+ auto cleanupAcl = [&]() {
104+ if (stream != nullptr) {
105+ aclrtDestroyStream(stream);
106+ stream = nullptr;
107+ }
108+ if (deviceSet) {
109+ aclrtResetDevice(deviceId);
110+ deviceSet = false;
111+ }
112+ if (aclInitialized) {
113+ aclFinalize();
114+ aclInitialized = false;
115+ }
116+ };
117+ try {
118+ MatmulTilingData tilingData;
119+ std::unique_ptr<MatmulTilingBase> tilingEngine = std::make_unique<MatmulTilingStreamK>();
120+ uint64_t dataTypeSize = sizeof(float);
121+ tilingEngine->GetTilingData(m, n, k, dataTypeSize, transA, transB, tilingData);
122+ CHECK_COND(aclInit(nullptr) == ACL_SUCCESS, "Failed to initialize ACL runtime.");
123+ aclInitialized = true;
124+ uint32_t deviceCount = 0;
125+ CHECK_COND(aclrtGetDeviceCount(&deviceCount) == ACL_SUCCESS, "Failed to query ACL device count.");
X
Xxubinglin6月12日

ACL runtime 要求先初始化 ACL，再调用 runtime 查询、设卡、建流等接口。当前顺序可能导致样例启动阶段失败。 https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/910beta1/API/runtimeapi/aclcppdevg_03_0022.html

likedislike
W
wuyufei
6月12日 评论：
126+ CHECK_COND(deviceCount > 0U, "No ACL devices are available.");
127+ CHECK_COND(aclrtSetDevice(deviceId) == ACL_SUCCESS, "Failed to set the ACL device.");
128+ deviceSet = true;
129+ CHECK_COND(aclrtCreateStream(&stream) == ACL_SUCCESS, "Failed to create the ACL stream.");
130+ size_t sizeA = m * k * dataTypeSize;
131+ size_t sizeB = k * n * dataTypeSize;
132+ size_t sizeC = m * n * dataTypeSize;
133+ size_t sizeWorkSpace = tilingEngine->GetWorkSpace();
134+ char exePath[PATH_MAX];
135+ ssize_t len = readlink("/proc/self/exe", exePath, sizeof(exePath) - 1);
136+ std::string baseDir = ".";
137+ if (len > 0) {
138+ exePath[len] = '\0';
139+ baseDir = exePath;
140+ size_t lastSlash = baseDir.find_last_of('/');
141+ if (lastSlash != std::string::npos) {
142+ baseDir.resize(lastSlash);
143+ }
144+ }
145+ std::string inputDir = baseDir + "/input";
146+ std::string outputDir = baseDir + "/output";
147+ struct stat st;
148+ if (stat((inputDir + "/input_a.bin").c_str(), &st) != 0) {
149+ std::cout << "Input files not found !" << std::endl;
150+ return 0;
151+ }
152+ float* hA = nullptr;
153+ float* hB = nullptr;
154+ float* hC = nullptr;
155+ float* hWorkSpace = nullptr;
156+ uint8_t* dA = nullptr;
157+ uint8_t* dB = nullptr;
158+ uint8_t* dC = nullptr;
159+ uint8_t* dWorkSpace = nullptr;
160+ CHECK_COND(aclrtMallocHost((void**)&hA, sizeA) == ACL_SUCCESS, "Failed to allocate host buffer for input A.");
161+ std::unique_ptr<void, aclError (*)(void*)> hostA(hA, aclrtFreeHost);
162+ CHECK_COND(aclrtMallocHost((void**)&hB, sizeB) == ACL_SUCCESS, "Failed to allocate host buffer for input B.");
163+ std::unique_ptr<void, aclError (*)(void*)> hostB(hB, aclrtFreeHost);
164+ CHECK_COND(aclrtMallocHost((void**)&hC, sizeC) == ACL_SUCCESS, "Failed to allocate host buffer for output C.");
165+ std::unique_ptr<void, aclError (*)(void*)> hostC(hC, aclrtFreeHost);
166+ CHECK_COND(
167+ aclrtMallocHost((void**)&hWorkSpace, sizeWorkSpace) == ACL_SUCCESS,
168+ "Failed to allocate host buffer for WorkSpace");
169+ std::unique_ptr<void, aclError (*)(void*)> hostWorkSpace(hWorkSpace, aclrtFreeHost);
170+ ReadFile(inputDir + "/input_a.bin", sizeA, hA, sizeA);
171+ ReadFile(inputDir + "/input_b.bin", sizeB, hB, sizeB);
172+ for (int32_t i = 0; i < EXAMPLE_KERNEL_RUN_COUNT; i++) {
173+ void* dAPtr = nullptr;
174+ CHECK_COND(
175+ aclrtMalloc(&dAPtr, sizeA, ACL_MEM_MALLOC_HUGE_ONLY) == ACL_SUCCESS,
176+ "Failed to allocate the device buffer for input A.");
177+ dA = reinterpret_cast<uint8_t*>(dAPtr);
178+ std::unique_ptr<void, aclError (*)(void*)> deviceA(dAPtr, aclrtFree);
179+ void* dBPtr = nullptr;
180+ CHECK_COND(
181+ aclrtMalloc(&dBPtr, sizeB, ACL_MEM_MALLOC_HUGE_ONLY) == ACL_SUCCESS,
182+ "Failed to allocate the device buffer for input B.");
183+ dB = reinterpret_cast<uint8_t*>(dBPtr);
184+ std::unique_ptr<void, aclError (*)(void*)> deviceB(dBPtr, aclrtFree);
185+ void* dCPtr = nullptr;
186+ CHECK_COND(
187+ aclrtMalloc(&dCPtr, sizeC, ACL_MEM_MALLOC_HUGE_ONLY) == ACL_SUCCESS,
188+ "Failed to allocate the device buffer for output C.");
189+ dC = reinterpret_cast<uint8_t*>(dCPtr);
190+ std::unique_ptr<void, aclError (*)(void*)> deviceC(dCPtr, aclrtFree);
191+ void* dWorkSpacePtr = nullptr;
192+ CHECK_COND(
193+ aclrtMalloc(&dWorkSpacePtr, sizeWorkSpace, ACL_MEM_MALLOC_HUGE_ONLY) == ACL_SUCCESS,
194+ "Failed to allocate the device buffer for WorkSpace.");
195+ dWorkSpace = reinterpret_cast<uint8_t*>(dWorkSpacePtr);
196+ std::unique_ptr<void, aclError (*)(void*)> deviceWorkSpace(dWorkSpacePtr, aclrtFree);
197+ CHECK_COND(
198+ aclrtMemcpyAsync(dA, sizeA, hA, sizeA, ACL_MEMCPY_HOST_TO_DEVICE, stream) == ACL_SUCCESS,
199+ "Failed to copy input A from host to device.");
200+ CHECK_COND(
201+ aclrtMemcpyAsync(dB, sizeB, hB, sizeB, ACL_MEMCPY_HOST_TO_DEVICE, stream) == ACL_SUCCESS,
202+ "Failed to copy input B from host to device.");
203+ if (transA && transB) {
204+ LaunchKernel<true, true>(dA, dB, dC, dWorkSpace, tilingData, stream);
205+ } else if (transA && !transB) {
206+ LaunchKernel<true, false>(dA, dB, dC, dWorkSpace, tilingData, stream);
207+ } else if (!transA && transB) {
208+ LaunchKernel<false, true>(dA, dB, dC, dWorkSpace, tilingData, stream);
209+ } else {
210+ LaunchKernel<false, false>(dA, dB, dC, dWorkSpace, tilingData, stream);
211+ }
212+ CHECK_COND(
213+ aclrtMemcpyAsync(hC, sizeC, dC, sizeC, ACL_MEMCPY_DEVICE_TO_HOST, stream) == ACL_SUCCESS,
214+ "Failed to copy output C from device to host.");
215+ CHECK_COND(
216+ aclrtSynchronizeStream(stream) == ACL_SUCCESS,
217+ "Failed to synchronize the ACL stream after kernel execution.");
218+ }
219+ WriteFile(outputDir + "/npu_out.bin", hC, sizeC);
220+ std::string cmd =
221+ "cd \"" + baseDir + "\" && python3 verify_result.py " + std::to_string(m) + " " + std::to_string(n);
222+ if (std::system(cmd.c_str()) != 0) {
223+ cleanupAcl();
224+ return 1;
225+ }
226+ cleanupAcl();
227+ return 0;
228+ } catch (const std::exception& e) {
229+ cleanupAcl();
230+ std::cerr << e.what() << std::endl;
231+ return 1;
232+ }
233+}
ASamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a32w32/matmul_a32w32_swat.asc+218-0
@@ -0,0 +1,218 @@
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 matmul_a32w32_swat.cpp
13+ * \brief Sample launcher for A32W32 MATMUL SWAT streaming example.
14+ */
15+#include <cstdint>
16+#include <cstdlib>
17+#include <exception>
18+#include <iomanip>
19+#include <iostream>
20+#include <limits.h>
21+#include <memory>
22+#include <stdexcept>
23+#include <string>
24+#include <sys/stat.h>
25+#include <unistd.h>
26+#include "acl/acl.h"
27+#include "kernel_operator.h"
28+#include "block/block_mmad.h"
29+#include "block/block_scheduler_policy.h"
30+#include "host_utils/common_utils.h"
31+#include "host_utils/io_utils.h"
32+#include "kernel/matmul_kernel_swat.h"
33+#include "tiling/matmul_tiling_swat.h"
34+#include "tiling/matmul_tiling_data.h"
35+ 
36+template <class LAYOUT_A, class LAYOUT_B, class LAYOUT_C>
37+__global__ __aicore__ __cube__ void MatMulA32W32SwatKernel(
38+ GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, MatmulTilingData matmulTilingData)
39+{
40+ using TypeA = float;
41+ using TypeB = float;
42+ using TypeC = float;
43+ using LayoutA = LAYOUT_A;
44+ using LayoutB = LAYOUT_B;
45+ using LayoutC = LAYOUT_C;
46+ 
47+ using BlockScheduler = MatmulSwatScheduler<NO_FULL_LOAD_MODE>;
48+ using DispatchPolicy = MatmulMultiBlockWithSwat<NO_FULL_LOAD_MODE>;
49+ using BlockMmad = Block::BlockMmad<DispatchPolicy, TypeA, LayoutA, TypeB, LayoutB, TypeC, LayoutC>;
50+ using ProblemShape = MatmulShape;
51+ using MatmulKernelImpl = Kernel::MatmulKernelSwat<ProblemShape, BlockMmad, BlockScheduler>;
52+ using Params = typename MatmulKernelImpl::Params;
53+ Params params = {
54+ {matmulTilingData.m, matmulTilingData.n, matmulTilingData.k, 1UL},
55+ {dA, dB, dC},
56+ {matmulTilingData.mL1, matmulTilingData.nL1, matmulTilingData.kL1,
57+ matmulTilingData.baseM, matmulTilingData.baseN, matmulTilingData.baseK,
58+ matmulTilingData.mBaseTailSplitCnt, matmulTilingData.nBaseTailSplitCnt,
59+ matmulTilingData.mTailMain, matmulTilingData.nTailMain, matmulTilingData.mTailCnt,
60+ matmulTilingData.nTailCnt},
61+ {matmulTilingData.mL1, matmulTilingData.nL1, matmulTilingData.kL1,
62+ matmulTilingData.baseM, matmulTilingData.baseN, matmulTilingData.baseK,
63+ matmulTilingData.isHf32, matmulTilingData.l0cDB}};
64+ MatmulKernelImpl matmulKernelImpl;
65+ matmulKernelImpl(params);
66+}
67+ 
68+namespace {
69+template <bool TransA, bool TransB>
70+void LaunchKernel(uint8_t* dA, uint8_t* dB, uint8_t* dC, const MatmulTilingData& tilingData, aclrtStream stream)
71+{
72+ using LAYOUT_A = std::conditional_t<
73+ TransA, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,
74+ AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>;
75+ using LAYOUT_B = std::conditional_t<
76+ TransB, AscendC::Te::FrameLayoutFormat<AscendC::Te::DNExtLayoutPtn>,
77+ AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>;
78+ MatMulA32W32SwatKernel<LAYOUT_A, LAYOUT_B, AscendC::Te::FrameLayoutFormat<AscendC::Te::NDExtLayoutPtn>>
79+ <<<tilingData.usedCoreNum, nullptr, stream>>>(dA, dB, dC, tilingData);
80+}
81+} // namespace
82+ 
83+int main(int argc, char* argv[])
84+{
85+ uint64_t m = 0;
86+ uint64_t k = 0;
87+ uint64_t n = 0;
88+ bool transA = false;
89+ bool transB = false;
90+ try {
91+ ParseArguments(argc, argv, m, k, n, transA, transB);
92+ } catch (const std::exception& e) {
93+ std::cerr << e.what() << std::endl;
94+ PrintUsage(argv[0]);
95+ return 1;
96+ }
97+ constexpr int32_t deviceId = 0;
98+ aclrtStream stream = nullptr;
99+ bool aclInitialized = false;
100+ bool deviceSet = false;
101+ auto cleanupAcl = [&]() {
102+ if (stream != nullptr) {
103+ aclrtDestroyStream(stream);
104+ stream = nullptr;
105+ }
106+ if (deviceSet) {
107+ aclrtResetDevice(deviceId);
108+ deviceSet = false;
109+ }
110+ if (aclInitialized) {
111+ aclFinalize();
112+ aclInitialized = false;
113+ }
114+ };
115+ try {
116+ MatmulTilingData tilingData;
117+ std::unique_ptr<MatmulTilingBase> tilingEngine = std::make_unique<MatmulTilingSwat>();
118+ uint64_t dataTypeSize = sizeof(float);
119+ tilingEngine->GetTilingData(m, n, k, dataTypeSize, transA, transB, tilingData);
120+ CHECK_COND(aclInit(nullptr) == ACL_SUCCESS, "Failed to initialize ACL runtime.");
121+ aclInitialized = true;
122+ uint32_t deviceCount = 0;
123+ CHECK_COND(aclrtGetDeviceCount(&deviceCount) == ACL_SUCCESS, "Failed to query ACL device count.");
124+ CHECK_COND(deviceCount > 0U, "No ACL devices are available.");
125+ CHECK_COND(aclrtSetDevice(deviceId) == ACL_SUCCESS, "Failed to set the ACL device.");
126+ deviceSet = true;
127+ CHECK_COND(aclrtCreateStream(&stream) == ACL_SUCCESS, "Failed to create the ACL stream.");
128+ size_t sizeA = m * k * dataTypeSize;
129+ size_t sizeB = k * n * dataTypeSize;
130+ size_t sizeC = m * n * dataTypeSize;
131+ char exePath[PATH_MAX];
132+ ssize_t len = readlink("/proc/self/exe", exePath, sizeof(exePath) - 1);
133+ std::string baseDir = ".";
134+ if (len > 0) {
135+ exePath[len] = '\0';
136+ baseDir = exePath;
137+ size_t lastSlash = baseDir.find_last_of('/');
138+ if (lastSlash != std::string::npos) {
139+ baseDir.resize(lastSlash);
140+ }
141+ }
142+ std::string inputDir = baseDir + "/input";
143+ std::string outputDir = baseDir + "/output";
144+ struct stat st;
145+ if (stat((inputDir + "/input_a.bin").c_str(), &st) != 0) {
146+ std::cout << "Input files not found !" << std::endl;
147+ return 0;
148+ }
149+ float* hA = nullptr;
150+ float* hB = nullptr;
151+ float* hC = nullptr;
152+ uint8_t* dA = nullptr;
153+ uint8_t* dB = nullptr;
154+ uint8_t* dC = nullptr;
155+ CHECK_COND(aclrtMallocHost((void**)&hA, sizeA) == ACL_SUCCESS, "Failed to allocate host buffer for input A.");
156+ std::unique_ptr<void, aclError (*)(void*)> hostA(hA, aclrtFreeHost);
157+ CHECK_COND(aclrtMallocHost((void**)&hB, sizeB) == ACL_SUCCESS, "Failed to allocate host buffer for input B.");
158+ std::unique_ptr<void, aclError (*)(void*)> hostB(hB, aclrtFreeHost);
159+ CHECK_COND(aclrtMallocHost((void**)&hC, sizeC) == ACL_SUCCESS, "Failed to allocate host buffer for output C.");
160+ std::unique_ptr<void, aclError (*)(void*)> hostC(hC, aclrtFreeHost);
161+ ReadFile(inputDir + "/input_a.bin", sizeA, hA, sizeA);
162+ ReadFile(inputDir + "/input_b.bin", sizeB, hB, sizeB);
163+ for (int32_t i = 0; i < EXAMPLE_KERNEL_RUN_COUNT; i++) {
164+ void* dAPtr = nullptr;
165+ CHECK_COND(
166+ aclrtMalloc(&dAPtr, sizeA, ACL_MEM_MALLOC_HUGE_ONLY) == ACL_SUCCESS,
167+ "Failed to allocate the device buffer for input A.");
168+ dA = reinterpret_cast<uint8_t*>(dAPtr);
169+ std::unique_ptr<void, aclError (*)(void*)> deviceA(dAPtr, aclrtFree);
170+ void* dBPtr = nullptr;
171+ CHECK_COND(
172+ aclrtMalloc(&dBPtr, sizeB, ACL_MEM_MALLOC_HUGE_ONLY) == ACL_SUCCESS,
173+ "Failed to allocate the device buffer for input B.");
174+ dB = reinterpret_cast<uint8_t*>(dBPtr);
175+ std::unique_ptr<void, aclError (*)(void*)> deviceB(dBPtr, aclrtFree);
176+ void* dCPtr = nullptr;
177+ CHECK_COND(
178+ aclrtMalloc(&dCPtr, sizeC, ACL_MEM_MALLOC_HUGE_ONLY) == ACL_SUCCESS,
179+ "Failed to allocate the device buffer for output C.");
180+ dC = reinterpret_cast<uint8_t*>(dCPtr);
181+ std::unique_ptr<void, aclError (*)(void*)> deviceC(dCPtr, aclrtFree);
182+ CHECK_COND(
183+ aclrtMemcpyAsync(dA, sizeA, hA, sizeA, ACL_MEMCPY_HOST_TO_DEVICE, stream) == ACL_SUCCESS,
184+ "Failed to copy input A from host to device.");
185+ CHECK_COND(
186+ aclrtMemcpyAsync(dB, sizeB, hB, sizeB, ACL_MEMCPY_HOST_TO_DEVICE, stream) == ACL_SUCCESS,
187+ "Failed to copy input B from host to device.");
188+ if (transA && transB) {
189+ LaunchKernel<true, true>(dA, dB, dC, tilingData, stream);
190+ } else if (transA && !transB) {
191+ LaunchKernel<true, false>(dA, dB, dC, tilingData, stream);
192+ } else if (!transA && transB) {
193+ LaunchKernel<false, true>(dA, dB, dC, tilingData, stream);
194+ } else {
195+ LaunchKernel<false, false>(dA, dB, dC, tilingData, stream);
196+ }
197+ CHECK_COND(
198+ aclrtMemcpyAsync(hC, sizeC, dC, sizeC, ACL_MEMCPY_DEVICE_TO_HOST, stream) == ACL_SUCCESS,
199+ "Failed to copy output C from device to host.");
200+ CHECK_COND(
201+ aclrtSynchronizeStream(stream) == ACL_SUCCESS,
202+ "Failed to synchronize the ACL stream after kernel execution.");
203+ }
204+ WriteFile(outputDir + "/npu_out.bin", hC, sizeC);
205+ std::string cmd =
206+ "cd \"" + baseDir + "\" && python3 verify_result.py " + std::to_string(m) + " " + std::to_string(n);
207+ if (std::system(cmd.c_str()) != 0) {
208+ cleanupAcl();
209+ return 1;
210+ }
211+ cleanupAcl();
212+ return 0;
213+ } catch (const std::exception& e) {
214+ cleanupAcl();
215+ std::cerr << e.what() << std::endl;
216+ return 1;
217+ }
218+}
ASamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a32w32/scripts/gen_data.py+82-0
@@ -0,0 +1,82 @@
1+#!/usr/bin/python3
2+# coding=utf-8
3+ 
4+# ----------------------------------------------------------------------------------------------------------
5+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
6+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
7+# CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
10+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11+# See LICENSE in the root of the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------------------------------------
13+ 
14+import os
15+import sys
16+ 
17+os.environ["TORCH_DEVICE_BACKEND_AUTOLOAD"] = "0"
18+ 
19+import numpy as np
20+import torch
21+ 
22+DATA_TYPE = torch.float32
23+ 
24+def write_artifacts(base_dir, a_data, b_data, out):
25+ input_dir = os.path.join(base_dir, "input")
26+ output_dir = os.path.join(base_dir, "output")
27+ os.makedirs(input_dir, exist_ok=True)
28+ os.makedirs(output_dir, exist_ok=True)
29+ 
30+ a_data.view(torch.uint32).numpy().tofile(os.path.join(input_dir, "input_a.bin"))
31+ b_data.view(torch.uint32).numpy().tofile(os.path.join(input_dir, "input_b.bin"))
32+ out.view(torch.uint32).numpy().tofile(os.path.join(output_dir, "cpu_output.bin"))
33+ 
34+ 
35+def gen_golden_data_simple(m, k, n, transpose_a, transpose_b):
36+ M = m
37+ K = k
38+ N = n
39+ 
40+ a_ori = (np.random.uniform(0, 1, (K, M)).astype(np.float32) if transpose_a
41+ else np.random.uniform(0, 1, (M, K)).astype(np.float32))
42+ b_ori = (np.random.uniform(0, 1, (N, K)).astype(np.float32) if transpose_b
43+ else np.random.uniform(0, 1, (K, N)).astype(np.float32))
44+ 
45+ a_cpu = torch.from_numpy(a_ori).to(DATA_TYPE)
46+ b_cpu = torch.from_numpy(b_ori).to(DATA_TYPE)
47+ 
48+ a_cpu_t = a_cpu.t() if transpose_a else a_cpu
49+ b_cpu_t = b_cpu.t() if transpose_b else b_cpu
50+ 
51+ out = torch.matmul(a_cpu_t, b_cpu_t).to(DATA_TYPE)
52+ 
53+ current_dir = os.getcwd()
54+ write_artifacts(current_dir, a_cpu, b_cpu, out)
55+ 
56+ script_dir = os.path.dirname(os.path.abspath(__file__))
57+ if os.path.normcase(os.path.abspath(script_dir)) != os.path.normcase(os.path.abspath(current_dir)):
58+ write_artifacts(script_dir, a_cpu, b_cpu, out)
59+ 
60+ print("Data generated successfully!")
61+ 
62+ 
63+if __name__ == "__main__":
64+ if len(sys.argv) != 4 and len(sys.argv) != 6:
65+ print("Usage: python3 gen_data.py m k n")
66+ print("Or")
67+ print("Usage: python3 gen_data.py m k n transA transB")
68+ print("Example1: python3 gen_data.py 100 50 200")
69+ print("Example2: python3 gen_data.py 100 50 200 false true")
70+ sys.exit(1)
71+ 
72+ m = int(sys.argv[1])
73+ k = int(sys.argv[2])
74+ n = int(sys.argv[3])
75+ if len(sys.argv) == 6:
76+ transpose_a = sys.argv[4].lower() == "true"
77+ transpose_b = sys.argv[5].lower() == "true"
78+ else:
79+ transpose_a = False
80+ transpose_b = True
81+ 
82+ gen_golden_data_simple(m, k, n, transpose_a, transpose_b)
ASamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a32w32/scripts/matmul_a32w32_algorithm_recommend.py+456-0
@@ -0,0 +1,456 @@
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+import os
12+import csv
13+import shutil
14+import subprocess
15+import sys
16+import tempfile
17+from dataclasses import dataclass
18+from pathlib import Path
19+from typing import List, Optional, Tuple
20+ 
21+ 
22+MSPROF_OUTPUT_DIR_NAME = "msprof_recommend"
23+MSPROF_PROF_DIR_PREFIX = "PROF_"
24+MSPROF_OP_SUMMARY_GLOB = "op_summary_*.csv"
25+PROFILE_METRIC_SPECS = (
26+ ("kernel_time_us", "kernel(us)", "Task Duration(us)"),
27+ ("mac_time_us", "mac(us)", "aic_mac_time(us)"),
28+ ("scalar_time_us", "scalar(us)", "aic_scalar_time(us)"),
29+ ("mte1_time_us", "mte1(us)", "aic_mte1_time(us)"),
30+ ("mte2_time_us", "mte2(us)", "aic_mte2_time(us)"),
31+ ("fixpipe_time_us", "fixpipe(us)", "aic_fixpipe_time(us)"),
32+ ("icache_miss_rate", "icache_miss(%)", "aic_icache_miss_rate"),
33+)
34+ 
35+ 
36+@dataclass(frozen=True)
37+class ProfileMetrics:
38+ """Performance fields extracted from one op_summary row."""
39+ 
40+ kernel_time_us: float
41+ mac_time_us: float
42+ scalar_time_us: float
43+ mte1_time_us: float
44+ mte2_time_us: float
45+ fixpipe_time_us: float
46+ icache_miss_rate: float
47+ 
48+ 
49+@dataclass(frozen=True)
50+class Candidate:
51+ """One installed executable that can participate in recommendation."""
52+ 
53+ label: str
54+ executable_name: str
55+ 
56+ 
57+@dataclass
58+class CandidateResult:
59+ """Execution record used for compatibility filtering and final ranking."""
60+ 
61+ label: str
62+ executable_path: Path
63+ kernel_time_us: Optional[float]
64+ profile_metrics: Optional[ProfileMetrics]
65+ return_code: int
66+ output: str
67+ 
68+ @property
69+ def succeeded(self) -> bool:
70+ return self.return_code == 0 and self.kernel_time_us is not None and self.profile_metrics is not None
71+ 
72+ 
73+def print_usage(program_name: str) -> None:
74+ print(f"Usage: {program_name} m k n [transA transB] [--print-target]")
75+ print("Args:")
76+ print(" m: row of matrix A")
77+ print(" k: col of matrix A")
78+ print(" n: col of matrix B")
79+ print(" transA: transdata of matrix A (optional, default: false)")
80+ print(" transB: transdata of matrix B (optional, default: true)")
81+ print("Options:")
82+ print(" --print-target: print only the recommended executable name")
83+ print(f"Example: {program_name} 100 50 200")
84+ print(f"Example: {program_name} 100 50 200 false true")
85+ 
86+ 
87+def parse_positive_uint64(arg: str, name: str) -> int:
88+ if not arg.isdigit():
89+ raise ValueError(f"{name} must be a positive integer")
90+ value = int(arg)
91+ if value <= 0:
92+ raise ValueError(f"{name} must be greater than 0")
93+ if value > 4294967295:
94+ raise ValueError(f"{name} must not exceed UINT32_MAX")
95+ return value
96+ 
97+ 
98+def parse_arguments(argv: List[str]) -> Tuple[int, int, int, bool, bool, bool]:
99+ if len(argv) >= 2 and argv[1] in ("-h", "--help"):
100+ print_usage(Path(argv[0]).name)
101+ raise SystemExit(0)
102+ 
103+ print_target_only = False
104+ positional: List[str] = []
105+ for arg in argv[1:]:
106+ if arg == "--print-target":
107+ print_target_only = True
108+ continue
109+ if arg.startswith("-"):
110+ raise ValueError(f"Unknown option: {arg}")
111+ positional.append(arg)
112+ 
113+ if len(positional) != 3 and len(positional) != 5:
114+ raise ValueError("Expected exactly 3 arguments: m k n, or 5 arguments: m k n transA transB")
115+ 
116+ m = parse_positive_uint64(positional[0], "m")
117+ k = parse_positive_uint64(positional[1], "k")
118+ n = parse_positive_uint64(positional[2], "n")
119+ 
120+ if len(positional) == 5:
121+ trans_a = positional[3].lower() == "true"
122+ trans_b = positional[4].lower() == "true"
123+ else:
124+ trans_a = False
125+ trans_b = True
126+ 
127+ return m, k, n, trans_a, trans_b, print_target_only
128+ 
129+ 
130+def get_ranked_results(results: List[CandidateResult]) -> List[CandidateResult]:
131+ return sorted(
132+ [item for item in results if item.succeeded],
133+ key=lambda item: item.kernel_time_us if item.kernel_time_us is not None else float("inf"),
134+ )
135+ 
136+ 
137+def resolve_executable(script_dir: Path, executable_name: str) -> Path:
138+ direct_path = script_dir / executable_name
139+ if direct_path.exists():
140+ return direct_path
141+ 
142+ windows_path = script_dir / f"{executable_name}.exe"
143+ if windows_path.exists():
144+ return windows_path
145+ 
146+ raise FileNotFoundError(f"Executable not found: {executable_name}")
147+ 
148+ 
149+def discover_candidates(script_dir: Path) -> List[Candidate]:
150+ candidates: List[Candidate] = []
151+ seen_names = set()
152+ script_stem = Path(__file__).stem
153+ 
154+ for entry in sorted(script_dir.iterdir(), key=lambda item: item.name):
155+ if not entry.is_file():
156+ continue
157+ 
158+ is_windows_executable = entry.suffix.lower() == ".exe"
159+ is_posix_executable = entry.suffix == "" and os.access(entry, os.X_OK)
160+ if not (is_windows_executable or is_posix_executable):
161+ continue
162+ 
163+ executable_name = entry.stem if is_windows_executable else entry.name
164+ if executable_name == script_stem:
165+ continue
166+ if executable_name in seen_names:
167+ continue
168+ 
169+ label = executable_name
170+ candidates.append(Candidate(label=label, executable_name=executable_name))
171+ seen_names.add(executable_name)
172+ 
173+ return candidates
174+ 
175+ 
176+def read_command_log(log_file) -> str:
177+ log_file.seek(0)
178+ return log_file.read().strip()
179+ 
180+ 
181+def format_command_output(prefix: str, raw_output: str) -> str:
182+ if not raw_output:
183+ return prefix
184+ return f"{prefix}\n{raw_output}"
185+ 
186+ 
187+def resolve_gen_data_script(script_dir: Path) -> Path:
188+ script_path = script_dir / "gen_data.py"
189+ if script_path.exists():
190+ return script_path
191+ 
192+ raise FileNotFoundError(f"gen_data.py was not found in {script_dir}")
193+ 
194+ 
195+def generate_input(script_dir: Path, m: int, k: int, n: int, trans_a: bool, trans_b: bool) -> None:
196+ script_path = resolve_gen_data_script(script_dir)
197+ args = [sys.executable, str(script_path), str(m), str(k), str(n)]
198+ args.extend([str(trans_a).lower(), str(trans_b).lower()])
199+ with tempfile.TemporaryFile(mode="w+t", encoding="utf-8") as log_file:
200+ result = subprocess.run(
201+ args,
202+ cwd=script_path.parent,
203+ text=True,
204+ stdout=log_file,
205+ stderr=subprocess.STDOUT,
206+ check=False,
207+ )
208+ if result.returncode != 0:
209+ output = read_command_log(log_file)
210+ raise RuntimeError(f"Failed to generate input data.\n{output}")
211+ 
212+ 
213+def cleanup_msprof_output_dir(msprof_output_dir: Path) -> None:
214+ if msprof_output_dir.exists():
215+ shutil.rmtree(msprof_output_dir, ignore_errors=True)
216+ 
217+ 
218+def list_prof_directories(msprof_output_dir: Path) -> set[Path]:
219+ if not msprof_output_dir.exists():
220+ return set()
221+ 
222+ return {
223+ entry.resolve()
224+ for entry in msprof_output_dir.iterdir()
225+ if entry.is_dir() and entry.name.startswith(MSPROF_PROF_DIR_PREFIX)
226+ }
227+ 
228+ 
229+def resolve_latest_prof_directory(msprof_output_dir: Path) -> Path:
230+ prof_dirs = list_prof_directories(msprof_output_dir)
231+ if not prof_dirs:
232+ raise FileNotFoundError(
233+ f"No {MSPROF_PROF_DIR_PREFIX}* directory was generated under {msprof_output_dir}"
234+ )
235+ 
236+ return max(prof_dirs, key=lambda entry: entry.stat().st_mtime_ns)
237+ 
238+ 
239+def resolve_op_summary_csv(prof_dir: Path) -> Path:
240+ profiler_output_dir = prof_dir / "mindstudio_profiler_output"
241+ if not profiler_output_dir.is_dir():
242+ raise FileNotFoundError(f"mindstudio_profiler_output was not found in {prof_dir}")
243+ 
244+ csv_files = sorted(
245+ profiler_output_dir.glob(MSPROF_OP_SUMMARY_GLOB),
246+ key=lambda entry: entry.stat().st_mtime_ns,
247+ reverse=True,
248+ )
249+ if not csv_files:
250+ raise FileNotFoundError(f"No {MSPROF_OP_SUMMARY_GLOB} file was found in {profiler_output_dir}")
251+ return csv_files[0]
252+ 
253+ 
254+def parse_metric_value(raw_value: Optional[str], column_name: str, csv_path: Path) -> float:
255+ if raw_value is None:
256+ raise ValueError(f"{column_name} column was not found in {csv_path}")
257+ 
258+ normalized_value = raw_value.strip().replace(",", "")
259+ if column_name == "aic_icache_miss_rate":
260+ normalized_value = normalized_value.rstrip("%")
261+ 
262+ if not normalized_value:
263+ raise ValueError(f"{column_name} is empty in {csv_path}")
264+ 
265+ try:
266+ return float(normalized_value)
267+ except ValueError as error:
268+ raise ValueError(f"Failed to parse {column_name} value '{raw_value}' from {csv_path}") from error
269+ 
270+ 
271+def parse_profile_metrics_from_csv(csv_path: Path) -> ProfileMetrics:
272+ with csv_path.open("r", encoding="utf-8-sig", newline="") as csv_file:
273+ reader = csv.DictReader(csv_file)
274+ header = reader.fieldnames
275+ first_row = next(reader, None)
276+ 
277+ if not header:
278+ raise ValueError(f"CSV header is missing in {csv_path}")
279+ if not first_row:
280+ raise ValueError(f"CSV data row is missing in {csv_path}")
281+ 
282+ metric_values = {
283+ field_name: parse_metric_value(first_row.get(column_name), column_name, csv_path)
284+ for field_name, _display_name, column_name in PROFILE_METRIC_SPECS
285+ }
286+ metric_values["icache_miss_rate"] *= 100.0
287+ return ProfileMetrics(**metric_values)
288+ 
289+ 
290+def resolve_candidate_msprof_output_dir(script_dir: Path, executable_path: Path) -> Path:
291+ return script_dir / MSPROF_OUTPUT_DIR_NAME / executable_path.stem
292+ 
293+ 
294+def run_candidate_with_msprof(
295+ script_dir: Path, executable_path: Path, m: int, k: int, n: int, trans_a: bool, trans_b: bool
296+) -> ProfileMetrics:
297+ msprof_output_dir = resolve_candidate_msprof_output_dir(script_dir, executable_path)
298+ cleanup_msprof_output_dir(msprof_output_dir)
299+ msprof_output_dir.parent.mkdir(parents=True, exist_ok=True)
300+ application = f"./{executable_path.name}"
301+ cmd_args = ["msprof", f"--output={msprof_output_dir}", application, str(m), str(k), str(n)]
302+ cmd_args.extend([str(trans_a).lower(), str(trans_b).lower()])
303+ with tempfile.TemporaryFile(mode="w+t", encoding="utf-8") as log_file:
304+ result = subprocess.run(
305+ cmd_args,
306+ cwd=script_dir,
307+ text=True,
308+ stdout=log_file,
309+ stderr=subprocess.STDOUT,
310+ check=False,
311+ )
312+ if result.returncode != 0:
313+ raise RuntimeError(format_command_output("[msprof]", read_command_log(log_file)))
314+ 
315+ try:
316+ prof_dir = resolve_latest_prof_directory(msprof_output_dir)
317+ op_summary_csv = resolve_op_summary_csv(prof_dir)
318+ return parse_profile_metrics_from_csv(op_summary_csv)
319+ except Exception as error:
320+ command_output = format_command_output("[msprof]", read_command_log(log_file))
321+ raise RuntimeError(f"{command_output}\n[msprof parse error]\n{error}") from error
322+ 
323+ 
324+def run_candidate(
325+ script_dir: Path, candidate: Candidate, m: int, k: int, n: int, trans_a: bool, trans_b: bool
326+) -> CandidateResult:
327+ executable_path = resolve_executable(script_dir, candidate.executable_name)
328+ try:
329+ profile_metrics = run_candidate_with_msprof(script_dir, executable_path, m, k, n, trans_a, trans_b)
330+ kernel_time_us = profile_metrics.kernel_time_us
331+ output = ""
332+ return_code = 0
333+ except Exception as error:
334+ kernel_time_us = None
335+ profile_metrics = None
336+ output = str(error)
337+ return_code = 1
338+ 
339+ return CandidateResult(
340+ label=candidate.label,
341+ executable_path=executable_path,
342+ kernel_time_us=kernel_time_us,
343+ profile_metrics=profile_metrics,
344+ return_code=return_code,
345+ output=output,
346+ )
347+ 
348+ 
349+def format_metric_cell(value: float) -> str:
350+ return f"{value:.3f}"
351+ 
352+ 
353+def build_ascii_table(headers: List[str], rows: List[List[str]], right_aligned_columns: set[int]) -> List[str]:
354+ widths = []
355+ for column_index, header in enumerate(headers):
356+ column_values = [row[column_index] for row in rows]
357+ widths.append(max(len(header), *(len(value) for value in column_values)))
358+ 
359+ def format_row(row: List[str]) -> str:
360+ cells = []
361+ for column_index, value in enumerate(row):
362+ width = widths[column_index]
363+ if column_index in right_aligned_columns:
364+ cells.append(f" {value.rjust(width)} ")
365+ else:
366+ cells.append(f" {value.ljust(width)} ")
367+ return "|" + "|".join(cells) + "|"
368+ 
369+ border = "+" + "+".join("-" * (width + 2) for width in widths) + "+"
370+ header_separator = "+" + "+".join("=" * (width + 2) for width in widths) + "+"
371+ lines = [border, format_row(headers), header_separator]
372+ for row in rows:
373+ lines.append(format_row(row))
374+ lines.append(border)
375+ return lines
376+ 
377+ 
378+def print_profile_table(results: List[CandidateResult]) -> None:
379+ headers = ["algorithm"] + [display_name for _field_name, display_name, _column_name in PROFILE_METRIC_SPECS]
380+ rows = []
381+ for result in results:
382+ if result.profile_metrics is None:
383+ raise ValueError(f"Profile metrics are missing for algorithm {result.label}")
384+ metric_row = [result.label]
385+ for field_name, _display_name, _column_name in PROFILE_METRIC_SPECS:
386+ metric_row.append(format_metric_cell(getattr(result.profile_metrics, field_name)))
387+ rows.append(metric_row)
388+ 
389+ print("\n[Profile Breakdown]")
390+ for line in build_ascii_table(headers, rows, right_aligned_columns=set(range(1, len(headers)))):
391+ print(line)
392+ 
393+ 
394+def print_ranking(results: List[CandidateResult]) -> None:
395+ ranked_results = sorted(
396+ [item for item in results if item.succeeded],
397+ key=lambda item: item.kernel_time_us if item.kernel_time_us is not None else float("inf"),
398+ )
399+ 
400+ print("\n[Recommended Algorithm Ranking]")
401+ 
402+ if not ranked_results:
403+ print(" No compatible algorithm found for the current shape.")
404+ return
405+ 
406+ for index, result in enumerate(ranked_results, start=1):
407+ print(f" {index}. {result.label}")
408+ 
409+ print_profile_table(ranked_results)
410+ print("Note: Only algorithms that support the current shape are listed.\n")
411+ 
412+ 
413+def main(argv: List[str]) -> int:
414+ try:
415+ m, k, n, trans_a, trans_b, print_target_only = parse_arguments(argv)
416+ except ValueError as error:
417+ print(f"ERROR: {error}")
418+ print_usage(Path(argv[0]).name)
419+ return 1
420+ 
421+ script_dir = Path(__file__).resolve().parent
422+ msprof_output_dir = script_dir / MSPROF_OUTPUT_DIR_NAME
423+ candidates = discover_candidates(script_dir)
424+ if not candidates:
425+ print(f"ERROR: No executable files were found in {script_dir}")
426+ return 1
427+ 
428+ try:
429+ try:
430+ generate_input(script_dir, m, k, n, trans_a, trans_b)
431+ except Exception as error:
432+ print(f"ERROR: {error}")
433+ return 1
434+ 
435+ results: List[CandidateResult] = []
436+ for candidate in candidates:
437+ candidate_result = run_candidate(script_dir, candidate, m, k, n, trans_a, trans_b)
438+ results.append(candidate_result)
439+ 
440+ ranked_results = get_ranked_results(results)
441+ if print_target_only:
442+ if not ranked_results:
443+ print("ERROR: No compatible algorithm found for the current shape.")
444+ return 1
445+ print(ranked_results[0].label)
446+ return 0
447+ 
448+ print_ranking(results)
449+ return 0 if ranked_results else 1
450+ finally:
451+ cleanup_msprof_output_dir(msprof_output_dir)
452+ 
453+ 
454+if __name__ == "__main__":
455+ sys.exit(main(sys.argv))
456+ 
ASamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a32w32/scripts/run.sh+156-0
@@ -0,0 +1,156 @@
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+set -euo pipefail
12+ 
13+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
14+ 
15+find_repo_root() {
16+ local dir="$SCRIPT_DIR"
17+ while [[ "$dir" != "/" ]]; do
18+ if [[ -f "$dir/.ci/build.sh" ]]; then
19+ echo "$dir"
20+ return 0
21+ fi
22+ dir="$(dirname "$dir")"
23+ done
24+ return 1
25+}
26+ 
27+REPO_ROOT="$(find_repo_root || true)"
28+if [[ -z "$REPO_ROOT" ]]; then
29+ echo "ERROR: cannot locate repo root containing .ci/build.sh"
30+ exit 1
31+fi
32+ 
33+INSTALL_DIR="${REPO_ROOT}/build_out/2_Performance/matmul_story/matmul_recipes/matmul_a32w32"
34+TARGET=""
35+SKIP_BUILD=false
36+M=""
37+K=""
38+N=""
39+TRANS_A=""
40+TRANS_B=""
41+ 
42+usage() {
43+ cat <<'EOF'
44+Usage: bash run.sh [OPTIONS] m k n [transA transB]
45+ 
46+Options:
47+ --target <name> Specify executable name to run.
48+ --skip-build Skip build/install stage.
49+ -h, --help Show this help.
50+ 
51+Arguments:
52+ m k n Matrix dimensions (required).
53+ transA Transpose matrix A (optional, default: false).
54+ transB Transpose matrix B (optional, default: true).
55+ 
56+When --target is omitted, run.sh auto-selects target via:
57+ python3 matmul_a32w32_algorithm_recommend.py --print-target m k n [transA transB]
58+ 
59+Examples:
60+ bash run.sh 32 4096 64
61+ bash run.sh 32 4096 64 false true
62+ bash run.sh --target matmul_a32w32_swat 32 4096 64
63+EOF
64+}
65+ 
66+while [[ $# -gt 0 ]]; do
67+ case "$1" in
68+ --target)
69+ [[ -z "${2:-}" ]] && { echo "ERROR: --target needs a value"; exit 1; }
70+ TARGET="$2"
71+ shift 2
72+ ;;
73+ --skip-build)
74+ SKIP_BUILD=true
75+ shift
76+ ;;
77+ -h|--help)
78+ usage
79+ exit 0
80+ ;;
81+ -*)
82+ echo "ERROR: unknown option: $1"
83+ usage
84+ exit 1
85+ ;;
86+ *)
87+ if [[ -z "$M" ]]; then
88+ M="$1"
89+ elif [[ -z "$K" ]]; then
90+ K="$1"
91+ elif [[ -z "$N" ]]; then
92+ N="$1"
93+ elif [[ -z "$TRANS_A" ]]; then
94+ TRANS_A="$1"
95+ elif [[ -z "$TRANS_B" ]]; then
96+ TRANS_B="$1"
97+ else
98+ echo "ERROR: unexpected argument: $1"
99+ usage
100+ exit 1
101+ fi
102+ shift
103+ ;;
104+ esac
105+done
106+ 
107+if [[ -z "$M" || -z "$K" || -z "$N" ]]; then
108+ echo "ERROR: m k n are required"
109+ usage
110+ exit 1
111+fi
112+ 
113+if [[ -z "$TRANS_A" ]]; then
114+ TRANS_A="false"
115+fi
116+if [[ -z "$TRANS_B" ]]; then
117+ TRANS_B="true"
118+fi
119+ 
120+if [[ "$SKIP_BUILD" != true ]]; then
121+ bash "${REPO_ROOT}/.ci/build.sh" dav-3510
122+fi
123+ 
124+if [[ ! -d "$INSTALL_DIR" ]]; then
125+ echo "ERROR: install dir not found: $INSTALL_DIR"
126+ echo "Hint: remove --skip-build for full build/install."
127+ exit 1
128+fi
129+ 
130+cd "$INSTALL_DIR"
131+ 
132+if [[ "$TRANS_A" == "false" && "$TRANS_B" == "true" ]]; then
133+ python3 gen_data.py "$M" "$K" "$N"
134+else
135+ python3 gen_data.py "$M" "$K" "$N" "$TRANS_A" "$TRANS_B"
136+fi
137+ 
138+if [[ -z "$TARGET" ]]; then
139+ if [[ "$TRANS_A" == "false" && "$TRANS_B" == "true" ]]; then
140+ TARGET="$(python3 matmul_a32w32_algorithm_recommend.py --print-target "$M" "$K" "$N")"
141+ else
142+ TARGET="$(python3 matmul_a32w32_algorithm_recommend.py --print-target "$M" "$K" "$N" "$TRANS_A" "$TRANS_B")"
143+ fi
144+fi
145+ 
146+if [[ -z "$TARGET" ]]; then
147+ echo "ERROR: failed to select a target executable"
148+ exit 1
149+fi
150+ 
151+if [[ ! -x "./$TARGET" ]]; then
152+ echo "ERROR: executable not found: $INSTALL_DIR/$TARGET"
153+ exit 1
154+fi
155+ 
156+"./$TARGET" "$M" "$K" "$N" "$TRANS_A" "$TRANS_B"
ASamples/2_Performance/matmul_story/matmul_recipes/examples/matmul_a32w32/scripts/verify_result.py+143-0
@@ -0,0 +1,143 @@
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+import os
12+import sys
13+ 
14+os.environ["TORCH_DEVICE_BACKEND_AUTOLOAD"] = "0"
15+ 
16+import numpy as np
17+import torch
18+ 
19+POINT_ERROR_TOL = 1e-1
20+RATIO_POINT_ERROR_TOL = 1e-4
21+ERROR_RATIO_TOL = 1e-3
22+DATA_TYPE = np.float32
23+ 
24+FULL_TENSOR_PRINT_MAX_ELEMENTS = 128
25+CORNER_ROWS = 4
26+CORNER_COLS = 4
27+ 
28+ 
29+def _print_large_tensor_summary(golden_tensor: torch.Tensor, npu_output_tensor: torch.Tensor, m: int, n: int) -> None:
30+ g = golden_tensor.float()
31+ p = npu_output_tensor.float()
32+ diff = p - g
33+ abs_err = diff.abs()
34+ denom = g.abs().clamp_min(1e-8)
35+ rel_err = abs_err / denom
36+ 
37+ numel = m * n
38+ max_ab = torch.maximum(g, p)
39+ ratio_rel_err = torch.where(
40+ max_ab > 0,
41+ abs_err / max_ab,
42+ torch.where(abs_err == 0, torch.zeros_like(abs_err), torch.full_like(abs_err, float("inf"))),
43+ )
44+ over_tol = (ratio_rel_err > RATIO_POINT_ERROR_TOL).sum().item()
45+ 
46+ print(f"\n[verify] shape=({m}, {n}), elements={numel} - summary (large matrix, full tensors omitted)")
47+ print(
48+ f" abs_err: max={abs_err.max().item():.6e}, mean={abs_err.mean().item():.6e}, "
49+ f"rmse={(diff.pow(2).mean().sqrt()).item():.6e}"
50+ )
51+ print(f" rel_err: max={rel_err.max().item():.6e}")
52+ print(f" count(rel_err(abs(a-b)/max(a,b)) > {RATIO_POINT_ERROR_TOL:g}): {over_tol} / {numel}")
53+ 
54+ cr = min(CORNER_ROWS, m)
55+ cc = min(CORNER_COLS, n)
56+ if cr > 0 and cc > 0:
57+ print(f" cpu golden (top-left {cr}x{cc}):\n{golden_tensor[:cr, :cc]}")
58+ print(f" npu output (top-left {cr}x{cc}):\n{npu_output_tensor[:cr, :cc]}")
59+ 
60+ 
61+def verify_result(m, n):
62+ output = np.fromfile("./output/npu_out.bin", dtype=DATA_TYPE)
63+ golden = np.fromfile("./output/cpu_output.bin", dtype=DATA_TYPE)
64+ 
65+ if output.size != golden.size:
66+ raise ValueError("npu output size != cpu output size")
67+ 
68+ npu_output_tensor = torch.from_numpy(output).view(torch.float32).reshape(m, n)
69+ golden_tensor = torch.from_numpy(golden).view(torch.float32).reshape(m, n)
70+ 
71+ numel = m * n
72+ if numel <= FULL_TENSOR_PRINT_MAX_ELEMENTS:
73+ print("\ncpu golden:\n", golden_tensor)
74+ print("npu output:\n", npu_output_tensor)
75+ else:
76+ _print_large_tensor_summary(golden_tensor, npu_output_tensor, m, n)
77+ golden_f32 = golden_tensor.to(torch.float32)
78+ npu_f32 = npu_output_tensor.to(torch.float32)
79+ abs_diff = torch.abs(golden_f32 - npu_f32)
80+ non_finite_mask = ~(torch.isfinite(golden_f32) & torch.isfinite(npu_f32) & torch.isfinite(abs_diff))
81+ abs_golden = torch.abs(golden_f32)
82+ rel_diff = torch.where(
83+ abs_golden > 0,
84+ abs_diff / abs_golden,
85+ torch.where(abs_diff == 0, torch.zeros_like(abs_diff), torch.full_like(abs_diff, float("inf"))),
86+ )
87+ max_ab = torch.maximum(golden_f32, npu_f32)
88+ ratio_rel_diff = torch.where(
89+ max_ab > 0,
90+ abs_diff / max_ab,
91+ torch.where(abs_diff == 0, torch.zeros_like(abs_diff), torch.full_like(abs_diff, float("inf"))),
92+ )
93+ point_error_mask = (rel_diff > POINT_ERROR_TOL) | non_finite_mask
94+ ratio_error_mask = (ratio_rel_diff > RATIO_POINT_ERROR_TOL) | non_finite_mask
95+ point_error_count = int(point_error_mask.sum().item())
96+ error_count = int(ratio_error_mask.sum().item())
97+ error_ratio = error_count / numel if numel else 0.0
98+ 
99+ print(f"max abs diff: {abs_diff.max().item() if numel else 0.0}")
100+ print(f"point error count(>{POINT_ERROR_TOL}): {point_error_count}/{numel}")
101+ if point_error_count > 0:
102+ point_error_indices = torch.nonzero(point_error_mask, as_tuple=False)
103+ print(f"point error details(rel diff > {POINT_ERROR_TOL} or non-finite):")
104+ for idx in point_error_indices:
105+ row = int(idx[0].item())
106+ col = int(idx[1].item())
107+ golden_val = float(golden_f32[row, col].item())
108+ npu_val = float(npu_f32[row, col].item())
109+ diff_val = float(abs_diff[row, col].item())
110+ rel_val = float(rel_diff[row, col].item())
111+ print(
112+ f" (row={row}, col={col}) "
113+ f"golden={golden_val}, npu={npu_val}, abs_diff={diff_val}, rel_diff={rel_val}"
114+ )
115+ print(
116+ f"ratio error count(rel_diff(abs(a-b)/max(a,b)) > {RATIO_POINT_ERROR_TOL}): {error_count}/{numel}, "
117+ f"error ratio: {error_ratio:.6f}"
118+ )
119+ 
120+ return point_error_count == 0 and error_ratio <= ERROR_RATIO_TOL
121+ 
122+ 
123+if __name__ == "__main__":
124+ if len(sys.argv) != 3:
125+ print("Usage: python3 verify_result.py m n")
126+ sys.exit(1)
127+ 
128+ m = int(sys.argv[1])
129+ n = int(sys.argv[2])
130+ try:
131+ res = verify_result(m, n)
132+ if not res:
133+ raise ValueError(
134+ f"[ERROR] NPU results differ from CPU. "
135+ f"Single-point relative error (abs_diff/abs(golden)) must be <= {POINT_ERROR_TOL}, "
136+ f"and the ratio of points with relative error abs(a-b)/max(a,b) > {RATIO_POINT_ERROR_TOL} "
137+ f"must be <= {ERROR_RATIO_TOL}.\n"
138+ )
139+ print("[PASS] NPU results are consistent with CPU.\n")
140+ 
141+ except Exception as e:
142+ print(e)
143+ sys.exit(1)
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/block_mmad.h+2-2
@@ -27,8 +27,8 @@ class BlockMmad {
27} // namespace Block27} // namespace Block
28 28 
29// Include all concrete BlockMmad specializations here.29// Include all concrete BlockMmad specializations here.
30-#include "matmul_a16w16_block_mmad_swat.h"30+#include "matmul_block_mmad_swat.h"
31-#include "matmul_a16w16_block_mmad_streamk.h"31+#include "matmul_block_mmad_streamk.h"
32#include "quant_matmul_mx_block_mmad_swat.h"32#include "quant_matmul_mx_block_mmad_swat.h"
33#include "quant_matmul_mx_block_mmad_swat_4_buffer.h"33#include "quant_matmul_mx_block_mmad_swat_4_buffer.h"
34#include "quant_matmul_mx_block_mmad_a_full_load.h"34#include "quant_matmul_mx_block_mmad_a_full_load.h"
MSamples/2_Performance/matmul_story/matmul_recipes/include/block/block_scheduler_policy.h+2-2
@@ -28,11 +28,11 @@ struct QuantMatmulMxSwatScheduler {
28};28};
29 29 
30template <uint64_t FULL_LOAD_MODE_>30template <uint64_t FULL_LOAD_MODE_>
31-struct MatmulA16W16SwatScheduler {31+struct MatmulSwatScheduler {
32 static constexpr uint64_t fullLoadMode = FULL_LOAD_MODE_;32 static constexpr uint64_t fullLoadMode = FULL_LOAD_MODE_;
33};33};
34 34 
35-struct MatmulA16W16StreamKScheduler {};35+struct MatmulStreamKScheduler {};
36 36 
37template <uint64_t FULL_LOAD_MODE_>37template <uint64_t FULL_LOAD_MODE_>
38struct QuantBatchMatmulV3Scheduler {38struct QuantBatchMatmulV3Scheduler {
RSamples/2_Performance/matmul_story/matmul_recipes/include/block/matmul_a16w16_block_mmad_streamk.h→Samples/2_Performance/matmul_story/matmul_recipes/include/block/matmul_block_mmad_streamk.h+3-3
@@ -9,8 +9,8 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_block_mmad_streamk.h12+ * \file matmul_block_mmad_streamk.h
13- * \brief Block-level A16W16 MMAD pipeline for StreamK path.13+ * \brief Block-level MMAD pipeline for StreamK path.
14 */14 */
15 15 
16#pragma once16#pragma once
@@ -27,7 +27,7 @@ template <
27 class DispatchPolicy_, class TypeA_, class LayoutA_, class TypeB_, class LayoutB_, class TypeC_, class LayoutC_>27 class DispatchPolicy_, class TypeA_, class LayoutA_, class TypeB_, class LayoutB_, class TypeC_, class LayoutC_>
28class BlockMmad<28class BlockMmad<
29 DispatchPolicy_, TypeA_, LayoutA_, TypeB_, LayoutB_, TypeC_, LayoutC_,29 DispatchPolicy_, TypeA_, LayoutA_, TypeB_, LayoutB_, TypeC_, LayoutC_,
30- AscendC::Std::enable_if_t<AscendC::Std::is_base_of_v<MatmulA16W16MultiBlockWithStreamK, DispatchPolicy_>>> {30+ AscendC::Std::enable_if_t<AscendC::Std::is_base_of_v<MatmulMultiBlockWithStreamK, DispatchPolicy_>>> {
31public:31public:
32 using TypeA = TypeA_;32 using TypeA = TypeA_;
33 using TypeB = TypeB_;33 using TypeB = TypeB_;
RSamples/2_Performance/matmul_story/matmul_recipes/include/block/matmul_a16w16_block_mmad_swat.h→Samples/2_Performance/matmul_story/matmul_recipes/include/block/matmul_block_mmad_swat.h+3-3
@@ -9,8 +9,8 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_block_mmad_swat.h12+ * \file matmul_block_mmad_swat.h
13- * \brief Block-level A16W16 MMAD pipeline for SWAT non-full-load path.13+ * \brief Block-level MMAD pipeline for SWAT non-full-load path.
14 */14 */
15 15 
16#pragma once16#pragma once
@@ -28,7 +28,7 @@ template <
28class BlockMmad<28class BlockMmad<
29 DispatchPolicy_, TypeA_, LayoutA_, TypeB_, LayoutB_, TypeC_, LayoutC_,29 DispatchPolicy_, TypeA_, LayoutA_, TypeB_, LayoutB_, TypeC_, LayoutC_,
30 AscendC::Std::enable_if_t<30 AscendC::Std::enable_if_t<
31- AscendC::Std::is_base_of_v<MatmulA16W16MultiBlockWithSwat<NO_FULL_LOAD_MODE>, DispatchPolicy_>>> {31+ AscendC::Std::is_base_of_v<MatmulMultiBlockWithSwat<NO_FULL_LOAD_MODE>, DispatchPolicy_>>> {
32public:32public:
33 using TypeA = TypeA_;33 using TypeA = TypeA_;
34 using TypeB = TypeB_;34 using TypeB = TypeB_;
RSamples/2_Performance/matmul_story/matmul_recipes/include/block/matmul_a16w16_block_scheduler_streamk.h→Samples/2_Performance/matmul_story/matmul_recipes/include/block/matmul_block_scheduler_streamk.h+7-7
@@ -9,21 +9,21 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_block_scheduler_streamk.h12+ * \file matmul_block_scheduler_streamk.h
13- * \brief StreamK block scheduler for A16W16 implementation.13+ * \brief StreamK block scheduler for implementation.
14 */14 */
15 15 
16#pragma once16#pragma once
17 17 
18#include "kernel_utils/common_utils.h"18#include "kernel_utils/common_utils.h"
19#include "utils/constant.h"19#include "utils/constant.h"
20-#include "tiling/matmul_a16w16_tiling_data.h"20+#include "tiling/matmul_tiling_data.h"
21#include "./block_scheduler_utils.h"21#include "./block_scheduler_utils.h"
22 22 
23namespace Block {23namespace Block {
24 24 
25template <class ProblemShape_>25template <class ProblemShape_>
26-class BlockSchedulerA16W16StreamK {26+class BlockSchedulerStreamK {
27public:27public:
28 int64_t usedCoreNum_{0};28 int64_t usedCoreNum_{0};
29 int64_t mTileNum_{0};29 int64_t mTileNum_{0};
@@ -65,7 +65,7 @@ public:
65 };65 };
66 66 
67public:67public:
68- __aicore__ inline BlockSchedulerA16W16StreamK(const ProblemShape& shape, const Params& params)68+ __aicore__ inline BlockSchedulerStreamK(const ProblemShape& shape, const Params& params)
69 {69 {
70 usedCoreNum_ = params.usedCoreNum;70 usedCoreNum_ = params.usedCoreNum;
71 m_ = shape.m;71 m_ = shape.m;
@@ -179,8 +179,8 @@ private:
179};179};
180 180 
181template <class ProblemShape_, bool TransA_, bool TransB_>181template <class ProblemShape_, bool TransA_, bool TransB_>
182-struct BlockSchedulerSelector<ProblemShape_, MatmulA16W16StreamKScheduler, TransA_, TransB_> {182+struct BlockSchedulerSelector<ProblemShape_, MatmulStreamKScheduler, TransA_, TransB_> {
183- using SchedulerOp = BlockSchedulerA16W16StreamK<ProblemShape_>;183+ using SchedulerOp = BlockSchedulerStreamK<ProblemShape_>;
184};184};
185 185 
186} // namespace Block186} // namespace Block
RSamples/2_Performance/matmul_story/matmul_recipes/include/block/matmul_a16w16_block_scheduler_swat.h→Samples/2_Performance/matmul_story/matmul_recipes/include/block/matmul_block_scheduler_swat.h+7-7
@@ -9,21 +9,21 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_block_scheduler_swat.h12+ * \file matmul_block_scheduler_swat.h
13- * \brief SWAT block scheduler for *A16W16 non-full-load path.13+ * \brief SWAT block scheduler for * non-full-load path.
14 */14 */
15 15 
16#pragma once16#pragma once
17 17 
18#include "kernel_utils/common_utils.h"18#include "kernel_utils/common_utils.h"
19#include "utils/constant.h"19#include "utils/constant.h"
20-#include "tiling/matmul_a16w16_tiling_data.h"20+#include "tiling/matmul_tiling_data.h"
21#include "./block_scheduler_utils.h"21#include "./block_scheduler_utils.h"
22 22 
23namespace Block {23namespace Block {
24 24 
25template <class ProblemShape_>25template <class ProblemShape_>
26-class BlockSchedulerA16W16Swat {26+class BlockSchedulerSwat {
27public:27public:
28 int64_t mTileNum_{0};28 int64_t mTileNum_{0};
29 int64_t nTileNum_{0};29 int64_t nTileNum_{0};
@@ -85,7 +85,7 @@ public:
85 };85 };
86 86 
87public:87public:
88- __aicore__ inline BlockSchedulerA16W16Swat(88+ __aicore__ inline BlockSchedulerSwat(
89 const ProblemShape& shape, int64_t blockIdx, int64_t blockNum, const Params& params)89 const ProblemShape& shape, int64_t blockIdx, int64_t blockNum, const Params& params)
90 : blockIdx_(blockIdx), blockNum_(blockNum)90 : blockIdx_(blockIdx), blockNum_(blockNum)
91 {91 {
@@ -225,7 +225,7 @@ private:
225};225};
226 226 
227template <class ProblemShape_, bool TransA_, bool TransB_>227template <class ProblemShape_, bool TransA_, bool TransB_>
228-struct BlockSchedulerSelector<ProblemShape_, MatmulA16W16SwatScheduler<NO_FULL_LOAD_MODE>, TransA_, TransB_> {228+struct BlockSchedulerSelector<ProblemShape_, MatmulSwatScheduler<NO_FULL_LOAD_MODE>, TransA_, TransB_> {
229- using SchedulerOp = BlockSchedulerA16W16Swat<ProblemShape_>;229+ using SchedulerOp = BlockSchedulerSwat<ProblemShape_>;
230};230};
231} // namespace Block231} // namespace Block
RSamples/2_Performance/matmul_story/matmul_recipes/include/epilogue/matmul_a16w16_block_epilogue_streamk.h→Samples/2_Performance/matmul_story/matmul_recipes/include/epilogue/matmul_block_epilogue_streamk.h+2-2
@@ -9,8 +9,8 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_block_epilogue_streamk.h12+ * \file matmul_block_epilogue_streamk.h
13- * \brief Block-level A16W16 accumulation operation in the Vector Core for StreamK path.13+ * \brief Block-level accumulation operation in the Vector Core for StreamK path.
14 */14 */
15 15 
16#pragma once16#pragma once
RSamples/2_Performance/matmul_story/matmul_recipes/include/kernel/matmul_a16w16_kernel_streamk.h→Samples/2_Performance/matmul_story/matmul_recipes/include/kernel/matmul_kernel_streamk.h+21-8
@@ -9,8 +9,8 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_kernel_streamk.h12+ * \file matmul_kernel_streamk.h
13- * \brief Kernel-side StreamK A16W16 implementation.13+ * \brief Kernel-side StreamK implementation.
14 */14 */
15 15 
16#pragma once16#pragma once
@@ -25,18 +25,18 @@
25#include "kernel_utils/common_utils.h"25#include "kernel_utils/common_utils.h"
26#include "include/tensor_api/tensor.h"26#include "include/tensor_api/tensor.h"
27 27 
28-#include "../block/matmul_a16w16_block_mmad_streamk.h"28+#include "../block/matmul_block_mmad_streamk.h"
29-#include "../block/matmul_a16w16_block_scheduler_streamk.h"29+#include "../block/matmul_block_scheduler_streamk.h"
30#include "../utils/constant.h"30#include "../utils/constant.h"
31 31 
32namespace Kernel {32namespace Kernel {
33 33 
34template <class ProblemShape, class BlockMmad, class BlockScheduler, class BlockEpilogue>34template <class ProblemShape, class BlockMmad, class BlockScheduler, class BlockEpilogue>
35-class MatmulA16W16KernelStreamK {35+class MatmulKernelStreamK {
36public:36public:
37- __aicore__ inline MatmulA16W16KernelStreamK()37+ __aicore__ inline MatmulKernelStreamK()
38 {}38 {}
39- __aicore__ inline ~MatmulA16W16KernelStreamK()39+ __aicore__ inline ~MatmulKernelStreamK()
40 {}40 {}
41 41 
42 static constexpr bool transA = BlockMmad::transA;42 static constexpr bool transA = BlockMmad::transA;
@@ -71,6 +71,7 @@ public:
71 uint32_t baseN;71 uint32_t baseN;
72 uint32_t baseK;72 uint32_t baseK;
73 uint32_t usedCoreNum;73 uint32_t usedCoreNum;
74+ bool isHf32;
74 };75 };
75 76 
76 struct Params {77 struct Params {
@@ -84,6 +85,12 @@ public:
84public:85public:
85 __aicore__ inline void operator()(const Params& params);86 __aicore__ inline void operator()(const Params& params);
86 87 
88+ __aicore__ inline void UnsetHf32(bool isHf32) {
89+ if (isHf32) {
90+ AscendC::SetHF32Mode(0);
91+ }
92+ }
93+ 
87private:94private:
88 __aicore__ inline TupleShape ToShapeTuple(const ProblemShape& problemShape)95 __aicore__ inline TupleShape ToShapeTuple(const ProblemShape& problemShape)
89 {96 {
@@ -97,7 +104,7 @@ private:
97};104};
98 105 
99template <class ProblemShape, class BlockMmad, class BlockScheduler, class BlockEpilogue>106template <class ProblemShape, class BlockMmad, class BlockScheduler, class BlockEpilogue>
100-__aicore__ inline void MatmulA16W16KernelStreamK<ProblemShape, BlockMmad, BlockScheduler, BlockEpilogue>::operator()(107+__aicore__ inline void MatmulKernelStreamK<ProblemShape, BlockMmad, BlockScheduler, BlockEpilogue>::operator()(
101 const Params& params)108 const Params& params)
102{109{
103 usedCoreNum_ = params.kernelParams.usedCoreNum;110 usedCoreNum_ = params.kernelParams.usedCoreNum;
@@ -125,6 +132,11 @@ __aicore__ inline void MatmulA16W16KernelStreamK<ProblemShape, BlockMmad, BlockS
125 AscendC::CrossCoreSetFlag<AIC_SYNC_AIV_MODE_4, PIPE_FIX>(AIC_SYNC_AIV_FLAG + FLAG_ID_MAX);132 AscendC::CrossCoreSetFlag<AIC_SYNC_AIV_MODE_4, PIPE_FIX>(AIC_SYNC_AIV_FLAG + FLAG_ID_MAX);
126 return;133 return;
127 }134 }
135+ // enbale Hf32
136+ if (params.kernelParams.isHf32) {
137+ AscendC::SetHF32Mode(1);
X
Xxubinglin6月12日

新增 isHf32 后，kernel 启用 SetHF32Mode(1) 和 SetHF32TransMode(1)，但退出时只恢复 SetHF32Mode(0)，没有恢复 SetHF32TransMode(0)。如果 HF32 trans mode 属于核内/线程上下文状态，后续计算可能继承错误转换模式。

likedislike
W
wuyufei
6月12日 评论：
138+ AscendC::SetHF32TransMode(1);
139+ }
128 AscendC::SetMMLayoutTransform(true);140 AscendC::SetMMLayoutTransform(true);
129 BlockMmad blockMmadOp(problemShape_, tileL1, tileL0);141 BlockMmad blockMmadOp(problemShape_, tileL1, tileL0);
130 int64_t tailSKTotalTileNum = static_cast<int64_t>(((mTileNum * nTileNum) % usedCoreNum_) * skKTileNum);142 int64_t tailSKTotalTileNum = static_cast<int64_t>(((mTileNum * nTileNum) % usedCoreNum_) * skKTileNum);
@@ -187,6 +199,7 @@ __aicore__ inline void MatmulA16W16KernelStreamK<ProblemShape, BlockMmad, BlockS
187 }199 }
188 }200 }
189 AscendC::SetMMLayoutTransform(false);201 AscendC::SetMMLayoutTransform(false);
202+ UnsetHf32(params.kernelParams.isHf32);
190 }203 }
191 if ASCEND_IS_AIV {204 if ASCEND_IS_AIV {
192 uint64_t lastLoopTotalCnt = (mTileNum * nTileNum % usedCoreNum_) * skKTileNum;205 uint64_t lastLoopTotalCnt = (mTileNum * nTileNum % usedCoreNum_) * skKTileNum;
RSamples/2_Performance/matmul_story/matmul_recipes/include/kernel/matmul_a16w16_kernel_swat.h→Samples/2_Performance/matmul_story/matmul_recipes/include/kernel/matmul_kernel_swat.h+22-9
@@ -9,8 +9,8 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_kernel_swat.h12+ * \file matmul_kernel_swat.h
13- * \brief Kernel-side SWAT A16W16 implementation for the non-full-load path.13+ * \brief Kernel-side SWAT implementation for the non-full-load path.
14 */14 */
15 15 
16#pragma once16#pragma once
@@ -25,18 +25,18 @@
25#include "kernel_utils/common_utils.h"25#include "kernel_utils/common_utils.h"
26#include "include/tensor_api/tensor.h"26#include "include/tensor_api/tensor.h"
27 27 
28-#include "../block/matmul_a16w16_block_mmad_swat.h"28+#include "../block/matmul_block_mmad_swat.h"
29-#include "../block/matmul_a16w16_block_scheduler_swat.h"29+#include "../block/matmul_block_scheduler_swat.h"
30#include "../utils/constant.h"30#include "../utils/constant.h"
31 31 
32namespace Kernel {32namespace Kernel {
33 33 
34template <class ProblemShape, class BlockMmad, class BlockScheduler>34template <class ProblemShape, class BlockMmad, class BlockScheduler>
35-class MatmulA16W16KernelSwat {35+class MatmulKernelSwat {
36public:36public:
37- __aicore__ inline MatmulA16W16KernelSwat()37+ __aicore__ inline MatmulKernelSwat()
38 {}38 {}
39- __aicore__ inline ~MatmulA16W16KernelSwat()39+ __aicore__ inline ~MatmulKernelSwat()
40 {}40 {}
41 41 
42 static constexpr bool transA = BlockMmad::transA;42 static constexpr bool transA = BlockMmad::transA;
@@ -67,6 +67,7 @@ public:
67 uint32_t baseM;67 uint32_t baseM;
68 uint32_t baseN;68 uint32_t baseN;
69 uint32_t baseK;69 uint32_t baseK;
70+ bool isHf32;
70 uint8_t dbL0C;71 uint8_t dbL0C;
71 };72 };
72 73 
@@ -80,6 +81,12 @@ public:
80public:81public:
81 __aicore__ inline void operator()(const Params& params);82 __aicore__ inline void operator()(const Params& params);
82 83 
84+ __aicore__ inline void UnsetHf32(bool isHf32) {
85+ if (isHf32) {
86+ AscendC::SetHF32Mode(0);
87+ }
88+ }
89+ 
83private:90private:
84 __aicore__ inline TupleShape ToShapeTuple(const ProblemShape& problemShape)91 __aicore__ inline TupleShape ToShapeTuple(const ProblemShape& problemShape)
85 {92 {
@@ -92,7 +99,7 @@ private:
92};99};
93 100 
94template <class ProblemShape, class BlockMmad, class BlockScheduler>101template <class ProblemShape, class BlockMmad, class BlockScheduler>
95-__aicore__ inline void MatmulA16W16KernelSwat<ProblemShape, BlockMmad, BlockScheduler>::operator()(const Params& params)102+__aicore__ inline void MatmulKernelSwat<ProblemShape, BlockMmad, BlockScheduler>::operator()(const Params& params)
96{103{
97 if ASCEND_IS_AIV {104 if ASCEND_IS_AIV {
98 return;105 return;
@@ -109,8 +116,13 @@ __aicore__ inline void MatmulA16W16KernelSwat<ProblemShape, BlockMmad, BlockSche
109 if (curBlockIdx >= realBlockNum) {116 if (curBlockIdx >= realBlockNum) {
110 return;117 return;
111 }118 }
112- AscendC::SetMMLayoutTransform(true);
113 bool l0cDB = params.kernelParams.dbL0C > 1;119 bool l0cDB = params.kernelParams.dbL0C > 1;
120+ // enbale Hf32
121+ if (params.kernelParams.isHf32) {
122+ AscendC::SetHF32Mode(1);
123+ AscendC::SetHF32TransMode(1);
124+ }
125+ AscendC::SetMMLayoutTransform(true);
114 // Instantiate mmadOp126 // Instantiate mmadOp
115 BlockMmad blockMmadOp(problemShape_, tileL1, tileL0, l0cDB);127 BlockMmad blockMmadOp(problemShape_, tileL1, tileL0, l0cDB);
116 128 
@@ -148,6 +160,7 @@ __aicore__ inline void MatmulA16W16KernelSwat<ProblemShape, BlockMmad, BlockSche
148 }160 }
149 161 
150 AscendC::SetMMLayoutTransform(false);162 AscendC::SetMMLayoutTransform(false);
163+ UnsetHf32(params.kernelParams.isHf32);
151}164}
152 165 
153} // namespace Kernel166} // namespace Kernel
MSamples/2_Performance/matmul_story/matmul_recipes/include/policy/dispatch_policy.h+4-4
@@ -30,22 +30,22 @@ struct QuantMatmulMxMultiBlockWithSwat {
30};30};
31 31 
32/**32/**
33- * @brief Dispatch tag for MatmulA16W16 matmul kernels that use the SWAT33+ * @brief Dispatch tag for Matmul matmul kernels that use the SWAT
34 * scheduling family.34 * scheduling family.
35 * @tparam SingleCoreShape Placeholder for the per-core tile shape recorded35 * @tparam SingleCoreShape Placeholder for the per-core tile shape recorded
36 * in the dispatch traits.36 * in the dispatch traits.
37 * @tparam FULL_LOAD_MODE_ Selects the SWAT variant: streaming or A-full-load or B-full-load.37 * @tparam FULL_LOAD_MODE_ Selects the SWAT variant: streaming or A-full-load or B-full-load.
38 */38 */
39template <uint64_t FULL_LOAD_MODE_>39template <uint64_t FULL_LOAD_MODE_>
40-struct MatmulA16W16MultiBlockWithSwat {40+struct MatmulMultiBlockWithSwat {
41 static constexpr uint64_t fullLoadMode = FULL_LOAD_MODE_;41 static constexpr uint64_t fullLoadMode = FULL_LOAD_MODE_;
42};42};
43 43 
44/**44/**
45- * @brief Dispatch tag for MatmulA16W16 matmul kernels that use the StreamK45+ * @brief Dispatch tag for Matmul matmul kernels that use the StreamK
46 * scheduling family.46 * scheduling family.
47 */47 */
48-struct MatmulA16W16MultiBlockWithStreamK {};48+struct MatmulMultiBlockWithStreamK {};
49 49 
50template <uint64_t FULL_LOAD_MODE_>50template <uint64_t FULL_LOAD_MODE_>
51struct MatmulWithScale {51struct MatmulWithScale {
RSamples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_a16w16_tiling_base.h→Samples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_tiling_base.h+18-15
@@ -9,29 +9,29 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_tiling_base.h12+ * \file matmul_tiling_base.h
13- * \brief Base tiling class for A16W16 matmul.13+ * \brief Base tiling class for matmul.
14 */14 */
15 15 
16#pragma once16#pragma once
17 17 
18#include <cstdint>18#include <cstdint>
19#include <memory>19#include <memory>
20-#include "tiling/matmul_a16w16_tiling_data.h"20+#include "tiling/matmul_tiling_data.h"
21-#include "tiling/matmul_a16w16_tiling_common.h"21+#include "tiling/matmul_tiling_common.h"
22#include "utils/constant.h"22#include "utils/constant.h"
23#include "platform/platform_ascendc.h"23#include "platform/platform_ascendc.h"
24 24 
25-class MatmulA16W16TilingBase {25+class MatmulTilingBase {
26public:26public:
27- MatmulA16W16TilingBase() = default;27+ MatmulTilingBase() = default;
28- virtual ~MatmulA16W16TilingBase() = default;28+ virtual ~MatmulTilingBase() = default;
29 29 
30 virtual void GetTilingData(30 virtual void GetTilingData(
31- uint64_t m, uint64_t n, uint64_t k, bool transA, bool transB, MatmulA16W16TilingData& tilingData)31+ uint64_t m, uint64_t n, uint64_t k, uint64_t dataTypeSize, bool transA, bool transB, MatmulTilingData& tilingData)
32 {32 {
33 InitCompileInfo();33 InitCompileInfo();
34- InitShapeArgs(m, n, k, transA, transB);34+ InitShapeArgs(m, n, k, dataTypeSize, transA, transB);
35 DoOpTiling(tilingData);35 DoOpTiling(tilingData);
36 PrintTilingData(tilingData);36 PrintTilingData(tilingData);
37 };37 };
@@ -43,12 +43,12 @@ public:
43 };43 };
44 44 
45protected:45protected:
46- MatmulA16W16PlatformInfo platformInfo_;46+ MatmulPlatformInfo platformInfo_;
47- MatmulA16W16Args args_;47+ MatmulArgs args_;
48- MatmulA16W16RunInfo runInfo_;48+ MatmulRunInfo runInfo_;
49 49 
50 virtual const char* TilingName() const = 0;50 virtual const char* TilingName() const = 0;
51- virtual void DoOpTiling(MatmulA16W16TilingData& tilingData) = 0;51+ virtual void DoOpTiling(MatmulTilingData& tilingData) = 0;
52 52 
53private:53private:
54 void InitCompileInfo()54 void InitCompileInfo()
@@ -66,17 +66,19 @@ private:
66 ascendcPlatform->GetCoreMemSize(platform_ascendc::CoreMemType::BT, platformInfo_.btSize);66 ascendcPlatform->GetCoreMemSize(platform_ascendc::CoreMemType::BT, platformInfo_.btSize);
67 }67 }
68 68 
69- void InitShapeArgs(uint64_t m, uint64_t n, uint64_t k, bool transA, bool transB, bool hasBias = false)69+ void InitShapeArgs(uint64_t m, uint64_t n, uint64_t k, uint64_t dataTypeSize, bool transA, bool transB, bool isHf32 = true, bool hasBias = false)
70 {70 {
71 args_.m = m;71 args_.m = m;
72 args_.n = n;72 args_.n = n;
73 args_.k = k;73 args_.k = k;
74 args_.isATrans = transA;74 args_.isATrans = transA;
75 args_.isBTrans = transB;75 args_.isBTrans = transB;
76+ args_.dataTypeSize = dataTypeSize;
77+ args_.isHf32 = dataTypeSize == DATA_SIZE_FP32 ? isHf32 : false;
76 args_.hasBias = hasBias;78 args_.hasBias = hasBias;
77 }79 }
78 80 
79- void PrintTilingData(const MatmulA16W16TilingData& tilingData) const81+ void PrintTilingData(const MatmulTilingData& tilingData) const
80 {82 {
81 printf("[Matmul Strategy]\n");83 printf("[Matmul Strategy]\n");
82 printf(" strategy : %s\n", TilingName());84 printf(" strategy : %s\n", TilingName());
@@ -98,6 +100,7 @@ private:
98 printf(" nBaseTailSplitCnt : %u\n", tilingData.nBaseTailSplitCnt);100 printf(" nBaseTailSplitCnt : %u\n", tilingData.nBaseTailSplitCnt);
99 printf(" mTailMain : %u\n", tilingData.mTailMain);101 printf(" mTailMain : %u\n", tilingData.mTailMain);
100 printf(" nTailMain : %u\n", tilingData.nTailMain);102 printf(" nTailMain : %u\n", tilingData.nTailMain);
103+ printf(" isHf32 : %u\n", tilingData.isHf32);
101 printf(" l1BufferNum : %u\n", tilingData.l1BufferNum);104 printf(" l1BufferNum : %u\n", tilingData.l1BufferNum);
102 printf(" l0cDB : %u\n", tilingData.l0cDB);105 printf(" l0cDB : %u\n", tilingData.l0cDB);
103 }106 }
RSamples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_a16w16_tiling_common.h→Samples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_tiling_common.h+5-3
@@ -21,7 +21,7 @@
21 21 
22#include "tiling/platform/platform_ascendc.h"22#include "tiling/platform/platform_ascendc.h"
23 23 
24-struct MatmulA16W16PlatformInfo {24+struct MatmulPlatformInfo {
25 uint32_t aicNum{0};25 uint32_t aicNum{0};
26 uint32_t aivNum{0};26 uint32_t aivNum{0};
27 uint64_t ubSize{0};27 uint64_t ubSize{0};
@@ -34,13 +34,15 @@ struct MatmulA16W16PlatformInfo {
34 platform_ascendc::SocVersion socVersion{0};34 platform_ascendc::SocVersion socVersion{0};
35};35};
36 36 
37-struct MatmulA16W16Args {37+struct MatmulArgs {
38 uint64_t m{0};38 uint64_t m{0};
39 uint64_t n{0};39 uint64_t n{0};
40 uint64_t k{0};40 uint64_t k{0};
41+ uint64_t dataTypeSize{0};
41 bool hasBias{false};42 bool hasBias{false};
42 bool isATrans{false};43 bool isATrans{false};
43 bool isBTrans{false};44 bool isBTrans{false};
45+ bool isHf32{false};
44};46};
45 47 
46struct MatMulV3TailInfo {48struct MatMulV3TailInfo {
@@ -51,7 +53,7 @@ struct MatMulV3TailInfo {
51 uint64_t nTailMain = 0UL;53 uint64_t nTailMain = 0UL;
52};54};
53 55 
54-struct MatmulA16W16RunInfo {56+struct MatmulRunInfo {
55 uint64_t baseM{1};57 uint64_t baseM{1};
56 uint64_t baseN{1};58 uint64_t baseN{1};
57 uint64_t baseK{1};59 uint64_t baseK{1};
RSamples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_a16w16_tiling_data.h→Samples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_tiling_data.h+3-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_tiling_data.h12+ * \file matmul_tiling_data.h
13 * \brief Serialized tiling data passed from host launcher to kernel.13 * \brief Serialized tiling data passed from host launcher to kernel.
14 */14 */
15 15 
@@ -22,7 +22,7 @@
22#include "kernel_tiling/kernel_tiling.h"22#include "kernel_tiling/kernel_tiling.h"
23 23 
24#pragma pack(push, 8)24#pragma pack(push, 8)
25-struct alignas(8) MatmulA16W16TilingData {25+struct alignas(8) MatmulTilingData {
26 uint32_t usedCoreNum{0};26 uint32_t usedCoreNum{0};
27 uint32_t m{0};27 uint32_t m{0};
28 uint32_t n{0};28 uint32_t n{0};
@@ -40,6 +40,7 @@ struct alignas(8) MatmulA16W16TilingData {
40 uint32_t nBaseTailSplitCnt{1};40 uint32_t nBaseTailSplitCnt{1};
41 uint32_t mTailMain{0};41 uint32_t mTailMain{0};
42 uint32_t nTailMain{0};42 uint32_t nTailMain{0};
43+ bool isHf32{0};
43 uint8_t l1BufferNum{0};44 uint8_t l1BufferNum{0};
44 uint8_t l0cDB{1};45 uint8_t l0cDB{1};
45};46};
RSamples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_a16w16_tiling_streamk.h→Samples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_tiling_streamk.h+24-19
@@ -9,8 +9,8 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_tiling_streamk.h12+ * \file matmul_tiling_streamk.h
13- * \brief StreamK tiling specialization for A16W16.13+ * \brief StreamK tiling specialization for .
14 */14 */
15 15 
16#pragma once16#pragma once
@@ -18,13 +18,13 @@
18#include <algorithm>18#include <algorithm>
19#include <cmath>19#include <cmath>
20#include <cstdlib>20#include <cstdlib>
21-#include "tiling/matmul_a16w16_tiling_base.h"21+#include "tiling/matmul_tiling_base.h"
22#include "host_utils/common_utils.h"22#include "host_utils/common_utils.h"
23 23 
24-class MatmulA16W16TilingStreamK : public MatmulA16W16TilingBase {24+class MatmulTilingStreamK : public MatmulTilingBase {
25public:25public:
26- MatmulA16W16TilingStreamK() = default;26+ MatmulTilingStreamK() = default;
27- ~MatmulA16W16TilingStreamK() override = default;27+ ~MatmulTilingStreamK() override = default;
28 28 
29protected:29protected:
30 const char* TilingName() const override30 const char* TilingName() const override
@@ -32,7 +32,7 @@ protected:
32 return "streamk";32 return "streamk";
33 }33 }
34 34 
35- void DoOpTiling(MatmulA16W16TilingData& tilingData) override35+ void DoOpTiling(MatmulTilingData& tilingData) override
36 {36 {
37 IsCapable();37 IsCapable();
38 ResetBase();38 ResetBase();
@@ -65,12 +65,16 @@ private:
65 {65 {
66 constexpr uint64_t STREAM_K_MIN_K_THRESHOLD = 8192UL;66 constexpr uint64_t STREAM_K_MIN_K_THRESHOLD = 8192UL;
67 uint64_t kThreshold =67 uint64_t kThreshold =
68- std::max(STREAM_K_MIN_K_THRESHOLD, platformInfo_.aicNum * BASIC_BLOCK_SIZE_256) / DATA_SIZE_FP16;68+ std::max(STREAM_K_MIN_K_THRESHOLD, platformInfo_.aicNum * BASIC_BLOCK_SIZE_256) / args_.dataTypeSize;
69 if (Align(args_.k, BASIC_BLOCK_SIZE_256) < kThreshold) {69 if (Align(args_.k, BASIC_BLOCK_SIZE_256) < kThreshold) {
70 return false;70 return false;
71 }71 }
72- uint64_t mCnt = CeilDiv(args_.m, BASIC_BLOCK_SIZE_256);72+ uint64_t alignValue = BASIC_BLOCK_SIZE_256;
73- uint64_t nCnt = CeilDiv(args_.n, BASIC_BLOCK_SIZE_256);73+ if (args_.dataTypeSize == DATA_SIZE_FP32 && !args_.isHf32) {
74+ alignValue = BLOCK_BYTE_SIZE;
75+ }
76+ uint64_t mCnt = CeilDiv(args_.m, alignValue);
77+ uint64_t nCnt = CeilDiv(args_.n, alignValue);
74 return (mCnt * nCnt <= platformInfo_.aicNum / NUM_TWO);78 return (mCnt * nCnt <= platformInfo_.aicNum / NUM_TWO);
75 }79 }
76 80 
@@ -81,7 +85,7 @@ private:
81 return false;85 return false;
82 }86 }
83 uint64_t kThreshold =87 uint64_t kThreshold =
84- std::max(STREAM_K_MIN_K_THRESHOLD, platformInfo_.aicNum * BASIC_BLOCK_SIZE_128) / DATA_SIZE_FP16;88+ std::max(STREAM_K_MIN_K_THRESHOLD, platformInfo_.aicNum * BASIC_BLOCK_SIZE_128) / args_.dataTypeSize;
85 if (args_.k < kThreshold) {89 if (args_.k < kThreshold) {
86 return false;90 return false;
87 }91 }
@@ -97,7 +101,7 @@ private:
97 runInfo_.usedCoreNum = platformInfo_.aicNum;101 runInfo_.usedCoreNum = platformInfo_.aicNum;
98 runInfo_.baseM = BASIC_BLOCK_SIZE_256;102 runInfo_.baseM = BASIC_BLOCK_SIZE_256;
99 runInfo_.baseN = BASIC_BLOCK_SIZE_256;103 runInfo_.baseN = BASIC_BLOCK_SIZE_256;
100- runInfo_.baseK = BASIC_BLOCK_SIZE_128 / DATA_SIZE_FP16;104+ runInfo_.baseK = BASIC_BLOCK_SIZE_128 / args_.dataTypeSize;
101 runInfo_.stepM = 1;105 runInfo_.stepM = 1;
102 runInfo_.stepN = 1;106 runInfo_.stepN = 1;
103 runInfo_.iterateOrder = 0;107 runInfo_.iterateOrder = 0;
@@ -140,9 +144,9 @@ private:
140 void CalBaseK()144 void CalBaseK()
141 {145 {
142 uint64_t baseKAlignValue =146 uint64_t baseKAlignValue =
143- (!args_.isATrans || args_.isBTrans) ? BASIC_BLOCK_SIZE_128 / DATA_SIZE_FP16 : BASIC_BLOCK_SIZE_16;147+ (!args_.isATrans || args_.isBTrans) ? BASIC_BLOCK_SIZE_128 / args_.dataTypeSize : BASIC_BLOCK_SIZE_16;
144 uint64_t kValueMax = FloorAlign(148 uint64_t kValueMax = FloorAlign(
145- platformInfo_.l0aSize / DB_SIZE / DATA_SIZE_FP16 / std::max(runInfo_.baseM, runInfo_.baseN),149+ platformInfo_.l0aSize / DB_SIZE / args_.dataTypeSize / std::max(runInfo_.baseM, runInfo_.baseN),
146 baseKAlignValue);150 baseKAlignValue);
147 runInfo_.baseK = std::min(runInfo_.singleCoreK, kValueMax);151 runInfo_.baseK = std::min(runInfo_.singleCoreK, kValueMax);
148 }152 }
@@ -151,11 +155,11 @@ private:
151 {155 {
152 uint64_t totalL1Size = platformInfo_.l1Size;156 uint64_t totalL1Size = platformInfo_.l1Size;
153 uint64_t reserveBTSize = args_.hasBias ? BASIC_BLOCK_SIZE_256 * DATA_SIZE_FP32 : 0UL;157 uint64_t reserveBTSize = args_.hasBias ? BASIC_BLOCK_SIZE_256 * DATA_SIZE_FP32 : 0UL;
154- runInfo_.depthA1 = totalL1Size / NUM_TWO / runInfo_.baseM / runInfo_.baseK / DATA_SIZE_FP16; // 2: half of l1158+ runInfo_.depthA1 = totalL1Size / NUM_TWO / runInfo_.baseM / runInfo_.baseK / args_.dataTypeSize; // 2: half of l1
155- runInfo_.depthB1 = totalL1Size / NUM_TWO / runInfo_.baseN / runInfo_.baseK / DATA_SIZE_FP16; // 2: half of l1159+ runInfo_.depthB1 = totalL1Size / NUM_TWO / runInfo_.baseN / runInfo_.baseK / args_.dataTypeSize; // 2: half of l1
156 160 
157- uint64_t depthASize = runInfo_.depthA1 * runInfo_.baseM * runInfo_.baseK * DATA_SIZE_FP16;161+ uint64_t depthASize = runInfo_.depthA1 * runInfo_.baseM * runInfo_.baseK * args_.dataTypeSize;
158- uint64_t depthBSize = runInfo_.depthB1 * runInfo_.baseN * runInfo_.baseK * DATA_SIZE_FP16;162+ uint64_t depthBSize = runInfo_.depthB1 * runInfo_.baseN * runInfo_.baseK * args_.dataTypeSize;
159 if (depthASize + depthBSize > totalL1Size - reserveBTSize) {163 if (depthASize + depthBSize > totalL1Size - reserveBTSize) {
160 if (runInfo_.baseM <= runInfo_.baseN) {164 if (runInfo_.baseM <= runInfo_.baseN) {
161 runInfo_.depthA1 = std::max(runInfo_.depthA1 / NUM_TWO, 1UL); // 2: adjust deptch for l1 buffer165 runInfo_.depthA1 = std::max(runInfo_.depthA1 / NUM_TWO, 1UL); // 2: adjust deptch for l1 buffer
@@ -205,7 +209,7 @@ private:
205 }209 }
206 }210 }
207 211 
208- void BuildTilingData(MatmulA16W16TilingData& tilingData) const212+ void BuildTilingData(MatmulTilingData& tilingData) const
209 {213 {
210 tilingData = {};214 tilingData = {};
211 tilingData.m = static_cast<uint32_t>(args_.m);215 tilingData.m = static_cast<uint32_t>(args_.m);
@@ -228,6 +232,7 @@ private:
228 tilingData.mTailMain = runInfo_.tailInfo.mTailMain;232 tilingData.mTailMain = runInfo_.tailInfo.mTailMain;
229 tilingData.nTailMain = runInfo_.tailInfo.nTailMain;233 tilingData.nTailMain = runInfo_.tailInfo.nTailMain;
230 tilingData.usedCoreNum = static_cast<uint32_t>(runInfo_.usedCoreNum);234 tilingData.usedCoreNum = static_cast<uint32_t>(runInfo_.usedCoreNum);
235+ tilingData.isHf32 = args_.isHf32;
231 tilingData.l1BufferNum = static_cast<uint8_t>(runInfo_.l1BufferNum);236 tilingData.l1BufferNum = static_cast<uint8_t>(runInfo_.l1BufferNum);
232 tilingData.l0cDB = static_cast<uint8_t>(runInfo_.dbL0c);237 tilingData.l0cDB = static_cast<uint8_t>(runInfo_.dbL0c);
233 }238 }
RSamples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_a16w16_tiling_swat.h→Samples/2_Performance/matmul_story/matmul_recipes/include/tiling/matmul_tiling_swat.h+55-37
@@ -9,8 +9,8 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_a16w16_tiling_swat.h12+ * \file matmul_tiling_swat.h
13- * \brief SWAT tiling specialization for A16W16 non-full-load path.13+ * \brief SWAT tiling specialization for non-full-load path.
14 */14 */
15 15 
16#pragma once16#pragma once
@@ -19,13 +19,13 @@
19#include <cmath>19#include <cmath>
20#include <tuple>20#include <tuple>
21#include <vector>21#include <vector>
22-#include "tiling/matmul_a16w16_tiling_base.h"22+#include "tiling/matmul_tiling_base.h"
23#include "host_utils/common_utils.h"23#include "host_utils/common_utils.h"
24 24 
25-class MatmulA16W16TilingSwat : public MatmulA16W16TilingBase {25+class MatmulTilingSwat : public MatmulTilingBase {
26public:26public:
27- MatmulA16W16TilingSwat() = default;27+ MatmulTilingSwat() = default;
28- ~MatmulA16W16TilingSwat() override = default;28+ ~MatmulTilingSwat() override = default;
29 29 
30protected:30protected:
31 const char* TilingName() const override31 const char* TilingName() const override
@@ -33,22 +33,19 @@ protected:
33 return "swat";33 return "swat";
34 }34 }
35 35 
36- void DoOpTiling(MatmulA16W16TilingData& tilingData) override36+ void DoOpTiling(MatmulTilingData& tilingData) override
37 {37 {
38- ResetBase();
39- FormulateBasicBlock();
40- OptimizeEdgeBasicBlock();
41- CalcTailBasicBlock();
42- CalL1Tiling();
43 ResetBase();38 ResetBase();
44 FormulateLoadBalanceBlock();39 FormulateLoadBalanceBlock();
45 if (runInfo_.baseM == BASIC_BLOCK_SIZE_256 && runInfo_.baseN == BASIC_BLOCK_SIZE_256) {40 if (runInfo_.baseM == BASIC_BLOCK_SIZE_256 && runInfo_.baseN == BASIC_BLOCK_SIZE_256) {
46 OptimizeEdgeBasicBlock();41 OptimizeEdgeBasicBlock();
47 }42 }
43+ CalcTailBasicBlock();
44+ CalL1Tiling();
48 uint64_t remainSizeForAL1BL1 =45 uint64_t remainSizeForAL1BL1 =
49 args_.hasBias ? (platformInfo_.l1Size - BIAS_TABLE_NUM * DATA_SIZE_FP32) : platformInfo_.l1Size;46 args_.hasBias ? (platformInfo_.l1Size - BIAS_TABLE_NUM * DATA_SIZE_FP32) : platformInfo_.l1Size;
50 runInfo_.stepKa =47 runInfo_.stepKa =
51- remainSizeForAL1BL1 / NUM_TWO / ((runInfo_.baseM + runInfo_.baseN) * runInfo_.baseK) / DATA_SIZE_FP16;48+ remainSizeForAL1BL1 / NUM_TWO / ((runInfo_.baseM + runInfo_.baseN) * runInfo_.baseK) / args_.dataTypeSize;
52 runInfo_.stepKb = runInfo_.stepKa; // has bias, adjust stepK to suitable value49 runInfo_.stepKb = runInfo_.stepKa; // has bias, adjust stepK to suitable value
53 runInfo_.depthA1 = runInfo_.stepKa * DB_SIZE;50 runInfo_.depthA1 = runInfo_.stepKa * DB_SIZE;
54 runInfo_.depthB1 = runInfo_.stepKb * DB_SIZE;51 runInfo_.depthB1 = runInfo_.stepKb * DB_SIZE;
@@ -62,7 +59,7 @@ private:
62 runInfo_.usedCoreNum = platformInfo_.aicNum;59 runInfo_.usedCoreNum = platformInfo_.aicNum;
63 runInfo_.baseM = BASIC_BLOCK_SIZE_256;60 runInfo_.baseM = BASIC_BLOCK_SIZE_256;
64 runInfo_.baseN = BASIC_BLOCK_SIZE_256; // 256 is better base61 runInfo_.baseN = BASIC_BLOCK_SIZE_256; // 256 is better base
65- runInfo_.baseK = BASIC_BLOCK_SIZE_128 / DATA_SIZE_FP16;62+ runInfo_.baseK = BASIC_BLOCK_SIZE_128 / args_.dataTypeSize;
66 runInfo_.stepM = 1;63 runInfo_.stepM = 1;
67 runInfo_.stepN = 1;64 runInfo_.stepN = 1;
68 runInfo_.iterateOrder = 0;65 runInfo_.iterateOrder = 0;
@@ -91,7 +88,7 @@ private:
91 runInfo_.usedCoreNum = mCore * nCore;88 runInfo_.usedCoreNum = mCore * nCore;
92 uint64_t kValueAlign = Align(args_.k, BASIC_BLOCK_SIZE_16);89 uint64_t kValueAlign = Align(args_.k, BASIC_BLOCK_SIZE_16);
93 uint64_t kValueMax = FloorAlign(90 uint64_t kValueMax = FloorAlign(
94- platformInfo_.l0aSize / DB_SIZE / DATA_SIZE_FP16 / std::max(runInfo_.baseM, runInfo_.baseN),91+ platformInfo_.l0aSize / DB_SIZE / args_.dataTypeSize / std::max(runInfo_.baseM, runInfo_.baseN),
95 BASIC_BLOCK_SIZE_16);92 BASIC_BLOCK_SIZE_16);
96 runInfo_.baseK = std::min(kValueAlign, kValueMax);93 runInfo_.baseK = std::min(kValueAlign, kValueMax);
97 }94 }
@@ -140,10 +137,13 @@ private:
140 mCore = CeilDiv(args_.m, runInfo_.baseM);137 mCore = CeilDiv(args_.m, runInfo_.baseM);
141 nCore = CeilDiv(args_.n, runInfo_.baseN);138 nCore = CeilDiv(args_.n, runInfo_.baseN);
142 runInfo_.usedCoreNum = std::min(mCore * nCore, static_cast<uint64_t>(platformInfo_.aicNum));139 runInfo_.usedCoreNum = std::min(mCore * nCore, static_cast<uint64_t>(platformInfo_.aicNum));
143- uint64_t kValueAlign = Align(args_.k, BASIC_BLOCK_SIZE_16);140+ uint64_t baseKAlignValue = !args_.isATrans && args_.isBTrans && args_.dataTypeSize == DATA_SIZE_FP32 ?
141+ BLOCK_BYTE_SIZE / args_.dataTypeSize :
142+ BASIC_BLOCK_SIZE_16;
143+ uint64_t kValueAlign = Align(args_.k, baseKAlignValue);
144 uint64_t kValueMax = FloorAlign(144 uint64_t kValueMax = FloorAlign(
145- platformInfo_.l0aSize / DB_SIZE / DATA_SIZE_FP16 / std::max(runInfo_.baseM, runInfo_.baseN),145+ platformInfo_.l0aSize / DB_SIZE / args_.dataTypeSize / std::max(runInfo_.baseM, runInfo_.baseN),
146- BASIC_BLOCK_SIZE_16);146+ baseKAlignValue);
147 runInfo_.baseK = std::min(kValueAlign, kValueMax);147 runInfo_.baseK = std::min(kValueAlign, kValueMax);
148 }148 }
149 149 
@@ -202,6 +202,17 @@ private:
202 }202 }
203 }203 }
204 204 
205+ uint64_t GetAswWindowLen() const
206+ {
207+ uint64_t sqrtNum = static_cast<uint64_t>(sqrt(platformInfo_.aicNum));
208+ for (uint64_t factor = sqrtNum; factor >= 1UL; --factor) {
209+ if (platformInfo_.aicNum % factor == 0UL) {
210+ return factor;
211+ }
212+ }
213+ return 1UL;
214+ }
215+ 
205 void GetOuterAxisTailCnt(bool nLoadBalance, uint32_t& baseTailSplitCnt, uint64_t& tailMain)216 void GetOuterAxisTailCnt(bool nLoadBalance, uint32_t& baseTailSplitCnt, uint64_t& tailMain)
206 {217 {
207 uint64_t aicNum = platformInfo_.aicNum;218 uint64_t aicNum = platformInfo_.aicNum;
@@ -220,9 +231,10 @@ private:
220 uint64_t yCnt = CeilDiv(y, baseY);231 uint64_t yCnt = CeilDiv(y, baseY);
221 uint64_t xTail = x % baseX;232 uint64_t xTail = x % baseX;
222 233 
234+ uint64_t aswWindowLen = GetAswWindowLen();
223 uint64_t totalWindows = CeilDiv(xCnt * yCnt, aicNum);235 uint64_t totalWindows = CeilDiv(xCnt * yCnt, aicNum);
224 uint64_t mainWindows = CeilDiv((xCnt - 1UL) * yCnt + yCnt % aicNum, aicNum);236 uint64_t mainWindows = CeilDiv((xCnt - 1UL) * yCnt + yCnt % aicNum, aicNum);
225- if (yCnt % aicNum == 0UL && (xCnt % WINDOW_LEN == 0UL || WINDOW_LEN % xCnt == 0UL)) {237+ if (yCnt % aicNum == 0UL && (xCnt % aswWindowLen == 0UL || aswWindowLen % xCnt == 0UL)) {
226 mainWindows = totalWindows;238 mainWindows = totalWindows;
227 }239 }
228 uint64_t tailWindows = totalWindows - mainWindows;240 uint64_t tailWindows = totalWindows - mainWindows;
@@ -274,11 +286,11 @@ private:
274 {286 {
275 uint64_t totalL1Size = platformInfo_.l1Size;287 uint64_t totalL1Size = platformInfo_.l1Size;
276 uint64_t reserveBTSize = args_.hasBias ? BASIC_BLOCK_SIZE_256 * DATA_SIZE_FP32 : 0UL;288 uint64_t reserveBTSize = args_.hasBias ? BASIC_BLOCK_SIZE_256 * DATA_SIZE_FP32 : 0UL;
277- runInfo_.depthA1 = totalL1Size / NUM_TWO / runInfo_.baseM / runInfo_.baseK / DATA_SIZE_FP16; // 2: half of l1289+ runInfo_.depthA1 = totalL1Size / NUM_TWO / runInfo_.baseM / runInfo_.baseK / args_.dataTypeSize; // 2: half of l1
278- runInfo_.depthB1 = totalL1Size / NUM_TWO / runInfo_.baseN / runInfo_.baseK / DATA_SIZE_FP16; // 2: half of l1290+ runInfo_.depthB1 = totalL1Size / NUM_TWO / runInfo_.baseN / runInfo_.baseK / args_.dataTypeSize; // 2: half of l1
279 291 
280- uint64_t depthASize = runInfo_.depthA1 * runInfo_.baseM * runInfo_.baseK * DATA_SIZE_FP16;292+ uint64_t depthASize = runInfo_.depthA1 * runInfo_.baseM * runInfo_.baseK * args_.dataTypeSize;
281- uint64_t depthBSize = runInfo_.depthB1 * runInfo_.baseN * runInfo_.baseK * DATA_SIZE_FP16;293+ uint64_t depthBSize = runInfo_.depthB1 * runInfo_.baseN * runInfo_.baseK * args_.dataTypeSize;
282 if (depthASize + depthBSize > totalL1Size - reserveBTSize) {294 if (depthASize + depthBSize > totalL1Size - reserveBTSize) {
283 if (runInfo_.baseM <= runInfo_.baseN) {295 if (runInfo_.baseM <= runInfo_.baseN) {
284 runInfo_.depthA1 = std::max(runInfo_.depthA1 / NUM_TWO, 1UL); // 2: adjust deptch for l1 buffer296 runInfo_.depthA1 = std::max(runInfo_.depthA1 / NUM_TWO, 1UL); // 2: adjust deptch for l1 buffer
@@ -398,18 +410,18 @@ private:
398 {410 {
399 if (isMLarger) {411 if (isMLarger) {
400 runInfo_.baseN = minMN;412 runInfo_.baseN = minMN;
401- runInfo_.baseM = platformInfo_.l0cSize / runInfo_.dbL0c / runInfo_.baseN;413+ runInfo_.baseM = platformInfo_.l0cSize / runInfo_.dbL0c / runInfo_.baseN / DATA_SIZE_FP32;
402 runInfo_.baseM = Floor(runInfo_.baseM, BASIC_BLOCK_SIZE_16);414 runInfo_.baseM = Floor(runInfo_.baseM, BASIC_BLOCK_SIZE_16);
403- CalcLargeSingleSide(maxMN, runInfo_.baseM, isMLarger);415+ CalcLargeSingleSide(minMN, maxMN, runInfo_.baseM, isMLarger);
404 } else {416 } else {
405 runInfo_.baseM = minMN;417 runInfo_.baseM = minMN;
406- runInfo_.baseN = platformInfo_.l0cSize / runInfo_.dbL0c / runInfo_.baseM;418+ runInfo_.baseN = platformInfo_.l0cSize / runInfo_.dbL0c / runInfo_.baseM / DATA_SIZE_FP32;
407 runInfo_.baseN = Floor(runInfo_.baseN, BLOCK_BYTE_SIZE);419 runInfo_.baseN = Floor(runInfo_.baseN, BLOCK_BYTE_SIZE);
408- CalcLargeSingleSide(maxMN, runInfo_.baseN, isMLarger);420+ CalcLargeSingleSide(minMN, maxMN, runInfo_.baseN, isMLarger);
409 }421 }
410 }422 }
411 423 
412- void CalcLargeSingleSide(uint64_t maxMN, uint64_t& targetBase, bool isMLarger)424+ void CalcLargeSingleSide(uint64_t minMN, uint64_t maxMN, uint64_t& targetBase, bool isMLarger)
413 {425 {
414 uint64_t minCoreNum = (platformInfo_.aicNum + 1UL) * NUM_NINE / NUM_TEN;426 uint64_t minCoreNum = (platformInfo_.aicNum + 1UL) * NUM_NINE / NUM_TEN;
415 for (uint64_t tmpCoreNum = platformInfo_.aicNum; tmpCoreNum >= minCoreNum; tmpCoreNum--) {427 for (uint64_t tmpCoreNum = platformInfo_.aicNum; tmpCoreNum >= minCoreNum; tmpCoreNum--) {
@@ -417,7 +429,8 @@ private:
417 while (loop <= MAX_LOOP_NUM) {429 while (loop <= MAX_LOOP_NUM) {
418 uint64_t baseBlock = CeilDiv(maxMN, tmpCoreNum * loop);430 uint64_t baseBlock = CeilDiv(maxMN, tmpCoreNum * loop);
419 baseBlock = UpdateBaseBlock(baseBlock, isMLarger);431 baseBlock = UpdateBaseBlock(baseBlock, isMLarger);
420- if (baseBlock >= MIN_BASE_BLOCK && baseBlock <= MAX_BASE_BLOCK) {432+ uint64_t tileSize = baseBlock * minMN * DATA_SIZE_FP32;
433+ if (baseBlock >= MIN_BASE_BLOCK && baseBlock <= MAX_BASE_BLOCK && tileSize <= platformInfo_.l0cSize) {
421 targetBase = baseBlock;434 targetBase = baseBlock;
422 return;435 return;
423 }436 }
@@ -430,7 +443,11 @@ private:
430 uint64_t UpdateBaseBlock(uint64_t baseBlock, bool isMLarger)443 uint64_t UpdateBaseBlock(uint64_t baseBlock, bool isMLarger)
431 {444 {
432 if (!isMLarger) {445 if (!isMLarger) {
433- return Align(baseBlock, BASIC_BLOCK_SIZE_128 / DATA_SIZE_FP16);446+ if (!args_.isBTrans || args_.k <= BASIC_BLOCK_SIZE_256) {
447+ return Align(baseBlock, BASIC_BLOCK_SIZE_128 / args_.dataTypeSize);
448+ } else {
449+ return Align(baseBlock, BASIC_BLOCK_SIZE_64 / args_.dataTypeSize);
450+ }
434 } else {451 } else {
435 return Align(baseBlock, BASIC_BLOCK_SIZE_16);452 return Align(baseBlock, BASIC_BLOCK_SIZE_16);
436 }453 }
@@ -445,7 +462,7 @@ private:
445 runInfo_.defaultBalance,462 runInfo_.defaultBalance,
446 BASIC_BLOCK_SIZE_256,463 BASIC_BLOCK_SIZE_256,
447 BASIC_BLOCK_SIZE_256,464 BASIC_BLOCK_SIZE_256,
448- BASIC_BLOCK_SIZE_128 / DATA_SIZE_FP16};465+ BASIC_BLOCK_SIZE_128 / args_.dataTypeSize};
449 if (CalcBestBalance(params1, isMLarger)) {466 if (CalcBestBalance(params1, isMLarger)) {
450 return;467 return;
451 }468 }
@@ -456,7 +473,7 @@ private:
456 runInfo_.defaultBalance,473 runInfo_.defaultBalance,
457 BASIC_BLOCK_SIZE_256,474 BASIC_BLOCK_SIZE_256,
458 BASIC_BLOCK_SIZE_256,475 BASIC_BLOCK_SIZE_256,
459- BASIC_BLOCK_SIZE_128 / DATA_SIZE_FP16};476+ BASIC_BLOCK_SIZE_128 / args_.dataTypeSize};
460 if (CalcBestBalance(params2, isMLarger)) {477 if (CalcBestBalance(params2, isMLarger)) {
461 return;478 return;
462 }479 }
@@ -508,10 +525,10 @@ private:
508 // If n is the inner axis or k<=256, align n to 128B, otherwise 64B; if m is the inner axis, align to525 // If n is the inner axis or k<=256, align n to 128B, otherwise 64B; if m is the inner axis, align to
509 // 128B, otherwise 64B526 // 128B, otherwise 64B
510 bool nNotAligned = !args_.isBTrans || args_.k <= BASIC_BLOCK_SIZE_256 ?527 bool nNotAligned = !args_.isBTrans || args_.k <= BASIC_BLOCK_SIZE_256 ?
511- currentBaseN * DATA_SIZE_FP16 % BASIC_BLOCK_SIZE_128 != 0UL :528+ currentBaseN * args_.dataTypeSize % BASIC_BLOCK_SIZE_128 != 0UL :
512- currentBaseN * DATA_SIZE_FP16 % BASIC_BLOCK_SIZE_64 != 0UL;529+ currentBaseN * args_.dataTypeSize % BASIC_BLOCK_SIZE_64 != 0UL;
513- bool mNotAligned = args_.isATrans ? currentBaseM * DATA_SIZE_FP16 % BASIC_BLOCK_SIZE_128 != 0UL :530+ bool mNotAligned = args_.isATrans ? currentBaseM * args_.dataTypeSize % BASIC_BLOCK_SIZE_128 != 0UL :
514- currentBaseM * DATA_SIZE_FP16 % BASIC_BLOCK_SIZE_64 != 0UL;531+ currentBaseM * args_.dataTypeSize % BASIC_BLOCK_SIZE_64 != 0UL;
515 if (mNotAligned || nNotAligned) {532 if (mNotAligned || nNotAligned) {
516 continue;533 continue;
517 }534 }
@@ -612,7 +629,7 @@ private:
612 {128, 272, 48, 1.242, 1.126}, {128, 256, 64, 1.267, 1}, {128, 240, 64, 1.295, 1},629 {128, 272, 48, 1.242, 1.126}, {128, 256, 64, 1.267, 1}, {128, 240, 64, 1.295, 1},
613 {112, 336, 48, 1.287, 1.126}};630 {112, 336, 48, 1.287, 1.126}};
614 631 
615- void BuildTilingData(MatmulA16W16TilingData& tilingData) const632+ void BuildTilingData(MatmulTilingData& tilingData) const
616 {633 {
617 tilingData = {};634 tilingData = {};
618 tilingData.m = static_cast<uint32_t>(args_.m);635 tilingData.m = static_cast<uint32_t>(args_.m);
@@ -634,6 +651,7 @@ private:
634 tilingData.mTailMain = runInfo_.tailInfo.mTailMain;651 tilingData.mTailMain = runInfo_.tailInfo.mTailMain;
635 tilingData.nTailMain = runInfo_.tailInfo.nTailMain;652 tilingData.nTailMain = runInfo_.tailInfo.nTailMain;
636 tilingData.usedCoreNum = static_cast<uint32_t>(runInfo_.usedCoreNum);653 tilingData.usedCoreNum = static_cast<uint32_t>(runInfo_.usedCoreNum);
654+ tilingData.isHf32 = args_.isHf32;
637 tilingData.l1BufferNum = static_cast<uint8_t>(runInfo_.l1BufferNum);655 tilingData.l1BufferNum = static_cast<uint8_t>(runInfo_.l1BufferNum);
638 tilingData.l0cDB = static_cast<uint8_t>(runInfo_.dbL0c);656 tilingData.l0cDB = static_cast<uint8_t>(runInfo_.dbL0c);
639 }657 }