diff --git a/examples/matrix/batch_matmul/host_tiling/batch_matmul_custom_tiling.h b/examples/matrix/batch_matmul/host_tiling/batch_matmul_custom_tiling.h index e88925d127b78cc039c3dc48d65177f6c921cb0f..7f950d02f4729997301e28d2d7b69630b09b6e9b 100644 --- a/examples/matrix/batch_matmul/host_tiling/batch_matmul_custom_tiling.h +++ b/examples/matrix/batch_matmul/host_tiling/batch_matmul_custom_tiling.h @@ -13,7 +13,6 @@ #include "register/tilingdata_base.h" #include "tiling/tiling_api.h" -using namespace matmul_tiling; namespace optiling { BEGIN_TILING_DATA_DEF(MatmulCustomTilingData) TILING_DATA_FIELD_DEF_STRUCT(TCubeTiling, cubeTilingData); @@ -22,7 +21,7 @@ END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(BatchMatmulCustom, MatmulCustomTilingData) } -bool ComputeTiling(optiling::TCubeTiling& tiling, MultiCoreMatmulTiling* cubeTiling, bool isBias) +bool ComputeTiling(optiling::TCubeTiling& tiling, matmul_tiling::MultiCoreMatmulTiling* cubeTiling, bool isBias) { int32_t M = 32; int32_t N = 256; @@ -30,10 +29,10 @@ bool ComputeTiling(optiling::TCubeTiling& tiling, MultiCoreMatmulTiling* cubeTil int32_t baseM = 32; int32_t baseN = 32; cubeTiling->SetDim(1); - cubeTiling->SetAType(TPosition::GM, CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16); - cubeTiling->SetBType(TPosition::GM, CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16, true); - cubeTiling->SetCType(TPosition::GM, CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); - cubeTiling->SetBiasType(TPosition::GM, CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); + cubeTiling->SetAType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16); + cubeTiling->SetBType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16, true); + cubeTiling->SetCType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); + cubeTiling->SetBiasType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); cubeTiling->SetShape(M, N, K); cubeTiling->SetOrgShape(M, N, K); cubeTiling->SetFixSplit(baseM, baseN, -1); diff --git a/examples/matrix/batch_matmul/kernel_impl/batch_matmul_custom_impl.h b/examples/matrix/batch_matmul/kernel_impl/batch_matmul_custom_impl.h index 0a5ad98ba01c45aa580e92564a85a70d8597ec41..065f151cbf3218348ab4d7dcc40ec59f5a3aada0 100644 --- a/examples/matrix/batch_matmul/kernel_impl/batch_matmul_custom_impl.h +++ b/examples/matrix/batch_matmul/kernel_impl/batch_matmul_custom_impl.h @@ -13,8 +13,6 @@ #include "kernel_operator.h" #include "lib/matmul_intf.h" -using namespace matmul; - constexpr int USED_CORE_NUM = 2; template @@ -24,7 +22,7 @@ class BatchMatmulKernel { __aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace, const TCubeTiling& tiling); template __aicore__ inline void Process(AscendC::TPipe* pipe, int32_t batchA, int32_t batchB); - Matmul matmulObj; + matmul::Matmul matmulObj; private: __aicore__ inline void CalcOffset(int32_t blockIdx, const TCubeTiling& tiling, int32_t& offsetA, int32_t& offsetB, int32_t& offsetC, int32_t& offsetBias); diff --git a/examples/matrix/batch_matmul/kernel_launch_method_by_direct/batch_matmul_custom.cpp b/examples/matrix/batch_matmul/kernel_launch_method_by_direct/batch_matmul_custom.cpp index b1abfcd9f25ea7af7fb3fc9c27add3b0d2dd7952..94206b3efc982492ca4661f50a8f7e299de3197e 100644 --- a/examples/matrix/batch_matmul/kernel_launch_method_by_direct/batch_matmul_custom.cpp +++ b/examples/matrix/batch_matmul/kernel_launch_method_by_direct/batch_matmul_custom.cpp @@ -12,8 +12,6 @@ #include "lib/matmul_intf.h" #include "../kernel_impl/batch_matmul_custom_impl.h" -using namespace matmul; - constexpr int32_t FULL_L1_SIZE = 512 * 1024; constexpr int32_t FULL_L0C_SIZE = 128 * 1024; @@ -33,10 +31,10 @@ extern "C" __global__ __aicore__ void batch_matmul_custom(GM_ADDR a, GM_ADDR b, { TCubeTiling tiling; CopyTiling(&tiling, tilingGm); - typedef MatmulType A_TYPE; - typedef MatmulType B_TYPE; - typedef MatmulType C_TYPE; - typedef MatmulType BIAS_TYPE; + typedef matmul::MatmulType A_TYPE; + typedef matmul::MatmulType B_TYPE; + typedef matmul::MatmulType C_TYPE; + typedef matmul::MatmulType BIAS_TYPE; BatchMatmulKernel batchMatmulKernel; AscendC::TPipe pipe; tiling.shareMode = 0; // 0, share mode diff --git a/examples/matrix/batch_matmul/kernel_launch_method_by_direct/batch_matmul_custom_tiling.cpp b/examples/matrix/batch_matmul/kernel_launch_method_by_direct/batch_matmul_custom_tiling.cpp index 0ea8dad7ff5e88f15dc970bd79291d5f8db3ce00..54108191b28d3d114ad7a90549728a79881a9f00 100644 --- a/examples/matrix/batch_matmul/kernel_launch_method_by_direct/batch_matmul_custom_tiling.cpp +++ b/examples/matrix/batch_matmul/kernel_launch_method_by_direct/batch_matmul_custom_tiling.cpp @@ -14,7 +14,6 @@ #include #include #include "../host_tiling/batch_matmul_custom_tiling.h" -using namespace matmul_tiling; using namespace std; uint8_t *GetTilingBuf(optiling::TCubeTiling *tilingData) @@ -37,7 +36,7 @@ uint8_t *GetTilingBuf(optiling::TCubeTiling *tilingData) uint8_t *GenerateTiling() { optiling::TCubeTiling tilingData; - MultiCoreMatmulTiling tilingApi; + matmul_tiling::MultiCoreMatmulTiling tilingApi; bool res = ComputeTiling(tilingData, &tilingApi, false); if (!res) { std::cout << "gen tiling failed" << std::endl; diff --git a/examples/matrix/batch_matmul/kernel_launch_method_by_framework/op_host/batch_matmul_custom.cpp b/examples/matrix/batch_matmul/kernel_launch_method_by_framework/op_host/batch_matmul_custom.cpp index ff053d118fa21741d401bbaab613f1c5516b4491..f36003929bfd4e941d367047af267675004d0c8b 100644 --- a/examples/matrix/batch_matmul/kernel_launch_method_by_framework/op_host/batch_matmul_custom.cpp +++ b/examples/matrix/batch_matmul/kernel_launch_method_by_framework/op_host/batch_matmul_custom.cpp @@ -12,13 +12,12 @@ #include "register/op_def_registry.h" #include "tiling/platform/platform_ascendc.h" #include "tiling/tiling_api.h" -using namespace matmul_tiling; namespace optiling { static ge::graphStatus TilingFunc(gert::TilingContext* context) { auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo()); - MultiCoreMatmulTiling cubeTiling(ascendcPlatform); + matmul_tiling::MultiCoreMatmulTiling cubeTiling(ascendcPlatform); MatmulCustomTilingData tiling; bool res = ComputeTiling(tiling.cubeTilingData, &cubeTiling, true); if (!res) { diff --git a/examples/matrix/batch_matmul/kernel_launch_method_by_framework/op_kernel/batch_matmul_custom.cpp b/examples/matrix/batch_matmul/kernel_launch_method_by_framework/op_kernel/batch_matmul_custom.cpp index a653a0f5b69e4b881cccd4f744a980576261c307..c55caba16e0cc356523fc44b7dc8c61a0f27c201 100644 --- a/examples/matrix/batch_matmul/kernel_launch_method_by_framework/op_kernel/batch_matmul_custom.cpp +++ b/examples/matrix/batch_matmul/kernel_launch_method_by_framework/op_kernel/batch_matmul_custom.cpp @@ -12,17 +12,15 @@ #include "lib/matmul_intf.h" #include "../../../../../../kernel_impl/batch_matmul_custom_impl.h" -using namespace matmul; - constexpr int32_t FULL_L1_SIZE = 512 * 1024; constexpr int32_t FULL_L0C_SIZE = 128 * 1024; extern "C" __global__ __aicore__ void batch_matmul_custom(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); - typedef MatmulType A_TYPE; - typedef MatmulType B_TYPE; - typedef MatmulType C_TYPE; - typedef MatmulType BIAS_TYPE; + typedef matmul::MatmulType A_TYPE; + typedef matmul::MatmulType B_TYPE; + typedef matmul::MatmulType C_TYPE; + typedef matmul::MatmulType BIAS_TYPE; BatchMatmulKernel batchMatmulKernel; AscendC::TPipe pipe; tilingData.cubeTilingData.shareMode = 0; // 0, share mode diff --git a/examples/matrix/matmul/host_tiling/matmul_custom_tiling.h b/examples/matrix/matmul/host_tiling/matmul_custom_tiling.h index 93658ca3f4f662f390cb2651906bace340fce349..89bcba1a6d63a8c96f1fa20d438e64f2210403c0 100644 --- a/examples/matrix/matmul/host_tiling/matmul_custom_tiling.h +++ b/examples/matrix/matmul/host_tiling/matmul_custom_tiling.h @@ -13,7 +13,6 @@ #include "register/tilingdata_base.h" #include "tiling/tiling_api.h" -using namespace matmul_tiling; namespace optiling { BEGIN_TILING_DATA_DEF(MatmulCustomTilingData) TILING_DATA_FIELD_DEF_STRUCT(TCubeTiling, cubeTilingData); @@ -22,16 +21,16 @@ END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(MatmulCustom, MatmulCustomTilingData) } -bool ComputeTiling(optiling::TCubeTiling& tiling, MultiCoreMatmulTiling* cubeTiling, int32_t M, int32_t N, int32_t K, +bool ComputeTiling(optiling::TCubeTiling& tiling, matmul_tiling::MultiCoreMatmulTiling* cubeTiling, int32_t M, int32_t N, int32_t K, int32_t blockDim, bool isBias) { cubeTiling->SetDim(blockDim); int32_t baseM = 128; int32_t baseN = 256; - cubeTiling->SetAType(TPosition::GM, CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16, true); // A is transposed - cubeTiling->SetBType(TPosition::GM, CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16); - cubeTiling->SetCType(TPosition::GM, CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); - cubeTiling->SetBiasType(TPosition::GM, CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); + cubeTiling->SetAType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16, true); // A is transposed + cubeTiling->SetBType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16); + cubeTiling->SetCType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); + cubeTiling->SetBiasType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); cubeTiling->SetShape(M, N, K); cubeTiling->SetOrgShape(M, N, K); cubeTiling->SetFixSplit(baseM, baseN, -1); diff --git a/examples/matrix/matmul/kernel_impl/matmul_custom_impl.h b/examples/matrix/matmul/kernel_impl/matmul_custom_impl.h index 40a764227e305e812d03a763ee0c8b8d83af08d9..6f85560549f5a1a7274538d03d0d32ba95aae833 100644 --- a/examples/matrix/matmul/kernel_impl/matmul_custom_impl.h +++ b/examples/matrix/matmul/kernel_impl/matmul_custom_impl.h @@ -13,8 +13,6 @@ #include "kernel_operator.h" #include "lib/matmul_intf.h" -using namespace matmul; - template class MatmulKernel { public: @@ -22,10 +20,10 @@ class MatmulKernel { __aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace, const TCubeTiling& tiling); template __aicore__ inline void Process(AscendC::TPipe* pipe); - Matmul, - MatmulType, - MatmulType, - MatmulType, CFG_MDL> matmulObj; + matmul::Matmul, + matmul::MatmulType, + matmul::MatmulType, + matmul::MatmulType, CFG_MDL> matmulObj; private: __aicore__ inline void CalcOffset(int32_t blockIdx, const TCubeTiling& tiling, int32_t& offsetA, int32_t& offsetB, @@ -71,8 +69,8 @@ __aicore__ inline void MatmulKernel::Process(Asce if constexpr (setTmpSpace) { AscendC::TBuf<> tmpMMFormatUb; AscendC::LocalTensor mmformatUb; - pipe->InitBuffer(tmpMMFormatUb, TOTAL_VEC_LOCAL_SIZE); - mmformatUb = tmpMMFormatUb.Get(TOTAL_VEC_LOCAL_SIZE); + pipe->InitBuffer(tmpMMFormatUb, AscendC::TOTAL_VEC_LOCAL_SIZE); + mmformatUb = tmpMMFormatUb.Get(AscendC::TOTAL_VEC_LOCAL_SIZE); matmulObj.SetLocalWorkspace(mmformatUb); } diff --git a/examples/matrix/matmul/kernel_launch_method_by_direct/matmul_custom.cpp b/examples/matrix/matmul/kernel_launch_method_by_direct/matmul_custom.cpp index a274b83f0b1dec6b23707382c9a27a29f8b61207..3c4040b092ab47b25dee1baff4263036f691f816 100644 --- a/examples/matrix/matmul/kernel_launch_method_by_direct/matmul_custom.cpp +++ b/examples/matrix/matmul/kernel_launch_method_by_direct/matmul_custom.cpp @@ -12,8 +12,6 @@ #include "lib/matmul_intf.h" #include "../kernel_impl/matmul_custom_impl.h" -using namespace matmul; - __aicore__ inline void CopyTiling(TCubeTiling* tiling, GM_ADDR tilingGM) { uint32_t* ptr = reinterpret_cast(tiling); diff --git a/examples/matrix/matmul/kernel_launch_method_by_direct/matmul_custom_tiling.cpp b/examples/matrix/matmul/kernel_launch_method_by_direct/matmul_custom_tiling.cpp index 21387bebb30fc75194c548a167eb98b2b6d63715..5128b7d3c89d0542a9a3f322e3533c875d4f9d96 100644 --- a/examples/matrix/matmul/kernel_launch_method_by_direct/matmul_custom_tiling.cpp +++ b/examples/matrix/matmul/kernel_launch_method_by_direct/matmul_custom_tiling.cpp @@ -16,7 +16,6 @@ #include "tiling/tiling_api.h" #include "tiling/platform/platform_ascendc.h" #include "../host_tiling/matmul_custom_tiling.h" -using namespace matmul_tiling; using namespace std; constexpr int32_t USED_CORE_NUM = 4; @@ -37,7 +36,7 @@ uint8_t *GenerateTiling() bool isBias = false; optiling::TCubeTiling tilingData; auto ascendcPlatform = platform_ascendc::PlatformAscendCManager::GetInstance(); - MultiCoreMatmulTiling cubeTiling(*ascendcPlatform); + matmul_tiling::MultiCoreMatmulTiling cubeTiling(*ascendcPlatform); bool res = ComputeTiling(tilingData, &cubeTiling, M, N, K, USED_CORE_NUM, false); if (!res) { cout << "gen tiling failed" << endl; diff --git a/examples/matrix/matmul/kernel_launch_method_by_framework/op_host/matmul_custom.cpp b/examples/matrix/matmul/kernel_launch_method_by_framework/op_host/matmul_custom.cpp index 95d12755765457605eff03212ce9b4ac050e13bb..97b9c7d56db7fee0adda4bcf4003fbdbbc954144 100644 --- a/examples/matrix/matmul/kernel_launch_method_by_framework/op_host/matmul_custom.cpp +++ b/examples/matrix/matmul/kernel_launch_method_by_framework/op_host/matmul_custom.cpp @@ -11,7 +11,6 @@ #include "matmul_custom_tiling.h" #include "register/op_def_registry.h" #include "tiling/platform/platform_ascendc.h" -using namespace matmul_tiling; constexpr int32_t USED_CORE_NUM = 20; constexpr int32_t USED_CORE_NUM_310P = 8; @@ -25,7 +24,7 @@ static ge::graphStatus TilingFunc(gert::TilingContext* context) int32_t N = shape_b.GetDim(1); int32_t K = shape_a.GetDim(1); bool isBias = true; - MultiCoreMatmulTiling cubeTiling(ascendcPlatform); + matmul_tiling::MultiCoreMatmulTiling cubeTiling(ascendcPlatform); bool isASCEND310P = ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCEND310P; int blockDim = USED_CORE_NUM * 2; if (isASCEND310P) { diff --git a/examples/matrix/matmul/kernel_launch_method_by_framework/op_kernel/matmul_custom.cpp b/examples/matrix/matmul/kernel_launch_method_by_framework/op_kernel/matmul_custom.cpp index af60b3d49c0b0d5d85f81757438834636294ef33..c2970e971d699582b96554a6ff806093ea18267b 100644 --- a/examples/matrix/matmul/kernel_launch_method_by_framework/op_kernel/matmul_custom.cpp +++ b/examples/matrix/matmul/kernel_launch_method_by_framework/op_kernel/matmul_custom.cpp @@ -12,8 +12,6 @@ #include "lib/matmul_intf.h" #include "../../../../../../kernel_impl/matmul_custom_impl.h" -using namespace matmul; - 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) { GET_TILING_DATA(tilingData, tiling); MatmulKernel matmulKernel;