diff --git a/operator/ascendc/2_features/14_matmul_api_constant/MatmulApiContantCustom.json b/operator/ascendc/2_features/14_matmul_api_constant/MatmulApiConstantCustom.json similarity index 100% rename from operator/ascendc/2_features/14_matmul_api_constant/MatmulApiContantCustom.json rename to operator/ascendc/2_features/14_matmul_api_constant/MatmulApiConstantCustom.json diff --git a/operator/ascendc/2_features/14_matmul_api_constant/MatmulApiConstantCustom/op_kernel/matmul_api_constant_custom.cpp b/operator/ascendc/2_features/14_matmul_api_constant/MatmulApiConstantCustom/op_kernel/matmul_api_constant_custom.cpp index c95abf2991036ca42d270092f497e0082aedff1f..4d1a8545579a3e42123333e3ac9f82a727799f38 100644 --- a/operator/ascendc/2_features/14_matmul_api_constant/MatmulApiConstantCustom/op_kernel/matmul_api_constant_custom.cpp +++ b/operator/ascendc/2_features/14_matmul_api_constant/MatmulApiConstantCustom/op_kernel/matmul_api_constant_custom.cpp @@ -17,13 +17,6 @@ __aicore__ inline uint32_t Ceiling(uint32_t a, uint32_t b) return (a + b - 1) / b; } -// The specified value remains consistent with the runtime tiling paramters. -// singleCoreM, singleCoreN, singleCoreK, baseM, baseN, baseK. -constexpr static MatmulShapeParams shapeParams = {512, 640, 256, 128, 128, 128}; -constexpr static MatmulConfig mmConfig = GetMMConfig(shapeParams); -// Get the fully constant template parameters. -constexpr static MatmulApiStaticTiling staticConfig = GetMatmulApiTiling(mmConfig); - template class MatmulApiConstantKernel { public: __aicore__ inline MatmulApiConstantKernel(){}; @@ -39,6 +32,13 @@ public: typedef MatmulType cMatmulType; typedef MatmulType biasMatmulType; + // The specified value remains consistent with the runtime tiling paramters. + // singleCoreM, singleCoreN, singleCoreK, baseM, baseN, baseK. + constexpr static MatmulShapeParams shapeParams = {512, 640, 256, 128, 128, 128}; + constexpr static MatmulConfig mmConfig = GetMMConfig(shapeParams); + // Get the fully constant template parameters. + constexpr static MatmulApiStaticTiling staticConfig = GetMatmulApiTiling(mmConfig); + Matmul matmulObj; AscendC::GlobalTensor aGlobal; @@ -111,7 +111,7 @@ extern "C" __global__ __aicore__ void matmul_api_constant_custom( AscendC::TPipe pipe; // With the fully constant template parameters, nullptr can be passed into REGIST_MATMUL_OBJ to replace tiling. - REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulApiConstantKernel.matmulObj, nullptr); + REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulApiConstantKernel.matmulObj, &tilingData.cubeTilingData); matmulApiConstantKernel.Init(a, b, bias, c, workspace, tilingData.cubeTilingData); matmulApiConstantKernel.Process(&pipe); } \ No newline at end of file