feat: MatMulV3 weightNz bf16/fp16 kernel support other templates-ops-nn-AtomGit

Mmatmul/common/op_host/op_api/matmul_util.cpp+27-2

	@@ -461,6 +461,13 @@ static bool CheckSupportInfoFormatNzNzNd(const MmOpInfo& mmOpInfo)
461	mmOpInfo.support_info.output_format == ge::FORMAT_ND;	461	mmOpInfo.support_info.output_format == ge::FORMAT_ND;
462	}	462	}
463		463
		464	+static bool CheckSupportInfoFormatNdNzNd(const MmOpInfo& mmOpInfo)
		唐唐唐超6月1日 WeightNz
		465	+{
		466	+ return mmOpInfo.support_info.self_format == ge::FORMAT_ND &&
		467	+ mmOpInfo.support_info.mat2_format == ge::FORMAT_FRACTAL_NZ &&
		468	+ mmOpInfo.support_info.output_format == ge::FORMAT_ND;
		469	+}
		470	+
464	static const aclTensor* GetMatMulOp(	471	static const aclTensor* GetMatMulOp(
465	const aclTensor* x1, const aclTensor* x2, const aclTensor* bias, MmOpInfo& mmOpInfo, const bool transposeX1,	472	const aclTensor* x1, const aclTensor* x2, const aclTensor* bias, MmOpInfo& mmOpInfo, const bool transposeX1,
466	const bool transposeX2, const bool offsetX, const int64_t opImplModeEnum, aclOpExecutor* executor)	473	const bool transposeX2, const bool offsetX, const int64_t opImplModeEnum, aclOpExecutor* executor)
	@@ -507,7 +514,10 @@ static const aclTensor* GetMatMulOp(
507	l0op::MatMulV3NzNzNd(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);	514	l0op::MatMulV3NzNzNd(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);
508	return mmOut;	515	return mmOut;
509	}	516	}
510	-	517	+ if (CheckSupportInfoFormatNdNzNd(mmOpInfo)) {
		518	+ x1 = l0op::ReFormat(x1, op::Format::FORMAT_ND);
		519	+ x2 = l0op::ReFormat(x2, op::Format::FORMAT_FRACTAL_NZ);
		520	+ }
511	const aclTensor* mmOut =	521	const aclTensor* mmOut =
512	l0op::MatMulV3Nd(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);	522	l0op::MatMulV3Nd(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);
513	return mmOut;	523	return mmOut;
	@@ -1888,6 +1898,16 @@ bool IsFormatSupportNd(const aclTensor* self, const aclTensor* mat2)
1888	return true;	1898	return true;
1889	}	1899	}
1890		1900
		1901	+bool CheckMMV3NdNzNdSupport(const aclTensor* self, const aclTensor* mat2, MmOpInfo& mmOpInfo) {
		1902	+ // 当前切换场景不支持输入self/mat2混精度场景，不支持self/out升精度输出场景，且输入类型只支持fp16/bf16/fp32
		1903	+ bool isDtypeSupport = mmOpInfo.support_info.self_dtype == mmOpInfo.support_info.mat2_dtype &&
		1904	+ mmOpInfo.support_info.self_dtype == mmOpInfo.support_info.output_dtype &&
		1905	+ ALIGN_UNIT_MAP.find(mmOpInfo.support_info.self_dtype) != ALIGN_UNIT_MAP.end();
		1906	+ auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
		1907	+ return isDtypeSupport && npuArch == NpuArch::DAV_2201 &&
		1908	+ self->GetStorageFormat() == Format::FORMAT_ND && mat2->GetStorageFormat() == Format::FORMAT_FRACTAL_NZ;
		1909	+}
		1910	+
