Star908
1826
代码
代码
Issues226
Pull Requests308
Actions
项目讨论
Wiki
分析
Star908
1826
已合并
aclnn的mm、bmm、addmm和baddbmm接口支持16in32Out场景 #5481
HKFLYE创建于 5月30日
aclnn的mm、bmm、addmm和baddbmm接口支持16in32Out场景 #5481
已合并
HKFLYE创建于 5月30日
从HuangKun8682/ops-nn_hkflyepersonal_addmm_support_16in32out_gemmV3_branch合入到cann/ops-nnmaster
共 27 个文件变更+3747-289
Mmatmul/batch_mat_mul_v3/docs/aclnnBaddbmm&aclnnInplaceBaddbmm.md+1-0
@@ -547,6 +547,7 @@ aclnnBaddbmm&aclnnInplaceBaddbmm默认确定性实现。
547 547 
548- <term>Atlas 训练系列产品</term>、<term>Atlas 推理系列产品</term>：Cube单元不支持FLOAT32计算。当输入为FLOAT32，可通过设置cubeMathType=1（ALLOW_FP32_DOWN_PRECISION）来允许接口内部cast到FLOAT16进行计算。548- <term>Atlas 训练系列产品</term>、<term>Atlas 推理系列产品</term>：Cube单元不支持FLOAT32计算。当输入为FLOAT32，可通过设置cubeMathType=1（ALLOW_FP32_DOWN_PRECISION）来允许接口内部cast到FLOAT16进行计算。
549- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：不支持batch1和batch2两输入其中一个输入为BFLOAT16, 另一个输入为FLOAT或FLOAT16的数据类型推导。549- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：不支持batch1和batch2两输入其中一个输入为BFLOAT16, 另一个输入为FLOAT或FLOAT16的数据类型推导。
550+- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：当输入self、mat2的数据类型都为FLOAT16或BFLOAT16，并且指定out数据类型为FLOAT32时，输出y的数据类型为FLOAT32。
G
Ggitee-duhuiping6月2日

接口中没有mat2、y参数啊？

