From bc9893166f224d19f87688a6019b81e4f28fa0a8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=9C=B1=E5=AE=87=E6=99=A8?= Date: Sun, 13 Jul 2025 21:35:47 +0800 Subject: [PATCH] change to new tiling definitions --- .../op_host/matmul_custom.cpp | 16 +++++++--------- .../op_kernel/matmul_custom.cpp | 4 +++- .../matmul_custom_tiling.h | 18 +++++++----------- .../10_matmul_frameworklaunch/README.md | 3 +++ 4 files changed, 20 insertions(+), 21 deletions(-) rename operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/{op_host => op_kernel}/matmul_custom_tiling.h (50%) diff --git a/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_host/matmul_custom.cpp b/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_host/matmul_custom.cpp index 49bc45d64..f1911480c 100644 --- a/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_host/matmul_custom.cpp +++ b/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_host/matmul_custom.cpp @@ -7,7 +7,7 @@ * but WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. */ -#include "matmul_custom_tiling.h" +#include "../op_kernel/matmul_custom_tiling.h" #include "register/op_def_registry.h" #include "tiling/platform/platform_ascendc.h" #include "tiling/tiling_api.h" @@ -45,27 +45,25 @@ static ge::graphStatus TilingFunc(gert::TilingContext *context) } cubeTiling.SetBias(true); cubeTiling.SetBufferSpace(-1, -1, -1); - MatmulCustomTilingData tiling; - if (cubeTiling.GetTiling(tiling.cubeTilingData) == -1) { + MatmulCustomTilingData *tiling = context->GetTilingData(); + if (cubeTiling.GetTiling(tiling->cubeTilingData) == -1) { return ge::GRAPH_FAILED; } uint64_t localMemSize; ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, localMemSize); - tiling.set_localMemSize(localMemSize); + tiling->localMemSize = localMemSize; if (ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCEND310P) { context->SetBlockDim(2); context->SetTilingKey(2); } else { - /* SetBlockDim here refers to the number of cube cores, so for separated arch(AIC:AIV=1:2), - vector cores number is set 48 by SetDim, cube core number need to be set 24 here.*/ - context->SetBlockDim(24); + /* SetBlockDim here refers to the number of cube cores, so for separated arch(AIC:AIV=1:2), + vector cores number is set 48 by SetDim, cube core number need to be set 24 here.*/ + context->SetBlockDim(24); context->SetTilingKey(1); } - tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); - context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); size_t userWorkspaceSize = 0; size_t systemWorkspaceSize = static_cast(ascendcPlatform.GetLibApiWorkSpaceSize()); size_t *currentWorkspace = context->GetWorkspaceSizes(1); diff --git a/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_kernel/matmul_custom.cpp b/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_kernel/matmul_custom.cpp index 77a323fca..d0d86f000 100644 --- a/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_kernel/matmul_custom.cpp +++ b/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_kernel/matmul_custom.cpp @@ -9,6 +9,7 @@ */ #include "kernel_operator.h" #include "lib/matmul_intf.h" +#include "matmul_custom_tiling.h" using namespace matmul; @@ -141,12 +142,13 @@ MatmulKernel::CalcOffset(int32_t blockIdx, const * @param bias: Bias gm addr. * @param c: C matrix gm addr. * @param workspace: Temporary gm space addr required by matmul calc. - * @param tiling: Tiling data addr. + * @param tiling: Tiling data addr. * @retval None */ extern "C" __global__ __aicore__ void matmul_custom(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace, GM_ADDR tiling) { + REGISTER_TILING_DEFAULT(MatmulCustomTilingData); GET_TILING_DATA(tilingData, tiling); MatmulKernel matmulKernel; AscendC::TPipe pipe; diff --git a/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_host/matmul_custom_tiling.h b/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_kernel/matmul_custom_tiling.h similarity index 50% rename from operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_host/matmul_custom_tiling.h rename to operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_kernel/matmul_custom_tiling.h index fd898cba9..8f32f3418 100644 --- a/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_host/matmul_custom_tiling.h +++ b/operator/ascendc/0_introduction/10_matmul_frameworklaunch/MatmulCustomMultiCore/op_kernel/matmul_custom_tiling.h @@ -10,16 +10,12 @@ #ifndef MATMUL_TILING_H #define MATMUL_TILING_H -#include "register/tilingdata_base.h" -#include "tiling/tiling_api.h" +#include +#include "kernel_tiling/kernel_tiling.h" -namespace optiling { -BEGIN_TILING_DATA_DEF(MatmulCustomTilingData) -TILING_DATA_FIELD_DEF(uint64_t, localMemSize); -TILING_DATA_FIELD_DEF_STRUCT(TCubeTiling, cubeTilingData); -END_TILING_DATA_DEF; +struct MatmulCustomTilingData { + uint64_t localMemSize; + TCubeTiling cubeTilingData; +}; -REGISTER_TILING_DATA_CLASS(MatmulCustom, MatmulCustomTilingData) -} // namespace optiling - -#endif \ No newline at end of file +#endif // MATMUL_TILING_H \ No newline at end of file diff --git a/operator/ascendc/0_introduction/10_matmul_frameworklaunch/README.md b/operator/ascendc/0_introduction/10_matmul_frameworklaunch/README.md index 3b58d140e..05aeaa0c3 100644 --- a/operator/ascendc/0_introduction/10_matmul_frameworklaunch/README.md +++ b/operator/ascendc/0_introduction/10_matmul_frameworklaunch/README.md @@ -46,6 +46,8 @@ C = A * B + Bias ## 算子工程介绍 本样例介绍了多核场景（[MatmulCustomMultiCore](./MatmulCustomMultiCore/)）和单核场景（[MatmulCustomSingleCore](./MatmulCustomSingleCore/)）两种MamMul算子实现。可以根据使用场景，自行选择多核算子工程或单核算子工程，并在编译算子工程时，进入选择的算子实现工程中完成编译和安装。 +其中[MatmulCustomMultiCore](./MatmulCustomMultiCore/)使用标准C++语法定义Tiling结构体，[MatmulCustomSingleCore](./MatmulCustomSingleCore/)使用宏定义方式定义Tiling结构体。相较于使用宏定义方式，标准C++语法定义Tiling结构体不仅更符合C++开发者的开发习惯，并且提供了强大的灵活性。 + 以单核算子工程为例，算子工程目录MatmulCustomSingleCore包含算子的实现文件，如下所示： ``` ├── MatmulCustomSingleCore // Matmul自定义算子工程 @@ -140,3 +142,4 @@ CANN软件包中提供了工程创建工具msOpGen，MatmulCustom算子工程可 | 2024/05/27 | 更新readme | | 2024/11/11 | 样例目录调整 | | 2024/11/18 | 算子工程改写为由msOpGen生成 | +| 2025/07/14 | MatmulCustomMultiCore使用标准C++语法定义Tiling结构体 | -- Gitee