1891	bool IsSupportNzNzNd(const aclTensor* self, const aclTensor* mat2)	1911	bool IsSupportNzNzNd(const aclTensor* self, const aclTensor* mat2)
1892	{	1912	{
1893	op::Shape selfShape = self->GetViewShape();	1913	op::Shape selfShape = self->GetViewShape();
	@@ -2034,7 +2054,12 @@ aclnnStatus SetMmSupportFormat(const aclTensor* self, const aclTensor* mat2, MmO
2034	mmOpInfo.support_info.mat2_format = Format::FORMAT_FRACTAL_NZ;	2054	mmOpInfo.support_info.mat2_format = Format::FORMAT_FRACTAL_NZ;
2035	return ACLNN_SUCCESS;	2055	return ACLNN_SUCCESS;
2036	}	2056	}
2037	-	2057	+ if (CheckMMV3NdNzNdSupport(self, mat2, mmOpInfo)) {
		2058	+ mmOpInfo.support_info.output_format = Format::FORMAT_ND;
		2059	+ mmOpInfo.support_info.self_format = Format::FORMAT_ND;
		2060	+ mmOpInfo.support_info.mat2_format = Format::FORMAT_FRACTAL_NZ;
		2061	+ return ACLNN_SUCCESS;
		2062	+ }
2038	mmOpInfo.support_info.output_format = Format::FORMAT_FRACTAL_NZ;	2063	mmOpInfo.support_info.output_format = Format::FORMAT_FRACTAL_NZ;
2039	mmOpInfo.support_info.self_format = Format::FORMAT_FRACTAL_NZ;	2064	mmOpInfo.support_info.self_format = Format::FORMAT_FRACTAL_NZ;
2040	mmOpInfo.support_info.mat2_format = Format::FORMAT_FRACTAL_NZ;	2065	mmOpInfo.support_info.mat2_format = Format::FORMAT_FRACTAL_NZ;

Mmatmul/mat_mul_v3/op_host/op_tiling/matmul_v3_base_tiling.cpp+12-12文件内容审核中，请稍后刷新重试

Mmatmul/mat_mul_v3/op_kernel/mat_mul_sc_splitk_kernel_gm_to_l1.h+4-4

Mmatmul/mat_mul_v3/op_kernel/mat_mul_v3.cpp+1-6

	@@ -61,11 +61,6 @@ constexpr CubeFormat format_y = CubeFormat::NZ;
61	constexpr CubeFormat format_y = CubeFormat::ND;	61	constexpr CubeFormat format_y = CubeFormat::ND;
62	#endif	62	#endif
63		63
64	-#if defined(FORMAT_X1) && FORMAT_X1 == FORMAT_ND && defined(ORIG_DTYPE_X2) && ORIG_DTYPE_X2 == DT_FLOAT
65	- #define IS_ND_NZ_FP32 1
66	-#else
67	- #define IS_ND_NZ_FP32 0
68	-#endif
69	// ND_NZ_FP32场景支持走入多个模板	64	// ND_NZ_FP32场景支持走入多个模板
70		65
71	#define MMV3_IMPL(templateFunc, cFormat, ...) \	66	#define MMV3_IMPL(templateFunc, cFormat, ...) \
	@@ -192,7 +187,7 @@ __global__ __aicore__ void mat_mul_v3(
192	MatmulBaseKernel, format_x1, MatmulBaseBlock, MM_CFG_VEC_ND2NZ	187	MatmulBaseKernel, format_x1, MatmulBaseBlock, MM_CFG_VEC_ND2NZ
193	);	188	);
194	}	189	}
195	-#elif defined(FORMAT_X2) && FORMAT_X2 == FRACTAL_NZ && !IS_ND_NZ_FP32	190	+#elif defined(FORMAT_X1) && FORMAT_X1 == FORMAT_FRACTAL_NZ && defined(FORMAT_X2) && FORMAT_X2 == FORMAT_FRACTAL_NZ
196	if constexpr (LOADMODE == MAT_MUL_V3_BASE_FULLLOAD && SPLITCOREMODE == MAT_MUL_V3_BASE_SPLIT_K &&	191	if constexpr (LOADMODE == MAT_MUL_V3_BASE_FULLLOAD && SPLITCOREMODE == MAT_MUL_V3_BASE_SPLIT_K &&
197	FIXOPTI == MAT_MUL_V3_BASE_FIXOPTI && MIXND2NZ == MAT_MUL_V3_MIXND2NZ_TRUE) {	192	FIXOPTI == MAT_MUL_V3_BASE_FIXOPTI && MIXND2NZ == MAT_MUL_V3_MIXND2NZ_TRUE) {
198	MMV3_IMPL_CLASS(	193	MMV3_IMPL_CLASS(

Mmatmul/mat_mul_v3/op_kernel/mat_mul_v3_tiling_key.h+3-19

	@@ -43,15 +43,6 @@
43	#define FP32_ADDMM_DISABLE 0	43	#define FP32_ADDMM_DISABLE 0
44	#define FP32_ADDMM_ENABLE 1	44	#define FP32_ADDMM_ENABLE 1
45		45
46	-#define FRACTAL_NZ 29
47	-
48	-#if defined(FORMAT_X1) && FORMAT_X1 == FORMAT_ND && defined(ORIG_DTYPE_X2) && ORIG_DTYPE_X2 == DT_FLOAT
49	- #define IS_ND_NZ_FP32 1
50	-#else
51	- #define IS_ND_NZ_FP32 0
52	-#endif
53	-// ND_NZ_FP32场景支持走入多个模板
54	-
55	ASCENDC_TPL_ARGS_DECL(	46	ASCENDC_TPL_ARGS_DECL(
56	MatMulV3,	47	MatMulV3,
57	ASCENDC_TPL_UINT_DECL(	48	ASCENDC_TPL_UINT_DECL(
	@@ -88,7 +79,8 @@ ASCENDC_TPL_SEL(
88	ASCENDC_TPL_UINT_SEL(SPECIALOPT, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_K_NOT_SHIFT),	79	ASCENDC_TPL_UINT_SEL(SPECIALOPT, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_K_NOT_SHIFT),
89	ASCENDC_TPL_UINT_SEL(FP32ADDMM, ASCENDC_TPL_UI_LIST, FP32_ADDMM_DISABLE)),	80	ASCENDC_TPL_UINT_SEL(FP32ADDMM, ASCENDC_TPL_UI_LIST, FP32_ADDMM_DISABLE)),
90	);	81	);
91	-#elif defined(FORMAT_X2) && FORMAT_X2 == FRACTAL_NZ && !IS_ND_NZ_FP32	82	+// NZ_NZ仅支持基础模板
		83	+#elif defined(FORMAT_X1) && FORMAT_X1 == FORMAT_FRACTAL_NZ && defined(FORMAT_X2) && FORMAT_X2 == FORMAT_FRACTAL_NZ
92	ASCENDC_TPL_SEL(	84	ASCENDC_TPL_SEL(
93	ASCENDC_TPL_ARGS_SEL(	85	ASCENDC_TPL_ARGS_SEL(
94	ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),	86	ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
	@@ -104,15 +96,7 @@ ASCENDC_TPL_SEL(
104	ASCENDC_TPL_UINT_SEL(SPLITCOREMODE, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_BASE_SPLIT_K),	96	ASCENDC_TPL_UINT_SEL(SPLITCOREMODE, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_BASE_SPLIT_K),
105	ASCENDC_TPL_UINT_SEL(FIXOPTI, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_BASE_FIXOPTI),	97	ASCENDC_TPL_UINT_SEL(FIXOPTI, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_BASE_FIXOPTI),
106	ASCENDC_TPL_UINT_SEL(MIXND2NZ, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_MIXND2NZ_FALSE),	98	ASCENDC_TPL_UINT_SEL(MIXND2NZ, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_MIXND2NZ_FALSE),
107	- ASCENDC_TPL_UINT_SEL(SPECIALOPT, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_K_NOT_SHIFT),	99	+ ASCENDC_TPL_UINT_SEL(SPECIALOPT, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_K_SHIFT, MAT_MUL_V3_K_NOT_SHIFT),
108	- ASCENDC_TPL_UINT_SEL(FP32ADDMM, ASCENDC_TPL_UI_LIST, FP32_ADDMM_DISABLE)),
109	- ASCENDC_TPL_ARGS_SEL(
110	- ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIC_ONLY),
111	- ASCENDC_TPL_UINT_SEL(LOADMODE, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_BASE_FULLLOAD),
112	- ASCENDC_TPL_UINT_SEL(SPLITCOREMODE, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_BASE_SPLIT_K),
113	- ASCENDC_TPL_UINT_SEL(FIXOPTI, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_BASE_FIXOPTI),
114	- ASCENDC_TPL_UINT_SEL(MIXND2NZ, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_MIXND2NZ_FALSE),
115	- ASCENDC_TPL_UINT_SEL(SPECIALOPT, ASCENDC_TPL_UI_LIST, MAT_MUL_V3_K_SHIFT),
116	ASCENDC_TPL_UINT_SEL(FP32ADDMM, ASCENDC_TPL_UI_LIST, FP32_ADDMM_DISABLE)),	100	ASCENDC_TPL_UINT_SEL(FP32ADDMM, ASCENDC_TPL_UI_LIST, FP32_ADDMM_DISABLE)),
117	);	101	);
118	#else	102	#else

Mmatmul/mat_mul_v3/tests/ut/op_host/test_matmul_v3_tiling.cpp+3-3

	@@ -737,8 +737,8 @@ static TilingTestParam ascend910B_cases_params[] = {
737	"block_dim":{"CORE_NUM":24, "vector_core_cnt": 48},"corerect_range_flag":null,"dynamic_mode":"dynamic_mkn", "fused_double_operand_num": 0,	737	"block_dim":{"CORE_NUM":24, "vector_core_cnt": 48},"corerect_range_flag":null,"dynamic_mode":"dynamic_mkn", "fused_double_operand_num": 0,
738	"hardware_info": {"BT_SIZE": 1024, "load3d_constraints": "unknown", "Intrinsic_fix_pipe_l0c2out": true, "Intrinsic_data_move_l12ub": false, "Intrinsic_data_move_l0c2ub": false, "Intrinsic_data_move_out2l1_nd2nz": true, "UB_SIZE": 196608, "L2_SIZE": 201326592, "L1_SIZE": 524288, "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM":24, "vector_core_cnt": 48, "socVersion": "Ascend910B" },	738	"hardware_info": {"BT_SIZE": 1024, "load3d_constraints": "unknown", "Intrinsic_fix_pipe_l0c2out": true, "Intrinsic_data_move_l12ub": false, "Intrinsic_data_move_l0c2ub": false, "Intrinsic_data_move_out2l1_nd2nz": true, "UB_SIZE": 196608, "L2_SIZE": 201326592, "L1_SIZE": 524288, "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM":24, "vector_core_cnt": 48, "socVersion": "Ascend910B" },
739	"format_a":"ND","format_b":"FRACTAL_NZ","repo_range":{},"repo_seeds":{}})",	739	"format_a":"ND","format_b":"FRACTAL_NZ","repo_range":{},"repo_seeds":{}})",
740	- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, true, false, 0, 0, {24576, 16}, {1, 1536, 16, 16}, {16, 16}, {24576, 16}, {24576, 16}, {16, 16}, false, 0, 0, 24, 65536,	740	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, true, false, 0, 0, {24576, 16}, {1, 1536, 16, 16}, {16, 16}, {24576, 16}, {24576, 16}, {16, 16}, false, 0, 0, 24, 65584,
741	- "24 16 16 24576 24576 16 1024 24576 16 1024 16 512 8 1 1 0 0 0 0 393216 1024 0 1 1 1 1 256 4 0 0 2 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 0 0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 "	741	+ "24 16 16 24576 24576 384 16 384 128 128 128 9 6 3 1 0 0 1 0 393216 1024 0 1 1 1 1 3 3 0 0 2 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 0 0 1 0 0 0 1 1 0 0 0 0 0 0 0 0 "
742	},	742	},
743	{	743	{
744	"MatMulV3_basic_test_weightNz_tiling_k_shift", "MatMulV3", R"({"_pattern": "MatMul", "attrs":{"transpose_a":false,"transpose_b":false,"offset_x":0,"opImplMode":0},	744	"MatMulV3_basic_test_weightNz_tiling_k_shift", "MatMulV3", R"({"_pattern": "MatMul", "attrs":{"transpose_a":false,"transpose_b":false,"offset_x":0,"opImplMode":0},
	@@ -747,7 +747,7 @@ static TilingTestParam ascend910B_cases_params[] = {
747	"hardware_info": {"BT_SIZE": 1024, "load3d_constraints": "unknown", "Intrinsic_fix_pipe_l0c2out": true, "Intrinsic_data_move_l12ub": false, "Intrinsic_data_move_l0c2ub": false, "Intrinsic_data_move_out2l1_nd2nz": true, "UB_SIZE": 196608, "L2_SIZE": 201326592, "L1_SIZE": 524288, "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM":24, "vector_core_cnt": 48, "socVersion": "Ascend910B" },	747	"hardware_info": {"BT_SIZE": 1024, "load3d_constraints": "unknown", "Intrinsic_fix_pipe_l0c2out": true, "Intrinsic_data_move_l12ub": false, "Intrinsic_data_move_l0c2ub": false, "Intrinsic_data_move_out2l1_nd2nz": true, "UB_SIZE": 196608, "L2_SIZE": 201326592, "L1_SIZE": 524288, "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM":24, "vector_core_cnt": 48, "socVersion": "Ascend910B" },
748	"format_a":"ND","format_b":"FRACTAL_NZ","repo_range":{},"repo_seeds":{}})",	748	"format_a":"ND","format_b":"FRACTAL_NZ","repo_range":{},"repo_seeds":{}})",
749	ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, false, false, 0, 0, {192, 6144}, {512, 384, 16, 16}, {192, 8192}, {192, 6144}, {6144, 8192}, {192, 8192}, false, 0, 0, 24, 65536,	749	ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, false, false, 0, 0, {192, 6144}, {512, 384, 16, 16}, {192, 8192}, {192, 6144}, {6144, 8192}, {192, 8192}, false, 0, 0, 24, 65536,
750	- "24 192 8192 6144 6144 128 128 6144 128 128 64 16 16 1 1 0 0 0 0 458752 98304 0 1 1 1 1 8 8 0 0 2 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 64 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 ", DT_BF16, DT_BF16	750	+ "24 192 8192 6144 6144 128 128 6144 128 128 64 16 16 1 1 0 0 0 0 458752 98304 0 1 1 1 1 8 8 0 0 2 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 11 2 6 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 ", DT_BF16, DT_BF16
751	},	751	},
752	{	752	{
753	"MatMulV3_basic_force_grp_acc_for_fp32", "MatMulV3", R"({"_pattern": "MatMul", "attrs":{"transpose_a":false,"transpose_b":false,"offset_x":0,"opImplMode":4},	753	"MatMulV3_basic_force_grp_acc_for_fp32", "MatMulV3", R"({"_pattern": "MatMul", "attrs":{"transpose_a":false,"transpose_b":false,"offset_x":0,"opImplMode":4},

	@@ -299,7 +299,7 @@ MatMulBaseKernelSingleCoreSplitKGmToL1<A_TYPE, B_TYPE, L0C_TYPE, OUTPUT_TYPE, BI
299	CopyInAl1ND(M_1, k0, K_INDEX_0, k0);	299	CopyInAl1ND(M_1, k0, K_INDEX_0, k0);
300	}	300	}
301	int32_t nPingReal = N_INDEX_0 == nnloop_ - 1 ? block_.params_.innerSingleCoreN : N0;	301	int32_t nPingReal = N_INDEX_0 == nnloop_ - 1 ? block_.params_.innerSingleCoreN : N0;
302	- CopyInBl1ND(bpingflag, k0, K_INDEX_0, nPingReal, N_INDEX_0);	302	+ CopyInBl1(bpingflag, k0, K_INDEX_0, nPingReal, N_INDEX_0);
303	for (uint64_t kIndex = 0; kIndex < block_.params_.loopK; ++kIndex) {	303	for (uint64_t kIndex = 0; kIndex < block_.params_.loopK; ++kIndex) {
304	bool lastK = kIndex == block_.params_.loopK - 1;	304	bool lastK = kIndex == block_.params_.loopK - 1;
305	int32_t realK = lastK ? block_.matmulTilingData_->matmulTiling.Ka - kIndex * k0 : k0;	305	int32_t realK = lastK ? block_.matmulTilingData_->matmulTiling.Ka - kIndex * k0 : k0;
	@@ -320,7 +320,7 @@ MatMulBaseKernelSingleCoreSplitKGmToL1<A_TYPE, B_TYPE, L0C_TYPE, OUTPUT_TYPE, BI
320	}	320	}
321	if (nnloop_ > 1) {	321	if (nnloop_ > 1) {
322	int32_t nPongReal = N_INDEX_0 + 1 == nnloop_ - 1 ? block_.params_.innerSingleCoreN - (N_INDEX_0 + 1) * N0: N0;	322	int32_t nPongReal = N_INDEX_0 + 1 == nnloop_ - 1 ? block_.params_.innerSingleCoreN - (N_INDEX_0 + 1) * N0: N0;
323	- CopyInBl1ND(!bpingflag, realK, K_INDEX_0, nPongReal, N_INDEX_0 + 1);	323	+ CopyInBl1(!bpingflag, realK, K_INDEX_0, nPongReal, N_INDEX_0 + 1);
324	}	324	}
325	ProcessBaseMNK(lastK, realK, nextRealK, kIndex);	325	ProcessBaseMNK(lastK, realK, nextRealK, kIndex);
326	}	326	}
	@@ -646,10 +646,10 @@ MatMulBaseKernelSingleCoreSplitKGmToL1<A_TYPE, B_TYPE, L0C_TYPE, OUTPUT_TYPE, BI
646	int32_t realN = lastN ? block_.params_.innerSingleCoreN - nloop * N0 : N0;	646	int32_t realN = lastN ? block_.params_.innerSingleCoreN - nloop * N0 : N0;
647	if (!lastN && !firstN) {	647	if (!lastN && !firstN) {
648	int32_t nextRealN = nloop == nnloop_ - 2 ? block_.params_.innerSingleCoreN - (nloop + 1) * N0 : N0;	648	int32_t nextRealN = nloop == nnloop_ - 2 ? block_.params_.innerSingleCoreN - (nloop + 1) * N0 : N0;
649	- CopyInBl1ND(!bpingflag, realK, K_INDEX_0, nextRealN, nloop + 1);	649	+ CopyInBl1(!bpingflag, realK, K_INDEX_0, nextRealN, nloop + 1);
650	}	650	}
651	if (lastN && !lastK) {	651	if (lastN && !lastK) {
652	- CopyInBl1ND(!bpingflag, nextRealK, K_INDEX_0 + 1, N0, N_INDEX_0); // B nextk	652	+ CopyInBl1(!bpingflag, nextRealK, K_INDEX_0 + 1, N0, N_INDEX_0); // B nextk
653	}	653	}
654	if (nmloop_ - 1 >= M_0) {	654	if (nmloop_ - 1 >= M_0) {
655	AMatMulB(bpingflag, nloop, realK, realN, kIndex, aL1Ping, eventIdAPingMte2Mte1, M_0);	655	AMatMulB(bpingflag, nloop, realK, realN, kIndex, aL1Ping, eventIdAPingMte2Mte1, M_0);