likedislike
HKFLYE
HKFLYE
6月2日 评论：
HKFLYE
HKFLYE
6月2日 评论：
550 551 
551## 调用示例552## 调用示例
552 553 
Mmatmul/batch_mat_mul_v3/op_host/config/ascend910_93/batch_mat_mul_v3_binary.json+480-0
@@ -721,6 +721,486 @@
721 "value": false721 "value": false
722 }722 }
723 ]723 ]
724+ },
725+ {
726+ "bin_filename": "BatchMatMulV3_ND_ND_ND_ND_FP16_FP16_FP32_FP32",
727+ "simplified_key": "diy,2/2/2/2/1/1/0/0",
728+ "inputs": [
729+ {
730+ "name": "x1",
731+ "index": 0,
732+ "dtype": "float16",
733+ "format": "ND",
734+ "paramType": "required",
735+ "shape": [
736+ -2
737+ ]
738+ },
739+ {
740+ "name": "x2",
741+ "index": 1,
742+ "dtype": "float16",
743+ "format": "ND",
744+ "paramType": "required",
745+ "shape": [
746+ -2
747+ ]
748+ },
749+ {
750+ "name": "bias",
751+ "index": 2,
752+ "dtype": "float32",
753+ "format": "ND",
754+ "paramType": "optional",
755+ "shape": [
756+ -2
757+ ]
758+ },
759+ {
760+ "name": "offset_w",
761+ "index": 3,
762+ "dtype": "int8",
763+ "format": "ND",
764+ "paramType": "optional",
765+ "shape": [
766+ -2
767+ ]
768+ }
769+ ],
770+ "outputs": [
771+ {
772+ "name": "y",
773+ "index": 0,
774+ "dtype": "float32",
775+ "format": "ND",
776+ "paramType": "required",
777+ "shape": [
778+ -2
779+ ]
780+ }
781+ ],
782+ "attrs": [
783+ {
784+ "name": "adj_x1",
785+ "dtype": "bool",
786+ "value": false
787+ },
788+ {
789+ "name": "adj_x2",
790+ "dtype": "bool",
791+ "value": false
792+ },
793+ {
794+ "name": "offset_x",
795+ "dtype": "int",
796+ "value": 0
797+ },
798+ {
799+ "name": "enable_hf32",
800+ "dtype": "bool",
801+ "value": false
802+ }
803+ ]
804+ },
805+ {
806+ "bin_filename": "BatchMatMulV3_ND_ND_ND_ND_BF16_BF16_FP32_FP32",
807+ "simplified_key": "diy,2/2/2/2/27/27/0/0",
808+ "inputs": [
809+ {
810+ "name": "x1",
811+ "index": 0,
812+ "dtype": "bfloat16",
813+ "format": "ND",
814+ "paramType": "required",
815+ "shape": [
816+ -2
817+ ]
818+ },
819+ {
820+ "name": "x2",
821+ "index": 1,
822+ "dtype": "bfloat16",
823+ "format": "ND",
824+ "paramType": "required",
825+ "shape": [
826+ -2
827+ ]
828+ },
829+ {
830+ "name": "bias",
831+ "index": 2,
832+ "dtype": "float32",
833+ "format": "ND",
834+ "paramType": "optional",
835+ "shape": [
836+ -2
837+ ]
838+ },
839+ {
840+ "name": "offset_w",
841+ "index": 3,
842+ "dtype": "int8",
843+ "format": "ND",
844+ "paramType": "optional",
845+ "shape": [
846+ -2
847+ ]
848+ }
849+ ],
850+ "outputs": [
851+ {
852+ "name": "y",
853+ "index": 0,
854+ "dtype": "float32",
855+ "format": "ND",
856+ "paramType": "required",
857+ "shape": [
858+ -2
859+ ]
860+ }
861+ ],
862+ "attrs": [
863+ {
864+ "name": "adj_x1",
865+ "dtype": "bool",
866+ "value": false
867+ },
868+ {
869+ "name": "adj_x2",
870+ "dtype": "bool",
871+ "value": false
872+ },
873+ {
874+ "name": "offset_x",
875+ "dtype": "int",
876+ "value": 0
877+ },
878+ {
879+ "name": "enable_hf32",
880+ "dtype": "bool",
881+ "value": false
882+ }
883+ ]
884+ },
885+ {
886+ "bin_filename": "BatchMatMulV3_ND_NZ_ND_ND_FP16_FP16_FP16_FP32",
887+ "simplified_key": "diy,2/29/2/2/1/1/1/0",
888+ "inputs": [
889+ {
890+ "name": "x1",
891+ "index": 0,
892+ "dtype": "float16",
893+ "format": "ND",
894+ "paramType": "required",
895+ "shape": [
896+ -2
897+ ]
898+ },
899+ {
900+ "name": "x2",
901+ "index": 1,
902+ "dtype": "float16",
903+ "format": "FRACTAL_NZ",
904+ "paramType": "required",
905+ "shape": [
906+ -2
907+ ]
908+ },
909+ {
910+ "name": "bias",
911+ "index": 2,
912+ "dtype": "float16",
913+ "format": "ND",
914+ "paramType": "optional",
915+ "shape": [
916+ -2
917+ ]
918+ },
919+ {
920+ "name": "offset_w",
921+ "index": 3,
922+ "dtype": "int8",
923+ "format": "ND",
924+ "paramType": "optional",
925+ "shape": [
926+ -2
927+ ]
928+ }
929+ ],
930+ "outputs": [
931+ {
932+ "name": "y",
933+ "index": 0,
934+ "dtype": "float32",
935+ "format": "ND",
936+ "paramType": "required",
937+ "shape": [
938+ -2
939+ ]
940+ }
941+ ],
942+ "attrs": [
943+ {
944+ "name": "adj_x1",
945+ "dtype": "bool",
946+ "value": false
947+ },
948+ {
949+ "name": "adj_x2",
950+ "dtype": "bool",
951+ "value": false
952+ },
953+ {
954+ "name": "offset_x",
955+ "dtype": "int",
956+ "value": 0
957+ },
958+ {
959+ "name": "enable_hf32",
960+ "dtype": "bool",
961+ "value": false
962+ }
963+ ]
964+ },
965+ {
966+ "bin_filename": "BatchMatMulV3_ND_NZ_ND_ND_FP16_FP16_FP32_FP32",
967+ "simplified_key": "diy,2/29/2/2/1/1/0/0",
968+ "inputs": [
969+ {
970+ "name": "x1",
971+ "index": 0,
972+ "dtype": "float16",
973+ "format": "ND",
974+ "paramType": "required",
975+ "shape": [
976+ -2
977+ ]
978+ },
979+ {
980+ "name": "x2",
981+ "index": 1,
982+ "dtype": "float16",
983+ "format": "FRACTAL_NZ",
984+ "paramType": "required",
985+ "shape": [
986+ -2
987+ ]
988+ },
989+ {
990+ "name": "bias",
991+ "index": 2,
992+ "dtype": "float32",
993+ "format": "ND",
994+ "paramType": "optional",
995+ "shape": [
996+ -2
997+ ]
998+ },
999+ {
1000+ "name": "offset_w",
1001+ "index": 3,
1002+ "dtype": "int8",
1003+ "format": "ND",
1004+ "paramType": "optional",
1005+ "shape": [
1006+ -2
1007+ ]
1008+ }
1009+ ],
1010+ "outputs": [
1011+ {
1012+ "name": "y",
1013+ "index": 0,
1014+ "dtype": "float32",
1015+ "format": "ND",
1016+ "paramType": "required",
1017+ "shape": [
1018+ -2
1019+ ]
1020+ }
1021+ ],
1022+ "attrs": [
1023+ {
1024+ "name": "adj_x1",
1025+ "dtype": "bool",
1026+ "value": false
1027+ },
1028+ {
1029+ "name": "adj_x2",
1030+ "dtype": "bool",
1031+ "value": false
1032+ },
1033+ {
1034+ "name": "offset_x",
1035+ "dtype": "int",
1036+ "value": 0
1037+ },
1038+ {
1039+ "name": "enable_hf32",
1040+ "dtype": "bool",
1041+ "value": false
1042+ }
1043+ ]
1044+ },
1045+ {
1046+ "bin_filename": "BatchMatMulV3_ND_NZ_ND_ND_BF16_BF16_BF16_FP32",
1047+ "simplified_key": "diy,2/29/2/2/27/27/27/0",
1048+ "inputs": [
1049+ {
1050+ "name": "x1",
1051+ "index": 0,
1052+ "dtype": "bfloat16",
1053+ "format": "ND",
1054+ "paramType": "required",
1055+ "shape": [
1056+ -2
1057+ ]
1058+ },
1059+ {
1060+ "name": "x2",
1061+ "index": 1,
1062+ "dtype": "bfloat16",
1063+ "format": "FRACTAL_NZ",
1064+ "paramType": "required",
1065+ "shape": [
1066+ -2
1067+ ]
1068+ },
1069+ {
1070+ "name": "bias",
1071+ "index": 2,
1072+ "dtype": "bfloat16",
1073+ "format": "ND",
1074+ "paramType": "optional",
1075+ "shape": [
1076+ -2
1077+ ]
1078+ },
1079+ {
1080+ "name": "offset_w",
1081+ "index": 3,
1082+ "dtype": "int8",
1083+ "format": "ND",
1084+ "paramType": "optional",
1085+ "shape": [
1086+ -2
1087+ ]
1088+ }
1089+ ],
1090+ "outputs": [
1091+ {
1092+ "name": "y",
1093+ "index": 0,
1094+ "dtype": "float32",
1095+ "format": "ND",
1096+ "paramType": "required",
1097+ "shape": [
1098+ -2
1099+ ]
1100+ }
1101+ ],
1102+ "attrs": [
1103+ {
1104+ "name": "adj_x1",
1105+ "dtype": "bool",
1106+ "value": false
1107+ },
1108+ {
1109+ "name": "adj_x2",
1110+ "dtype": "bool",
1111+ "value": false
1112+ },
1113+ {
1114+ "name": "offset_x",
1115+ "dtype": "int",
1116+ "value": 0
1117+ },
1118+ {
1119+ "name": "enable_hf32",
1120+ "dtype": "bool",
1121+ "value": false
1122+ }
1123+ ]
1124+ },
1125+ {
1126+ "bin_filename": "BatchMatMulV3_ND_NZ_ND_ND_BF16_BF16_FP32_FP32",
1127+ "simplified_key": "diy,2/29/2/2/27/27/0/0",
1128+ "inputs": [
1129+ {
1130+ "name": "x1",
1131+ "index": 0,
1132+ "dtype": "bfloat16",
1133+ "format": "ND",
1134+ "paramType": "required",
1135+ "shape": [
1136+ -2
1137+ ]
1138+ },
1139+ {
1140+ "name": "x2",
1141+ "index": 1,
1142+ "dtype": "bfloat16",
1143+ "format": "FRACTAL_NZ",
1144+ "paramType": "required",
1145+ "shape": [
1146+ -2
1147+ ]
1148+ },
1149+ {
1150+ "name": "bias",
1151+ "index": 2,
1152+ "dtype": "float32",
1153+ "format": "ND",
1154+ "paramType": "optional",
1155+ "shape": [
1156+ -2
1157+ ]
1158+ },
1159+ {
1160+ "name": "offset_w",
1161+ "index": 3,
1162+ "dtype": "int8",
1163+ "format": "ND",
1164+ "paramType": "optional",
1165+ "shape": [
1166+ -2
1167+ ]
1168+ }
1169+ ],
1170+ "outputs": [
1171+ {
1172+ "name": "y",
1173+ "index": 0,
1174+ "dtype": "float32",
1175+ "format": "ND",
1176+ "paramType": "required",
1177+ "shape": [
1178+ -2
1179+ ]
1180+ }
1181+ ],
1182+ "attrs": [
1183+ {
1184+ "name": "adj_x1",
1185+ "dtype": "bool",
1186+ "value": false
1187+ },
1188+ {
1189+ "name": "adj_x2",
1190+ "dtype": "bool",
1191+ "value": false
1192+ },
1193+ {
1194+ "name": "offset_x",
1195+ "dtype": "int",
1196+ "value": 0
1197+ },
1198+ {
1199+ "name": "enable_hf32",
1200+ "dtype": "bool",
1201+ "value": false
1202+ }
1203+ ]
724 }1204 }
725 ]1205 ]
726}1206}
Mmatmul/batch_mat_mul_v3/op_host/config/ascend910b/batch_mat_mul_v3_binary.json+480-0
@@ -721,6 +721,486 @@
721 "value": false721 "value": false
722 }722 }
723 ]723 ]
724+ },
725+ {
726+ "bin_filename": "BatchMatMulV3_ND_ND_ND_ND_FP16_FP16_FP32_FP32",
727+ "simplified_key": "diy,2/2/2/2/1/1/0/0",
728+ "inputs": [
729+ {
730+ "name": "x1",
731+ "index": 0,
732+ "dtype": "float16",
733+ "format": "ND",
734+ "paramType": "required",
735+ "shape": [
736+ -2
737+ ]
738+ },
739+ {
740+ "name": "x2",
741+ "index": 1,
742+ "dtype": "float16",
743+ "format": "ND",
744+ "paramType": "required",
745+ "shape": [
746+ -2
747+ ]
748+ },
749+ {
750+ "name": "bias",
751+ "index": 2,
752+ "dtype": "float32",
753+ "format": "ND",
754+ "paramType": "optional",
755+ "shape": [
756+ -2
757+ ]
758+ },
759+ {
760+ "name": "offset_w",
761+ "index": 3,
762+ "dtype": "int8",
763+ "format": "ND",
764+ "paramType": "optional",
765+ "shape": [
766+ -2
767+ ]
768+ }
769+ ],
770+ "outputs": [
771+ {
772+ "name": "y",
773+ "index": 0,
774+ "dtype": "float32",
775+ "format": "ND",
776+ "paramType": "required",
777+ "shape": [
778+ -2
779+ ]
780+ }
781+ ],
782+ "attrs": [
783+ {
784+ "name": "adj_x1",
785+ "dtype": "bool",
786+ "value": false
787+ },
788+ {
789+ "name": "adj_x2",
790+ "dtype": "bool",
791+ "value": false
792+ },
793+ {
794+ "name": "offset_x",
795+ "dtype": "int",
796+ "value": 0
797+ },
798+ {
799+ "name": "enable_hf32",
800+ "dtype": "bool",
801+ "value": false
802+ }
803+ ]
804+ },
805+ {
806+ "bin_filename": "BatchMatMulV3_ND_ND_ND_ND_BF16_BF16_FP32_FP32",
807+ "simplified_key": "diy,2/2/2/2/27/27/0/0",
808+ "inputs": [
809+ {
810+ "name": "x1",
811+ "index": 0,
812+ "dtype": "bfloat16",
813+ "format": "ND",
814+ "paramType": "required",
815+ "shape": [
816+ -2
817+ ]
818+ },
819+ {
820+ "name": "x2",
821+ "index": 1,
822+ "dtype": "bfloat16",
823+ "format": "ND",
824+ "paramType": "required",
825+ "shape": [
826+ -2
827+ ]
828+ },
829+ {
830+ "name": "bias",
831+ "index": 2,
832+ "dtype": "float32",
833+ "format": "ND",
834+ "paramType": "optional",
835+ "shape": [
836+ -2
837+ ]
838+ },
839+ {
840+ "name": "offset_w",
841+ "index": 3,
842+ "dtype": "int8",
843+ "format": "ND",
844+ "paramType": "optional",
845+ "shape": [
846+ -2
847+ ]
848+ }
849+ ],
850+ "outputs": [
851+ {
852+ "name": "y",
853+ "index": 0,
854+ "dtype": "float32",
855+ "format": "ND",
856+ "paramType": "required",
857+ "shape": [
858+ -2
859+ ]
860+ }
861+ ],
862+ "attrs": [
863+ {
864+ "name": "adj_x1",
865+ "dtype": "bool",
866+ "value": false
867+ },
868+ {
869+ "name": "adj_x2",
870+ "dtype": "bool",
871+ "value": false
872+ },
873+ {
874+ "name": "offset_x",
875+ "dtype": "int",
876+ "value": 0
877+ },
878+ {
879+ "name": "enable_hf32",
880+ "dtype": "bool",
881+ "value": false
882+ }
883+ ]
884+ },
885+ {
886+ "bin_filename": "BatchMatMulV3_ND_NZ_ND_ND_FP16_FP16_FP16_FP32",
887+ "simplified_key": "diy,2/29/2/2/1/1/1/0",
888+ "inputs": [
889+ {
890+ "name": "x1",
891+ "index": 0,
892+ "dtype": "float16",
893+ "format": "ND",
894+ "paramType": "required",
895+ "shape": [
896+ -2
897+ ]
898+ },
899+ {
900+ "name": "x2",
901+ "index": 1,
902+ "dtype": "float16",
903+ "format": "FRACTAL_NZ",
904+ "paramType": "required",
905+ "shape": [
906+ -2
907+ ]
908+ },
909+ {
910+ "name": "bias",
911+ "index": 2,
912+ "dtype": "float16",
913+ "format": "ND",
914+ "paramType": "optional",
915+ "shape": [
916+ -2
917+ ]
918+ },
919+ {
920+ "name": "offset_w",
921+ "index": 3,
922+ "dtype": "int8",
923+ "format": "ND",
924+ "paramType": "optional",
925+ "shape": [
926+ -2
927+ ]
928+ }
929+ ],
930+ "outputs": [
931+ {
932+ "name": "y",
933+ "index": 0,
934+ "dtype": "float32",
935+ "format": "ND",
936+ "paramType": "required",
937+ "shape": [
938+ -2
939+ ]
940+ }
941+ ],
942+ "attrs": [
943+ {
944+ "name": "adj_x1",
945+ "dtype": "bool",
946+ "value": false
947+ },
948+ {
949+ "name": "adj_x2",
950+ "dtype": "bool",
951+ "value": false
952+ },
953+ {
954+ "name": "offset_x",
955+ "dtype": "int",
956+ "value": 0
957+ },
958+ {
959+ "name": "enable_hf32",
960+ "dtype": "bool",
961+ "value": false
962+ }
963+ ]
964+ },
965+ {
966+ "bin_filename": "BatchMatMulV3_ND_NZ_ND_ND_FP16_FP16_FP32_FP32",
967+ "simplified_key": "diy,2/29/2/2/1/1/0/0",
968+ "inputs": [
969+ {
970+ "name": "x1",
971+ "index": 0,
972+ "dtype": "float16",
973+ "format": "ND",
974+ "paramType": "required",
975+ "shape": [
976+ -2
977+ ]
978+ },
979+ {
980+ "name": "x2",
981+ "index": 1,
982+ "dtype": "float16",
983+ "format": "FRACTAL_NZ",
984+ "paramType": "required",
985+ "shape": [
986+ -2
987+ ]
988+ },
989+ {
990+ "name": "bias",
991+ "index": 2,
992+ "dtype": "float32",
993+ "format": "ND",
994+ "paramType": "optional",
995+ "shape": [
996+ -2
997+ ]
998+ },
999+ {
1000+ "name": "offset_w",
1001+ "index": 3,
1002+ "dtype": "int8",
1003+ "format": "ND",
1004+ "paramType": "optional",
1005+ "shape": [
1006+ -2
1007+ ]
1008+ }
1009+ ],
1010+ "outputs": [
1011+ {
1012+ "name": "y",
1013+ "index": 0,
1014+ "dtype": "float32",
1015+ "format": "ND",
1016+ "paramType": "required",
1017+ "shape": [
1018+ -2
1019+ ]
1020+ }
1021+ ],
1022+ "attrs": [
1023+ {
1024+ "name": "adj_x1",
1025+ "dtype": "bool",
1026+ "value": false
1027+ },
1028+ {
1029+ "name": "adj_x2",
1030+ "dtype": "bool",
1031+ "value": false
1032+ },
1033+ {
1034+ "name": "offset_x",
1035+ "dtype": "int",
1036+ "value": 0
1037+ },
1038+ {
1039+ "name": "enable_hf32",
1040+ "dtype": "bool",
1041+ "value": false
1042+ }
1043+ ]
1044+ },
1045+ {
1046+ "bin_filename": "BatchMatMulV3_ND_NZ_ND_ND_BF16_BF16_BF16_FP32",
1047+ "simplified_key": "diy,2/29/2/2/27/27/27/0",
1048+ "inputs": [
1049+ {
1050+ "name": "x1",
1051+ "index": 0,
1052+ "dtype": "bfloat16",
1053+ "format": "ND",
1054+ "paramType": "required",
1055+ "shape": [
1056+ -2
1057+ ]
1058+ },
1059+ {
1060+ "name": "x2",
1061+ "index": 1,
1062+ "dtype": "bfloat16",
1063+ "format": "FRACTAL_NZ",
1064+ "paramType": "required",
1065+ "shape": [
1066+ -2
1067+ ]
1068+ },
1069+ {
1070+ "name": "bias",
1071+ "index": 2,
1072+ "dtype": "bfloat16",
1073+ "format": "ND",
1074+ "paramType": "optional",
1075+ "shape": [
1076+ -2
1077+ ]
1078+ },
1079+ {
1080+ "name": "offset_w",
1081+ "index": 3,
1082+ "dtype": "int8",
1083+ "format": "ND",
1084+ "paramType": "optional",
1085+ "shape": [
1086+ -2
1087+ ]
1088+ }
1089+ ],
1090+ "outputs": [
1091+ {
1092+ "name": "y",
1093+ "index": 0,
1094+ "dtype": "float32",
1095+ "format": "ND",
1096+ "paramType": "required",
1097+ "shape": [
1098+ -2
1099+ ]
1100+ }
1101+ ],
1102+ "attrs": [
1103+ {
1104+ "name": "adj_x1",
1105+ "dtype": "bool",
1106+ "value": false
1107+ },
1108+ {
1109+ "name": "adj_x2",
1110+ "dtype": "bool",
1111+ "value": false
1112+ },
1113+ {
1114+ "name": "offset_x",
1115+ "dtype": "int",
1116+ "value": 0
1117+ },
1118+ {
1119+ "name": "enable_hf32",
1120+ "dtype": "bool",
1121+ "value": false
1122+ }
1123+ ]
1124+ },
1125+ {
1126+ "bin_filename": "BatchMatMulV3_ND_NZ_ND_ND_BF16_BF16_FP32_FP32",
1127+ "simplified_key": "diy,2/29/2/2/27/27/0/0",
1128+ "inputs": [
1129+ {
1130+ "name": "x1",
1131+ "index": 0,
1132+ "dtype": "bfloat16",
1133+ "format": "ND",
1134+ "paramType": "required",
1135+ "shape": [
1136+ -2
1137+ ]
1138+ },
1139+ {
1140+ "name": "x2",
1141+ "index": 1,
1142+ "dtype": "bfloat16",
1143+ "format": "FRACTAL_NZ",
1144+ "paramType": "required",
1145+ "shape": [
1146+ -2
1147+ ]
1148+ },
1149+ {
1150+ "name": "bias",
1151+ "index": 2,
1152+ "dtype": "float32",
1153+ "format": "ND",
1154+ "paramType": "optional",
1155+ "shape": [
1156+ -2
1157+ ]
1158+ },
1159+ {
1160+ "name": "offset_w",
1161+ "index": 3,
1162+ "dtype": "int8",
1163+ "format": "ND",
1164+ "paramType": "optional",
1165+ "shape": [
1166+ -2
1167+ ]
1168+ }
1169+ ],
1170+ "outputs": [
1171+ {
1172+ "name": "y",
1173+ "index": 0,
1174+ "dtype": "float32",
1175+ "format": "ND",
1176+ "paramType": "required",
1177+ "shape": [
1178+ -2
1179+ ]
1180+ }
1181+ ],
1182+ "attrs": [
1183+ {
1184+ "name": "adj_x1",
1185+ "dtype": "bool",
1186+ "value": false
1187+ },
1188+ {
1189+ "name": "adj_x2",
1190+ "dtype": "bool",
1191+ "value": false
1192+ },
1193+ {
1194+ "name": "offset_x",
1195+ "dtype": "int",
1196+ "value": 0
1197+ },
1198+ {
1199+ "name": "enable_hf32",
1200+ "dtype": "bool",
1201+ "value": false
1202+ }
1203+ ]
724 }1204 }
725 ]1205 ]
726}1206}
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_baddbmm.cpp+15-3
@@ -270,7 +270,11 @@ public:
270 convOut = bmmOut;270 convOut = bmmOut;
271 return ACLNN_SUCCESS;271 return ACLNN_SUCCESS;
272 }272 }
273- if (CheckGemmV3WithAlphaBeta(bias, matA, matB, cubeMathType)) {273+ bool enable16In32Out = NeedEnableFp32Output(
274+ matA->GetDataType(), matB->GetDataType(), output->GetDataType(), cubeMathType);
275+ bool needBroadcast = CheckAddmmTensorShapeNeedBroadcast(matA, matB, bias);
276+ // A2/A3上对于 16in32out,且不需要broadcast场景 直接走gemmV3
277+ if (CheckGemmV3WithAlphaBeta(bias, matA, matB, cubeMathType) || (enable16In32Out && !needBroadcast)) {
274 const aclTensor* bmmOut = ExecGemmV3WithAlphaBetaOp(bias, matA, matB, alpha, beta, executor);278 const aclTensor* bmmOut = ExecGemmV3WithAlphaBetaOp(bias, matA, matB, alpha, beta, executor);
275 CHECK_RET(bmmOut != nullptr, ACLNN_ERR_INNER_NULLPTR);279 CHECK_RET(bmmOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
276 convOut = bmmOut;280 convOut = bmmOut;
@@ -279,7 +283,12 @@ public:
279 // self(bias) * beta283 // self(bias) * beta
280 const aclTensor* selfContiguous = l0op::Contiguous(bias, executor);284 const aclTensor* selfContiguous = l0op::Contiguous(bias, executor);
281 CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);285 CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
282- const aclTensor* mulOut = l0op::Muls(selfContiguous, beta->ToFloat(), executor);286+ const aclTensor* selfContiguousCast = selfContiguous;
287+ if (enable16In32Out) {
288+ selfContiguousCast = l0op::Cast(selfContiguous, output->GetDataType(), executor);
289+ CHECK_RET(selfContiguousCast != nullptr, ACLNN_ERR_INNER_NULLPTR);
290+ }
291+ const aclTensor* mulOut = l0op::Muls(selfContiguousCast, beta->ToFloat(), executor);
283 CHECK_RET(mulOut != nullptr, ACLNN_ERR_INNER_NULLPTR);292 CHECK_RET(mulOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
284 293 
285 // bmmOut = batch1(matA) @ batch2(matB)294 // bmmOut = batch1(matA) @ batch2(matB)
@@ -357,17 +366,20 @@ aclnnStatus aclnnBaddbmmGetWorkspaceSize(
357 // 参数检查366 // 参数检查
358 auto ret = CheckInputParams(self, batch1, batch2, beta, alpha, out, cubeMathType);367 auto ret = CheckInputParams(self, batch1, batch2, beta, alpha, out, cubeMathType);
359 CHECK_RET(ret == ACLNN_SUCCESS, ret);368 CHECK_RET(ret == ACLNN_SUCCESS, ret);
369+ CHECK_RET(Check16In32OutWithBiasValid(batch1->GetDataType(), batch2->GetDataType(), out->GetDataType(), self), ACLNN_ERR_PARAM_INVALID);
360 370 
361 // 创建OpExecutor371 // 创建OpExecutor
362 auto uniqueExecutor = CREATE_EXECUTOR();372 auto uniqueExecutor = CREATE_EXECUTOR();
363 CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);373 CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
364 374 
375+ const aclTensor* castOut = nullptr;
365 // 输入空Tensor处理方法，直接返回了一个空 Tensor，未做计算376 // 输入空Tensor处理方法，直接返回了一个空 Tensor，未做计算
366 if (isProcessEmptyTensor(batch1, batch2)) {377 if (isProcessEmptyTensor(batch1, batch2)) {
367 auto emptyOut = bmmProcessEmptyTensor(batch1, batch2, uniqueExecutor.get());378 auto emptyOut = bmmProcessEmptyTensor(batch1, batch2, uniqueExecutor.get());
368 CHECK_RET(emptyOut != nullptr, ACLNN_ERR_INNER_NULLPTR);379 CHECK_RET(emptyOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
369 380 
370- auto viewCopyResult = l0op::ViewCopy(emptyOut, out, uniqueExecutor.get());381+ castOut = l0op::Cast(emptyOut, out->GetDataType(), uniqueExecutor.get());
382+ auto viewCopyResult = l0op::ViewCopy(castOut, out, uniqueExecutor.get());
371 CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);383 CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
372 384 
373 // 固定写法，获取计算过程中需要使用的workspace大小385 // 固定写法，获取计算过程中需要使用的workspace大小
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_batch_matmul.cpp+2-2
@@ -204,11 +204,11 @@ bool CheckDtypeValidWeightNz(const aclTensor* self, const aclTensor* mat2, const
204 "batchmatmulweightnz is unsupported in this npu arch");204 "batchmatmulweightnz is unsupported in this npu arch");
205 return false;205 return false;
206 }206 }
207- bool enable3510Fp32Output = NeedEnable3510Fp32Output(207+ bool enable16In32Out = NeedEnableFp32Output(
208 self->GetDataType(), mat2->GetDataType(), out->GetDataType(), cubeMathType);208 self->GetDataType(), mat2->GetDataType(), out->GetDataType(), cubeMathType);
209 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);209 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);
210 OP_CHECK_DTYPE_NOT_SUPPORT(mat2, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);210 OP_CHECK_DTYPE_NOT_SUPPORT(mat2, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);
211- if (enable3510Fp32Output) {211+ if (enable16In32Out) {
212 OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST, return false);212 OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST, return false);
213 } else {213 } else {
214 OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);214 OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);
Mmatmul/batch_mat_mul_v3/tests/ut/op_host/test_aclnn_baddbmm.cpp+55-1
@@ -1040,4 +1040,58 @@ TEST_F(l2_baddbmm_test, ascend950_test_bmm_transpose_add2)
1040 auto ut_false_false = OP_API_UT(aclnnBaddbmm, INPUT(self, batch1, batch2, beta, alpha), OUTPUT(out), math_type);1040 auto ut_false_false = OP_API_UT(aclnnBaddbmm, INPUT(self, batch1, batch2, beta, alpha), OUTPUT(out), math_type);
1041 aclRet = ut_false_false.TestGetWorkspaceSize(&workspace_size);1041 aclRet = ut_false_false.TestGetWorkspaceSize(&workspace_size);
1042 EXPECT_EQ(aclRet, ACL_SUCCESS);1042 EXPECT_EQ(aclRet, ACL_SUCCESS);
1043-}1043+}
1044+ 
1045+TEST_F(l2_baddbmm_test, baddbmm_16in32out_fp16_fp16_fp32_use_hf32)
1046+{
1047+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1048+ auto self = TensorDesc({4, 16, 16}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(0, 2);
1049+ auto batch1 = TensorDesc({4, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(0, 2);
1050+ auto batch2 = TensorDesc({4, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(0, 2);
1051+ auto out = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1052+ auto beta = ScalarDesc(1.0f);
1053+ auto alpha = ScalarDesc(1.0f);
1054+ int8_t cubeMathType = USE_HF32;
1055+ 
1056+ auto ut = OP_API_UT(aclnnBaddbmm, INPUT(self, batch1, batch2, beta, alpha), OUTPUT(out), cubeMathType);
1057+ 
1058+ uint64_t workspace_size = 0;
1059+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1060+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1061+}
1062+ 
1063+TEST_F(l2_baddbmm_test, baddbmm_16in32out_bf16_bf16_fp32_use_hf32)
1064+{
1065+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1066+ auto self = TensorDesc({4, 16, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(0, 2);
1067+ auto batch1 = TensorDesc({4, 16, 32}, ACL_BF16, ACL_FORMAT_ND).ValueRange(0, 2);
1068+ auto batch2 = TensorDesc({4, 32, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(0, 2);
1069+ auto out = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1070+ auto beta = ScalarDesc(1.0f);
1071+ auto alpha = ScalarDesc(1.0f);
1072+ int8_t cubeMathType = USE_HF32;
1073+ 
1074+ auto ut = OP_API_UT(aclnnBaddbmm, INPUT(self, batch1, batch2, beta, alpha), OUTPUT(out), cubeMathType);
1075+ 
1076+ uint64_t workspace_size = 0;
1077+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1078+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1079+}
1080+ 
1081+TEST_F(l2_baddbmm_test, baddbmm_16in32out_fp16_fp16_fp32_use_fp32_add_bf16_out)
1082+{
1083+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1084+ auto self = TensorDesc({4, 16, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(0, 2);
1085+ auto batch1 = TensorDesc({4, 16, 32}, ACL_BF16, ACL_FORMAT_ND).ValueRange(0, 2);
1086+ auto batch2 = TensorDesc({4, 32, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(0, 2);
1087+ auto out = TensorDesc({4, 16, 16}, ACL_BF16, ACL_FORMAT_ND).Precision(0.001, 0.001);
1088+ auto beta = ScalarDesc(1.0f);
1089+ auto alpha = ScalarDesc(1.0f);
1090+ int8_t cubeMathType = USE_FP32_ADD;
1091+ 
1092+ auto ut = OP_API_UT(aclnnBaddbmm, INPUT(self, batch1, batch2, beta, alpha), OUTPUT(out), cubeMathType);
1093+ 
1094+ uint64_t workspace_size = 0;
1095+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1096+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1097+}
Mmatmul/batch_mat_mul_v3/tests/ut/op_host/test_aclnn_batch_matmul.cpp+273-0
@@ -14,6 +14,7 @@
14 14 
15#include "../../../op_host/op_api/aclnn_batch_matmul.h"15#include "../../../op_host/op_api/aclnn_batch_matmul.h"
16#include "opdev/platform.h"16#include "opdev/platform.h"
17+#include "op_api/op_api_def.h"
17#include "op_api_ut_common/op_api_ut.h"18#include "op_api_ut_common/op_api_ut.h"
18#include "op_api_ut_common/scalar_desc.h"19#include "op_api_ut_common/scalar_desc.h"
19#include "op_api_ut_common/tensor_desc.h"20#include "op_api_ut_common/tensor_desc.h"
@@ -1391,4 +1392,276 @@ TEST_F(l2_batch_matmul_test, batch_matmul_310_FP32_FP16_FP16FP32_KEEP_DTYPE)
1391 OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);1392 OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1392 aclRet = ut_false_false.TestGetWorkspaceSize(&workspace_size);1393 aclRet = ut_false_false.TestGetWorkspaceSize(&workspace_size);
1393 EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);1394 EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
1395+}
1396+ 
1397+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_fp16_fp16_fp32_keep_dtype_aligned)
1398+{
1399+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1400+ auto tensor_1_desc = TensorDesc({8, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1401+ auto tensor_2_desc = TensorDesc({8, 32, 64}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1402+ auto out_tensor_desc = TensorDesc({8, 16, 64}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.005, 0.005);
1403+ int8_t cube_math_type = KEEP_DTYPE;
1404+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1405+ 
1406+ uint64_t workspace_size = 0;
1407+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1408+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1409+}
1410+ 
1411+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_bf16_bf16_fp32_keep_dtype_aligned)
1412+{
1413+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1414+ auto tensor_1_desc = TensorDesc({8, 16, 32}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1415+ auto tensor_2_desc = TensorDesc({8, 32, 64}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1416+ auto out_tensor_desc = TensorDesc({8, 16, 64}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.005, 0.005);
1417+ int8_t cube_math_type = KEEP_DTYPE;
1418+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1419+ 
1420+ uint64_t workspace_size = 0;
1421+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1422+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1423+}
1424+ 
1425+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_fp16_fp16_fp32_keep_dtype_not_aligned)
1426+{
1427+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1428+ auto tensor_1_desc = TensorDesc({3, 5, 7}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1429+ auto tensor_2_desc = TensorDesc({3, 7, 9}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1430+ auto out_tensor_desc = TensorDesc({3, 5, 9}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1431+ int8_t cube_math_type = KEEP_DTYPE;
1432+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1433+ 
1434+ uint64_t workspace_size = 0;
1435+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1436+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1437+}
1438+ 
1439+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_bf16_bf16_fp32_keep_dtype_not_aligned)
1440+{
1441+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1442+ auto tensor_1_desc = TensorDesc({3, 5, 7}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1443+ auto tensor_2_desc = TensorDesc({3, 7, 9}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1444+ auto out_tensor_desc = TensorDesc({3, 5, 9}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1445+ int8_t cube_math_type = KEEP_DTYPE;
1446+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1447+ 
1448+ uint64_t workspace_size = 0;
1449+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1450+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1451+}
1452+ 
1453+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_fp16_fp16_fp32_use_hf32)
1454+{
1455+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1456+ auto tensor_1_desc = TensorDesc({4, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1457+ auto tensor_2_desc = TensorDesc({4, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1458+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1459+ int8_t cube_math_type = USE_HF32;
1460+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1461+ 
1462+ uint64_t workspace_size = 0;
1463+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1464+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1465+}
1466+ 
1467+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_bf16_bf16_fp32_use_hf32)
1468+{
1469+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1470+ auto tensor_1_desc = TensorDesc({4, 16, 32}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1471+ auto tensor_2_desc = TensorDesc({4, 32, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1472+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1473+ int8_t cube_math_type = USE_HF32;
1474+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1475+ 
1476+ uint64_t workspace_size = 0;
1477+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1478+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1479+}
1480+ 
1481+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_fp16_fp16_fp32_allow_fp32_down_precision)
1482+{
1483+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1484+ auto tensor_1_desc = TensorDesc({4, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1485+ auto tensor_2_desc = TensorDesc({4, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1486+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1487+ int8_t cube_math_type = ALLOW_FP32_DOWN_PRECISION;
1488+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1489+ 
1490+ uint64_t workspace_size = 0;
1491+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1492+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1493+}
1494+ 
1495+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_bf16_bf16_fp32_allow_fp32_down_precision)
1496+{
1497+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1498+ auto tensor_1_desc = TensorDesc({4, 16, 32}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1499+ auto tensor_2_desc = TensorDesc({4, 32, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1500+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1501+ int8_t cube_math_type = ALLOW_FP32_DOWN_PRECISION;
1502+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1503+ 
1504+ uint64_t workspace_size = 0;
1505+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1506+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1507+}
1508+ 
1509+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_fp16_fp16_fp32_large_batch)
1510+{
1511+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1512+ auto tensor_1_desc = TensorDesc({128, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1513+ auto tensor_2_desc = TensorDesc({128, 32, 64}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1514+ auto out_tensor_desc = TensorDesc({128, 16, 64}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.005, 0.005);
1515+ int8_t cube_math_type = KEEP_DTYPE;
1516+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1517+ 
1518+ uint64_t workspace_size = 0;
1519+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1520+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1521+}
1522+ 
1523+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_bf16_bf16_fp32_large_batch)
1524+{
1525+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1526+ auto tensor_1_desc = TensorDesc({128, 16, 32}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1527+ auto tensor_2_desc = TensorDesc({128, 32, 64}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1528+ auto out_tensor_desc = TensorDesc({128, 16, 64}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.005, 0.005);
1529+ int8_t cube_math_type = KEEP_DTYPE;
1530+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1531+ 
1532+ uint64_t workspace_size = 0;
1533+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1534+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1535+}
1536+ 
1537+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_fp16_fp16_fp32_batch_broadcast)
1538+{
1539+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1540+ auto tensor_1_desc = TensorDesc({4, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1541+ auto tensor_2_desc = TensorDesc({1, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1542+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1543+ int8_t cube_math_type = KEEP_DTYPE;
1544+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1545+ 
1546+ uint64_t workspace_size = 0;
1547+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1548+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1549+}
1550+ 
1551+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_bf16_bf16_fp32_batch_broadcast)
1552+{
1553+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1554+ auto tensor_1_desc = TensorDesc({4, 16, 32}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1555+ auto tensor_2_desc = TensorDesc({1, 32, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1556+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1557+ int8_t cube_math_type = KEEP_DTYPE;
1558+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1559+ 
1560+ uint64_t workspace_size = 0;
1561+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1562+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1563+}
1564+ 
1565+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_fp16_fp16_fp32_k1_special)
1566+{
1567+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1568+ auto tensor_1_desc = TensorDesc({4, 16, 1}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1569+ auto tensor_2_desc = TensorDesc({4, 1, 16}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1570+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1571+ int8_t cube_math_type = KEEP_DTYPE;
1572+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1573+ 
1574+ uint64_t workspace_size = 0;
1575+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1576+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1577+}
1578+ 
1579+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_bf16_bf16_fp32_k1_special)
1580+{
1581+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1582+ auto tensor_1_desc = TensorDesc({4, 16, 1}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1583+ auto tensor_2_desc = TensorDesc({4, 1, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1584+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1585+ int8_t cube_math_type = KEEP_DTYPE;
1586+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1587+ 
1588+ uint64_t workspace_size = 0;
1589+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1590+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1591+}
1592+ 
1593+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_fp16_fp16_fp32_empty_tensor)
1594+{
1595+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1596+ auto tensor_1_desc = TensorDesc({2, 0, 4}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1597+ auto tensor_2_desc = TensorDesc({2, 4, 5}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1598+ auto out_tensor_desc = TensorDesc({2, 0, 5}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.001, 0.001);
1599+ int8_t cube_math_type = KEEP_DTYPE;
1600+ auto ut = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc), OUTPUT(out_tensor_desc), cube_math_type);
1601+ 
1602+ uint64_t workspace_size = 0;
1603+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1604+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1605+}
1606+ 
1607+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_fp16_fp16_fp32_transpose)
1608+{
1609+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1610+ auto tensor_1_desc = TensorDesc({4, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1611+ auto tensor_1_desc_t = TensorDesc({4, 16, 32}, ACL_FLOAT16, ACL_FORMAT_ND,
1612+ {512, 1, 16}, 0, {4, 32, 16}).ValueRange(-2, 2);
1613+ auto tensor_2_desc = TensorDesc({4, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-2, 2);
1614+ auto tensor_2_desc_t = TensorDesc({4, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND,
1615+ {64, 1, 32}, 0, {4, 16, 32}).ValueRange(-2, 2);
1616+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.005, 0.005);
1617+ 
1618+ int8_t math_type = KEEP_DTYPE;
1619+ uint64_t workspace_size = 0;
1620+ aclnnStatus aclRet = 0;
1621+ auto ut_false_false = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc),
1622+ OUTPUT(out_tensor_desc), math_type);
1623+ aclRet = ut_false_false.TestGetWorkspaceSize(&workspace_size);
1624+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1625+ auto ut_true_false = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc_t, tensor_2_desc),
1626+ OUTPUT(out_tensor_desc), math_type);
1627+ aclRet = ut_true_false.TestGetWorkspaceSize(&workspace_size);
1628+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1629+ auto ut_false_true = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc_t),
1630+ OUTPUT(out_tensor_desc), math_type);
1631+ aclRet = ut_false_true.TestGetWorkspaceSize(&workspace_size);
1632+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1633+ auto ut_true_true = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc_t, tensor_2_desc_t),
1634+ OUTPUT(out_tensor_desc), math_type);
1635+ aclRet = ut_true_true.TestGetWorkspaceSize(&workspace_size);
1636+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1637+}
1638+ 
1639+TEST_F(l2_batch_matmul_test, batch_matmul_16in32out_bf16_bf16_fp32_transpose)
1640+{
1641+ op::SocVersionManager versionManager(op::SocVersion::ASCEND910B);
1642+ auto tensor_1_desc = TensorDesc({4, 16, 32}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1643+ auto tensor_1_desc_t = TensorDesc({4, 16, 32}, ACL_BF16, ACL_FORMAT_ND, {512, 1, 16}, 0, {4, 32, 16}).ValueRange(-2, 2);
1644+ auto tensor_2_desc = TensorDesc({4, 32, 16}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-2, 2);
1645+ auto tensor_2_desc_t = TensorDesc({4, 32, 16}, ACL_BF16, ACL_FORMAT_ND, {64, 1, 32}, 0, {4, 16, 32}).ValueRange(-2, 2);
1646+ auto out_tensor_desc = TensorDesc({4, 16, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.005, 0.005);
1647+ 
1648+ int8_t math_type = KEEP_DTYPE;
1649+ uint64_t workspace_size = 0;
1650+ aclnnStatus aclRet = 0;
1651+ auto ut_false_false = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc),
1652+ OUTPUT(out_tensor_desc), math_type);
1653+ aclRet = ut_false_false.TestGetWorkspaceSize(&workspace_size);
1654+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1655+ auto ut_true_false = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc_t, tensor_2_desc),
1656+ OUTPUT(out_tensor_desc), math_type);
1657+ aclRet = ut_true_false.TestGetWorkspaceSize(&workspace_size);
1658+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1659+ auto ut_false_true = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc, tensor_2_desc_t),
1660+ OUTPUT(out_tensor_desc), math_type);
1661+ aclRet = ut_false_true.TestGetWorkspaceSize(&workspace_size);
1662+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1663+ auto ut_true_true = OP_API_UT(aclnnBatchMatMul, INPUT(tensor_1_desc_t, tensor_2_desc_t),
1664+ OUTPUT(out_tensor_desc), math_type);
1665+ aclRet = ut_true_true.TestGetWorkspaceSize(&workspace_size);
1666+ EXPECT_EQ(aclRet, ACL_SUCCESS);
1394}1667}
Mmatmul/common/op_host/op_api/batch_matmul_util.cpp+40-18
@@ -386,30 +386,29 @@ const aclTensor* GetBatchMatmulOp(
386{386{
387 auto bmmOpOut = selfTransdata;387 auto bmmOpOut = selfTransdata;
388 auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();388 auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
389- bool enable3510Fp32Output = NeedEnable3510Fp32Output(389+ bool enable16In32Out = NeedEnableFp32Output(
390 matmulOpInfo.support_info.self_dtype, matmulOpInfo.support_info.mat2_dtype,390 matmulOpInfo.support_info.self_dtype, matmulOpInfo.support_info.mat2_dtype,
391 matmulOpInfo.support_info.output_dtype, KEEP_DTYPE, bias);391 matmulOpInfo.support_info.output_dtype, KEEP_DTYPE, bias);
392- if (CheckAscendCScenario(selfTransdata, mat2Transdata, bias, matmulOpInfo, adjX1, adjX2)) {392+ bool isNdNd = matmulOpInfo.support_info.self_format == ge::FORMAT_ND &&
393+ matmulOpInfo.support_info.mat2_format == ge::FORMAT_ND;
394+ bool bmm16In32OutRorA2 = npuArch == NpuArch::DAV_2201 && enable16In32Out && isNdNd;
395+ bool enable2201Baddbmm16In32Out = (matmulOpInfo.support_info.self_dtype == op::DataType::DT_FLOAT16 ||
396+ matmulOpInfo.support_info.self_dtype == op::DataType::DT_BF16) && isBaddbmm &&
397+ (npuArch == NpuArch::DAV_2201);
398+ if (CheckAscendCScenario(selfTransdata, mat2Transdata, bias, matmulOpInfo, adjX1, adjX2) || bmm16In32OutRorA2) {
393 if ((npuArch == NpuArch::DAV_3510) && // 1.多维*2维(左非转置)2.多维*多维batch为1399 if ((npuArch == NpuArch::DAV_3510) && // 1.多维*2维(左非转置)2.多维*多维batch为1
394 (GetBatchDimAll(mat2Transdata) <= 1 &&400 (GetBatchDimAll(mat2Transdata) <= 1 &&
395 (!adjX1 || GetBatchDimAll(selfTransdata) <= 1))) {401 (!adjX1 || GetBatchDimAll(selfTransdata) <= 1))) {
396 int64_t opImplModeEnumV3 = matmulOpInfo.enableHf32 ? 0x40 : (matmulOpInfo.enableForceGrpAccForFp32 ? 0x4 : 0x1);402 int64_t opImplModeEnumV3 = matmulOpInfo.enableHf32 ? 0x40 : (matmulOpInfo.enableForceGrpAccForFp32 ? 0x4 : 0x1);
397 return TransBmm2Mm(403 return TransBmm2Mm(
398 selfTransdata, mat2Transdata, bias, opImplModeEnumV3, adjX1, adjX2, offsetX,404 selfTransdata, mat2Transdata, bias, opImplModeEnumV3, adjX1, adjX2, offsetX,
399- enable3510Fp32Output, executor);405+ enable16In32Out, executor);
400 }406 }
401 OP_LOGI("Hit batch_mat_mul_v3 scenario.");407 OP_LOGI("Hit batch_mat_mul_v3 scenario.");
402- if ((matmulOpInfo.support_info.self_dtype == op::DataType::DT_FLOAT16 ||408+ if (enable16In32Out) {
403- matmulOpInfo.support_info.self_dtype == op::DataType::DT_BF16) &&
404- isBaddbmm &&
405- (npuArch == NpuArch::DAV_2201)) {
406 OP_LOGI("Hit batch_mat_mul_v3 fp16/bf16 in - fp32 out scenario.");409 OP_LOGI("Hit batch_mat_mul_v3 fp16/bf16 in - fp32 out scenario.");
407 bmmOpOut = l0op::BatchMatMulV3NdFp16Bf162Fp32(410 bmmOpOut = l0op::BatchMatMulV3NdFp16Bf162Fp32(
408 selfTransdata, mat2Transdata, bias, nullptr, adjX1, adjX2, offsetX, matmulOpInfo.enableHf32, executor);411 selfTransdata, mat2Transdata, bias, nullptr, adjX1, adjX2, offsetX, matmulOpInfo.enableHf32, executor);
409- } else if (enable3510Fp32Output) {
410- OP_LOGI("Hit batch_mat_mul_v3 fp16/bf16 in - fp32 out 3510 scenario.");
411- bmmOpOut = l0op::BatchMatMulV3NdFp16Bf162Fp32(
412- selfTransdata, mat2Transdata, bias, nullptr, adjX1, adjX2, offsetX, matmulOpInfo.enableHf32, executor);
413 } else {412 } else {
414 OP_LOGI("Hit BatchMatMulV3Nd scenario.");413 OP_LOGI("Hit BatchMatMulV3Nd scenario.");
415 bmmOpOut = l0op::BatchMatMulV3Nd(414 bmmOpOut = l0op::BatchMatMulV3Nd(
@@ -418,8 +417,9 @@ const aclTensor* GetBatchMatmulOp(
418 return bmmOpOut;417 return bmmOpOut;
419 }418 }
420 // 输入是FP16的场景419 // 输入是FP16的场景
421- if (matmulOpInfo.support_info.self_dtype == op::DataType::DT_FLOAT16) {420+ if (matmulOpInfo.support_info.self_dtype == op::DataType::DT_FLOAT16 ||
422- if (matmulOpInfo.support_info.output_dtype == op::DataType::DT_FLOAT16) {421+ matmulOpInfo.support_info.self_dtype == op::DataType::DT_BF16) {
422+ if (matmulOpInfo.support_info.output_dtype != op::DataType::DT_FLOAT) {
423 // 输入是FP16, 输出是FP16的场景423 // 输入是FP16, 输出是FP16的场景
424 if (matmulOpInfo.support_info.self_format == op::Format::FORMAT_ND) {424 if (matmulOpInfo.support_info.self_format == op::Format::FORMAT_ND) {
425 bmmOpOut = l0op::BatchMatMulNd(425 bmmOpOut = l0op::BatchMatMulNd(
@@ -443,7 +443,7 @@ const aclTensor* GetBatchMatmulOp(
443 }443 }
444 }444 }
445 } else {445 } else {
446- // 输入是FP32/BF16,输出是FP32/BF16的场景446+ // 输入是FP32,输出是FP32/BF16的场景
447 bmmOpOut = l0op::BatchMatMulNd(447 bmmOpOut = l0op::BatchMatMulNd(
448 selfTransdata, mat2Transdata, bias, nullptr, adjX1, adjX2, offsetX, matmulOpInfo.opImplModeEnum, executor);448 selfTransdata, mat2Transdata, bias, nullptr, adjX1, adjX2, offsetX, matmulOpInfo.opImplModeEnum, executor);
449 }449 }
@@ -820,11 +820,11 @@ bool CheckDtypeValidWeightNz(const aclTensor* self, const aclTensor* mat2, const
820 "batchmatmulweightnz is unsupported in this npu arch");820 "batchmatmulweightnz is unsupported in this npu arch");
821 return false;821 return false;
822 }822 }
823- bool enable3510Fp32Output = NeedEnable3510Fp32Output(823+ bool enable16In32Out = NeedEnableFp32Output(
824 self->GetDataType(), mat2->GetDataType(), out->GetDataType(), cubeMathType);824 self->GetDataType(), mat2->GetDataType(), out->GetDataType(), cubeMathType);
825 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);825 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);
826 OP_CHECK_DTYPE_NOT_SUPPORT(mat2, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);826 OP_CHECK_DTYPE_NOT_SUPPORT(mat2, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);
827- if (enable3510Fp32Output) {827+ if (enable16In32Out) {
828 OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST, return false);828 OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST, return false);
829 } else {829 } else {
830 OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);830 OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);
@@ -909,11 +909,22 @@ const aclTensor* ExecBatchMatmulOpWithBiasAndAttrs(
909 }909 }
910 910 
911 const aclTensor* bmmOpOut = nullptr;911 const aclTensor* bmmOpOut = nullptr;
912+ bool enable16In32Out = NeedEnableFp32Output(
913+ self->GetDataType(), mat2->GetDataType(), out->GetDataType(), cubeMathType, bias);
914+ auto selfCastfp32 = selfTransdata;
915+ auto mat2Castfp32 = mat2Transdata;
912 if (isTransposeMat2Contiguous) {916 if (isTransposeMat2Contiguous) {
913 bmmOpOut = GetBatchMatmulOp(917 bmmOpOut = GetBatchMatmulOp(
914 selfTransdata, mat2Transdata, bias, matmulOpInfo, adjX1, adjX2, 0, executor, isBaddbmm);918 selfTransdata, mat2Transdata, bias, matmulOpInfo, adjX1, adjX2, 0, executor, isBaddbmm);
915 } else if (ifKEqual1) {919 } else if (ifKEqual1) {
916- bmmOpOut = l0op::Mul(selfTransdata, mat2Transdata, executor);920+ if (enable16In32Out) {
921+ // 16in32out场景下升精度处理
922+ selfCastfp32 = l0op::Cast(selfTransdata, op::DataType::DT_FLOAT, executor);
923+ CHECK_RET(selfCastfp32 != nullptr, nullptr);
924+ mat2Castfp32 = l0op::Cast(mat2Transdata, op::DataType::DT_FLOAT, executor);
925+ CHECK_RET(mat2Castfp32 != nullptr, nullptr);
926+ }
927+ bmmOpOut = l0op::Mul(selfCastfp32, mat2Castfp32, executor);
917 } else {928 } else {
918 bmmOpOut = GetBatchMatmulOp(929 bmmOpOut = GetBatchMatmulOp(
919 selfTransdata, mat2Transdata, bias, matmulOpInfo, adjX1, adjX2, 0, executor, isBaddbmm);930 selfTransdata, mat2Transdata, bias, matmulOpInfo, adjX1, adjX2, 0, executor, isBaddbmm);
@@ -1213,8 +1224,19 @@ const aclTensor* ExecBatchMatmulOpWithBiasAndAttrsV2(
1213 }1224 }
1214 1225 
1215 const aclTensor* bmmOpOut = nullptr;1226 const aclTensor* bmmOpOut = nullptr;
1227+ bool enable16In32Out = NeedEnableFp32Output(
1228+ self->GetDataType(), mat2->GetDataType(), out->GetDataType(), cubeMathType, bias);
1229+ auto selfCastfp32 = selfTransdata;
1230+ auto mat2Castfp32 = mat2Transdata;
1216 if (ifKEqual1) {1231 if (ifKEqual1) {
1217- bmmOpOut = l0op::Mul(selfTransdata, mat2Transdata, executor);1232+ if (enable16In32Out) {
1233+ // 16in32out场景下升精度处理
1234+ selfCastfp32 = l0op::Cast(selfTransdata, op::DataType::DT_FLOAT, executor);
1235+ CHECK_RET(selfCastfp32 != nullptr, nullptr);
1236+ mat2Castfp32 = l0op::Cast(mat2Transdata, op::DataType::DT_FLOAT, executor);
1237+ CHECK_RET(mat2Castfp32 != nullptr, nullptr);
1238+ }
1239+ bmmOpOut = l0op::Mul(selfCastfp32, mat2Castfp32, executor);
1218 } else {1240 } else {
1219 bmmOpOut = GetBatchMatmulOp(1241 bmmOpOut = GetBatchMatmulOp(
1220 selfTransdata, mat2Transdata, contiguousBias, matmulOpInfo, adjX1, adjX2, 0, executor, isBaddbmm);1242 selfTransdata, mat2Transdata, contiguousBias, matmulOpInfo, adjX1, adjX2, 0, executor, isBaddbmm);
Mmatmul/common/op_host/op_api/cube_util.cpp+36-29
@@ -134,6 +134,30 @@ bool CheckCubeMathTypeForMm(const op::DataType cubeTensorDtype, int8_t cubeMathT
134 }134 }
135}135}
136 136 
137+bool CheckAddmmTensorShapeNeedBroadcast(const aclTensor* mat1, const aclTensor* mat2, const aclTensor* self)
138+{
139+ uint64_t dimNum = mat1->GetViewShape().GetDimNum();
140+ uint64_t selfDimNum = self->GetViewShape().GetDimNum();
141+ if (dimNum != selfDimNum) {
142+ OP_LOGI("self's dimnum != matmul out's dimnum.");
143+ return true;
144+ }
145+ const op::Shape selfShape = self->GetViewShape();
146+ const op::Shape mat1Shape = mat1->GetViewShape();
147+ const op::Shape mat2Shape = mat2->GetViewShape();
148+ if (dimNum == 3UL) {
149+ OP_CHECK(selfShape[0] == mat1Shape[0] && selfShape[1] == mat1Shape[1] && selfShape[2] == mat2Shape[2],
150+ OP_LOGI("self shape not equal to matmul out shape."),
151+ return true);
152+ } else if (dimNum == 2UL) {
153+ OP_CHECK(selfShape[0] == mat1Shape[0] && selfShape[1] == mat2Shape[1],
154+ OP_LOGI("self shape not equal to matmul out shape."),
155+ return true);
156+ }
157+ return false;
158+}
159+ 
160+ 
137bool CheckCubeMathTypeForAddMm(161bool CheckCubeMathTypeForAddMm(
138 const aclTensor* mat1, const aclTensor* mat2, const aclTensor* self, const aclTensor* out, int8_t cubeMathType)162 const aclTensor* mat1, const aclTensor* mat2, const aclTensor* self, const aclTensor* out, int8_t cubeMathType)
139{163{
@@ -146,43 +170,26 @@ bool CheckCubeMathTypeForAddMm(
146 cubeMathType);170 cubeMathType);
147 return false;171 return false;
148 }172 }
173+ 
174+ auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
149 if (cubeMathType != USE_FP32_ADD) {175 if (cubeMathType != USE_FP32_ADD) {
150 return true;176 return true;
151 }177 }
152 // 平台校验178 // 平台校验
153- if (GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_2201 &&179+ if (npuArch != NpuArch::DAV_2201 && npuArch != NpuArch::DAV_3510) {
154- GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_3510) {180+ OP_LOGE(
155- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "current platform not support cubeMathType = 4: USE_FP32_ADD.");181+ ACLNN_ERR_PARAM_INVALID,
182+ "current platform not support cubeMathType = 4: USE_FP32_ADD.");
156 return false;183 return false;
157 }184 }
158 // A2平台上，当cubeMathType=USE_FP32_ADD时，当前不支持self与mmout broadcast185 // A2平台上，当cubeMathType=USE_FP32_ADD时，当前不支持self与mmout broadcast
159- uint64_t dimNum = mat1->GetViewShape().GetDimNum();186+ if (npuArch == NpuArch::DAV_2201) {
160- uint64_t selfDimNum = self->GetViewShape().GetDimNum();187+ bool needBroadcast = CheckAddmmTensorShapeNeedBroadcast(mat1, mat2, self);
161- if (GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201) {188+ OP_CHECK(!needBroadcast,
162- if (dimNum != selfDimNum) {189+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
163- OP_LOGE(190+ "when cubeMathType = 4:USE_FP32_ADD, do not support broadcast between self and mmout."),
164- ACLNN_ERR_PARAM_INVALID,
165- "when cubeMathType = 4:USE_FP32_ADD, do not support self's dimnum != matmul out's dimnum.");
166 return false;191 return false;
167- }192+ );
168- const op::Shape selfShape = self->GetViewShape();
169- const op::Shape mat1Shape = mat1->GetViewShape();
170- const op::Shape mat2Shape = mat2->GetViewShape();
171- if (dimNum == 3UL) {
172- OP_CHECK(
173- selfShape[0] == mat1Shape[0] && selfShape[1] == mat1Shape[1] && selfShape[2] == mat2Shape[2],
174- OP_LOGE(
175- ACLNN_ERR_PARAM_INVALID,
176- "when cubeMathType = 4:USE_FP32_ADD, self shape should equal to matmul out shape."),
177- return false);
178- } else if (dimNum == 2UL) {
179- OP_CHECK(
180- selfShape[0] == mat1Shape[0] && selfShape[1] == mat2Shape[1],
181- OP_LOGE(
182- ACLNN_ERR_PARAM_INVALID,
183- "when cubeMathType = 4:USE_FP32_ADD, self shape should equal to matmul out shape."),
184- return false);
185- }
186 }193 }
187 return true;194 return true;
188}195}
Mmatmul/common/op_host/op_api/cube_util.h+3-0
@@ -26,6 +26,9 @@ bool CheckCubeMathType(const op::DataType cubeTensorDtype, int8_t cubeMathType);
26// 校验针对mm算子 tensor的dtype，cubeMathType的值是否符合预期26// 校验针对mm算子 tensor的dtype，cubeMathType的值是否符合预期
27bool CheckCubeMathTypeForMm(const op::DataType cubeTensorDtype, int8_t cubeMathType);27bool CheckCubeMathTypeForMm(const op::DataType cubeTensorDtype, int8_t cubeMathType);
28 28 
29+// 校验针对addmm算子的输入shape检验其是否需要广播，是否可以广播需要预先校验
30+bool CheckAddmmTensorShapeNeedBroadcast(const aclTensor* mat1, const aclTensor* mat2, const aclTensor* self);
31+ 
29// 校验针对Addmm算子的cubeMathType和平台是否符合预期32// 校验针对Addmm算子的cubeMathType和平台是否符合预期
30bool CheckCubeMathTypeForAddMm(const aclTensor* mat1, const aclTensor* mat2, const aclTensor* self,33bool CheckCubeMathTypeForAddMm(const aclTensor* mat1, const aclTensor* mat2, const aclTensor* self,
31 const aclTensor* out, int8_t cubeMathType);34 const aclTensor* out, int8_t cubeMathType);
Mmatmul/common/op_host/op_api/matmul.cpp+6-4
@@ -190,7 +190,7 @@ const aclTensor* GemmV3Nd(
190 return gemmOut;190 return gemmOut;
191};191};
192 192 
193-const aclTensor* GemmV3NdWithAlphaBeta(const aclTensor* x1,193+const aclTensor* GemmV3NdNzWithAlphaBeta(const aclTensor* x1,
194 const aclTensor* x2,194 const aclTensor* x2,
195 const aclTensor* self,195 const aclTensor* self,
196 float alpha,196 float alpha,
@@ -198,10 +198,12 @@ const aclTensor* GemmV3NdWithAlphaBeta(const aclTensor* x1,
198 bool transposeX1,198 bool transposeX1,
199 bool transposeX2,199 bool transposeX2,
200 bool enableHf32,200 bool enableHf32,
201- aclOpExecutor* executor)201+ aclOpExecutor* executor,
202+ bool enable16In32Out)
202{203{
203- L0_DFX(GemmV3NdWithAlphaBeta, x1, x2, self, alpha, beta, transposeX1, transposeX2, enableHf32);204+ L0_DFX(GemmV3NdNzWithAlphaBeta, x1, x2, self, alpha, beta, transposeX1, transposeX2, enableHf32);
204- aclTensor* out = executor->AllocTensor(self->GetDataType(), Format::FORMAT_ND, Format::FORMAT_ND);205+ op::DataType outDtype = enable16In32Out ? DataType::DT_FLOAT : self->GetDataType();
206+ aclTensor* out = executor->AllocTensor(outDtype, Format::FORMAT_ND, Format::FORMAT_ND);
205 auto ret = INFER_SHAPE(207 auto ret = INFER_SHAPE(
206 GemmV3, OP_INPUT(x1, x2, self), OP_OUTPUT(out), OP_ATTR(alpha, beta, transposeX1, transposeX2, enableHf32));208 GemmV3, OP_INPUT(x1, x2, self), OP_OUTPUT(out), OP_ATTR(alpha, beta, transposeX1, transposeX2, enableHf32));
207 if (ret != ACLNN_SUCCESS) {209 if (ret != ACLNN_SUCCESS) {
Mmatmul/common/op_host/op_api/matmul.h+11-11
@@ -63,16 +63,16 @@ const aclTensor* GemmV3Nd(
63 const aclTensor* x1, const aclTensor* x2, const aclTensor* c, bool transposeX1, bool transposeX2, bool enableHf32,63 const aclTensor* x1, const aclTensor* x2, const aclTensor* c, bool transposeX1, bool transposeX2, bool enableHf32,
64 aclOpExecutor* executor);64 aclOpExecutor* executor);
65 65 
66-// 输入ND66+// 输入ND/ NDNZ
67-const aclTensor* GemmV3NdWithAlphaBeta(const aclTensor* x1,67+const aclTensor* GemmV3NdNzWithAlphaBeta(const aclTensor* x1,
68- const aclTensor* x2,68+ const aclTensor* x2,
69- const aclTensor* self,69+ const aclTensor* self,
70- float alpha,70+ float alpha,
71- float beta,71+ float beta,
72- bool transposeX1,72+ bool transposeX1,
73- bool transposeX2,73+ bool transposeX2,
74- bool enableHf32,74+ bool enableHf32,
75- aclOpExecutor* executor);75+ aclOpExecutor* executor,
76- 76+ bool enable16In32Out = false);
77} // namespace l0op77} // namespace l0op
78 78 
Mmatmul/common/op_host/op_api/matmul_util.cpp+103-63
@@ -323,12 +323,12 @@ static aclnnStatus SetMatmulOpSupportInfo(
323 bool isNdNzIn =323 bool isNdNzIn =
324 self->GetStorageFormat() == Format::FORMAT_ND && mat2->GetStorageFormat() == Format::FORMAT_FRACTAL_NZ;324 self->GetStorageFormat() == Format::FORMAT_ND && mat2->GetStorageFormat() == Format::FORMAT_FRACTAL_NZ;
325 // 支持weightnz16进32出325 // 支持weightnz16进32出
326- bool enable3510Fp32Output = NeedEnable3510Fp32Output(326+ bool enable16In32Out = NeedEnableFp32Output(
327 self->GetDataType(), mat2->GetDataType(), mmOpInfo.support_info.output_dtype, cubeMathType);327 self->GetDataType(), mat2->GetDataType(), mmOpInfo.support_info.output_dtype, cubeMathType);
328 mmOpInfo.support_info.mat2_format = isNdNzIn ? Format::FORMAT_FRACTAL_NZ : mmOpInfo.support_info.mat2_format;328 mmOpInfo.support_info.mat2_format = isNdNzIn ? Format::FORMAT_FRACTAL_NZ : mmOpInfo.support_info.mat2_format;
329 mmOpInfo.support_info.output_dtype =329 mmOpInfo.support_info.output_dtype =
330 isNdNzIn ? mmOpInfo.support_info.mat2_dtype : mmOpInfo.support_info.output_dtype;330 isNdNzIn ? mmOpInfo.support_info.mat2_dtype : mmOpInfo.support_info.output_dtype;
331- if (enable3510Fp32Output) {331+ if (enable16In32Out) {
332 mmOpInfo.support_info.output_dtype = op::DataType::DT_FLOAT;332 mmOpInfo.support_info.output_dtype = op::DataType::DT_FLOAT;
333 }333 }
334 return ACLNN_SUCCESS;334 return ACLNN_SUCCESS;
@@ -472,25 +472,23 @@ static const aclTensor* GetMatMulOp(
472 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,
473 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)
474{474{
475- bool enableFp32Output = NeedEnable3510Fp32Output(475+ auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
476+ bool enable16In32Out = NeedEnableFp32Output(
476 mmOpInfo.support_info.self_dtype, mmOpInfo.support_info.mat2_dtype, mmOpInfo.support_info.output_dtype,477 mmOpInfo.support_info.self_dtype, mmOpInfo.support_info.mat2_dtype, mmOpInfo.support_info.output_dtype,
477 KEEP_DTYPE, bias);478 KEEP_DTYPE, bias);
479+ bool supportNdNz = mmOpInfo.support_info.self_format == ge::FORMAT_ND &&
480+ mmOpInfo.support_info.mat2_format == ge::FORMAT_FRACTAL_NZ;
481+ bool addmm16In32Out = enable16In32Out && (bias != nullptr || supportNdNz);
478 if (CheckMatmulV3Support(x1, x2, bias, mmOpInfo, transposeX1, transposeX2, opImplModeEnum) ||482 if (CheckMatmulV3Support(x1, x2, bias, mmOpInfo, transposeX1, transposeX2, opImplModeEnum) ||
479- (CheckMMV3NzNzNdSupport(mmOpInfo) && CheckSupportInfoFormatNzNzNd(mmOpInfo))) {483+ (CheckMMV3NzNzNdSupport(mmOpInfo) && CheckSupportInfoFormatNzNzNd(mmOpInfo)) || addmm16In32Out) {
480 OP_LOGI("Hit matmul_v3 scenario.");484 OP_LOGI("Hit matmul_v3 scenario.");
481- 485+
482- if (enableFp32Output) {486+ if ((enable16In32Out && npuArch == NpuArch::DAV_3510)) {
483 const aclTensor* mmOut =487 const aclTensor* mmOut =
484 l0op::MatMulV3NdFp162Fp32(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);488 l0op::MatMulV3NdFp162Fp32(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);
485 return mmOut;489 return mmOut;
486- } else if (490+ } else if (enable16In32Out) {
487- mmOpInfo.support_info.output_dtype == DataType::DT_FLOAT &&491+ if (CheckSupportInfoFormatNzNzNd(mmOpInfo) && bias == nullptr) {
488- ((mmOpInfo.support_info.mat2_dtype == DataType::DT_FLOAT16 &&
489- mmOpInfo.support_info.self_dtype == DataType::DT_FLOAT16) ||
490- (mmOpInfo.support_info.mat2_dtype == DataType::DT_BF16 &&
491- mmOpInfo.support_info.self_dtype == DataType::DT_BF16)) &&
492- bias == nullptr) {
493- if (CheckSupportInfoFormatNzNzNd(mmOpInfo)) {
494 OP_LOGD("check SocVersion, call MatMulV3NzNzNdFp162Fp32.");492 OP_LOGD("check SocVersion, call MatMulV3NzNzNdFp162Fp32.");
495 x1 = l0op::ReFormat(x1, op::Format::FORMAT_FRACTAL_NZ);493 x1 = l0op::ReFormat(x1, op::Format::FORMAT_FRACTAL_NZ);
496 x2 = l0op::ReFormat(x2, op::Format::FORMAT_FRACTAL_NZ);494 x2 = l0op::ReFormat(x2, op::Format::FORMAT_FRACTAL_NZ);
@@ -502,6 +500,7 @@ static const aclTensor* GetMatMulOp(
502 500 
503 const aclTensor* mmOut =501 const aclTensor* mmOut =
504 l0op::MatMulV3NdFp162Fp32(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);502 l0op::MatMulV3NdFp162Fp32(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);
503+ OP_LOGI("hit matmulv3 fp16/bp16 in fp32 out case.");
505 return mmOut;504 return mmOut;
506 }505 }
507 506 
@@ -521,14 +520,9 @@ static const aclTensor* GetMatMulOp(
521 const aclTensor* mmOut =520 const aclTensor* mmOut =
522 l0op::MatMulV3Nd(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);521 l0op::MatMulV3Nd(x1, x2, bias, transposeX1, transposeX2, offsetX, opImplModeEnum, executor);
523 return mmOut;522 return mmOut;
524- } else if (523+ } else if (enable16In32Out && bias == nullptr) {
525- mmOpInfo.support_info.output_dtype == DataType::DT_FLOAT &&
526- ((mmOpInfo.support_info.mat2_dtype == DataType::DT_FLOAT16 &&
527- mmOpInfo.support_info.self_dtype == DataType::DT_FLOAT16) ||
528- (mmOpInfo.support_info.mat2_dtype == DataType::DT_BF16 &&
529- mmOpInfo.support_info.self_dtype == DataType::DT_BF16)) &&
530- bias == nullptr) {
531 // This is Split K Mode; Check if MatMul using Nd in Nd Out524 // This is Split K Mode; Check if MatMul using Nd in Nd Out
525+ OP_LOGI("hit matmulv2 fp16/bp16 in fp32 out case.");
532 const aclTensor* mmOut =526 const aclTensor* mmOut =
533 (mmOpInfo.support_info.self_format == ge::FORMAT_ND &&527 (mmOpInfo.support_info.self_format == ge::FORMAT_ND &&
534 mmOpInfo.support_info.output_format == ge::FORMAT_ND) ?528 mmOpInfo.support_info.output_format == ge::FORMAT_ND) ?
@@ -791,11 +785,30 @@ static aclnnStatus SetMatmulOpSupportFormat(const aclTensor* self, const aclTens
791 785 
792namespace Ops {786namespace Ops {
793namespace NN {787namespace NN {
794-bool NeedEnable3510Fp32Output(788+bool Check16In32OutWithBiasValid(
789+ op::DataType selfDtype, op::DataType mat2Dtype, op::DataType outputDtype, const aclTensor* bias)
790+{
791+ bool isFp32Output = outputDtype == DataType::DT_FLOAT;
792+ if (!isFp32Output) {
793+ return true;
794+ }
795+ bool isLowPrecisionInput =
796+ (selfDtype == DataType::DT_FLOAT16 && mat2Dtype == DataType::DT_FLOAT16) ||
797+ (selfDtype == DataType::DT_BF16 && mat2Dtype == DataType::DT_BF16);
798+ if (isLowPrecisionInput && bias != nullptr) {
799+ // bias Dtype should match mat Dtype, or be DT_FLOAT
800+ op::DataType biasDataType = bias->GetDataType();
801+ return biasDataType == selfDtype || biasDataType == DataType::DT_FLOAT;
802+ }
803+ return true;
804+}
805+ 
806+bool NeedEnableFp32Output(
795 op::DataType selfDtype, op::DataType mat2Dtype, op::DataType outputDtype, int8_t cubeMathType,807 op::DataType selfDtype, op::DataType mat2Dtype, op::DataType outputDtype, int8_t cubeMathType,
796 const aclTensor* bias, bool isFusion)808 const aclTensor* bias, bool isFusion)
797{809{
798- if (op::GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_3510) {810+ auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
811+ if (npuArch != NpuArch::DAV_2201 && npuArch != NpuArch::DAV_3510) {
799 return false;812 return false;
800 }813 }
801 bool isSameLowPrecisionInput = (selfDtype == DataType::DT_FLOAT16 && mat2Dtype == DataType::DT_FLOAT16) ||814 bool isSameLowPrecisionInput = (selfDtype == DataType::DT_FLOAT16 && mat2Dtype == DataType::DT_FLOAT16) ||
@@ -803,9 +816,9 @@ bool NeedEnable3510Fp32Output(
803 if (!isSameLowPrecisionInput) {816 if (!isSameLowPrecisionInput) {
804 return false;817 return false;
805 }818 }
806- bool isLowPrecisionFp32Output = outputDtype == DataType::DT_FLOAT;819+ bool isFp32Output = outputDtype == DataType::DT_FLOAT;
807 bool isFusionAddMatmulFp32Output = isFusion && bias == nullptr && cubeMathType == USE_FP32_ADD;820 bool isFusionAddMatmulFp32Output = isFusion && bias == nullptr && cubeMathType == USE_FP32_ADD;
808- return isLowPrecisionFp32Output || isFusionAddMatmulFp32Output;821+ return isFp32Output || isFusionAddMatmulFp32Output;
809}822}
810 823 
811// 非连续条件的shape范围限制824// 非连续条件的shape范围限制
@@ -1028,7 +1041,8 @@ MmOpInfo GetMatmulOpInfo(
1028 mmOpInfo.ori_info.output_dtype = self->GetDataType();1041 mmOpInfo.ori_info.output_dtype = self->GetDataType();
1029 op::DataType outDtype = out == nullptr ? DataType::DT_UNDEFINED : out->GetDataType();1042 op::DataType outDtype = out == nullptr ? DataType::DT_UNDEFINED : out->GetDataType();
1030 if (FP16FP32_KEEP_DTYPE == cubeMathType ||1043 if (FP16FP32_KEEP_DTYPE == cubeMathType ||
1031- NeedEnable3510Fp32Output(self->GetDataType(), mat2->GetDataType(), outDtype, cubeMathType, bias, isFusion)) {1044+ NeedEnableFp32Output(
1045+ self->GetDataType(), mat2->GetDataType(), outDtype, cubeMathType, bias, isFusion)) {
1032 mmOpInfo.ori_info.output_dtype = DataType::DT_FLOAT;1046 mmOpInfo.ori_info.output_dtype = DataType::DT_FLOAT;
1033 }1047 }
1034 mmOpInfo.ori_info.output_format = op::Format::FORMAT_ND;1048 mmOpInfo.ori_info.output_format = op::Format::FORMAT_ND;
@@ -1348,12 +1362,24 @@ const aclTensor* ExecMmOpWithBias(
1348 OP_LOGI("Format of mat2 is mat2TransdataOut [%s].", op::ToString(mat2TransdataOut->GetStorageShape()).GetString());1362 OP_LOGI("Format of mat2 is mat2TransdataOut [%s].", op::ToString(mat2TransdataOut->GetStorageShape()).GetString());
1349 1363 
1350 const aclTensor* mmOut = nullptr;1364 const aclTensor* mmOut = nullptr;
1365+ op::DataType outDtype = out == nullptr ? DataType::DT_UNDEFINED : out->GetDataType();
1366+ bool enable16In32Out = NeedEnableFp32Output(
1367+ self->GetDataType(), mat2->GetDataType(), outDtype, cubeMathType, bias);
1368+ auto selfCastfp32 = selfTransdataOut;
1369+ auto mat2Castfp32 = mat2TransdataOut;
1351 if (isSelfSlice) {1370 if (isSelfSlice) {
1352 mmOut = GetMatMulOp(1371 mmOut = GetMatMulOp(
1353 selfTransdataOut, mat2TransdataOut, contiguousBias, mmOpInfo, mmOpInfo.shapeInfo.transposeX1,1372 selfTransdataOut, mat2TransdataOut, contiguousBias, mmOpInfo, mmOpInfo.shapeInfo.transposeX1,
1354 mmOpInfo.shapeInfo.transposeX2, 0, mmOpInfo.opImplModeEnum, executor);1373 mmOpInfo.shapeInfo.transposeX2, 0, mmOpInfo.opImplModeEnum, executor);
1355 } else if (ifKEqual1) {1374 } else if (ifKEqual1) {
1356- mmOut = l0op::Mul(selfTransdataOut, mat2TransdataOut, executor);1375+ if (enable16In32Out) {
1376+ // 16in32out场景下升精度处理
1377+ selfCastfp32 = l0op::Cast(selfTransdataOut, op::DataType::DT_FLOAT, executor);
1378+ CHECK_RET(selfCastfp32 != nullptr, nullptr);
1379+ mat2Castfp32 = l0op::Cast(mat2TransdataOut, op::DataType::DT_FLOAT, executor);
1380+ CHECK_RET(mat2Castfp32 != nullptr, nullptr);
1381+ }
1382+ mmOut = l0op::Mul(selfCastfp32, mat2Castfp32, executor);
1357 } else {1383 } else {
1358 mmOut = GetMatMulOp(1384 mmOut = GetMatMulOp(
1359 selfTransdataOut, mat2TransdataOut, contiguousBias, mmOpInfo, mmOpInfo.shapeInfo.transposeX1,1385 selfTransdataOut, mat2TransdataOut, contiguousBias, mmOpInfo, mmOpInfo.shapeInfo.transposeX1,
@@ -2024,7 +2050,7 @@ aclnnStatus SetMmSupportDType(MmOpInfo& mmOpInfo, int8_t cubeMathType)
2024 } else if (IsInputSupportFp32() && cubeMathType == USE_FP32_ADD && lowPrecisionInput) {2050 } else if (IsInputSupportFp32() && cubeMathType == USE_FP32_ADD && lowPrecisionInput) {
2025 mmOpInfo.support_info.output_dtype = DataType::DT_FLOAT;2051 mmOpInfo.support_info.output_dtype = DataType::DT_FLOAT;
2026 } else if (2052 } else if (
2027- npuArch == NpuArch::DAV_3510 && lowPrecisionInput && (mmOpInfo.ori_info.output_dtype == DataType::DT_FLOAT)) {2053+ IsInputSupportFp32()&& lowPrecisionInput && (mmOpInfo.ori_info.output_dtype == DataType::DT_FLOAT)) {
2028 mmOpInfo.support_info.output_dtype = DataType::DT_FLOAT;2054 mmOpInfo.support_info.output_dtype = DataType::DT_FLOAT;
2029 }2055 }
2030 return ACLNN_SUCCESS;2056 return ACLNN_SUCCESS;
@@ -2403,12 +2429,12 @@ aclnnStatus Dav2201MatMulRule::PromoteDtype(
2403 // 更新 kernel support outputDtype2429 // 更新 kernel support outputDtype
2404 mmOpInfo.support_info.output_dtype = UpdateOutputDtype(upperDtype, out->GetDataType(), cubeMathType);2430 mmOpInfo.support_info.output_dtype = UpdateOutputDtype(upperDtype, out->GetDataType(), cubeMathType);
2405 2431 
2406- bool enable3510Fp32Output = NeedEnable3510Fp32Output(2432+ bool enable16In32Out = NeedEnableFp32Output(
2407- matA->GetDataType(), matB->GetDataType(), out->GetDataType(), cubeMathType, bias, isFusion);2433+ matA->GetDataType(), matB->GetDataType(), out->GetDataType(), cubeMathType, bias, isFusion);
2408- if (enable3510Fp32Output) {2434+ if (enable16In32Out) {
2409- mmOpInfo.ori_info.output_dtype = op::DataType::DT_FLOAT;2435+ mmOpInfo.ori_info.output_dtype = op::DataType::DT_FLOAT;
2410- mmOpInfo.support_info.output_dtype = op::DataType::DT_FLOAT;2436+ mmOpInfo.support_info.output_dtype = op::DataType::DT_FLOAT;
2411- }2437+ }
2412 2438 
2413 // 更新 biasDtype2439 // 更新 biasDtype
2414 if (bias != nullptr) {2440 if (bias != nullptr) {
@@ -2614,38 +2640,51 @@ op::DataType DefaultMatMulRule::PromoteOutputAndBiasDtype(op::DataType outputDty
2614 return op::DataType::DT_FLOAT16;2640 return op::DataType::DT_FLOAT16;
2615}2641}
2616 2642 
2643+const aclTensor* TransposeAndContiguousMat(const aclTensor* mat, aclOpExecutor* executor)
2644+{
2645+ CHECK_RET(mat != nullptr, nullptr);
2646+ auto transpose = Ops::NN::IsTransposeLastTwoDims(mat);
2647+ // 转置并转连续
2648+ auto contiguousMat = mat;
2649+ auto matStorageShape = mat->GetStorageShape();
2650+ if (transpose) {
2651+ contiguousMat = executor->CreateView(mat, SwapLastTwoDimValue(mat->GetViewShape()), mat->GetViewOffset());
2652+ } else {
2653+ contiguousMat = l0op::Contiguous(mat, executor);
2654+ }
2655+ if (mat->GetStorageFormat() == op::Format::FORMAT_FRACTAL_NZ) {
2656+ OP_LOGI("mat GetStorageFormat FORMAT_FRACTAL_NZ.");
2657+ aclTensor* matShapeSet = const_cast<aclTensor*>(contiguousMat);
2658+ matShapeSet->SetStorageShape(matStorageShape); // 对NZ的场景用原来的stroageShape刷新
2659+ }
2660+ OP_LOGI("mat storage shape is [%s].", op::ToString(matStorageShape).GetString());
2661+ CHECK_RET(contiguousMat != nullptr, nullptr);
2662+ return contiguousMat;
2663+}
2664+ 
2617const aclTensor* ExecGemmV3WithAlphaBetaOp(const aclTensor* bias,2665const aclTensor* ExecGemmV3WithAlphaBetaOp(const aclTensor* bias,
2618 const aclTensor* self,2666 const aclTensor* self,
2619 const aclTensor* mat2,2667 const aclTensor* mat2,
2620 const aclScalar* alpha,2668 const aclScalar* alpha,
2621 const aclScalar* beta,2669 const aclScalar* beta,
2622- aclOpExecutor* executor)2670+ aclOpExecutor* executor,
2671+ bool enable16In32Out)
2623{2672{
2624- auto transposeSelf = Ops::NN::IsTransposeLastTwoDims(self);2673+ bool isNdNzInput = self->GetStorageFormat() == Format::FORMAT_ND &&
2625- auto transposeMat2 = Ops::NN::IsTransposeLastTwoDims(mat2);2674+ mat2->GetStorageFormat() == Format::FORMAT_FRACTAL_NZ;
2626 // reformat, 转成ND2675 // reformat, 转成ND
2627 auto reformatSelf = self;2676 auto reformatSelf = self;
2628 reformatSelf = l0op::ReFormat(self, op::Format::FORMAT_ND);2677 reformatSelf = l0op::ReFormat(self, op::Format::FORMAT_ND);
2629 // 左输入非连续转连续2678 // 左输入非连续转连续
2630- auto contiguousSelf = reformatSelf;2679+ auto contiguousSelf = TransposeAndContiguousMat(reformatSelf, executor);
2631- if (transposeSelf) {
2632- contiguousSelf = executor->CreateView(
2633- reformatSelf, SwapLastTwoDimValue(reformatSelf->GetViewShape()), reformatSelf->GetViewOffset());
2634- } else {
2635- contiguousSelf = l0op::Contiguous(reformatSelf, executor);
2636- }
2637 CHECK_RET(contiguousSelf != nullptr, nullptr);2680 CHECK_RET(contiguousSelf != nullptr, nullptr);
2638 // reformat, 转成ND2681 // reformat, 转成ND
2639 auto reformatMat2 = mat2;2682 auto reformatMat2 = mat2;
2640- reformatMat2 = l0op::ReFormat(mat2, op::Format::FORMAT_ND);2683+ if (!isNdNzInput) {
2641- // 右输入非连续转连续2684+ reformatMat2 = l0op::ReFormat(mat2, op::Format::FORMAT_ND);
2642- auto contiguousMat2 = reformatMat2;
2643- if (transposeMat2) {
2644- contiguousMat2 = executor->CreateView(
2645- reformatMat2, SwapLastTwoDimValue(reformatMat2->GetViewShape()), reformatMat2->GetViewOffset());
2646- } else {
2647- contiguousMat2 = l0op::Contiguous(reformatMat2, executor);
2648 }2685 }
2686+ // 右输入非连续转连续
2687+ auto contiguousMat2 = TransposeAndContiguousMat(reformatMat2, executor);;
2649 CHECK_RET(contiguousMat2 != nullptr, nullptr);2688 CHECK_RET(contiguousMat2 != nullptr, nullptr);
2650 2689 
2651 // bias非连续转连续2690 // bias非连续转连续
@@ -2656,17 +2695,18 @@ const aclTensor* ExecGemmV3WithAlphaBetaOp(const aclTensor* bias,
2656 contiguousBias = l0op::Contiguous(reformatBias, executor);2695 contiguousBias = l0op::Contiguous(reformatBias, executor);
2657 CHECK_RET(contiguousBias != nullptr, nullptr);2696 CHECK_RET(contiguousBias != nullptr, nullptr);
2658 2697 
2659- // 执行GemmV3NdWithAlphaBeta l0接口2698+ // 执行GemmV3WithAlphaBeta l0接口
2660- OP_LOGI("Entering l0op::GemmV3NdWithAlphaBeta.");2699+ OP_LOGI("Entering l0op::GemmV3NdNzWithAlphaBeta.");
2661- const aclTensor* gemmV3OpOut = l0op::GemmV3NdWithAlphaBeta(contiguousSelf,2700+ const aclTensor* gemmV3OpOut = l0op::GemmV3NdNzWithAlphaBeta(contiguousSelf,
2662- contiguousMat2,2701+ contiguousMat2,
2663- contiguousBias,2702+ contiguousBias,
2664- alpha->ToFloat(),2703+ alpha->ToFloat(),
2665- beta->ToFloat(),2704+ beta->ToFloat(),
2666- transposeSelf,2705+ false,
2667- transposeMat2,2706+ false,
2668- false,2707+ false,
2669- executor);2708+ executor,
2709+ enable16In32Out);
2670 return gemmV3OpOut;2710 return gemmV3OpOut;
2671}2711}
2672} // namespace NN2712} // namespace NN
Mmatmul/common/op_host/op_api/matmul_util.h+6-2
@@ -101,7 +101,10 @@ bool IsTransposeLastTwoDims(const aclTensor* tensor);
101bool CheckGemmV3Support(const aclTensor* mat1, const aclTensor* mat2, MmOpInfo& mmOpInfo,101bool CheckGemmV3Support(const aclTensor* mat1, const aclTensor* mat2, MmOpInfo& mmOpInfo,
102 int8_t cubeMathType);102 int8_t cubeMathType);
103 103 
104-bool NeedEnable3510Fp32Output(104+bool Check16In32OutWithBiasValid(
105+ op::DataType selfDtype, op::DataType mat2Dtype, op::DataType outputDtype, const aclTensor* bias);
106+ 
107+bool NeedEnableFp32Output(
105 op::DataType selfDtype, op::DataType mat2Dtype, op::DataType outputDtype, int8_t cubeMathType,108 op::DataType selfDtype, op::DataType mat2Dtype, op::DataType outputDtype, int8_t cubeMathType,
106 const aclTensor* bias = nullptr, bool isFusion = false);109 const aclTensor* bias = nullptr, bool isFusion = false);
107 110 
@@ -121,7 +124,8 @@ const aclTensor* ExecGemmV3WithAlphaBetaOp(const aclTensor* bias,
121 const aclTensor* mat2,124 const aclTensor* mat2,
122 const aclScalar* alpha,125 const aclScalar* alpha,
123 const aclScalar* beta,126 const aclScalar* beta,
124- aclOpExecutor* executor);127+ aclOpExecutor* executor,
128+ bool enable16In32Out = false);
125 129 
126const aclTensor* ExecMmOp(130const aclTensor* ExecMmOp(
127 const aclTensor* self, const aclTensor* mat2, const aclTensor* out, int8_t cubeMathType, aclOpExecutor* executor,131 const aclTensor* self, const aclTensor* mat2, const aclTensor* out, int8_t cubeMathType, aclOpExecutor* executor,
Mmatmul/gemm_v3/op_host/config/ascend910_93/gemm_v3_binary.json+450-0
@@ -227,6 +227,156 @@
227 }227 }
228 ]228 ]
229 },229 },
230+ {
231+ "bin_filename": "GemmV3_ND_ND_ND_ND_FP16_FP16_FP16_FP32",
232+ "simplified_key": "diy,2/2/2/2/1/1/1/0",
233+ "inputs": [
234+ {
235+ "name": "a",
236+ "index": 0,
237+ "dtype": "float16",
238+ "format": "ND",
239+ "paramType": "required",
240+ "shape": [
241+ -2
242+ ]
243+ },
244+ {
245+ "name": "b",
246+ "index": 1,
247+ "dtype": "float16",
248+ "format": "ND",
249+ "paramType": "required",
250+ "shape": [
251+ -2
252+ ]
253+ },
254+ {
255+ "name": "c",
256+ "index": 2,
257+ "dtype": "float16",
258+ "format": "ND",
259+ "paramType": "optional",
260+ "shape": [
261+ -2
262+ ]
263+ }
264+ ],
265+ "outputs": [
266+ {
267+ "name": "y",
268+ "index": 0,
269+ "dtype": "float32",
270+ "format": "ND",
271+ "paramType": "required",
272+ "shape": [
273+ -2
274+ ]
275+ }
276+ ],
277+ "attrs": [
278+ {
279+ "name": "alpha",
280+ "dtype": "float",
281+ "value": 1.0
282+ },
283+ {
284+ "name": "beta",
285+ "dtype": "float",
286+ "value": 1.0
287+ },
288+ {
289+ "name": "transpose_a",
290+ "dtype": "bool",
291+ "value": false
292+ },
293+ {
294+ "name": "transpose_b",
295+ "dtype": "bool",
296+ "value": false
297+ },
298+ {
299+ "name": "enable_hf32",
300+ "dtype": "bool",
301+ "value": false
302+ }
303+ ]
304+ },
305+ {
306+ "bin_filename": "GemmV3_ND_ND_ND_ND_BF16_BF16_BF16_FP32",
307+ "simplified_key": "diy,2/2/2/2/27/27/27/0",
308+ "inputs": [
309+ {
310+ "name": "a",
311+ "index": 0,
312+ "dtype": "bfloat16",
313+ "format": "ND",
314+ "paramType": "required",
315+ "shape": [
316+ -2
317+ ]
318+ },
319+ {
320+ "name": "b",
321+ "index": 1,
322+ "dtype": "bfloat16",
323+ "format": "ND",
324+ "paramType": "required",
325+ "shape": [
326+ -2
327+ ]
328+ },
329+ {
330+ "name": "c",
331+ "index": 2,
332+ "dtype": "bfloat16",
333+ "format": "ND",
334+ "paramType": "optional",
335+ "shape": [
336+ -2
337+ ]
338+ }
339+ ],
340+ "outputs": [
341+ {
342+ "name": "y",
343+ "index": 0,
344+ "dtype": "float32",
345+ "format": "ND",
346+ "paramType": "required",
347+ "shape": [
348+ -2
349+ ]
350+ }
351+ ],
352+ "attrs": [
353+ {
354+ "name": "alpha",
355+ "dtype": "float",
356+ "value": 1.0
357+ },
358+ {
359+ "name": "beta",
360+ "dtype": "float",
361+ "value": 1.0
362+ },
363+ {
364+ "name": "transpose_a",
365+ "dtype": "bool",
366+ "value": false
367+ },
368+ {
369+ "name": "transpose_b",
370+ "dtype": "bool",
371+ "value": false
372+ },
373+ {
374+ "name": "enable_hf32",
375+ "dtype": "bool",
376+ "value": false
377+ }
378+ ]
379+ },
230 {380 {
231 "bin_filename": "GemmV3_ND_ND_ND_ND_FP16_FP16_FP16_FP16",381 "bin_filename": "GemmV3_ND_ND_ND_ND_FP16_FP16_FP16_FP16",
232 "simplified_key": "diy,2/2/2/2/1/1/1/1",382 "simplified_key": "diy,2/2/2/2/1/1/1/1",
@@ -301,6 +451,306 @@
301 "value": false451 "value": false
302 }452 }
303 ]453 ]
454+ },
455+ {
456+ "bin_filename": "GemmV3_ND_NZ_ND_ND_FP16_FP16_FP32_FP32",
457+ "simplified_key": "diy,2/29/2/2/1/1/0/0",
458+ "inputs": [
459+ {
460+ "name": "a",
461+ "index": 0,
462+ "dtype": "float16",
463+ "format": "ND",
464+ "paramType": "required",
465+ "shape": [
466+ -2
467+ ]
468+ },
469+ {
470+ "name": "b",
471+ "index": 1,
472+ "dtype": "float16",
473+ "format": "FRACTAL_NZ",
474+ "paramType": "required",
475+ "shape": [
476+ -2
477+ ]
478+ },
479+ {
480+ "name": "c",
481+ "index": 2,
482+ "dtype": "float32",
483+ "format": "ND",
484+ "paramType": "optional",
485+ "shape": [
486+ -2
487+ ]
488+ }
489+ ],
490+ "outputs": [
491+ {
492+ "name": "y",
493+ "index": 0,
494+ "dtype": "float32",
495+ "format": "ND",
496+ "paramType": "required",
497+ "shape": [
498+ -2
499+ ]
500+ }
501+ ],
502+ "attrs": [
503+ {
504+ "name": "alpha",
505+ "dtype": "float",
506+ "value": 1.0
507+ },
508+ {
509+ "name": "beta",
510+ "dtype": "float",
511+ "value": 1.0
512+ },
513+ {
514+ "name": "transpose_a",
515+ "dtype": "bool",
516+ "value": false
517+ },
518+ {
519+ "name": "transpose_b",
520+ "dtype": "bool",
521+ "value": false
522+ },
523+ {
524+ "name": "enable_hf32",
525+ "dtype": "bool",
526+ "value": false
527+ }
528+ ]
529+ },
530+ {
531+ "bin_filename": "GemmV3_ND_NZ_ND_ND_BF16_BF16_FP32_FP32",
532+ "simplified_key": "diy,2/29/2/2/27/27/0/0",
533+ "inputs": [
534+ {
535+ "name": "a",
536+ "index": 0,
537+ "dtype": "bfloat16",
538+ "format": "ND",
539+ "paramType": "required",
540+ "shape": [
541+ -2
542+ ]
543+ },
544+ {
545+ "name": "b",
546+ "index": 1,
547+ "dtype": "bfloat16",
548+ "format": "FRACTAL_NZ",
549+ "paramType": "required",
550+ "shape": [
551+ -2
552+ ]
553+ },
554+ {
555+ "name": "c",
556+ "index": 2,
557+ "dtype": "float32",
558+ "format": "ND",
559+ "paramType": "optional",
560+ "shape": [
561+ -2
562+ ]
563+ }
564+ ],
565+ "outputs": [
566+ {
567+ "name": "y",
568+ "index": 0,
569+ "dtype": "float32",
570+ "format": "ND",
571+ "paramType": "required",
572+ "shape": [
573+ -2
574+ ]
575+ }
576+ ],
577+ "attrs": [
578+ {
579+ "name": "alpha",
580+ "dtype": "float",
581+ "value": 1.0
582+ },
583+ {
584+ "name": "beta",
585+ "dtype": "float",
586+ "value": 1.0
587+ },
588+ {
589+ "name": "transpose_a",
590+ "dtype": "bool",
591+ "value": false
592+ },
593+ {
594+ "name": "transpose_b",
595+ "dtype": "bool",
596+ "value": false
597+ },
598+ {
599+ "name": "enable_hf32",
600+ "dtype": "bool",
601+ "value": false
602+ }
603+ ]
604+ },
605+ {
606+ "bin_filename": "GemmV3_ND_NZ_ND_ND_FP16_FP16_FP16_FP32",
607+ "simplified_key": "diy,2/29/2/2/1/1/1/0",
608+ "inputs": [
609+ {
610+ "name": "a",
611+ "index": 0,
612+ "dtype": "float16",
613+ "format": "ND",
614+ "paramType": "required",
615+ "shape": [
616+ -2
617+ ]
618+ },
619+ {
620+ "name": "b",
621+ "index": 1,
622+ "dtype": "float16",
623+ "format": "FRACTAL_NZ",
624+ "paramType": "required",
625+ "shape": [
626+ -2
627+ ]
628+ },
629+ {
630+ "name": "c",
631+ "index": 2,
632+ "dtype": "float16",
633+ "format": "ND",
634+ "paramType": "optional",
635+ "shape": [
636+ -2
637+ ]
638+ }
639+ ],
640+ "outputs": [
641+ {
642+ "name": "y",
643+ "index": 0,
644+ "dtype": "float32",
645+ "format": "ND",
646+ "paramType": "required",
647+ "shape": [
648+ -2
649+ ]
650+ }
651+ ],
652+ "attrs": [
653+ {
654+ "name": "alpha",
655+ "dtype": "float",
656+ "value": 1.0
657+ },
658+ {
659+ "name": "beta",
660+ "dtype": "float",
661+ "value": 1.0
662+ },
663+ {
664+ "name": "transpose_a",
665+ "dtype": "bool",
666+ "value": false
667+ },
668+ {
669+ "name": "transpose_b",
670+ "dtype": "bool",
671+ "value": false
672+ },
673+ {
674+ "name": "enable_hf32",
675+ "dtype": "bool",
676+ "value": false
677+ }
678+ ]
679+ },
680+ {
681+ "bin_filename": "GemmV3_ND_NZ_ND_ND_BF16_BF16_BF16_FP32",
682+ "simplified_key": "diy,2/29/2/2/27/27/27/0",
683+ "inputs": [
684+ {
685+ "name": "a",
686+ "index": 0,
687+ "dtype": "bfloat16",
688+ "format": "ND",
689+ "paramType": "required",
690+ "shape": [
691+ -2
692+ ]
693+ },
694+ {
695+ "name": "b",
696+ "index": 1,
697+ "dtype": "bfloat16",
698+ "format": "FRACTAL_NZ",
699+ "paramType": "required",
700+ "shape": [
701+ -2
702+ ]
703+ },
704+ {
705+ "name": "c",
706+ "index": 2,
707+ "dtype": "bfloat16",
708+ "format": "ND",
709+ "paramType": "optional",
710+ "shape": [
711+ -2
712+ ]
713+ }
714+ ],
715+ "outputs": [
716+ {
717+ "name": "y",
718+ "index": 0,
719+ "dtype": "float32",
720+ "format": "ND",
721+ "paramType": "required",
722+ "shape": [
723+ -2
724+ ]
725+ }
726+ ],
727+ "attrs": [
728+ {
729+ "name": "alpha",
730+ "dtype": "float",
731+ "value": 1.0
732+ },
733+ {
734+ "name": "beta",
735+ "dtype": "float",
736+ "value": 1.0
737+ },
738+ {
739+ "name": "transpose_a",
740+ "dtype": "bool",
741+ "value": false
742+ },
743+ {
744+ "name": "transpose_b",
745+ "dtype": "bool",
746+ "value": false
747+ },
748+ {
749+ "name": "enable_hf32",
750+ "dtype": "bool",
751+ "value": false
752+ }
753+ ]
304 }754 }
305 ]755 ]
306}756}
Mmatmul/gemm_v3/op_host/config/ascend910b/gemm_v3_binary.json+450-0
@@ -227,6 +227,156 @@
227 }227 }
228 ]228 ]
229 },229 },
230+ {
231+ "bin_filename": "GemmV3_ND_ND_ND_ND_FP16_FP16_FP16_FP32",
232+ "simplified_key": "diy,2/2/2/2/1/1/1/0",
233+ "inputs": [
234+ {
235+ "name": "a",
236+ "index": 0,
237+ "dtype": "float16",
238+ "format": "ND",
239+ "paramType": "required",
240+ "shape": [
241+ -2
242+ ]
243+ },
244+ {
245+ "name": "b",
246+ "index": 1,
247+ "dtype": "float16",
248+ "format": "ND",
249+ "paramType": "required",
250+ "shape": [
251+ -2
252+ ]
253+ },
254+ {
255+ "name": "c",
256+ "index": 2,
257+ "dtype": "float16",
258+ "format": "ND",
259+ "paramType": "optional",
260+ "shape": [
261+ -2
262+ ]
263+ }
264+ ],
265+ "outputs": [
266+ {
267+ "name": "y",
268+ "index": 0,
269+ "dtype": "float32",
270+ "format": "ND",
271+ "paramType": "required",
272+ "shape": [
273+ -2
274+ ]
275+ }
276+ ],
277+ "attrs": [
278+ {
279+ "name": "alpha",
280+ "dtype": "float",
281+ "value": 1.0
282+ },
283+ {
284+ "name": "beta",
285+ "dtype": "float",
286+ "value": 1.0
287+ },
288+ {
289+ "name": "transpose_a",
290+ "dtype": "bool",
291+ "value": false
292+ },
293+ {
294+ "name": "transpose_b",
295+ "dtype": "bool",
296+ "value": false
297+ },
298+ {
299+ "name": "enable_hf32",
300+ "dtype": "bool",
301+ "value": false
302+ }
303+ ]
304+ },
305+ {
306+ "bin_filename": "GemmV3_ND_ND_ND_ND_BF16_BF16_BF16_FP32",
307+ "simplified_key": "diy,2/2/2/2/27/27/27/0",
308+ "inputs": [
309+ {
310+ "name": "a",
311+ "index": 0,
312+ "dtype": "bfloat16",
313+ "format": "ND",
314+ "paramType": "required",
315+ "shape": [
316+ -2
317+ ]
318+ },
319+ {
320+ "name": "b",
321+ "index": 1,
322+ "dtype": "bfloat16",
323+ "format": "ND",
324+ "paramType": "required",
325+ "shape": [
326+ -2
327+ ]
328+ },
329+ {
330+ "name": "c",
331+ "index": 2,
332+ "dtype": "bfloat16",
333+ "format": "ND",
334+ "paramType": "optional",
335+ "shape": [
336+ -2
337+ ]
338+ }
339+ ],
340+ "outputs": [
341+ {
342+ "name": "y",
343+ "index": 0,
344+ "dtype": "float32",
345+ "format": "ND",
346+ "paramType": "required",
347+ "shape": [
348+ -2
349+ ]
350+ }
351+ ],
352+ "attrs": [
353+ {
354+ "name": "alpha",
355+ "dtype": "float",
356+ "value": 1.0
357+ },
358+ {
359+ "name": "beta",
360+ "dtype": "float",
361+ "value": 1.0
362+ },
363+ {
364+ "name": "transpose_a",
365+ "dtype": "bool",
366+ "value": false
367+ },
368+ {
369+ "name": "transpose_b",
370+ "dtype": "bool",
371+ "value": false
372+ },
373+ {
374+ "name": "enable_hf32",
375+ "dtype": "bool",
376+ "value": false
377+ }
378+ ]
379+ },
230 {380 {
231 "bin_filename": "GemmV3_ND_ND_ND_ND_FP16_FP16_FP16_FP16",381 "bin_filename": "GemmV3_ND_ND_ND_ND_FP16_FP16_FP16_FP16",
232 "simplified_key": "diy,2/2/2/2/1/1/1/1",382 "simplified_key": "diy,2/2/2/2/1/1/1/1",
@@ -301,6 +451,306 @@
301 "value": false451 "value": false
302 }452 }
303 ]453 ]
454+ },
455+ {
456+ "bin_filename": "GemmV3_ND_NZ_ND_ND_FP16_FP16_FP32_FP32",
457+ "simplified_key": "diy,2/29/2/2/1/1/0/0",
458+ "inputs": [
459+ {
460+ "name": "a",
461+ "index": 0,
462+ "dtype": "float16",
463+ "format": "ND",
464+ "paramType": "required",
465+ "shape": [
466+ -2
467+ ]
468+ },
469+ {
470+ "name": "b",
471+ "index": 1,
472+ "dtype": "float16",
473+ "format": "FRACTAL_NZ",
474+ "paramType": "required",
475+ "shape": [
476+ -2
477+ ]
478+ },
479+ {
480+ "name": "c",
481+ "index": 2,
482+ "dtype": "float32",
483+ "format": "ND",
484+ "paramType": "optional",
485+ "shape": [
486+ -2
487+ ]
488+ }
489+ ],
490+ "outputs": [
491+ {
492+ "name": "y",
493+ "index": 0,
494+ "dtype": "float32",
495+ "format": "ND",
496+ "paramType": "required",
497+ "shape": [
498+ -2
499+ ]
500+ }
501+ ],
502+ "attrs": [
503+ {
504+ "name": "alpha",
505+ "dtype": "float",
506+ "value": 1.0
507+ },
508+ {
509+ "name": "beta",
510+ "dtype": "float",
511+ "value": 1.0
512+ },
513+ {
514+ "name": "transpose_a",
515+ "dtype": "bool",
516+ "value": false
517+ },
518+ {
519+ "name": "transpose_b",
520+ "dtype": "bool",
521+ "value": false
522+ },
523+ {
524+ "name": "enable_hf32",
525+ "dtype": "bool",
526+ "value": false
527+ }
528+ ]
529+ },
530+ {
531+ "bin_filename": "GemmV3_ND_NZ_ND_ND_BF16_BF16_FP32_FP32",
532+ "simplified_key": "diy,2/29/2/2/27/27/0/0",
533+ "inputs": [
534+ {
535+ "name": "a",
536+ "index": 0,
537+ "dtype": "bfloat16",
538+ "format": "ND",
539+ "paramType": "required",
540+ "shape": [
541+ -2
542+ ]
543+ },
544+ {
545+ "name": "b",
546+ "index": 1,
547+ "dtype": "bfloat16",
548+ "format": "FRACTAL_NZ",
549+ "paramType": "required",
550+ "shape": [
551+ -2
552+ ]
553+ },
554+ {
555+ "name": "c",
556+ "index": 2,
557+ "dtype": "float32",
558+ "format": "ND",
559+ "paramType": "optional",
560+ "shape": [
561+ -2
562+ ]
563+ }
564+ ],
565+ "outputs": [
566+ {
567+ "name": "y",
568+ "index": 0,
569+ "dtype": "float32",
570+ "format": "ND",
571+ "paramType": "required",
572+ "shape": [
573+ -2
574+ ]
575+ }
576+ ],
577+ "attrs": [
578+ {
579+ "name": "alpha",
580+ "dtype": "float",
581+ "value": 1.0
582+ },
583+ {
584+ "name": "beta",
585+ "dtype": "float",
586+ "value": 1.0
587+ },
588+ {
589+ "name": "transpose_a",
590+ "dtype": "bool",
591+ "value": false
592+ },
593+ {
594+ "name": "transpose_b",
595+ "dtype": "bool",
596+ "value": false
597+ },
598+ {
599+ "name": "enable_hf32",
600+ "dtype": "bool",
601+ "value": false
602+ }
603+ ]
604+ },
605+ {
606+ "bin_filename": "GemmV3_ND_NZ_ND_ND_FP16_FP16_FP16_FP32",
607+ "simplified_key": "diy,2/29/2/2/1/1/1/0",
608+ "inputs": [
609+ {
610+ "name": "a",
611+ "index": 0,
612+ "dtype": "float16",
613+ "format": "ND",
614+ "paramType": "required",
615+ "shape": [
616+ -2
617+ ]
618+ },
619+ {
620+ "name": "b",
621+ "index": 1,
622+ "dtype": "float16",
623+ "format": "FRACTAL_NZ",
624+ "paramType": "required",
625+ "shape": [
626+ -2
627+ ]
628+ },
629+ {
630+ "name": "c",
631+ "index": 2,
632+ "dtype": "float16",
633+ "format": "ND",
634+ "paramType": "optional",
635+ "shape": [
636+ -2
637+ ]
638+ }
639+ ],
640+ "outputs": [
641+ {
642+ "name": "y",
643+ "index": 0,
644+ "dtype": "float32",
645+ "format": "ND",
646+ "paramType": "required",
647+ "shape": [
648+ -2
649+ ]
650+ }
651+ ],
652+ "attrs": [
653+ {
654+ "name": "alpha",
655+ "dtype": "float",
656+ "value": 1.0
657+ },
658+ {
659+ "name": "beta",
660+ "dtype": "float",
661+ "value": 1.0
662+ },
663+ {
664+ "name": "transpose_a",
665+ "dtype": "bool",
666+ "value": false
667+ },
668+ {
669+ "name": "transpose_b",
670+ "dtype": "bool",
671+ "value": false
672+ },
673+ {
674+ "name": "enable_hf32",
675+ "dtype": "bool",
676+ "value": false
677+ }
678+ ]
679+ },
680+ {
681+ "bin_filename": "GemmV3_ND_NZ_ND_ND_BF16_BF16_BF16_FP32",
682+ "simplified_key": "diy,2/29/2/2/27/27/27/0",
683+ "inputs": [
684+ {
685+ "name": "a",
686+ "index": 0,
687+ "dtype": "bfloat16",
688+ "format": "ND",
689+ "paramType": "required",
690+ "shape": [
691+ -2
692+ ]
693+ },
694+ {
695+ "name": "b",
696+ "index": 1,
697+ "dtype": "bfloat16",
698+ "format": "FRACTAL_NZ",
699+ "paramType": "required",
700+ "shape": [
701+ -2
702+ ]
703+ },
704+ {
705+ "name": "c",
706+ "index": 2,
707+ "dtype": "bfloat16",
708+ "format": "ND",
709+ "paramType": "optional",
710+ "shape": [
711+ -2
712+ ]
713+ }
714+ ],
715+ "outputs": [
716+ {
717+ "name": "y",
718+ "index": 0,
719+ "dtype": "float32",
720+ "format": "ND",
721+ "paramType": "required",
722+ "shape": [
723+ -2
724+ ]
725+ }
726+ ],
727+ "attrs": [
728+ {
729+ "name": "alpha",
730+ "dtype": "float",
731+ "value": 1.0
732+ },
733+ {
734+ "name": "beta",
735+ "dtype": "float",
736+ "value": 1.0
737+ },
738+ {
739+ "name": "transpose_a",
740+ "dtype": "bool",
741+ "value": false
742+ },
743+ {
744+ "name": "transpose_b",
745+ "dtype": "bool",
746+ "value": false
747+ },
748+ {
749+ "name": "enable_hf32",
750+ "dtype": "bool",
751+ "value": false
752+ }
753+ ]
304 }754 }
305 ]755 ]
306}756}
Mmatmul/gemm_v3/op_host/gemm_v3_def.cpp+8-8
@@ -21,20 +21,20 @@ public:
21 {21 {
22 this->Input("a")22 this->Input("a")
23 .ParamType(REQUIRED)23 .ParamType(REQUIRED)
24- .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16})24+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_BF16})
25- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
26 this->Input("b")26 this->Input("b")
27 .ParamType(REQUIRED)27 .ParamType(REQUIRED)
28- .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16})28+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_BF16})
29- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});29+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ});
30 this->Input("c")30 this->Input("c")
31 .ParamType(OPTIONAL)31 .ParamType(OPTIONAL)
32- .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})32+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT})
33- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});33+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
34 this->Output("y")34 this->Output("y")
35 .ParamType(REQUIRED)35 .ParamType(REQUIRED)
36- .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})36+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
37- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});37+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
38 this->Attr("alpha").AttrType(REQUIRED).Float(1.0);38 this->Attr("alpha").AttrType(REQUIRED).Float(1.0);
39 this->Attr("beta").AttrType(REQUIRED).Float(1.0);39 this->Attr("beta").AttrType(REQUIRED).Float(1.0);
40 this->Attr("transpose_a").AttrType(REQUIRED).Bool(false);40 this->Attr("transpose_a").AttrType(REQUIRED).Bool(false);
Mmatmul/gemm_v3/op_host/gemm_v3_infershape.cpp+16-31
@@ -38,6 +38,7 @@ ge::graphStatus InferShapeForGemmV3(InferShapeContext* context)
38 auto shape_out = context->GetOutputShape(kOutputIdx);38 auto shape_out = context->GetOutputShape(kOutputIdx);
39 auto attrs = context->GetAttrs();39 auto attrs = context->GetAttrs();
40 auto dtype_c = context->GetInputDesc(kMatmulV2MinShapeSize)->GetDataType();40 auto dtype_c = context->GetInputDesc(kMatmulV2MinShapeSize)->GetDataType();
41+ auto dtype_out = context->GetOutputDesc(kOutputIdx)->GetDataType();
41 // 当前仅支持累加场景，shape_c不支持为空42 // 当前仅支持累加场景，shape_c不支持为空
42 OP_CHECK_IF(43 OP_CHECK_IF(
43 shape_a == nullptr || shape_b == nullptr || shape_c == nullptr || shape_out == nullptr || attrs == nullptr,44 shape_a == nullptr || shape_b == nullptr || shape_c == nullptr || shape_out == nullptr || attrs == nullptr,
@@ -48,12 +49,6 @@ ge::graphStatus InferShapeForGemmV3(InferShapeContext* context)
48 OP_CHECK_IF(49 OP_CHECK_IF(
49 alpha == nullptr || beta == nullptr, CUBE_INNER_ERR_REPORT(op_name, "gemmv3 attribute is null"),50 alpha == nullptr || beta == nullptr, CUBE_INNER_ERR_REPORT(op_name, "gemmv3 attribute is null"),
50 return ge::GRAPH_FAILED);51 return ge::GRAPH_FAILED);
51- if (dtype_c == ge::DT_FLOAT) {
52- // 校验alpha和beta是否等于1.0f
53- OP_CHECK_IF(std::abs(*alpha - 1.0f) > epsilon || std::abs(*beta - 1.0f) > epsilon,
54- CUBE_INNER_ERR_REPORT(op_name, "alpha and beta must be 1.0"),
55- return ge::GRAPH_FAILED);
56- }
57 const bool* trans_a = attrs->GetAttrPointer<bool>(2);52 const bool* trans_a = attrs->GetAttrPointer<bool>(2);
58 const bool* trans_b = attrs->GetAttrPointer<bool>(3);53 const bool* trans_b = attrs->GetAttrPointer<bool>(3);
59 const bool* enable_hf32 = attrs->GetAttrPointer<bool>(4);54 const bool* enable_hf32 = attrs->GetAttrPointer<bool>(4);
@@ -65,38 +60,28 @@ ge::graphStatus InferShapeForGemmV3(InferShapeContext* context)
65 "gemmv3 a_shape: %s, b_shape: %s, c_shape:%s, transpose_a: %d, transpose_b: %d, enable_hf32: %d",60 "gemmv3 a_shape: %s, b_shape: %s, c_shape:%s, transpose_a: %d, transpose_b: %d, enable_hf32: %d",
66 Shape2String(*shape_a).c_str(), Shape2String(*shape_b).c_str(), Shape2String(*shape_c).c_str(), *trans_a,61 Shape2String(*shape_a).c_str(), Shape2String(*shape_b).c_str(), Shape2String(*shape_c).c_str(), *trans_a,
67 *trans_b, *enable_hf32);62 *trans_b, *enable_hf32);
68- if (dtype_c == ge::DT_FLOAT) {63+ if (shape_a->GetDimNum() == kMatmulV2MinShapeSize) { // addmm校验
69 OP_CHECK_IF(64 OP_CHECK_IF(
70 (shape_a->GetDimNum() != kMatmulV2MinShapeSize || shape_b->GetDimNum() != kMatmulV2MinShapeSize ||65 (shape_a->GetDimNum() != kMatmulV2MinShapeSize || shape_b->GetDimNum() != kMatmulV2MinShapeSize ||
71 shape_c->GetDimNum() != kMatmulV2MinShapeSize),66 shape_c->GetDimNum() != kMatmulV2MinShapeSize),
72 CUBE_INNER_ERR_REPORT(67 CUBE_INNER_ERR_REPORT(
73- op_name, "gemmv3 input dim num[%zu] [%zu] [%zu]is not 2!", shape_a->GetDimNum(), shape_b->GetDimNum(),68+ op_name, "addmm input dim num[%zu] [%zu] [%zu]is not 2!", shape_a->GetDimNum(), shape_b->GetDimNum(),
69+ shape_c->GetDimNum()),
70+ return ge::GRAPH_FAILED);
71+ } else if (shape_a->GetDimNum() == kAddmmMaxShapeSize) { // baddbmm校验
72+ OP_CHECK_IF(
73+ (shape_a->GetDimNum() != kAddmmMaxShapeSize || shape_b->GetDimNum() != kAddmmMaxShapeSize ||
74+ shape_c->GetDimNum() != kAddmmMaxShapeSize),
75+ CUBE_INNER_ERR_REPORT(
76+ op_name, "baddbmm input x1、x2、self dim num[%zu] [%zu] [%zu]is not 3!", shape_a->GetDimNum(), shape_b->GetDimNum(),
74 shape_c->GetDimNum()),77 shape_c->GetDimNum()),
75 return ge::GRAPH_FAILED);78 return ge::GRAPH_FAILED);
76 } else {79 } else {
77- if (shape_a->GetDimNum() == kMatmulV2MinShapeSize) { // addmm校验80+ OP_CHECK_IF(
78- OP_CHECK_IF(81+ shape_a->GetDimNum() != kAddmmMaxShapeSize && shape_a->GetDimNum() != kMatmulV2MinShapeSize,
79- (shape_a->GetDimNum() != kMatmulV2MinShapeSize || shape_b->GetDimNum() != kMatmulV2MinShapeSize ||82+ CUBE_INNER_ERR_REPORT(
80- shape_c->GetDimNum() != kMatmulV2MinShapeSize),83+ op_name, "unsupport addmm/baddbmm input x1 dim num[%zu], it should be 2 or 3!", shape_a->GetDimNum()),
81- CUBE_INNER_ERR_REPORT(84+ return ge::GRAPH_FAILED);
82- op_name, "addmm input dim num[%zu] [%zu] [%zu]is not 2!", shape_a->GetDimNum(), shape_b->GetDimNum(),
83- shape_c->GetDimNum()),
84- return ge::GRAPH_FAILED);
85- } else if (shape_a->GetDimNum() == kAddmmMaxShapeSize) { // baddbmm校验
86- OP_CHECK_IF(
87- (shape_a->GetDimNum() != kAddmmMaxShapeSize || shape_b->GetDimNum() != kAddmmMaxShapeSize ||
88- shape_c->GetDimNum() != kAddmmMaxShapeSize),
89- CUBE_INNER_ERR_REPORT(
90- op_name, "baddbmm input x1、x2、self dim num[%zu] [%zu] [%zu]is not 3!", shape_a->GetDimNum(), shape_b->GetDimNum(),
91- shape_c->GetDimNum()),
92- return ge::GRAPH_FAILED);
93- } else {
94- OP_CHECK_IF(
95- shape_a->GetDimNum() != kAddmmMaxShapeSize && shape_a->GetDimNum() != kMatmulV2MinShapeSize,
96- CUBE_INNER_ERR_REPORT(
97- op_name, "unsupport addmm/baddbmm input x1 dim num[%zu], it should be 2 or 3!", shape_a->GetDimNum()),
98- return ge::GRAPH_FAILED);
99- }
100 }85 }
101 86
102 size_t idx_m = static_cast<size_t>(*trans_a);87 size_t idx_m = static_cast<size_t>(*trans_a);
Mmatmul/gemm_v3/op_host/op_tiling/gemm_v3_base_tiling.cpp+10-4
@@ -81,18 +81,24 @@ bool GemmV3BaseTiling::IsCapable()
81 hwInfo_.socVersion != platform_ascendc::SocVersion::ASCEND910_93) {81 hwInfo_.socVersion != platform_ascendc::SocVersion::ASCEND910_93) {
82 return false;82 return false;
83 }83 }
84+ ge::DataType dtypeA = context_->GetInputDesc(INDEX_A)->GetDataType();
84 ge::DataType dtypeC = context_->GetInputDesc(INDEX_C)->GetDataType();85 ge::DataType dtypeC = context_->GetInputDesc(INDEX_C)->GetDataType();
85 ge::DataType dtypeY = context_->GetOutputDesc(INDEX_Y)->GetDataType();86 ge::DataType dtypeY = context_->GetOutputDesc(INDEX_Y)->GetDataType();
86- if (dtypeC != dtypeY) {87+ bool isLowPrecInput = (dtypeA == ge::DT_FLOAT16 || dtypeA == ge::DT_BF16);
88+ bool is16In32Out = dtypeY == ge::DT_FLOAT && isLowPrecInput;
89+ if ((!is16In32Out && dtypeC != dtypeY) || (is16In32Out && dtypeC != dtypeA && dtypeC != dtypeY)) {
87 OP_LOGW(params_.opName,90 OP_LOGW(params_.opName,
88- "self dtype (%s) and out dtype (%s) are not the same.",91+ "Expected self dtype(%s) to be equal to mat dtype or out dtype(%s) in 16in32out scenario. or"
92+ "Expected self dtype(%s) to be equal to out dtype(%s).",
89 Ops::Base::ToString(dtypeC).c_str(),93 Ops::Base::ToString(dtypeC).c_str(),
90 Ops::Base::ToString(dtypeY).c_str());94 Ops::Base::ToString(dtypeY).c_str());
91 return false;95 return false;
92 }96 }
93- if (dtypeY != ge::DT_FLOAT16 && dtypeY != ge::DT_BF16) {97+ bool isValidOutput = (dtypeY == ge::DT_FLOAT16 || dtypeY == ge::DT_BF16 ||
98+ (dtypeY == ge::DT_FLOAT && isLowPrecInput));
99+ if (!isValidOutput) {
94 OP_LOGW(params_.opName,100 OP_LOGW(params_.opName,
95- "invalid out dtype (%s), only support half or bfloat16 output.",101+ "invalid out dtype (%s), only support half, bfloat16 or float(with low-precision input) output.",
96 Ops::Base::ToString(dtypeY).c_str());102 Ops::Base::ToString(dtypeY).c_str());
97 return false;103 return false;
98 }104 }
Mmatmul/gemm_v3/op_kernel/gemm_v3.cpp+98-68
@@ -69,6 +69,36 @@ constexpr CubeFormat format_y = CubeFormat::NZ;
69constexpr CubeFormat format_y = CubeFormat::ND;69constexpr CubeFormat format_y = CubeFormat::ND;
70#endif70#endif
71 71 
72+#if defined (FORMAT_A) && (FORMAT_A == FORMAT_FRACTAL_NZ)
73+constexpr PpMatMulNS::DataFormat formatA = PpMatMulNS::DataFormat::NZ;
74+#else
75+constexpr PpMatMulNS::DataFormat formatA = PpMatMulNS::DataFormat::ND;
76+#endif
77+ 
78+#if defined (FORMAT_B) && (FORMAT_B == FORMAT_FRACTAL_NZ)
79+constexpr PpMatMulNS::DataFormat formatB = PpMatMulNS::DataFormat::NZ;
80+#else
81+constexpr PpMatMulNS::DataFormat formatB = PpMatMulNS::DataFormat::ND;
82+#endif
83+ 
84+#if defined(ORIG_DTYPE_Y) && (ORIG_DTYPE_Y == DT_FLOAT)
85+#define IS_FP32_OUT 1
86+#else
87+#define IS_FP32_OUT 0
88+#endif
89+ 
90+#if defined(ORIG_DTYPE_C) && (ORIG_DTYPE_C == DT_FLOAT)
91+ #define IS_FP32_BIAS 1
92+#else
93+ #define IS_FP32_BIAS 0
94+#endif
95+ 
96+#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220
97+ #define IS_A2_A3 1
98+#else
99+ #define IS_A2_A3 0
100+#endif
101+ 
72#define MMV3_IMPL_ACCU(templateFunc, cFormat, ...) \102#define MMV3_IMPL_ACCU(templateFunc, cFormat, ...) \
73 do { \103 do { \
74 using cType = MatmulType<AscendC::TPosition::GM, cFormat, DTYPE_Y>; \104 using cType = MatmulType<AscendC::TPosition::GM, cFormat, DTYPE_Y>; \
@@ -190,80 +220,80 @@ constexpr CubeFormat format_y = CubeFormat::ND;
190 } \220 } \
191 } while (0)221 } while (0)
192 222 
193-#define INVOKE_GEMM_V3_BASE_KERNEL(transA, transB, swizzleDirect...) \223+#define INVOKE_GEMM_V3_BASE_KERNEL(transA, transB, swizzleDirect...) \
194- do { \224+ do { \
195- using PpMatMulNS::GemmV3BaseKernel; \225+ using PpMatMulNS::GemmV3BaseKernel; \
196- AscendC::SetAtomicNone(); \226+ AscendC::SetAtomicNone(); \
197- if ASCEND_IS_AIV { \227+ if ASCEND_IS_AIV { \
198- AscendC::SetMaskNorm(); \228+ AscendC::SetMaskNorm(); \
199- AscendC::SetVectorMask<uint8_t>((uint64_t)-1, (uint64_t)-1); \229+ AscendC::SetVectorMask<uint8_t>((uint64_t)-1, (uint64_t)-1); \
200- AscendC::CrossCoreSetFlag<0x2, PIPE_MTE2>(2); \230+ AscendC::CrossCoreSetFlag<0x2, PIPE_MTE2>(2); \
201- AscendC::CrossCoreSetFlag<0x2, PIPE_MTE2>(3); \231+ AscendC::CrossCoreSetFlag<0x2, PIPE_MTE2>(3); \
202- } \232+ } \
203- if ASCEND_IS_AIC { \233+ if ASCEND_IS_AIC { \
204- AscendC::SetLoadDataPaddingValue<uint64_t>((uint64_t)0); \234+ AscendC::SetLoadDataPaddingValue<uint64_t>((uint64_t)0); \
205- AscendC::SetFixpipeNz2ndFlag(1, 0, 0); \235+ AscendC::SetFixpipeNz2ndFlag(1, 0, 0); \
206- } \236+ } \
207- if (transA == 0 && transB == 0) { \237+ if (transA == 0 && transB == 0) { \
208- if (swizzleDirect == 0) { \238+ if (swizzleDirect == 0) { \
209- GemmV3BaseKernel<0, false, false, DTYPE_A, DTYPE_A, float> op; \239+ GemmV3BaseKernel<0, false, false, DTYPE_A, DTYPE_C, DTYPE_Y, float, formatA, formatB> op; \
210- op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \240+ op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \
211- op.RunVector(); \241+ op.RunVector(); \
212- op.RunCube(); \242+ op.RunCube(); \
213- } else if (swizzleDirect == 1) { \243+ } else if (swizzleDirect == 1) { \
214- GemmV3BaseKernel<1, false, false, DTYPE_A, DTYPE_A, float> op; \244+ GemmV3BaseKernel<1, false, false, DTYPE_A, DTYPE_C, DTYPE_Y, float, formatA, formatB> op; \
215- op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \245+ op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \
216- op.RunVector(); \246+ op.RunVector(); \
217- op.RunCube(); \247+ op.RunCube(); \
218- } \248+ } \
219- } else if (transA == 0 && transB == 1) { \249+ } else if (transA == 0 && transB == 1) { \
220- if (swizzleDirect == 0) { \250+ if (swizzleDirect == 0) { \
221- GemmV3BaseKernel<0, false, true, DTYPE_A, DTYPE_A, float> op; \251+ GemmV3BaseKernel<0, false, true, DTYPE_A, DTYPE_C, DTYPE_Y, float, formatA, formatB> op; \
222- op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \252+ op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \
223- op.RunVector(); \253+ op.RunVector(); \
224- op.RunCube(); \254+ op.RunCube(); \
225- } else if (swizzleDirect == 1) { \255+ } else if (swizzleDirect == 1) { \
226- GemmV3BaseKernel<1, false, true, DTYPE_A, DTYPE_A, float> op; \256+ GemmV3BaseKernel<1, false, true, DTYPE_A, DTYPE_C, DTYPE_Y, float, formatA, formatB> op; \
227- op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \257+ op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \
228- op.RunVector(); \258+ op.RunVector(); \
229- op.RunCube(); \259+ op.RunCube(); \
230- } \260+ } \
231- } else if (transA == 1 && transB == 0) { \261+ } else if (transA == 1 && transB == 0) { \
232- if (swizzleDirect == 0) { \262+ if (swizzleDirect == 0) { \
233- GemmV3BaseKernel<0, true, false, DTYPE_A, DTYPE_A, float> op; \263+ GemmV3BaseKernel<0, true, false, DTYPE_A, DTYPE_C, DTYPE_Y, float, formatA, formatB> op; \
234- op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \264+ op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \
235- op.RunVector(); \265+ op.RunVector(); \
236- op.RunCube(); \266+ op.RunCube(); \
237- } else if (swizzleDirect == 1) { \267+ } else if (swizzleDirect == 1) { \
238- GemmV3BaseKernel<1, true, false, DTYPE_A, DTYPE_A, float> op; \268+ GemmV3BaseKernel<1, true, false, DTYPE_A, DTYPE_C, DTYPE_Y, float, formatA, formatB> op; \
239- op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \269+ op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \
240- op.RunVector(); \270+ op.RunVector(); \
241- op.RunCube(); \271+ op.RunCube(); \
242- } \272+ } \
243- } else if (transA == 1 && transB == 1) { \273+ } else if (transA == 1 && transB == 1) { \
244- if (swizzleDirect == 0) { \274+ if (swizzleDirect == 0) { \
245- GemmV3BaseKernel<0, true, true, DTYPE_A, DTYPE_A, float> op; \275+ GemmV3BaseKernel<0, true, true, DTYPE_A, DTYPE_C, DTYPE_Y, float, formatA, formatB> op; \
246- op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \276+ op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \
247- op.RunVector(); \277+ op.RunVector(); \
248- op.RunCube(); \278+ op.RunCube(); \
249- } else if (swizzleDirect == 1) { \279+ } else if (swizzleDirect == 1) { \
250- GemmV3BaseKernel<1, true, true, DTYPE_A, DTYPE_A, float> op; \280+ GemmV3BaseKernel<1, true, true, DTYPE_A, DTYPE_C, DTYPE_Y, float, formatA, formatB> op; \
251- op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \281+ op.Init(aGM, bGM, cGM, yGM, gmUserWorkspace, tilingData); \
252- op.RunVector(); \282+ op.RunVector(); \
253- op.RunCube(); \283+ op.RunCube(); \
254- } \284+ } \
255- } \285+ } \
256- if ASCEND_IS_AIC { \286+ if ASCEND_IS_AIC { \
257- AscendC::CrossCoreWaitFlag(2); \287+ AscendC::CrossCoreWaitFlag(2); \
258- AscendC::CrossCoreWaitFlag(3); \288+ AscendC::CrossCoreWaitFlag(3); \
259- } \289+ } \
260 } while (0)290 } while (0)
261 291 
262template <int LOADMODE, int SPLITCOREMODE, int FIXOPTI, int MIXND2NZ, int SPECIALOPT, int FP32ADDMM>292template <int LOADMODE, int SPLITCOREMODE, int FIXOPTI, int MIXND2NZ, int SPECIALOPT, int FP32ADDMM>
263__global__ __aicore__ void gemm_v3(293__global__ __aicore__ void gemm_v3(
264 GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)294 GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM)
265{295{
266-#if (defined(ORIG_DTYPE_C) && (ORIG_DTYPE_C == DT_FLOAT))296+#if ((IS_FP32_BIAS && !IS_A2_A3) || (IS_A2_A3 && IS_FP32_BIAS && !IS_FP32_OUT))
267 REGISTER_TILING_DEFAULT(MatmulTilingData);297 REGISTER_TILING_DEFAULT(MatmulTilingData);
268 __gm__ uint8_t *user = GetUserWorkspace(workspaceGM);298 __gm__ uint8_t *user = GetUserWorkspace(workspaceGM);
269 GM_ADDR biasGM = nullptr;299 GM_ADDR biasGM = nullptr;
Mmatmul/gemm_v3/op_kernel/gemm_v3_base_kernel.h+85-38
@@ -30,6 +30,7 @@
30#include "../transpose_batch_mat_mul/pp_matmul_common.h"30#include "../transpose_batch_mat_mul/pp_matmul_common.h"
31#include "gemm_v3_tiling_data.h"31#include "gemm_v3_tiling_data.h"
32#include "kernel_tensor.h"32#include "kernel_tensor.h"
33+#include <type_traits>
33 34 
34namespace PpMatMulNS {35namespace PpMatMulNS {
35 36 
@@ -42,6 +43,7 @@ template <uint32_t swizzleDirect,
42 bool transA,43 bool transA,
43 bool transB,44 bool transB,
44 typename InDtype,45 typename InDtype,
46+ typename BiasDtype,
45 typename OutDtype,47 typename OutDtype,
46 typename AccumDtype,48 typename AccumDtype,
47 DataFormat formatA = DataFormat::ND,49 DataFormat formatA = DataFormat::ND,
@@ -186,7 +188,7 @@ private:
186private:188private:
187 AscendC::GlobalTensor<InDtype> gmA_;189 AscendC::GlobalTensor<InDtype> gmA_;
188 AscendC::GlobalTensor<InDtype> gmB_;190 AscendC::GlobalTensor<InDtype> gmB_;
189- AscendC::GlobalTensor<OutDtype> gmC_;191+ AscendC::GlobalTensor<BiasDtype> gmC_;
190 AscendC::GlobalTensor<OutDtype> gmY_;192 AscendC::GlobalTensor<OutDtype> gmY_;
191 AscendC::GlobalTensor<AccumDtype> gmWorkspace_;193 AscendC::GlobalTensor<AccumDtype> gmWorkspace_;
192 AscendC::LocalTensor<InDtype> l1BaseA_;194 AscendC::LocalTensor<InDtype> l1BaseA_;
@@ -195,7 +197,7 @@ private:
195 AscendC::LocalTensor<InDtype> l0BaseB_;197 AscendC::LocalTensor<InDtype> l0BaseB_;
196 AscendC::LocalTensor<AccumDtype> l0BaseC_;198 AscendC::LocalTensor<AccumDtype> l0BaseC_;
197 AscendC::LocalTensor<AccumDtype> ubSum_;199 AscendC::LocalTensor<AccumDtype> ubSum_;
198- AscendC::LocalTensor<OutDtype> ubC_;200+ AscendC::LocalTensor<BiasDtype> ubC_;
199 AscendC::LocalTensor<OutDtype> ubY_;201 AscendC::LocalTensor<OutDtype> ubY_;
200 AscendC::LocalTensor<AccumDtype> ubProd_;202 AscendC::LocalTensor<AccumDtype> ubProd_;
201 203 
@@ -226,12 +228,13 @@ template <uint32_t swizzleDirect,
226 bool transA,228 bool transA,
227 bool transB,229 bool transB,
228 typename InDtype,230 typename InDtype,
231+ typename BiasDtype,
229 typename OutDtype,232 typename OutDtype,
230 typename AccumDtype,233 typename AccumDtype,
231 DataFormat formatA,234 DataFormat formatA,
232 DataFormat formatB>235 DataFormat formatB>
233__aicore__ FORCE_INLINE void236__aicore__ FORCE_INLINE void
234-GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, formatA, formatB>::Init(237+GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, BiasDtype, OutDtype, AccumDtype, formatA, formatB>::Init(
235 __gm__ uint8_t* __restrict__ a,238 __gm__ uint8_t* __restrict__ a,
236 __gm__ uint8_t* __restrict__ b,239 __gm__ uint8_t* __restrict__ b,
237 __gm__ uint8_t* __restrict__ c,240 __gm__ uint8_t* __restrict__ c,
@@ -269,7 +272,7 @@ GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, f
269 272 
270 gmA_.SetGlobalBuffer(reinterpret_cast<__gm__ InDtype*>(a));273 gmA_.SetGlobalBuffer(reinterpret_cast<__gm__ InDtype*>(a));
271 gmB_.SetGlobalBuffer(reinterpret_cast<__gm__ InDtype*>(b));274 gmB_.SetGlobalBuffer(reinterpret_cast<__gm__ InDtype*>(b));
272- gmC_.SetGlobalBuffer(reinterpret_cast<__gm__ OutDtype*>(c));275+ gmC_.SetGlobalBuffer(reinterpret_cast<__gm__ BiasDtype*>(c));
273 gmY_.SetGlobalBuffer(reinterpret_cast<__gm__ OutDtype*>(y));276 gmY_.SetGlobalBuffer(reinterpret_cast<__gm__ OutDtype*>(y));
274 gmWorkspace_.SetGlobalBuffer(reinterpret_cast<__gm__ AccumDtype*>(workspace) +277 gmWorkspace_.SetGlobalBuffer(reinterpret_cast<__gm__ AccumDtype*>(workspace) +
275 coreIdx_ * mTile_ * nTile_ * NUM_BUFFER);278 coreIdx_ * mTile_ * nTile_ * NUM_BUFFER);
@@ -283,12 +286,21 @@ template <uint32_t swizzleDirect,
283 bool transA,286 bool transA,
284 bool transB,287 bool transB,
285 typename InDtype,288 typename InDtype,
289+ typename BiasDtype,
286 typename OutDtype,290 typename OutDtype,
287 typename AccumDtype,291 typename AccumDtype,
288 DataFormat formatA,292 DataFormat formatA,
289 DataFormat formatB>293 DataFormat formatB>
290__aicore__ FORCE_INLINE void294__aicore__ FORCE_INLINE void
291-GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, formatA, formatB>::GetBlockIdx(295+GemmV3BaseKernel<swizzleDirect,
296+ transA,
297+ transB,
298+ InDtype,
299+ BiasDtype,
300+ OutDtype,
301+ AccumDtype,
302+ formatA,
303+ formatB>::GetBlockIdx(
292 const uint64_t index, MatCoord& tidx)304 const uint64_t index, MatCoord& tidx)
293{305{
294 uint64_t in_batch_idx = index % (tileDim_.m * tileDim_.n);306 uint64_t in_batch_idx = index % (tileDim_.m * tileDim_.n);
@@ -327,12 +339,13 @@ template <uint32_t swizzleDirect,
327 bool transA,339 bool transA,
328 bool transB,340 bool transB,
329 typename InDtype,341 typename InDtype,
342+ typename BiasDtype,
330 typename OutDtype,343 typename OutDtype,
331 typename AccumDtype,344 typename AccumDtype,
332 DataFormat formatA,345 DataFormat formatA,
333 DataFormat formatB>346 DataFormat formatB>
334__aicore__ FORCE_INLINE void347__aicore__ FORCE_INLINE void
335-GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, formatA, formatB>::RunCube()348+GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, BiasDtype, OutDtype, AccumDtype, formatA, formatB>::RunCube()
336{349{
337#ifdef __DAV_C220_CUBE__350#ifdef __DAV_C220_CUBE__
338 using LocalTensor = AscendC::LocalTensor<InDtype>;351 using LocalTensor = AscendC::LocalTensor<InDtype>;
@@ -571,16 +584,18 @@ GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, f
571#endif584#endif
572}585}
573 586 
587+// vector 数据搬运
574template <uint32_t swizzleDirect,588template <uint32_t swizzleDirect,
575 bool transA,589 bool transA,
576 bool transB,590 bool transB,
577 typename InDtype,591 typename InDtype,
592+ typename BiasDtype,
578 typename OutDtype,593 typename OutDtype,
579 typename AccumDtype,594 typename AccumDtype,
580 DataFormat formatA,595 DataFormat formatA,
581 DataFormat formatB>596 DataFormat formatB>
582__aicore__ FORCE_INLINE void597__aicore__ FORCE_INLINE void
583-GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, formatA, formatB>::RunVector()598+GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, BiasDtype, OutDtype, AccumDtype, formatA, formatB>::RunVector()
584{599{
585#ifdef __DAV_C220_VEC__600#ifdef __DAV_C220_VEC__
586 using namespace AscendC;601 using namespace AscendC;
@@ -592,8 +607,10 @@ GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, f
592 GetBlockIdx(loopIdx, tileIdx);607 GetBlockIdx(loopIdx, tileIdx);
593 uint64_t mTileActual = (tileIdx.m == (tileDim_.m - 1)) ? (m_ - tileIdx.m * mTile_) : mTile_;608 uint64_t mTileActual = (tileIdx.m == (tileDim_.m - 1)) ? (m_ - tileIdx.m * mTile_) : mTile_;
594 uint64_t nTileActual = (tileIdx.n == (tileDim_.n - 1)) ? (n_ - tileIdx.n * nTile_) : nTile_;609 uint64_t nTileActual = (tileIdx.n == (tileDim_.n - 1)) ? (n_ - tileIdx.n * nTile_) : nTile_;
595- uint64_t nRound = RoundUp<CONST_16>(nTileActual);610+ constexpr uint32_t roundNum = CONST_32 / sizeof(BiasDtype);
611+ uint64_t nRound = RoundUp<roundNum>(nTileActual);
596 uint32_t prodDstGap = (nRound - nTileActual) * sizeof(AccumDtype) >= CONST_32 ? 1 : 0;612 uint32_t prodDstGap = (nRound - nTileActual) * sizeof(AccumDtype) >= CONST_32 ? 1 : 0;
613+ uint32_t sumSrcGap = std::is_same_v<OutDtype, BiasDtype> ? 0 : prodDstGap;
597 uint64_t mTileHalf = (mTileActual + 1) / 2;614 uint64_t mTileHalf = (mTileActual + 1) / 2;
598 uint64_t mTileHalfActual = (AscendC::GetSubBlockIdx() == 0) ? mTileHalf : (mTileActual - mTileHalf);615 uint64_t mTileHalfActual = (AscendC::GetSubBlockIdx() == 0) ? mTileHalf : (mTileActual - mTileHalf);
599 uint64_t offsetCy = batchIdx * m_ * n_ + tileIdx.m * mTile_ * n_ + tileIdx.n * nTile_ +616 uint64_t offsetCy = batchIdx * m_ * n_ + tileIdx.m * mTile_ * n_ + tileIdx.n * nTile_ +
@@ -602,25 +619,32 @@ GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, f
602 uint64_t count = CeilDiv<VEC_ITER_NUMEL>(numel);619 uint64_t count = CeilDiv<VEC_ITER_NUMEL>(numel);
603 if (mTileHalfActual != 0) {620 if (mTileHalfActual != 0) {
604 WaitFlag<HardEvent::MTE3_MTE2>(EVENT_ID0);621 WaitFlag<HardEvent::MTE3_MTE2>(EVENT_ID0);
605- CopyGmToUbufAlign<OutDtype>(ubC_, // dst622+ AscendC::LocalTensor<BiasDtype> dstUbC = ubC_;
606- gmC_[offsetCy], // src623+ if constexpr (std::is_same_v<BiasDtype, AccumDtype>) {
607- 0, // sid624+ dstUbC = ubSum_;
608- mTileHalfActual, // nBurst625+ }
609- nTileActual * sizeof(OutDtype), // lenBurst626+ CopyGmToUbufAlign<BiasDtype>(dstUbC, // dst
610- 0, // leftPaddingNum627+ gmC_[offsetCy], // src
611- 0, // rightPaddingNum628+ 0, // sid
612- (n_ - nTileActual) * sizeof(OutDtype), // srcGap629+ mTileHalfActual, // nBurst
613- 0); // dstGap630+ nTileActual * sizeof(BiasDtype), // lenBurst
631+ 0, // leftPaddingNum
632+ 0, // rightPaddingNum
633+ (n_ - nTileActual) * sizeof(BiasDtype), // srcGap
634+ 0); // dstGap
614 SetFlag<HardEvent::MTE2_V>(EVENT_ID0);635 SetFlag<HardEvent::MTE2_V>(EVENT_ID0);
615 WaitFlag<HardEvent::MTE2_V>(EVENT_ID0);636 WaitFlag<HardEvent::MTE2_V>(EVENT_ID0);
616 for (uint32_t i = 0; i < count; ++i) {637 for (uint32_t i = 0; i < count; ++i) {
617- AscendC::Cast<AccumDtype, OutDtype, false>(ubSum_[i * VEC_ITER_NUMEL],638+ if constexpr (!std::is_same_v<BiasDtype, AccumDtype>) {
618- ubC_[i * VEC_ITER_NUMEL],639+ AscendC::printf("need cast ubc dtype.\n");
619- AscendC::RoundMode::CAST_NONE,640+ AscendC::Cast<AccumDtype, BiasDtype, false>(ubSum_[i * VEC_ITER_NUMEL], // dst
620- (uint64_t)0,641+ ubC_[i * VEC_ITER_NUMEL], // src
621- VEC_ITER_REPEAT,642+ AscendC::RoundMode::CAST_NONE, // mode
622- AscendC::UnaryRepeatParams(1, 1, 8, 4));643+ (uint64_t)0, // mask
623- AscendC::PipeBarrier<PIPE_V>();644+ VEC_ITER_REPEAT, // repeat
645+ AscendC::UnaryRepeatParams(1, 1, 8, 4));
646+ AscendC::PipeBarrier<PIPE_V>();
647+ }
624 AscendC::Muls<AccumDtype, false>(ubSum_[i * VEC_ITER_NUMEL], // dst648 AscendC::Muls<AccumDtype, false>(ubSum_[i * VEC_ITER_NUMEL], // dst
625 ubSum_[i * VEC_ITER_NUMEL], // src649 ubSum_[i * VEC_ITER_NUMEL], // src
626 beta_, // scalar650 beta_, // scalar
@@ -654,24 +678,31 @@ GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, f
654 VEC_ITER_REPEAT, // repeat678 VEC_ITER_REPEAT, // repeat
655 AscendC::UnaryRepeatParams(1, 1, 8, 8));679 AscendC::UnaryRepeatParams(1, 1, 8, 8));
656 AscendC::PipeBarrier<PIPE_V>();680 AscendC::PipeBarrier<PIPE_V>();
657- AscendC::Cast<OutDtype, AccumDtype, false>(ubY_[i * VEC_ITER_NUMEL], // dst681+ if constexpr (!std::is_same_v<OutDtype, AccumDtype>) {
658- ubSum_[i * VEC_ITER_NUMEL], // src682+ AscendC::printf("need cast uby dtype.\n");
659- AscendC::RoundMode::CAST_RINT, // mode683+ AscendC::Cast<OutDtype, AccumDtype, false>(ubY_[i * VEC_ITER_NUMEL], // dst
660- (uint64_t)0, // mask (disabled)684+ ubSum_[i * VEC_ITER_NUMEL], // src
661- VEC_ITER_REPEAT, // repeat685+ AscendC::RoundMode::CAST_RINT, // mode
662- AscendC::UnaryRepeatParams(1, 1, 4, 8));686+ (uint64_t)0, // mask (disabled)
663- AscendC::PipeBarrier<PIPE_V>();687+ VEC_ITER_REPEAT, // repeat
688+ AscendC::UnaryRepeatParams(1, 1, 4, 8));
689+ AscendC::PipeBarrier<PIPE_V>();
690+ }
664 }691 }
665 SetFlag<HardEvent::V_MTE3>(EVENT_ID0);692 SetFlag<HardEvent::V_MTE3>(EVENT_ID0);
666 WaitFlag<HardEvent::V_MTE3>(EVENT_ID0);693 WaitFlag<HardEvent::V_MTE3>(EVENT_ID0);
694+ AscendC::LocalTensor<OutDtype> SrcUbY = ubY_;
695+ if constexpr (std::is_same_v<OutDtype, AccumDtype>) {
696+ SrcUbY = ubSum_;
697+ }
667 CopyUbufToGmAlign(gmY_[offsetCy], // dst698 CopyUbufToGmAlign(gmY_[offsetCy], // dst
668- ubY_, // src699+ SrcUbY, // src
669 0, // sid700 0, // sid
670 mTileHalfActual, // nBurst701 mTileHalfActual, // nBurst
671 nTileActual * sizeof(OutDtype), // lenBurst702 nTileActual * sizeof(OutDtype), // lenBurst
672 0, // leftPaddingNum703 0, // leftPaddingNum
673 0, // rightPaddingNum704 0, // rightPaddingNum
674- 0, // srcGap705+ sumSrcGap, // srcGap
675 (n_ - nTileActual) * sizeof(OutDtype)); // dstGap706 (n_ - nTileActual) * sizeof(OutDtype)); // dstGap
676 SetFlag<HardEvent::MTE3_MTE2>(EVENT_ID0);707 SetFlag<HardEvent::MTE3_MTE2>(EVENT_ID0);
677 } else {708 } else {
@@ -687,13 +718,21 @@ template <uint32_t swizzleDirect,
687 bool transA,718 bool transA,
688 bool transB,719 bool transB,
689 typename InDtype,720 typename InDtype,
721+ typename BiasDtype,
690 typename OutDtype,722 typename OutDtype,
691 typename AccumDtype,723 typename AccumDtype,
692 DataFormat formatA,724 DataFormat formatA,
693 DataFormat formatB>725 DataFormat formatB>
694__aicore__ FORCE_INLINE void726__aicore__ FORCE_INLINE void
695-GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, formatA, formatB>::InitBufferCube(727+GemmV3BaseKernel<swizzleDirect,
696- const OcBuffer& buf)728+ transA,
729+ transB,
730+ InDtype,
731+ BiasDtype,
732+ OutDtype,
733+ AccumDtype,
734+ formatA,
735+ formatB>::InitBufferCube(const OcBuffer& buf)
697{736{
698#ifdef __DAV_C220_CUBE__737#ifdef __DAV_C220_CUBE__
699 l1BaseA_ = buf.template GetBuffer<BufferType::ASCEND_CB, InDtype>(0);738 l1BaseA_ = buf.template GetBuffer<BufferType::ASCEND_CB, InDtype>(0);
@@ -707,13 +746,21 @@ template <uint32_t swizzleDirect,
707 bool transA,746 bool transA,
708 bool transB,747 bool transB,
709 typename InDtype,748 typename InDtype,
749+ typename BiasDtype,
710 typename OutDtype,750 typename OutDtype,
711 typename AccumDtype,751 typename AccumDtype,
712 DataFormat formatA,752 DataFormat formatA,
713 DataFormat formatB>753 DataFormat formatB>
714__aicore__ FORCE_INLINE void754__aicore__ FORCE_INLINE void
715-GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, formatA, formatB>::InitBufferVector(755+GemmV3BaseKernel<swizzleDirect,
716- const OcBuffer& buf)756+ transA,
757+ transB,
758+ InDtype,
759+ BiasDtype,
760+ OutDtype,
761+ AccumDtype,
762+ formatA,
763+ formatB>::InitBufferVector(const OcBuffer& buf)
717{764{
718#ifdef __DAV_C220_VEC__765#ifdef __DAV_C220_VEC__
719 uint64_t sizeUbProd = 65536;766 uint64_t sizeUbProd = 65536;
@@ -721,7 +768,7 @@ GemmV3BaseKernel<swizzleDirect, transA, transB, InDtype, OutDtype, AccumDtype, f
721 uint64_t sizeUbC = 32768;768 uint64_t sizeUbC = 32768;
722 ubProd_ = buf.template GetBuffer<BufferType::ASCEND_UB, AccumDtype>(0);769 ubProd_ = buf.template GetBuffer<BufferType::ASCEND_UB, AccumDtype>(0);
723 ubSum_ = buf.template GetBuffer<BufferType::ASCEND_UB, AccumDtype>(sizeUbProd);770 ubSum_ = buf.template GetBuffer<BufferType::ASCEND_UB, AccumDtype>(sizeUbProd);
724- ubC_ = buf.template GetBuffer<BufferType::ASCEND_UB, OutDtype>(sizeUbProd + sizeUbSum);771+ ubC_ = buf.template GetBuffer<BufferType::ASCEND_UB, BiasDtype>(sizeUbProd + sizeUbSum);
725 ubY_ = buf.template GetBuffer<BufferType::ASCEND_UB, OutDtype>(sizeUbProd + sizeUbSum + sizeUbC);772 ubY_ = buf.template GetBuffer<BufferType::ASCEND_UB, OutDtype>(sizeUbProd + sizeUbSum + sizeUbC);
726#endif773#endif
727}774}