Star909
1826
代码
代码
Issues226
Pull Requests308
Actions
项目讨论
Wiki
分析
Star909
1826
已合并
log_fix #8828
Alpha123_创建于 20 天前
log_fix #8828
已合并
Alpha123_创建于 20 天前
从Alpha123_/ops-nn_new2master合入到cann/ops-nnmaster
共 109 个文件变更+422-424
Mmatmul/addmv/op_host/op_api/aclnn_addmv.cpp+12-11
@@ -96,7 +96,7 @@ static bool CheckPromoteType(const aclTensor* self, const aclTensor* mat, const
96 // 检查mat和vec能否做数据类型推导96 // 检查mat和vec能否做数据类型推导
97 DataType promoteType = PromoteType(mat->GetDataType(), vec->GetDataType());97 DataType promoteType = PromoteType(mat->GetDataType(), vec->GetDataType());
98 if (promoteType == DataType::DT_UNDEFINED) {98 if (promoteType == DataType::DT_UNDEFINED) {
99- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Mat dtype %s and Vec dtype %s can not promote dtype.",99+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Mat dtype %s and Vec dtype %s cannot be promoted.",
100 ToString(mat->GetDataType()).GetString(), ToString(vec->GetDataType()).GetString());100 ToString(mat->GetDataType()).GetString(), ToString(vec->GetDataType()).GetString());
101 return false;101 return false;
102 }102 }
@@ -106,7 +106,7 @@ static bool CheckPromoteType(const aclTensor* self, const aclTensor* mat, const
106 if (std::abs(alpha->ToFloat() - 1.0f) > std::numeric_limits<float>::epsilon()) {106 if (std::abs(alpha->ToFloat() - 1.0f) > std::numeric_limits<float>::epsilon()) {
107 promoteType2 = PromoteType(alpha->GetDataType(), promoteType);107 promoteType2 = PromoteType(alpha->GetDataType(), promoteType);
108 if (promoteType2 == DataType::DT_UNDEFINED) {108 if (promoteType2 == DataType::DT_UNDEFINED) {
109- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Alpha dtype %s and Mat/Vec dtype %s can not promote dtype.",109+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Alpha dtype %s and Mat/Vec dtype %s cannot be promoted.",
110 ToString(alpha->GetDataType()).GetString(), ToString(promoteType).GetString());110 ToString(alpha->GetDataType()).GetString(), ToString(promoteType).GetString());
111 return false;111 return false;
112 }112 }
@@ -117,7 +117,7 @@ static bool CheckPromoteType(const aclTensor* self, const aclTensor* mat, const
117 if (std::abs(beta->ToFloat() - 1.0f) > std::numeric_limits<float>::epsilon()) {117 if (std::abs(beta->ToFloat() - 1.0f) > std::numeric_limits<float>::epsilon()) {
118 promoteType3 = PromoteType(beta->GetDataType(), self->GetDataType());118 promoteType3 = PromoteType(beta->GetDataType(), self->GetDataType());
119 if (promoteType3 == DataType::DT_UNDEFINED) {119 if (promoteType3 == DataType::DT_UNDEFINED) {
120- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self dtype %s and Beta dtype %s can not promote dtype.",120+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self dtype %s and Beta dtype %s cannot be promoted.",
121 ToString(self->GetDataType()).GetString(), ToString(beta->GetDataType()).GetString());121 ToString(self->GetDataType()).GetString(), ToString(beta->GetDataType()).GetString());
122 return false;122 return false;
123 }123 }
@@ -127,7 +127,7 @@ static bool CheckPromoteType(const aclTensor* self, const aclTensor* mat, const
127 DataType promoteTypeFinal = PromoteType(promoteType3, promoteType2);127 DataType promoteTypeFinal = PromoteType(promoteType3, promoteType2);
128 if (promoteTypeFinal == DataType::DT_UNDEFINED) {128 if (promoteTypeFinal == DataType::DT_UNDEFINED) {
129 OP_LOGE(ACLNN_ERR_PARAM_INVALID,129 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
130- "part1(beta*self) dtype %s and part2(alpha*(mat@vec)) dtype %s can not promote dtype.",130+ "part1(beta*self) dtype %s and part2(alpha*(mat@vec)) dtype %s cannot be promoted.",
131 ToString(promoteType3).GetString(), ToString(promoteType2).GetString());131 ToString(promoteType3).GetString(), ToString(promoteType2).GetString());
132 return false;132 return false;
133 }133 }
@@ -149,7 +149,7 @@ static bool CheckTensorDimAndSize(const aclTensor* self, const aclTensor* mat, c
149 if ((selfShape.GetDimNum() > 1) || (matShape.GetDimNum() != MATRIX_DIM) || (vecShape.GetDimNum() != 1) ||149 if ((selfShape.GetDimNum() > 1) || (matShape.GetDimNum() != MATRIX_DIM) || (vecShape.GetDimNum() != 1) ||
150 (outShape.GetDimNum() != 1)) {150 (outShape.GetDimNum() != 1)) {
151 OP_LOGE(ACLNN_ERR_PARAM_INVALID,151 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
152- "Expect input tensor dim [0/1,2,1,1], but receive self [%zu], mat [%zu], vec [%zu], out [%zu].",152+ "Expect input tensor dim [0/1,2,1,1], but received self [%zu], mat [%zu], vec [%zu], out [%zu].",
153 selfShape.GetDimNum(), matShape.GetDimNum(), vecShape.GetDimNum(), outShape.GetDimNum());153 selfShape.GetDimNum(), matShape.GetDimNum(), vecShape.GetDimNum(), outShape.GetDimNum());
154 return false;154 return false;
155 }155 }
@@ -158,7 +158,7 @@ static bool CheckTensorDimAndSize(const aclTensor* self, const aclTensor* mat, c
158 (selfShape.GetDimNum() != 0 && matShape.GetDim(0) != selfShape.GetDim(0) &&158 (selfShape.GetDimNum() != 0 && matShape.GetDim(0) != selfShape.GetDim(0) &&
159 (selfShape.GetDim(0) != 1 || std::abs(alpha->ToFloat() - 0.0f) <= std::numeric_limits<float>::epsilon()))) {159 (selfShape.GetDim(0) != 1 || std::abs(alpha->ToFloat() - 0.0f) <= std::numeric_limits<float>::epsilon()))) {
160 OP_LOGE(ACLNN_ERR_PARAM_INVALID,160 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
161- "Input tensor shape not satisfy, current shape : self [%s], mat [%s], vec [%s], out [%s].",161+ "Input tensor shape does not satisfy, current shape: self [%s], mat [%s], vec [%s], out [%s].",
162 op::ToString(selfShape).GetString(), op::ToString(matShape).GetString(),162 op::ToString(selfShape).GetString(), op::ToString(matShape).GetString(),
163 op::ToString(vecShape).GetString(), op::ToString(outShape).GetString());163 op::ToString(vecShape).GetString(), op::ToString(outShape).GetString());
164 return false;164 return false;
@@ -176,7 +176,8 @@ static bool CheckFormat(const aclTensor* self, const aclTensor* mat, const aclTe
176 bool noSupportFormat = ((selfFormat == Format::FORMAT_FRACTAL_NZ) || (matFormat == Format::FORMAT_FRACTAL_NZ) ||176 bool noSupportFormat = ((selfFormat == Format::FORMAT_FRACTAL_NZ) || (matFormat == Format::FORMAT_FRACTAL_NZ) ||
177 (vecFormat == Format::FORMAT_FRACTAL_NZ) || (outTensorFormat == Format::FORMAT_FRACTAL_NZ));177 (vecFormat == Format::FORMAT_FRACTAL_NZ) || (outTensorFormat == Format::FORMAT_FRACTAL_NZ));
178 if (noSupportFormat) {178 if (noSupportFormat) {
179- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "One of the tensors ('self', 'mat', 'vec', 'out') does not support NZ format");179+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
180+ "One of the tensors ('self', 'mat', 'vec', 'out') does not support NZ format.");
180 return false;181 return false;
181 }182 }
182 return true;183 return true;
@@ -254,11 +255,11 @@ const aclTensor* GetMatMulResult(const aclTensor* mat, const aclTensor* vec, int
254 // 调用matmul之前需要对mat和vec做数据类型转换，以满足运算条件255 // 调用matmul之前需要对mat和vec做数据类型转换，以满足运算条件
255 auto promoteType = PromoteType(matContiguous->GetDataType(), mat2->GetDataType());256 auto promoteType = PromoteType(matContiguous->GetDataType(), mat2->GetDataType());
256 257 
257- // 将输入self的数据类型转换成隐式数据类型，根据具体算子语义按需调用258+ // 将输入 mat 的数据类型转换成隐式数据类型，根据具体算子语义按需调用
258 auto matCasted = l0op::Cast(matContiguous, promoteType, executor);259 auto matCasted = l0op::Cast(matContiguous, promoteType, executor);
259 CHECK_RET(matCasted != nullptr, nullptr);260 CHECK_RET(matCasted != nullptr, nullptr);
260 261 
261- // 将输入other的数据类型转换成隐式数据类型，根据具体算子语义按需调用262+ // 将输入 vec 的数据类型转换成隐式数据类型，根据具体算子语义按需调用
262 auto mat2Casted = l0op::Cast(mat2, promoteType, executor);263 auto mat2Casted = l0op::Cast(mat2, promoteType, executor);
263 CHECK_RET(mat2Casted != nullptr, nullptr);264 CHECK_RET(mat2Casted != nullptr, nullptr);
264 265 
@@ -290,11 +291,11 @@ const aclTensor* GetMulResult(const aclTensor* mat, const aclTensor* vec, const
290 // 调用matmul之前需要对mat和vec做数据类型转换，以满足运算条件291 // 调用matmul之前需要对mat和vec做数据类型转换，以满足运算条件
291 auto promoteType = PromoteType(matContiguous->GetDataType(), vecContiguous->GetDataType());292 auto promoteType = PromoteType(matContiguous->GetDataType(), vecContiguous->GetDataType());
292 293 
293- // 将输入self的数据类型转换成隐式数据类型，根据具体算子语义按需调用294+ // 将输入 mat 的数据类型转换成隐式数据类型，根据具体算子语义按需调用
294 auto matCasted = l0op::Cast(matContiguous, promoteType, executor);295 auto matCasted = l0op::Cast(matContiguous, promoteType, executor);
295 CHECK_RET(matCasted != nullptr, nullptr);296 CHECK_RET(matCasted != nullptr, nullptr);
296 297 
297- // 将输入other的数据类型转换成隐式数据类型，根据具体算子语义按需调用298+ // 将输入 vec 的数据类型转换成隐式数据类型，根据具体算子语义按需调用
298 auto vecCasted = l0op::Cast(vecContiguous, promoteType, executor);299 auto vecCasted = l0op::Cast(vecContiguous, promoteType, executor);
299 CHECK_RET(vecCasted != nullptr, nullptr);300 CHECK_RET(vecCasted != nullptr, nullptr);
300 301 
Mmatmul/addmv/op_host/op_api/aclnn_addmv.h+1-1
@@ -35,4 +35,4 @@ ACLNN_API aclnnStatus aclnnAddmv(void* workspace, uint64_t workspaceSize, aclOpE
35}35}
36#endif36#endif
37 37 
38-#endif // OP_API_INC_ADD_H_38+#endif // OP_API_INC_ADDMV_H_
Mmatmul/addmv/tests/ut/op_api/test_aclnn_addmv.cpp+4-4
@@ -132,10 +132,10 @@ TEST_F(l2_addmv_test, input_nullptr)
132// 输入数据类型遍历132// 输入数据类型遍历
133TEST_F(l2_addmv_test, input_alldtype_in_list)133TEST_F(l2_addmv_test, input_alldtype_in_list)
134{134{
135- vector<aclDataType> vaild_dtype_list{ACL_FLOAT, ACL_FLOAT16, ACL_DOUBLE, ACL_INT8, ACL_INT32,135+ vector<aclDataType> valid_dtype_list{ACL_FLOAT, ACL_FLOAT16, ACL_DOUBLE, ACL_INT8, ACL_INT32,
136 ACL_UINT8, ACL_INT16, ACL_INT64, ACL_BOOL};136 ACL_UINT8, ACL_INT16, ACL_INT64, ACL_BOOL};
137- vector<aclDataType> invaild_dtype_list{ACL_COMPLEX64, ACL_COMPLEX128};137+ vector<aclDataType> invalid_dtype_list{ACL_COMPLEX64, ACL_COMPLEX128};
138- for (auto dtype : vaild_dtype_list) {138+ for (auto dtype : valid_dtype_list) {
139 auto input_tensor_desc = TensorDesc({10}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);139 auto input_tensor_desc = TensorDesc({10}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);
140 auto mat_tensor_desc = TensorDesc({10, 5}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);140 auto mat_tensor_desc = TensorDesc({10, 5}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);
141 auto vec_tensor_desc = TensorDesc({5}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);141 auto vec_tensor_desc = TensorDesc({5}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);
@@ -154,7 +154,7 @@ TEST_F(l2_addmv_test, input_alldtype_in_list)
154 // SAMPLE: precision simulate154 // SAMPLE: precision simulate
155 // ut.TestPrecision();155 // ut.TestPrecision();
156 }156 }
157- for (auto dtype : invaild_dtype_list) {157+ for (auto dtype : invalid_dtype_list) {
158 auto input_tensor_desc = TensorDesc({10}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);158 auto input_tensor_desc = TensorDesc({10}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);
159 auto mat_tensor_desc = TensorDesc({10, 5}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);159 auto mat_tensor_desc = TensorDesc({10, 5}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);
160 auto vec_tensor_desc = TensorDesc({5}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);160 auto vec_tensor_desc = TensorDesc({5}, dtype, ACL_FORMAT_ND).ValueRange(-2, 2);
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_addbmm.cpp+5-5
@@ -110,7 +110,7 @@ static bool CheckShape(const aclTensor* selfTensor, const aclTensor* batch1Tenso
110 // check batch1 last dim and batch2 penultimate dim is equal or not110 // check batch1 last dim and batch2 penultimate dim is equal or not
111 if (batch1[THIRD_DIM] != batch2[SECOND_DIM]) {111 if (batch1[THIRD_DIM] != batch2[SECOND_DIM]) {
112 OP_LOGE(ACLNN_ERR_PARAM_INVALID,112 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
113- "batch1's last dim and batch2's penultimate dim shoule be same, batch1 [%ld], batch2 [%ld].",113+ "batch1's last dim and batch2's penultimate dim should be same, batch1 [%ld], batch2 [%ld].",
114 batch1[THIRD_DIM], batch2[SECOND_DIM]);114 batch1[THIRD_DIM], batch2[SECOND_DIM]);
115 return false;115 return false;
116 }116 }
@@ -118,7 +118,7 @@ static bool CheckShape(const aclTensor* selfTensor, const aclTensor* batch1Tenso
118 auto batch1DimNum = batch1Tensor->GetViewShape().GetDimNum();118 auto batch1DimNum = batch1Tensor->GetViewShape().GetDimNum();
119 auto batch2DimNum = batch2Tensor->GetViewShape().GetDimNum();119 auto batch2DimNum = batch2Tensor->GetViewShape().GetDimNum();
120 if (!CheckBatchDimBroadcast(batch1DimNum, batch2DimNum, batch1, batch2)) {120 if (!CheckBatchDimBroadcast(batch1DimNum, batch2DimNum, batch1, batch2)) {
121- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "self's batch dim and mat2's batch dim can not broadcast");121+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "self's batch dim and mat2's batch dim cannot broadcast");
122 return false;122 return false;
123 }123 }
124 124 
@@ -417,10 +417,10 @@ public:
417 // 进行Add或Axpy计算417 // 进行Add或Axpy计算
418 const aclTensor* addOut = nullptr;418 const aclTensor* addOut = nullptr;
419 if (std::abs(alpha->ToFloat() - 1.0f) <= std::numeric_limits<float>::epsilon()) {419 if (std::abs(alpha->ToFloat() - 1.0f) <= std::numeric_limits<float>::epsilon()) {
420- // alpha == 0 addOut = mulOutCasted + bmmOutCasted420+ // alpha == 1, addOut = mulOutCasted + reduceSumOutCasted
421 addOut = l0op::Add(mulOutCasted, reduceSumOutCasted, executor);421 addOut = l0op::Add(mulOutCasted, reduceSumOutCasted, executor);
422 } else {422 } else {
423- // alpha != 0 addOut = mulOutCasted + bmmOutCasted * alpha423+ // alpha != 1, addOut = mulOutCasted + reduceSumOutCasted * alpha
424 addOut = l0op::Axpy(mulOutCasted, reduceSumOutCasted, alpha->ToFloat(), executor);424 addOut = l0op::Axpy(mulOutCasted, reduceSumOutCasted, alpha->ToFloat(), executor);
425 }425 }
426 CHECK_RET(addOut != nullptr, ACLNN_ERR_INNER_NULLPTR);426 CHECK_RET(addOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
@@ -556,4 +556,4 @@ aclnnStatus aclnnInplaceAddbmm(void* workspace, uint64_t workspaceSize, aclOpExe
556 L2_DFX_PHASE_2(aclnnInplaceAddbmm);556 L2_DFX_PHASE_2(aclnnInplaceAddbmm);
557 // 固定写法，调用框架能力，完成计算557 // 固定写法，调用框架能力，完成计算
558 return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);558 return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
559-}559+}
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_addbmm.h+7-7
@@ -38,7 +38,7 @@ extern "C" {
38 * @param [in] beta: host侧的aclScalar，数据类型需要可转换成self与batch1@batch2推导后的数据类型。38 * @param [in] beta: host侧的aclScalar，数据类型需要可转换成self与batch1@batch2推导后的数据类型。
39 * @param [in] alpha: host侧的aclScalar，数据类型需要可转换成self与batch1@batch2推导后的数据类型。39 * @param [in] alpha: host侧的aclScalar，数据类型需要可转换成self与batch1@batch2推导后的数据类型。
40 * @param [in] cubeMathType:40 * @param [in] cubeMathType:
41- * INT8类型的枚举值，用于判断Cube单元应该使用那种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能41+ * INT8类型的枚举值，用于判断Cube单元应该使用哪种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能
42 * @param [in] out: npu42 * @param [in] out: npu
43 * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16类型，且数据类型需要与self保持一致，shape要求与batch1@batch2的后两维保持一致。43 * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16类型，且数据类型需要与self保持一致，shape要求与batch1@batch2的后两维保持一致。
44 * 支持非连续的Tensor，支持空Tensor传入，数据格式支持ND。44 * 支持非连续的Tensor，支持空Tensor传入，数据格式支持ND。
@@ -59,8 +59,8 @@ ACLNN_API aclnnStatus aclnnAddbmmGetWorkspaceSize(const aclTensor* self, const a
59 * 计算公式：59 * 计算公式：
60 * $$ out = βself+α(\sum_{i=0}^{b-1}batch1_{i}@batch2_{i}) $$60 * $$ out = βself+α(\sum_{i=0}^{b-1}batch1_{i}@batch2_{i}) $$
61 *61 *
62- * @param [in] workspace: 在npu device侧申请的workspace内存起址。62+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
63- * @param [in] workspace_size: 在npu device侧申请的workspace大小，由第一段接口aclnnAddbmmGetWorkspaceSize获取。63+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnAddbmmGetWorkspaceSize获取。
64 * @param [in] executor: op执行器，包含了算子计算流程。64 * @param [in] executor: op执行器，包含了算子计算流程。
65 * @param [in] stream: acl stream流。65 * @param [in] stream: acl stream流。
66 * @return aclnnStatus: 返回状态码。66 * @return aclnnStatus: 返回状态码。
@@ -88,7 +88,7 @@ ACLNN_API aclnnStatus aclnnAddbmm(void* workspace, uint64_t workspaceSize, aclOp
88 * @param [in] beta: host侧的aclScalar，数据类型需要可转换成self与batch1@batch2推导后的数据类型。88 * @param [in] beta: host侧的aclScalar，数据类型需要可转换成self与batch1@batch2推导后的数据类型。
89 * @param [in] alpha: host侧的aclScalar，数据类型需要可转换成self与batch1@batch2推导后的数据类型。89 * @param [in] alpha: host侧的aclScalar，数据类型需要可转换成self与batch1@batch2推导后的数据类型。
90 * @param [in] cubeMathType:90 * @param [in] cubeMathType:
91- * INT8类型的枚举值，用于判断Cube单元应该使用那种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能91+ * INT8类型的枚举值，用于判断Cube单元应该使用哪种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能
92 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。92 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
93 * @param [out] executor: 返回op执行器，包含算子计算流程。93 * @param [out] executor: 返回op执行器，包含算子计算流程。
94 * @return aclnnStatus: 返回状态码。94 * @return aclnnStatus: 返回状态码。
@@ -106,8 +106,8 @@ ACLNN_API aclnnStatus aclnnInplaceAddbmmGetWorkspaceSize(aclTensor* selfRef, con
106 * 计算公式：106 * 计算公式：
107 * $$ selfRef = βselfRef+α(\sum_{i=0}^{b-1}batch1_{i}@batch2_{i}) $$107 * $$ selfRef = βselfRef+α(\sum_{i=0}^{b-1}batch1_{i}@batch2_{i}) $$
108 *108 *
109- * @param [in] workspace: 在npu device侧申请的workspace内存起址。109+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
110- * @param [in] workspace_size: 在npu device侧申请的workspace大小，由第一段接口aclnnInplaceAddbmmGetWorkspaceSize获取。110+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnInplaceAddbmmGetWorkspaceSize获取。
111 * @param [in] executor: op执行器，包含了算子计算流程。111 * @param [in] executor: op执行器，包含了算子计算流程。
112 * @param [in] stream: acl stream流。112 * @param [in] stream: acl stream流。
113 * @return aclnnStatus: 返回状态码。113 * @return aclnnStatus: 返回状态码。
@@ -119,4 +119,4 @@ ACLNN_API aclnnStatus aclnnInplaceAddbmm(void* workspace, uint64_t workspaceSize
119}119}
120#endif120#endif
121 121 
122-#endif // OP_API_INC_BADDBMM_H_122+#endif // OP_API_INC_ADDBMM_H_
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_baddbmm.cpp+2-2
@@ -81,13 +81,13 @@ static bool CheckShape(const aclTensor* selfTensor, const aclTensor* batch1Tenso
81 // batch1DimNum - LAST_DIM means the last element, batch2DimNum - PENULTIMATE_DIM means the penultimate element81 // batch1DimNum - LAST_DIM means the last element, batch2DimNum - PENULTIMATE_DIM means the penultimate element
82 if (batch1[batch1DimNum - LAST_DIM] != batch2[batch2DimNum - PENULTIMATE_DIM]) {82 if (batch1[batch1DimNum - LAST_DIM] != batch2[batch2DimNum - PENULTIMATE_DIM]) {
83 OP_LOGE(ACLNN_ERR_PARAM_INVALID,83 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
84- "batch1's last dim and batch2's penultimate dim shoule be same, batch1 [%ld], batch2 [%ld].",84+ "batch1's last dim and batch2's penultimate dim should be same, batch1 [%ld], batch2 [%ld].",
85 batch1[batch1DimNum - LAST_DIM], batch2[batch2DimNum - PENULTIMATE_DIM]);85 batch1[batch1DimNum - LAST_DIM], batch2[batch2DimNum - PENULTIMATE_DIM]);
86 return false;86 return false;
87 }87 }
88 88 
89 if (!CheckBatchDimBroadcast(batch1DimNum, batch2DimNum, batch1, batch2)) {89 if (!CheckBatchDimBroadcast(batch1DimNum, batch2DimNum, batch1, batch2)) {
90- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "mat1's batch dim and mat2's batch dim can not broadcast");90+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "mat1's batch dim and mat2's batch dim cannot broadcast");
91 return false;91 return false;
92 }92 }
93 93 
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_baddbmm.h+6-6
@@ -37,7 +37,7 @@ extern "C" {
37 * @param [in] beta: host侧的aclScalar，默认为137 * @param [in] beta: host侧的aclScalar，默认为1
38 * @param [in] alpha: host侧的aclScalar，默认为138 * @param [in] alpha: host侧的aclScalar，默认为1
39 * @param [in] cubeMathType:39 * @param [in] cubeMathType:
40- * INT8类型的枚举值，用于判断Cube单元应该使用那种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能40+ * INT8类型的枚举值，用于判断Cube单元应该使用哪种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能
41 * @param [in] out: npu41 * @param [in] out: npu
42 * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16类型，dtype和format均需要与self、batch1@batch2保持一致。42 * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16类型，dtype和format均需要与self、batch1@batch2保持一致。
43 * 支持非连续的Tensor，数据格式支持ND。输出的shape要求与batch1@batch2的shape保持一致。43 * 支持非连续的Tensor，数据格式支持ND。输出的shape要求与batch1@batch2的shape保持一致。
@@ -57,8 +57,8 @@ ACLNN_API aclnnStatus aclnnBaddbmmGetWorkspaceSize(const aclTensor* self, const
57 * 计算公式：计算α与batch1、batch2的矩阵乘结果的乘积，再与β和self的乘积求和57 * 计算公式：计算α与batch1、batch2的矩阵乘结果的乘积，再与β和self的乘积求和
58 * $$ out = βself+α(batch1@batch2) $$58 * $$ out = βself+α(batch1@batch2) $$
59 *59 *
60- * @param [in] workspace: 在npu device侧申请的workspace内存起址。60+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
61- * @param [in] workspace_size: 在npu device侧申请的workspace大小，由第一段接口aclnnBaddbmmGetWorkspaceSize获取。61+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnBaddbmmGetWorkspaceSize获取。
62 * @param [in] executor: op执行器，包含了算子计算流程。62 * @param [in] executor: op执行器，包含了算子计算流程。
63 * @param [in] stream: acl stream流。63 * @param [in] stream: acl stream流。
64 * @return aclnnStatus: 返回状态码。64 * @return aclnnStatus: 返回状态码。
@@ -86,7 +86,7 @@ ACLNN_API aclnnStatus aclnnBaddbmm(void* workspace, uint64_t workspaceSize, aclO
86 * @param [in] beta: host侧的aclScalar，默认为186 * @param [in] beta: host侧的aclScalar，默认为1
87 * @param [in] alpha: host侧的aclScalar，默认为187 * @param [in] alpha: host侧的aclScalar，默认为1
88 * @param [in] cubeMathType:88 * @param [in] cubeMathType:
89- * INT8类型的枚举值，用于判断Cube单元应该使用那种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能89+ * INT8类型的枚举值，用于判断Cube单元应该使用哪种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能
90 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。90 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
91 * @param [out] executor: 返回op执行器，包含算子计算流程。91 * @param [out] executor: 返回op执行器，包含算子计算流程。
92 * @return aclnnStatus: 返回状态码。92 * @return aclnnStatus: 返回状态码。
@@ -103,8 +103,8 @@ ACLNN_API aclnnStatus aclnnInplaceBaddbmmGetWorkspaceSize(const aclTensor* selfR
103 * 计算公式：计算α与batch1、batch2的矩阵乘结果的乘积，再与β和self的乘积求和103 * 计算公式：计算α与batch1、batch2的矩阵乘结果的乘积，再与β和self的乘积求和
104 * $$ out = βself+α(batch1@batch2) $$104 * $$ out = βself+α(batch1@batch2) $$
105 *105 *
106- * @param [in] workspace: 在npu device侧申请的workspace内存起址。106+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
107- * @param [in] workspace_size: 在npu device侧申请的workspace大小，由第一段接口aclnnInplaceBaddbmmGetWorkspaceSize获取。107+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnInplaceBaddbmmGetWorkspaceSize获取。
108 * @param [in] executor: op执行器，包含了算子计算流程。108 * @param [in] executor: op执行器，包含了算子计算流程。
109 * @param [in] stream: acl stream流。109 * @param [in] stream: acl stream流。
110 * @return aclnnStatus: 返回状态码。110 * @return aclnnStatus: 返回状态码。
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_batch_matmul.cpp+9-10
@@ -102,22 +102,21 @@ static bool CheckShape(const aclTensor* selfTensor, const aclTensor* otherTensor
102 }102 }
103 // selfDimNum - 1 means self's last dim, and otherDimNum - 2 means mat2's penultimate dim103 // selfDimNum - 1 means self's last dim, and otherDimNum - 2 means mat2's penultimate dim
104 if (selfDimNum < 2 || otherDimNum < 2 || outDimNum < 2) {104 if (selfDimNum < 2 || otherDimNum < 2 || outDimNum < 2) {
105- OP_LOGE(105+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
106- ACLNN_ERR_PARAM_INVALID,106+ "dim of self, other or out must be > 2, actual selfDimNum [%zu], otherDimNum [%zu], outDimNum [%zu].",
107- "shapedim of self, other or out must > 2, actual selfshapeDim [%zu], otherDimNum [%zu] , outDimNum [%zu].",107+ selfDimNum, otherDimNum, outDimNum);
108- selfDimNum, otherDimNum, outDimNum);
109 return false;108 return false;
110 }109 }
111 if (self[selfDimNum - 1] != other[otherDimNum - PENULTIMATE_DIM]) {110 if (self[selfDimNum - 1] != other[otherDimNum - PENULTIMATE_DIM]) {
112 OP_LOGE(ACLNN_ERR_PARAM_INVALID,111 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
113- "self's last dim and mat2's penultimate dim shoule be same, self [%ld], mat2 [%ld].",112+ "self's last dim and mat2's penultimate dim should be same, self [%ld], mat2 [%ld].",
114 self[selfDimNum - LAST_DIM], other[otherDimNum - PENULTIMATE_DIM]);113 self[selfDimNum - LAST_DIM], other[otherDimNum - PENULTIMATE_DIM]);
115 return false;114 return false;
116 }115 }
117 if (self[FIRST_DIM] != other[FIRST_DIM] && self[FIRST_DIM] != 1 && other[FIRST_DIM] != 1) {116 if (self[FIRST_DIM] != other[FIRST_DIM] && self[FIRST_DIM] != 1 && other[FIRST_DIM] != 1) {
118 OP_LOGE(ACLNN_ERR_PARAM_INVALID,117 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
119- "self's first dim and mat2's first dim shoule be same, or at least one of the self's first dim and "118+ "self's first dim and mat2's first dim should be same, or at least one of the self's first dim and "
120- "mat2's first dim is 1.Now self [%ld], mat2 [%ld].",119+ "mat2's first dim is 1. Now self [%ld], mat2 [%ld].",
121 self[FIRST_DIM], other[FIRST_DIM]);120 self[FIRST_DIM], other[FIRST_DIM]);
122 return false;121 return false;
123 }122 }
@@ -125,7 +124,7 @@ static bool CheckShape(const aclTensor* selfTensor, const aclTensor* otherTensor
125 if (out[outDimNum - PENULTIMATE_DIM] != self[selfDimNum - PENULTIMATE_DIM] ||124 if (out[outDimNum - PENULTIMATE_DIM] != self[selfDimNum - PENULTIMATE_DIM] ||
126 out[outDimNum - LAST_DIM] != other[otherDimNum - LAST_DIM] || out[FIRST_DIM] != firstDim) {125 out[outDimNum - LAST_DIM] != other[otherDimNum - LAST_DIM] || out[FIRST_DIM] != firstDim) {
127 OP_LOGE(ACLNN_ERR_PARAM_INVALID,126 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
128- "output's shape is not match input, out_m[%ld] must be same with self_m[%ld], "127+ "output's shape does not match input, out_m[%ld] must be same with self_m[%ld], "
129 "out_n[%ld] must be same with other_n[%ld], out_batch[%ld] must be same with input_batch[%ld].",128 "out_n[%ld] must be same with other_n[%ld], out_batch[%ld] must be same with input_batch[%ld].",
130 out[outDimNum - PENULTIMATE_DIM], self[selfDimNum - PENULTIMATE_DIM], out[outDimNum - LAST_DIM],129 out[outDimNum - PENULTIMATE_DIM], self[selfDimNum - PENULTIMATE_DIM], out[outDimNum - LAST_DIM],
131 other[otherDimNum - LAST_DIM], out[FIRST_DIM], firstDim);130 other[otherDimNum - LAST_DIM], out[FIRST_DIM], firstDim);
@@ -179,7 +178,7 @@ static inline bool CheckStorageShape(const aclTensor* otherTensor)
179 auto storageShapeDim = storageShape.GetDimNum();178 auto storageShapeDim = storageShape.GetDimNum();
180 OP_CHECK(179 OP_CHECK(
181 storageShapeDim == 5,180 storageShapeDim == 5,
182- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Only support mat2 storageShapeDim is 5, which are [%zu].", storageShapeDim),181+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Only support mat2 storageShapeDim is 5, which is [%zu].", storageShapeDim),
183 return false);182 return false);
184 return true;183 return true;
185}184}
@@ -196,7 +195,7 @@ bool CheckDtypeValidWeightNz(const aclTensor* self, const aclTensor* mat2, const
196{195{
197 auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();196 auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
198 if ((npuArch != NpuArch::DAV_2201) && !IsNpuArch3510Series()) {197 if ((npuArch != NpuArch::DAV_2201) && !IsNpuArch3510Series()) {
199- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "batchmatmulweightnz is unsupported in this npu arch");198+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "BatchMatMulWeightNz is unsupported in this NPU arch.");
200 return false;199 return false;
201 }200 }
202 bool enable16In32Out = NeedEnableFp32Output(self->GetDataType(), mat2->GetDataType(), out->GetDataType(),201 bool enable16In32Out = NeedEnableFp32Output(self->GetDataType(), mat2->GetDataType(), out->GetDataType(),
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_batch_matmul.h+4-4
@@ -47,7 +47,7 @@ ACLNN_API aclnnStatus aclnnBatchMatMul(void* workspace, uint64_t workspaceSize,
47 * device侧的aclTensor，数据类型支持FLOAT16、BFLOAT16类型，且数据类型需要与self保持一致，shape要求与self@mat2的后两维保持一致。47 * device侧的aclTensor，数据类型支持FLOAT16、BFLOAT16类型，且数据类型需要与self保持一致，shape要求与self@mat2的后两维保持一致。
48 * 支持非连续的Tensor，支持空Tensor传入，数据格式支持ND。48 * 支持非连续的Tensor，支持空Tensor传入，数据格式支持ND。
49 * @param [in] cubeMathType:49 * @param [in] cubeMathType:
50- * INT8类型的枚举值，用于判断Cube单元应该使用那种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能50+ * INT8类型的枚举值，用于判断Cube单元应该使用哪种计算逻辑进行运算，可通过此开关使能如HFLOAT32等功能
51 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。51 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
52 * @param [out] executor: 返回op执行器，包含算子计算流程。52 * @param [out] executor: 返回op执行器，包含算子计算流程。
53 * @return aclnnStatus: 返回状态码。53 * @return aclnnStatus: 返回状态码。
@@ -58,8 +58,8 @@ ACLNN_API aclnnStatus aclnnBatchMatMulWeightNzGetWorkspaceSize(const aclTensor*
58/**58/**
59 * @brief aclnnBatchMatMulWeightNz的第二段接口，用于执行计算。59 * @brief aclnnBatchMatMulWeightNz的第二段接口，用于执行计算。
60 *60 *
61- * @param [in] workspace: 在npu device侧申请的workspace内存起址。61+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
62- * @param [in] workspace_size: 在npu62+ * @param [in] workspaceSize: 在npu
63 * device侧申请的workspace大小，由第一段接口aclnnBatchMatMulWeightNzGetWorkspaceSize获取。63 * device侧申请的workspace大小，由第一段接口aclnnBatchMatMulWeightNzGetWorkspaceSize获取。
64 * @param [in] executor: op执行器，包含了算子计算流程。64 * @param [in] executor: op执行器，包含了算子计算流程。
65 * @param [in] stream: acl stream流。65 * @param [in] stream: acl stream流。
@@ -72,4 +72,4 @@ ACLNN_API aclnnStatus aclnnBatchMatMulWeightNz(void* workspace, uint64_t workspa
72}72}
73#endif73#endif
74 74 
75-#endif // OP_API_ACLNN_BATCHMATMUL_H75+#endif // OP_API_ACLNN_BATCHMATMUL_H
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_einsum.cpp+18-18
@@ -142,24 +142,24 @@ aclnnStatus HandleABCDxABCED2ABCE(const aclTensorList* tensors, aclTensor* outpu
142 0 :142 0 :
143 g_useFP16;143 g_useFP16;
144 144 
145- auto tensor0Contigous = l0op::Contiguous((*tensors)[0], uniqueExecutor.get());145+ auto tensor0Contiguous = l0op::Contiguous((*tensors)[0], uniqueExecutor.get());
146- auto tensor1Contigous = l0op::Contiguous((*tensors)[1], uniqueExecutor.get());146+ auto tensor1Contiguous = l0op::Contiguous((*tensors)[1], uniqueExecutor.get());
147- CHECK_RET(tensor0Contigous != nullptr, ACLNN_ERR_INNER_NULLPTR);147+ CHECK_RET(tensor0Contiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
148- CHECK_RET(tensor1Contigous != nullptr, ACLNN_ERR_INNER_NULLPTR);148+ CHECK_RET(tensor1Contiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
149 149 
150- const aclTensor* tensor0Cast = tensor0Contigous;150+ const aclTensor* tensor0Cast = tensor0Contiguous;
151- const aclTensor* tensor1Cast = tensor1Contigous;151+ const aclTensor* tensor1Cast = tensor1Contiguous;
152 const aclTensor* outputCast = output;152 const aclTensor* outputCast = output;
153 153 
154 auto inputDtype = (*tensors)[0]->GetDataType();154 auto inputDtype = (*tensors)[0]->GetDataType();
155 if (inputDtype != op::DataType::DT_FLOAT && inputDtype != op::DataType::DT_FLOAT16) {155 if (inputDtype != op::DataType::DT_FLOAT && inputDtype != op::DataType::DT_FLOAT16) {
156- tensor0Cast = l0op::Cast(tensor0Contigous, op::DataType::DT_FLOAT16, uniqueExecutor.get());156+ tensor0Cast = l0op::Cast(tensor0Contiguous, op::DataType::DT_FLOAT16, uniqueExecutor.get());
157 CHECK_RET(tensor0Cast != nullptr, ACLNN_ERR_INNER_NULLPTR);157 CHECK_RET(tensor0Cast != nullptr, ACLNN_ERR_INNER_NULLPTR);
158- tensor1Cast = l0op::Cast(tensor1Contigous, op::DataType::DT_FLOAT16, uniqueExecutor.get());158+ tensor1Cast = l0op::Cast(tensor1Contiguous, op::DataType::DT_FLOAT16, uniqueExecutor.get());
159 CHECK_RET(tensor1Cast != nullptr, ACLNN_ERR_INNER_NULLPTR);159 CHECK_RET(tensor1Cast != nullptr, ACLNN_ERR_INNER_NULLPTR);
160- auto outputContigous = l0op::Contiguous(output, uniqueExecutor.get());160+ auto outputContiguous = l0op::Contiguous(output, uniqueExecutor.get());
161- CHECK_RET(outputContigous != nullptr, ACLNN_ERR_INNER_NULLPTR);161+ CHECK_RET(outputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
162- outputCast = l0op::Cast(outputContigous, op::DataType::DT_FLOAT16, uniqueExecutor.get());162+ outputCast = l0op::Cast(outputContiguous, op::DataType::DT_FLOAT16, uniqueExecutor.get());
163 CHECK_RET(outputCast != nullptr, ACLNN_ERR_INNER_NULLPTR);163 CHECK_RET(outputCast != nullptr, ACLNN_ERR_INNER_NULLPTR);
164 }164 }
165 auto expandA = l0op::UnsqueezeNd(tensor0Cast, DIM_FOUR, uniqueExecutor.get());165 auto expandA = l0op::UnsqueezeNd(tensor0Cast, DIM_FOUR, uniqueExecutor.get());
@@ -197,13 +197,13 @@ aclnnStatus HandleAxB2AB(const aclTensorList* tensors, aclTensor* output, uint64
197 return ACLNN_ERR_PARAM_INVALID); // 校验tensorList中第2个Tensor的dimNum为1197 return ACLNN_ERR_PARAM_INVALID); // 校验tensorList中第2个Tensor的dimNum为1
198 OP_CHECK_WRONG_DIMENSION(output, 2, return ACLNN_ERR_PARAM_INVALID); // 校验Tensor output的dimNum为2198 OP_CHECK_WRONG_DIMENSION(output, 2, return ACLNN_ERR_PARAM_INVALID); // 校验Tensor output的dimNum为2
199 199 
200- auto tensor0Contigous = l0op::Contiguous((*tensors)[0], uniqueExecutor.get());200+ auto tensor0Contiguous = l0op::Contiguous((*tensors)[0], uniqueExecutor.get());
201- CHECK_RET(tensor0Contigous != nullptr, ACLNN_ERR_INNER_NULLPTR);201+ CHECK_RET(tensor0Contiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
202- auto tensor1Contigous = l0op::Contiguous((*tensors)[1], uniqueExecutor.get());202+ auto tensor1Contiguous = l0op::Contiguous((*tensors)[1], uniqueExecutor.get());
203- CHECK_RET(tensor1Contigous != nullptr, ACLNN_ERR_INNER_NULLPTR);203+ CHECK_RET(tensor1Contiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
204 204 
205- auto leftMatrix = l0op::UnsqueezeNd(tensor0Contigous, DIM_ONE, uniqueExecutor.get());205+ auto leftMatrix = l0op::UnsqueezeNd(tensor0Contiguous, DIM_ONE, uniqueExecutor.get());
206- auto rightMatrix = l0op::UnsqueezeNd(tensor1Contigous, DIM_ZERO, uniqueExecutor.get());206+ auto rightMatrix = l0op::UnsqueezeNd(tensor1Contiguous, DIM_ZERO, uniqueExecutor.get());
207 auto result = l0op::Mul(leftMatrix, rightMatrix, uniqueExecutor.get());207 auto result = l0op::Mul(leftMatrix, rightMatrix, uniqueExecutor.get());
208 208 
209 // 固定写法，将计算结果拷贝到输出 output output可能是非连续的tensor209 // 固定写法，将计算结果拷贝到输出 output output可能是非连续的tensor
@@ -290,7 +290,7 @@ static inline bool CheckTensorValid(const aclTensorList* tensors, const aclTenso
290 auto input1Dtype = (*tensors)[1]->GetDataType();290 auto input1Dtype = (*tensors)[1]->GetDataType();
291 auto outputDtype = output->GetDataType();291 auto outputDtype = output->GetDataType();
292 if ((input0Dtype != input1Dtype) || (input0Dtype != outputDtype)) { // 校验三个tensor的数据类型一致292 if ((input0Dtype != input1Dtype) || (input0Dtype != outputDtype)) { // 校验三个tensor的数据类型一致
293- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "all inputs dtype is not equal, please check.");293+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "all inputs dtype are not equal, please check.");
294 return false;294 return false;
295 }295 }
296 return true;296 return true;
Mmatmul/batch_mat_mul_v3/op_host/op_api/aclnn_einsum.h+3-3
@@ -28,7 +28,7 @@ extern "C" {
28 * @param [in] equation: 输入字符串，数据类型支持const char *。28 * @param [in] equation: 输入字符串，数据类型支持const char *。
29 * @param [out] output:29 * @param [out] output:
30 * 输出Tensor，数据类型支持FLOAT16、FLOAT、INT16、UINT16、INT32、UINT32、INT64、UINT64。支持非连续Tensor，数据格式支持ND。30 * 输出Tensor，数据类型支持FLOAT16、FLOAT、INT16、UINT16、INT32、UINT32、INT64、UINT64。支持非连续Tensor，数据格式支持ND。
31- * @param [out] workspaceSize：返回用户需要在npu device侧申的的workspace大小。31+ * @param [out] workspaceSize：返回用户需要在npu device侧申请的workspace大小。
32 * @param [out] executor: 返回op执行器，包含了算子计算流程。32 * @param [out] executor: 返回op执行器，包含了算子计算流程。
33 * @return aclnnStatus: 返回状态码33 * @return aclnnStatus: 返回状态码
34 */34 */
@@ -37,7 +37,7 @@ ACLNN_API aclnnStatus aclnnEinsumGetWorkspaceSize(const aclTensorList* tensors,
37 37 
38/**38/**
39 * @brief aclnnEinsum的第二段接口，用于执行计算。39 * @brief aclnnEinsum的第二段接口，用于执行计算。
40- * @param [in] workspace: 在npu device侧申请的workspace内存起址。40+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
41 * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnEinsumGetWorkspaceSize获取。41 * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnEinsumGetWorkspaceSize获取。
42 * @param [in] stream: acl stream流。42 * @param [in] stream: acl stream流。
43 * @param [in] executor: op执行器，包含了算子计算流程。43 * @param [in] executor: op执行器，包含了算子计算流程。
@@ -49,4 +49,4 @@ ACLNN_API aclnnStatus aclnnEinsum(void* workspace, uint64_t workspaceSize, aclOp
49}49}
50#endif50#endif
51 51 
52-#endif // OP_API_INC_EINSUM_H_52+#endif // OP_API_INC_EINSUM_H_
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_asw_basic_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_asw_basic_tiling.cc12+ * \file batch_matmul_v3_asw_basic_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_asw_bl1_full_load_basic_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_asw_bl1_full_load_basic_tiling.cc12+ * \file batch_matmul_v3_asw_bl1_full_load_basic_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_asw_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_asw_tiling.cc12+ * \file batch_matmul_v3_asw_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_iterbatch_basicapi_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_iterbatch_basicapi_tiling.cc12+ * \file batch_matmul_v3_iterbatch_basicapi_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_iterbatch_basicapi_tiling.h+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_iterbatch_tiling.h12+ * \file batch_matmul_v3_iterbatch_basicapi_tiling.h
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_iterbatch_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_iterbatch_tiling.cc12+ * \file batch_matmul_v3_iterbatch_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_k_equal_zero_tiling.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_k_equal_zero_tiling.cc12+ * \file batch_matmul_v3_k_equal_zero_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
@@ -66,4 +66,4 @@ ge::graphStatus BatchMatMulV3KEqZeroTiling::GetTilingData(TilingResult& tiling)
66 return GetTilingDataImpl<MatMulV3KEqZeroBasicTilingData>(tiling);66 return GetTilingDataImpl<MatMulV3KEqZeroBasicTilingData>(tiling);
67}67}
68} // namespace batch_matmul_v3_advanced68} // namespace batch_matmul_v3_advanced
69-} // namespace optiling69+} // namespace optiling
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_matmul2mul_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_matmul2mul_tiling.cc12+ * \file batch_matmul_v3_matmul2mul_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_mergebatch_basicapi_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_mergebatch_basicapi_tiling.cc12+ * \file batch_matmul_v3_mergebatch_basicapi_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15#include <cmath>15#include <cmath>
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_mergebatch_basicapi_tiling.h+3-3
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_mergebatch_tiling.h12+ * \file batch_matmul_v3_mergebatch_basicapi_tiling.h
13 * \brief13 * \brief
14 */14 */
15 15 
@@ -23,9 +23,9 @@ using namespace matmul_v3_advanced;
23class BatchMatMulV3MergeBatchBasicApiTiling : public MatMulV3BaseTiling {23class BatchMatMulV3MergeBatchBasicApiTiling : public MatMulV3BaseTiling {
24public:24public:
25 BatchMatMulV3MergeBatchBasicApiTiling(gert::TilingContext* context, MatMulTilingCfg& cfg)25 BatchMatMulV3MergeBatchBasicApiTiling(gert::TilingContext* context, MatMulTilingCfg& cfg)
26- : MatMulV3BaseTiling(context, cfg){};26+ : MatMulV3BaseTiling(context, cfg) {};
27 27 
28- ~BatchMatMulV3MergeBatchBasicApiTiling() override{};28+ ~BatchMatMulV3MergeBatchBasicApiTiling() override {};
29 29 
30protected:30protected:
31 bool IsCapable() override;31 bool IsCapable() override;
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/arch35/batch_matmul_v3_tiling_key.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_tiling_key.cc12+ * \file batch_matmul_v3_tiling_key.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
@@ -24,4 +24,4 @@ uint64_t BatchMatMulV3TilingKey::GetTilingKey() const
24 static_cast<uint64_t>(btrans_), static_cast<uint64_t>(batchModel_),24 static_cast<uint64_t>(btrans_), static_cast<uint64_t>(batchModel_),
25 static_cast<uint64_t>(model_), static_cast<uint64_t>(fullLoad_),25 static_cast<uint64_t>(model_), static_cast<uint64_t>(fullLoad_),
26 static_cast<uint64_t>(out_));26 static_cast<uint64_t>(out_));
27-}27+}
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/batch_mat_mul_v3_base_tiling.cpp+12-12
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file batch_mat_mul_v3_base_tiling.cc12+ * \file batch_mat_mul_v3_base_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15#include "batch_mat_mul_v3_base_tiling.h"15#include "batch_mat_mul_v3_base_tiling.h"
@@ -209,7 +209,7 @@ void BatchMatmulV3BaseTiling::MergeBatchAndMAxis()
209 return;209 return;
210}210}
211 211 
212-bool BatchMatmulV3BaseTiling::CheckBMMTilingDataIsVaild() const212+bool BatchMatmulV3BaseTiling::CheckBMMTilingDataIsValid() const
213{213{
214 return (optiling::matmul_v3::CheckNumberIsValid(batchInfo_.batchA3, args_.opName, "batchInfo_.batchA3") ||214 return (optiling::matmul_v3::CheckNumberIsValid(batchInfo_.batchA3, args_.opName, "batchInfo_.batchA3") ||
215 optiling::matmul_v3::CheckNumberIsValid(batchInfo_.batchA2, args_.opName, "batchInfo_.batchA2") ||215 optiling::matmul_v3::CheckNumberIsValid(batchInfo_.batchA2, args_.opName, "batchInfo_.batchA2") ||
@@ -357,7 +357,7 @@ ge::graphStatus BatchMatmulV3BaseTiling::DoLibApiTiling()
357 auto ret = MatmulV3BaseTiling::DoLibApiTiling();357 auto ret = MatmulV3BaseTiling::DoLibApiTiling();
358 SetBatchDimInfo();358 SetBatchDimInfo();
359 CalcBatchDimAll();359 CalcBatchDimAll();
360- if (CheckBMMTilingDataIsVaild()) {360+ if (CheckBMMTilingDataIsValid()) {
361 return ge::GRAPH_FAILED;361 return ge::GRAPH_FAILED;
362 }362 }
363 bmmTilingData_.multiBatchInfo.biasWithBatch = static_cast<uint32_t>(batchInfo_.biasWithBatch);363 bmmTilingData_.multiBatchInfo.biasWithBatch = static_cast<uint32_t>(batchInfo_.biasWithBatch);
@@ -386,7 +386,7 @@ ge::graphStatus BatchMatmulV3BaseTiling::DoLibApiTiling()
386/*386/*
387 * Algorithm to calculate the best (baseM, baseN) that gives even workload amongst iterations.387 * Algorithm to calculate the best (baseM, baseN) that gives even workload amongst iterations.
388 * Parameter `divisor` is used to control the starting point of the algorithm.388 * Parameter `divisor` is used to control the starting point of the algorithm.
389- * Choosing different starting point can sometimes get better performence.389+ * Choosing different starting point can sometimes get better performance.
390 * The starting point of `divisor = 2` is half of that of `divisor = 1`390 * The starting point of `divisor = 2` is half of that of `divisor = 1`
391 */391 */
392static void CalcBaseMN(uint64_t& baseM, uint64_t& baseN, const matmul_v3::MatmulV3Args& args, uint64_t divisor = 1UL)392static void CalcBaseMN(uint64_t& baseM, uint64_t& baseN, const matmul_v3::MatmulV3Args& args, uint64_t divisor = 1UL)
@@ -845,9 +845,9 @@ void BatchMatmulV3BaseTiling::CalculateNd2nzWorkspaceSize()
845 * Func: tune down parameter x until either y(x) is below target value, or x has reached its minimum845 * Func: tune down parameter x until either y(x) is below target value, or x has reached its minimum
846 * Args: `target` - target value846 * Args: `target` - target value
847 * `y` - current value of y(x)847 * `y` - current value of y(x)
848- * `x` - the paramter being tuned848+ * `x` - the parameter being tuned
849 * `dydx` - the slope, dy/dx849 * `dydx` - the slope, dy/dx
850- * `step` - the step length (i.e. the minimul variation value) of x. Default is 1.850+ * `step` - the step length (i.e. the minimal variation value) of x. Default is 1.
851 * Note: `x` is assumed to be aligned to `step`, that makes `step` the minimum of `x`.851 * Note: `x` is assumed to be aligned to `step`, that makes `step` the minimum of `x`.
852 */852 */
853static void TuneDownParam(uint64_t target, uint64_t& y, uint64_t& x, uint64_t dydx, uint64_t step = 1UL)853static void TuneDownParam(uint64_t target, uint64_t& y, uint64_t& x, uint64_t dydx, uint64_t step = 1UL)
@@ -919,9 +919,9 @@ static void AL1FullLoadTiling(const matmul_v3::MatmulV3Args& args, uint64_t l1Si
919 const uint64_t biasSizePerStepN = args.hasBias ? baseN * ge::GetSizeByDataType(args.biasType) * NUM_TWO : 0UL;919 const uint64_t biasSizePerStepN = args.hasBias ? baseN * ge::GetSizeByDataType(args.biasType) * NUM_TWO : 0UL;
920 uint64_t loadSize = aSize + (bSizePerStepN + biasSizePerStepN) * stepN;920 uint64_t loadSize = aSize + (bSizePerStepN + biasSizePerStepN) * stepN;
921 // Tune down loadSize until it is fully loaded in L1921 // Tune down loadSize until it is fully loaded in L1
922- // Stage 1: try tunning stepN922+ // Stage 1: try tuning stepN
923 TuneDownParam(l1Size, loadSize, stepN, bSizePerStepN + biasSizePerStepN);923 TuneDownParam(l1Size, loadSize, stepN, bSizePerStepN + biasSizePerStepN);
924- // Stage 2: stepN has reached 1 yet loadSize's still too big for L1, tune down stepKb924+ // Stage 2: stepN has reached 1 yet loadSize is still too big for L1, tune down stepKb
925 const uint64_t bSizePerStepKb = bSizePerStepN / stepKb;925 const uint64_t bSizePerStepKb = bSizePerStepN / stepKb;
926 TuneDownParam(l1Size, loadSize, stepKb, bSizePerStepKb);926 TuneDownParam(l1Size, loadSize, stepKb, bSizePerStepKb);
927 // Stage 3: stepN & stepKb have both reached 1, tune down baseN927 // Stage 3: stepN & stepKb have both reached 1, tune down baseN
@@ -959,10 +959,10 @@ static void BL1FullLoadTiling(const matmul_v3::MatmulV3Args& args, uint64_t l1Si
959 const uint64_t biasSize = args.hasBias ? baseN * stepN * ge::GetSizeByDataType(args.biasType) * NUM_TWO : 0UL;959 const uint64_t biasSize = args.hasBias ? baseN * stepN * ge::GetSizeByDataType(args.biasType) * NUM_TWO : 0UL;
960 const uint64_t aSizePerStepM = stepKa * baseK * baseM * ge::GetSizeByDataType(args.aType) * NUM_TWO;960 const uint64_t aSizePerStepM = stepKa * baseK * baseM * ge::GetSizeByDataType(args.aType) * NUM_TWO;
961 uint64_t loadSize = bSize + biasSize + aSizePerStepM * stepM;961 uint64_t loadSize = bSize + biasSize + aSizePerStepM * stepM;
962- // Tune down loadSize util it is fully loaded in L1962+ // Tune down loadSize until it is fully loaded in L1
963- // Stage 1: try tunning stepM963+ // Stage 1: try tuning stepM
964 TuneDownParam(l1Size, loadSize, stepM, aSizePerStepM);964 TuneDownParam(l1Size, loadSize, stepM, aSizePerStepM);
965- // Stage 2: stepM has reached 1 yet loadSize's still too big for L1, tune down stepKa965+ // Stage 2: stepM has reached 1 yet loadSize is still too big for L1, tune down stepKa
966 const uint64_t aSizePerStepKa = aSizePerStepM / stepKa;966 const uint64_t aSizePerStepKa = aSizePerStepM / stepKa;
967 TuneDownParam(l1Size, loadSize, stepKa, aSizePerStepKa);967 TuneDownParam(l1Size, loadSize, stepKa, aSizePerStepKa);
968 // Stage 3: stepM & stepKa have both reached 1, tune down baseM968 // Stage 3: stepM & stepKa have both reached 1, tune down baseM
@@ -1341,4 +1341,4 @@ void BatchMatmulV3BaseTiling::DoTilingKeyCustom()
1341}1341}
1342 1342 
1343} // namespace optiling1343} // namespace optiling
1344-}1344+} // namespace optiling
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/batch_mat_mul_v3_base_tiling.h+1-1
@@ -125,7 +125,7 @@ protected:
125 void CalcBatchDimAll();125 void CalcBatchDimAll();
126 bool CheckNd2NzOnTheFlyLimit();126 bool CheckNd2NzOnTheFlyLimit();
127 void DoMultiBatchAndL1FullLoadTiling();127 void DoMultiBatchAndL1FullLoadTiling();
128- bool CheckBMMTilingDataIsVaild() const;128+ bool CheckBMMTilingDataIsValid() const;
129 void DoTilingKeyCustom();129 void DoTilingKeyCustom();
130 void DoUnAlignCommonTiling();130 void DoUnAlignCommonTiling();
131 void DoMultiBatchTiling();131 void DoMultiBatchTiling();
Mmatmul/batch_mat_mul_v3/op_host/op_tiling/batch_mat_mul_v3_tiling.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file batch_mat_mul_v3_tiling.cc12+ * \file batch_mat_mul_v3_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15#include "batch_mat_mul_v3_tiling.h"15#include "batch_mat_mul_v3_tiling.h"
@@ -99,4 +99,4 @@ IMPL_OP_OPTILING(BatchMatMulV3)
99 .Tiling(BatchMatMulV3TilingFunc)99 .Tiling(BatchMatMulV3TilingFunc)
100 .TilingParse<MatmulV3CompileInfo>(TilingPrepareForBatchMatMulV3)100 .TilingParse<MatmulV3CompileInfo>(TilingPrepareForBatchMatMulV3)
101 .GenSimplifiedKey(GenSimplifiedKey);101 .GenSimplifiedKey(GenSimplifiedKey);
102-} // namespace optiling102+} // namespace optiling
Mmatmul/batch_mat_mul_v3/op_kernel/arch35/batch_mat_mul_v3.cpp+5-5
@@ -156,7 +156,7 @@ constexpr CubeFormat format_y = CubeFormat::ND;
156 } \156 } \
157 } while (0)157 } while (0)
158 158 
159-#define BMMV3_IMPL_CLASS_COMMON_TRNAS(transA, transB, templateClass, ...) \159+#define BMMV3_IMPL_CLASS_COMMON_TRANS(transA, transB, templateClass, ...) \
160 do { \160 do { \
161 GET_TILING_DATA(tilingData, tilingGM); \161 GET_TILING_DATA(tilingData, tilingGM); \
162 using cType = MatmulType<AscendC::TPosition::GM, format_y, DTYPE_Y>; \162 using cType = MatmulType<AscendC::TPosition::GM, format_y, DTYPE_Y>; \
@@ -200,7 +200,7 @@ __global__ __aicore__ void batch_mat_mul_v3(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR bi
200 if constexpr (BATCH_API_LEVEL == MAT_MUL_HIGH_LEVEL && BMODEL == MAT_MUL_BASIC &&200 if constexpr (BATCH_API_LEVEL == MAT_MUL_HIGH_LEVEL && BMODEL == MAT_MUL_BASIC &&
201 BATCH_FULL_LOAD == MAT_MUL_NO_FULL_LOAD && BATCH_L0C2OUT_MODEL == MAT_MUL_ON_THE_FLY &&201 BATCH_FULL_LOAD == MAT_MUL_NO_FULL_LOAD && BATCH_L0C2OUT_MODEL == MAT_MUL_ON_THE_FLY &&
202 BATCH_ITER_MODEL == MAT_MUL_FOR_BATCH) {202 BATCH_ITER_MODEL == MAT_MUL_FOR_BATCH) {
203- BMMV3_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,203+ BMMV3_IMPL_CLASS_COMMON_TRANS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,
204 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD);204 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD);
205 } else if constexpr (BATCH_API_LEVEL == MAT_MUL_BASIC_LEVEL && BMODEL == MAT_MUL_BASIC &&205 } else if constexpr (BATCH_API_LEVEL == MAT_MUL_BASIC_LEVEL && BMODEL == MAT_MUL_BASIC &&
206 BATCH_FULL_LOAD == MAT_MUL_NO_FULL_LOAD && BATCH_L0C2OUT_MODEL == MAT_MUL_ON_THE_FLY &&206 BATCH_FULL_LOAD == MAT_MUL_NO_FULL_LOAD && BATCH_L0C2OUT_MODEL == MAT_MUL_ON_THE_FLY &&
@@ -289,7 +289,7 @@ __global__ __aicore__ void batch_mat_mul_v3(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR bi
289 GET_TILING_DATA_WITH_STRUCT(BatchMatMulV3TilingData, tilingData, tilingGM);289 GET_TILING_DATA_WITH_STRUCT(BatchMatMulV3TilingData, tilingData, tilingGM);
290#if IS_BLAZE290#if IS_BLAZE
291 if constexpr (format_x2 == CubeFormat::NZ) {291 if constexpr (format_x2 == CubeFormat::NZ) {
292- BMMV3_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,292+ BMMV3_IMPL_CLASS_COMMON_TRANS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,
293 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD);293 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD);
294 } else {294 } else {
295 BatchMatMulIterBatchBroadcastKernel<DTYPE_X1, DTYPE_X2, DTYPE_Y, DTYPE_BIAS, layoutA, layoutB, layoutC,295 BatchMatMulIterBatchBroadcastKernel<DTYPE_X1, DTYPE_X2, DTYPE_Y, DTYPE_BIAS, layoutA, layoutB, layoutC,
@@ -298,7 +298,7 @@ __global__ __aicore__ void batch_mat_mul_v3(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR bi
298 aGM, bGM, biasGM, cGM, nullptr, tilingData);298 aGM, bGM, biasGM, cGM, nullptr, tilingData);
299 }299 }
300#else300#else
301- BMMV3_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,301+ BMMV3_IMPL_CLASS_COMMON_TRANS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,
302 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD);302 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD);
303#endif303#endif
304 } else if constexpr (BATCH_API_LEVEL == MAT_MUL_TENSOR_LEVEL && BMODEL == MAT_MUL_BASIC &&304 } else if constexpr (BATCH_API_LEVEL == MAT_MUL_TENSOR_LEVEL && BMODEL == MAT_MUL_BASIC &&
@@ -309,7 +309,7 @@ __global__ __aicore__ void batch_mat_mul_v3(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR bi
309 BatchMatMulV3Advanced::BatchMatMulBroadcastKernel<DTYPE_X1, DTYPE_X2, DTYPE_Y, DTYPE_BIAS, layoutA, layoutB,309 BatchMatMulV3Advanced::BatchMatMulBroadcastKernel<DTYPE_X1, DTYPE_X2, DTYPE_Y, DTYPE_BIAS, layoutA, layoutB,
310 layoutC>(aGM, bGM, biasGM, cGM, nullptr, tilingData);310 layoutC>(aGM, bGM, biasGM, cGM, nullptr, tilingData);
311#else311#else
312- BMMV3_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,312+ BMMV3_IMPL_CLASS_COMMON_TRANS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,
313 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD);313 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD);
314#endif314#endif
315 } else if constexpr (BATCH_API_LEVEL == MAT_MUL_TENSOR_LEVEL && BMODEL == MAT_MUL_BASIC &&315 } else if constexpr (BATCH_API_LEVEL == MAT_MUL_TENSOR_LEVEL && BMODEL == MAT_MUL_BASIC &&
Mmatmul/batch_mat_mul_v3/op_kernel/arch35/batch_mat_mul_v3_matmul2mul_block_scheduler.h+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file batch_matmul_v3_matmul2mul_tiling.h12+ * \file batch_mat_mul_v3_matmul2mul_block_scheduler.h
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/batch_mat_mul_v3/op_kernel/batch_mat_mul_v3.cpp+2-2
@@ -127,7 +127,7 @@ constexpr CubeFormat format_y = CubeFormat::ND;
127 } \127 } \
128 } while (0)128 } while (0)
129 129 
130-#define BMMV3_IMPL_CLASS_COMMON_TRNAS(transA, transB, templateClass, ...) \130+#define BMMV3_IMPL_CLASS_COMMON_TRANS(transA, transB, templateClass, ...) \
131 do { \131 do { \
132 GET_TILING_DATA(tilingData, tilingGM); \132 GET_TILING_DATA(tilingData, tilingGM); \
133 using cType = MatmulType<AscendC::TPosition::GM, format_y, DTYPE_Y>; \133 using cType = MatmulType<AscendC::TPosition::GM, format_y, DTYPE_Y>; \
@@ -237,4 +237,4 @@ __global__ __aicore__ void batch_mat_mul_v3(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR bi
237#endif237#endif
238 }238 }
239#endif239#endif
240-}240+}
Mmatmul/common/op_host/op_api/cube_util.cpp+2-2
@@ -73,7 +73,7 @@ bool CheckUnSupportDtype(const aclTensor* input, const aclTensor* weight)
73int8_t routeCubeMathType4ToCubeMathType0DAV_2201(int8_t cubeMathType)73int8_t routeCubeMathType4ToCubeMathType0DAV_2201(int8_t cubeMathType)
74{74{
75 if (GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201 && cubeMathType == USE_FP32_ADD) {75 if (GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201 && cubeMathType == USE_FP32_ADD) {
76- OP_LOGD("The cubeMathType is USE_FP32_ADD for which the API does not support temporarily, "76+ OP_LOGD("The cubeMathType is USE_FP32_ADD for which the API is not temporarily supported, "
77 "route to KEEP_DTYPE instead.");77 "route to KEEP_DTYPE instead.");
78 return KEEP_DTYPE;78 return KEEP_DTYPE;
79 }79 }
@@ -343,4 +343,4 @@ bool NeedCubeGoHF32(const DataType cubeTensorPromoteType, int8_t cubeMathType)
343}343}
344 344 
345} // namespace NN345} // namespace NN
346-} // namespace Ops346+} // namespace Ops
Mmatmul/common/op_host/op_api/matmul_util.cpp+15-15
@@ -259,7 +259,7 @@ static const aclTensor* ProcessEmptyTensorWithTrans(const aclTensor* self, const
259static bool CheckSupportSingleSplitKFp16Bf16(const aclTensor* self, const aclTensor* mat2, const DataType selfDtype,259static bool CheckSupportSingleSplitKFp16Bf16(const aclTensor* self, const aclTensor* mat2, const DataType selfDtype,
260 const DataType mat2Dtype)260 const DataType mat2Dtype)
261{261{
262- // 判决门限262+ // 判断门限
263 // 1. 输入数据类型为fp16/bf16263 // 1. 输入数据类型为fp16/bf16
264 // 2. 在K轴非256字节对齐场景下，输入数据大小不超过INT32最大值264 // 2. 在K轴非256字节对齐场景下，输入数据大小不超过INT32最大值
265 // 3. K轴大于27392265 // 3. K轴大于27392
@@ -299,10 +299,10 @@ static aclnnStatus SetMatmulOpSupportInfo(const aclTensor* self, const aclTensor
299 // 判断当前Shape是否支持使用ND输入输出299 // 判断当前Shape是否支持使用ND输入输出
300 SetMmSupportFormat(self, mat2, mmOpInfo);300 SetMmSupportFormat(self, mat2, mmOpInfo);
301 301 
302- TensorInfo SpTensor_sefl = {self, mmOpInfo.support_info.self_dtype, mmOpInfo.support_info.self_format};302+ TensorInfo SpTensor_self = {self, mmOpInfo.support_info.self_dtype, mmOpInfo.support_info.self_format};
303 TensorInfo SpTensor_mat2 = {mat2, mmOpInfo.support_info.mat2_dtype, mmOpInfo.support_info.output_format};303 TensorInfo SpTensor_mat2 = {mat2, mmOpInfo.support_info.mat2_dtype, mmOpInfo.support_info.output_format};
304 304 
305- if (IsSplitk(&SpTensor_sefl, &SpTensor_mat2)) {305+ if (IsSplitk(&SpTensor_self, &SpTensor_mat2)) {
306 mmOpInfo.supporSplitK = true;306 mmOpInfo.supporSplitK = true;
307 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();307 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
308 if (npuArch == NpuArch::DAV_2002) {308 if (npuArch == NpuArch::DAV_2002) {
@@ -521,8 +521,8 @@ static const aclTensor* GetMatMulV2Op(const aclTensor* x1, const aclTensor* x2,
521 (mmOpInfo.support_info.self_dtype == DataType::DT_BF16 &&521 (mmOpInfo.support_info.self_dtype == DataType::DT_BF16 &&
522 mmOpInfo.support_info.mat2_dtype == DataType::DT_BF16);522 mmOpInfo.support_info.mat2_dtype == DataType::DT_BF16);
523 if ((enable16In32Out || (isFp32Out && bothMatFp16Bf16)) && bias == nullptr) {523 if ((enable16In32Out || (isFp32Out && bothMatFp16Bf16)) && bias == nullptr) {
524- // 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.");525+ OP_LOGI("hit matmulv2 fp16/bf16 in fp32 out case.");
526 const aclTensor* mmOut = (mmOpInfo.support_info.self_format == ge::FORMAT_ND &&526 const aclTensor* mmOut = (mmOpInfo.support_info.self_format == ge::FORMAT_ND &&
527 mmOpInfo.support_info.output_format == ge::FORMAT_ND) ?527 mmOpInfo.support_info.output_format == ge::FORMAT_ND) ?
528 l0op::MatMulNdFp162Fp32(x1, x2, nullptr, nullptr, transposeX1, transposeX2,528 l0op::MatMulNdFp162Fp32(x1, x2, nullptr, nullptr, transposeX1, transposeX2,
@@ -581,7 +581,7 @@ static const aclTensor* GetMatMulOp(const aclTensor* x1, const aclTensor* x2, co
581 581 
582 const aclTensor* mmOut = l0op::MatMulV3NdFp162Fp32(x1, x2, bias, transposeX1, transposeX2, offsetX,582 const aclTensor* mmOut = l0op::MatMulV3NdFp162Fp32(x1, x2, bias, transposeX1, transposeX2, offsetX,
583 opImplModeEnum, executor);583 opImplModeEnum, executor);
584- OP_LOGI("hit matmulv3 fp16/bp16 in fp32 out case.");584+ OP_LOGI("hit matmulv3 fp16/bf16 in fp32 out case.");
585 return mmOut;585 return mmOut;
586 }586 }
587 587 
@@ -808,10 +808,10 @@ static aclnnStatus SetMatmulOpSupportFormat(const aclTensor* self, const aclTens
808 // 判断当前Shape是否支持使用ND输入输出808 // 判断当前Shape是否支持使用ND输入输出
809 SetMmSupportFormat(self, mat2, mmOpInfo);809 SetMmSupportFormat(self, mat2, mmOpInfo);
810 810 
811- TensorInfo SpTensor_sefl = {self, mmOpInfo.support_info.self_dtype, mmOpInfo.support_info.self_format};811+ TensorInfo SpTensor_self = {self, mmOpInfo.support_info.self_dtype, mmOpInfo.support_info.self_format};
812 TensorInfo SpTensor_mat2 = {mat2, mmOpInfo.support_info.mat2_dtype, mmOpInfo.support_info.output_format};812 TensorInfo SpTensor_mat2 = {mat2, mmOpInfo.support_info.mat2_dtype, mmOpInfo.support_info.output_format};
813 813 
814- if (IsSplitk(&SpTensor_sefl, &SpTensor_mat2)) {814+ if (IsSplitk(&SpTensor_self, &SpTensor_mat2)) {
815 mmOpInfo.supporSplitK = true;815 mmOpInfo.supporSplitK = true;
816 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();816 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
817 if (npuArch == NpuArch::DAV_2002) {817 if (npuArch == NpuArch::DAV_2002) {
@@ -1265,7 +1265,7 @@ const aclTensor* ExecMmOp(const aclTensor* self, const aclTensor* mat2, const ac
1265}1265}
1266 1266 
1267/*1267/*
1268- 注意：虽然申明为const指针引用，但selfReshapeOutput和mat2ReshapeOutput会重新赋值修改指针指向新的const aclTensor1268+ 注意：虽然声明为const指针引用，但selfReshapeOutput和mat2ReshapeOutput会重新赋值修改指针指向新的const aclTensor
1269*/1269*/
1270int64_t ProcessSpecialCases(const aclTensor*& selfCastOut, const aclTensor*& mat2CastOut, MmOpInfo& mmOpInfo,1270int64_t ProcessSpecialCases(const aclTensor*& selfCastOut, const aclTensor*& mat2CastOut, MmOpInfo& mmOpInfo,
1271 const aclTensor*& bias, const aclTensor*& selfReshapeOutput,1271 const aclTensor*& bias, const aclTensor*& selfReshapeOutput,
@@ -1299,7 +1299,7 @@ int64_t ProcessSpecialCases(const aclTensor*& selfCastOut, const aclTensor*& mat
1299}1299}
1300 1300 
1301/*1301/*
1302-计算MatMul的workSize， 内涵MatMul算子的构图流程1302+计算MatMul的workSize， 包含MatMul算子的构图流程
1303 self mat21303 self mat2
1304 | |1304 | |
1305 contiguous contiguous1305 contiguous contiguous
@@ -1354,7 +1354,7 @@ const aclTensor* ExecMmOpWithBias(const aclTensor* self, const aclTensor* mat2,
1354 mat2 = l0op::ReFormat(mat2, op::Format::FORMAT_ND);1354 mat2 = l0op::ReFormat(mat2, op::Format::FORMAT_ND);
1355 CHECK_RET(mat2 != nullptr, nullptr);1355 CHECK_RET(mat2 != nullptr, nullptr);
1356 }1356 }
1357- OP_LOGI("mat2 origin storage shape is [%s].", op::ToString(mat2->GetStorageShape()).GetString());1357+ OP_LOGI("mat2 origin storage shape is [%s].", op::ToString(mat2->GetStorageShape()).GetString());
1358 // bias非连续转连续以及转换dtype1358 // bias非连续转连续以及转换dtype
1359 auto contiguousBias = bias;1359 auto contiguousBias = bias;
1360 if (contiguousBias != nullptr) {1360 if (contiguousBias != nullptr) {
@@ -1371,7 +1371,7 @@ const aclTensor* ExecMmOpWithBias(const aclTensor* self, const aclTensor* mat2,
1371 if (mat2->GetStorageFormat() == op::Format::FORMAT_FRACTAL_NZ) {1371 if (mat2->GetStorageFormat() == op::Format::FORMAT_FRACTAL_NZ) {
1372 OP_LOGI("mat2 GetStorageFormat FORMAT_FRACTAL_NZ.");1372 OP_LOGI("mat2 GetStorageFormat FORMAT_FRACTAL_NZ.");
1373 aclTensor* mat2ShapeSet = const_cast<aclTensor*>(mat2CastOut);1373 aclTensor* mat2ShapeSet = const_cast<aclTensor*>(mat2CastOut);
1374- mat2ShapeSet->SetStorageShape(mat2StorageShape); // 对NZ的场景用原来的stroageShape刷新1374+ mat2ShapeSet->SetStorageShape(mat2StorageShape); // 对NZ的场景用原来的storageShape刷新
1375 }1375 }
1376 OP_LOGI("mat2 storage shape is [%s].", op::ToString(mat2StorageShape).GetString());1376 OP_LOGI("mat2 storage shape is [%s].", op::ToString(mat2StorageShape).GetString());
1377 1377 
@@ -1507,7 +1507,7 @@ const aclTensor* MatmulCommonProcess(const aclTensor* self, const aclTensor* mat
1507 CHECK_RET(mat2 != nullptr, nullptr);1507 CHECK_RET(mat2 != nullptr, nullptr);
1508 }1508 }
1509 }1509 }
1510- OP_LOGI("mat2 origin storage shape is [%s].", op::ToString(mat2->GetStorageShape()).GetString());1510+ OP_LOGI("mat2 origin storage shape is [%s].", op::ToString(mat2->GetStorageShape()).GetString());
1511 1511 
1512 auto mat2CastOut = mat2;1512 auto mat2CastOut = mat2;
1513 auto mat2StorageShape = mat2->GetStorageShape();1513 auto mat2StorageShape = mat2->GetStorageShape();
@@ -1517,7 +1517,7 @@ const aclTensor* MatmulCommonProcess(const aclTensor* self, const aclTensor* mat
1517 if (mat2->GetStorageFormat() == op::Format::FORMAT_FRACTAL_NZ) {1517 if (mat2->GetStorageFormat() == op::Format::FORMAT_FRACTAL_NZ) {
1518 OP_LOGI("mat2 GetStorageFormat FORMAT_FRACTAL_NZ.");1518 OP_LOGI("mat2 GetStorageFormat FORMAT_FRACTAL_NZ.");
1519 aclTensor* mat2ShapeSet = const_cast<aclTensor*>(mat2CastOut);1519 aclTensor* mat2ShapeSet = const_cast<aclTensor*>(mat2CastOut);
1520- mat2ShapeSet->SetStorageShape(mat2StorageShape); // 对NZ的场景用原来的stroageShape刷新1520+ mat2ShapeSet->SetStorageShape(mat2StorageShape); // 对NZ的场景用原来的storageShape刷新
1521 }1521 }
1522 OP_LOGI("mat2 storage shape is [%s].", op::ToString(mat2StorageShape).GetString());1522 OP_LOGI("mat2 storage shape is [%s].", op::ToString(mat2StorageShape).GetString());
1523 // bias非连续转连续以及转换dtype1523 // bias非连续转连续以及转换dtype
@@ -2692,7 +2692,7 @@ const aclTensor* TransposeAndContiguousMat(const aclTensor* mat, aclOpExecutor*
2692 if (mat->GetStorageFormat() == op::Format::FORMAT_FRACTAL_NZ) {2692 if (mat->GetStorageFormat() == op::Format::FORMAT_FRACTAL_NZ) {
2693 OP_LOGI("mat GetStorageFormat FORMAT_FRACTAL_NZ.");2693 OP_LOGI("mat GetStorageFormat FORMAT_FRACTAL_NZ.");
2694 aclTensor* matShapeSet = const_cast<aclTensor*>(contiguousMat);2694 aclTensor* matShapeSet = const_cast<aclTensor*>(contiguousMat);
2695- matShapeSet->SetStorageShape(matStorageShape); // 对NZ的场景用原来的stroageShape刷新2695+ matShapeSet->SetStorageShape(matStorageShape); // 对NZ的场景用原来的storageShape刷新
2696 }2696 }
2697 OP_LOGI("mat storage shape is [%s].", op::ToString(matStorageShape).GetString());2697 OP_LOGI("mat storage shape is [%s].", op::ToString(matStorageShape).GetString());
2698 CHECK_RET(contiguousMat != nullptr, nullptr);2698 CHECK_RET(contiguousMat != nullptr, nullptr);
Mmatmul/common/op_host/op_tiling/debug_tiling.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file debug_tiilng.cpp12+ * \file debug_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
@@ -69,4 +69,4 @@ std::string DebugTilingData(gert::TilingContext* context)
69 return oss.str();69 return oss.str();
70}70}
71} // namespace NN71} // namespace NN
72-} // namespace Ops72+} // namespace Ops
Mmatmul/fused_mat_mul/op_host/op_api/aclnn_fused_matmul.cpp+4-4
@@ -105,7 +105,7 @@ bool CheckNotNull(const aclTensor* x, const aclTensor* x2, const aclTensor* bias
105 OP_CHECK_NULL(x, return false);105 OP_CHECK_NULL(x, return false);
106 OP_CHECK_NULL(x2, return false);106 OP_CHECK_NULL(x2, return false);
107 if (bias != nullptr && !IsInSupportedOpTypes(fusedOpType, kSupportedBiasOpTypes)) {107 if (bias != nullptr && !IsInSupportedOpTypes(fusedOpType, kSupportedBiasOpTypes)) {
108- OP_LOGE(ACLNN_ERR_PARAM_NULLPTR, "bias is not supported right now");108+ OP_LOGE(ACLNN_ERR_PARAM_NULLPTR, "bias is not supported for the current fusedOpType");
109 return false;109 return false;
110 }110 }
111 if (IsInSupportedOpTypes(fusedOpType, kSupportedX3OpTypes)) {111 if (IsInSupportedOpTypes(fusedOpType, kSupportedX3OpTypes)) {
@@ -121,7 +121,7 @@ static inline bool CheckMathType(const aclTensor* self, const aclTensor* mat2, i
121 bool mat2Float = mat2->GetDataType() == DataType::DT_FLOAT;121 bool mat2Float = mat2->GetDataType() == DataType::DT_FLOAT;
122 auto promoteType = selfFloat || mat2Float ? DataType::DT_FLOAT : self->GetDataType();122 auto promoteType = selfFloat || mat2Float ? DataType::DT_FLOAT : self->GetDataType();
123 if (cubeMathType != USE_HF32 && promoteType == DataType::DT_FLOAT) {123 if (cubeMathType != USE_HF32 && promoteType == DataType::DT_FLOAT) {
124- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "fusedmatmul is only supported bf16/fp16/hf32, does not support fp32.");124+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "fusedmatmul only supports bf16/fp16/hf32, does not support fp32.");
125 return false;125 return false;
126 }126 }
127 return CheckCubeMathTypeForMm(promoteType, cubeMathType);127 return CheckCubeMathTypeForMm(promoteType, cubeMathType);
@@ -266,7 +266,7 @@ static bool CheckBiasShape(const aclTensor* bias)
266 }266 }
267 size_t biasDimNum = bias->GetViewShape().GetDimNum();267 size_t biasDimNum = bias->GetViewShape().GetDimNum();
268 if (biasDimNum != 1 && biasDimNum != DIM_LEN_MIN) {268 if (biasDimNum != 1 && biasDimNum != DIM_LEN_MIN) {
269- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Input dim of bias cannot be larger than 2");269+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Input dimension of bias cannot be larger than 2");
270 return false;270 return false;
271 }271 }
272 return true;272 return true;
@@ -548,7 +548,7 @@ static const aclTensor* BuildFusedMatMulGraph(const aclTensor* x, const aclTenso
548 // 空tensor 处理，对于非16Cast32放开空tensor548 // 空tensor 处理，对于非16Cast32放开空tensor
549 bool allowEmptyTensor = IsInSupportedOpTypes(fusedOpType, kSupportedEmptyTensorOpTypes);549 bool allowEmptyTensor = IsInSupportedOpTypes(fusedOpType, kSupportedEmptyTensorOpTypes);
550 if (!allowEmptyTensor && (x->IsEmpty() || x2->IsEmpty())) {550 if (!allowEmptyTensor && (x->IsEmpty() || x2->IsEmpty())) {
551- OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "fused matmul is not supported empty tensor handle");551+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "FusedMatmul does not support empty tensor for this fusedOpType");
552 return nullptr;552 return nullptr;
553 }553 }
554 // 解析当前规格matmulop支持的dtype、format能力554 // 解析当前规格matmulop支持的dtype、format能力
Mmatmul/fused_mat_mul/op_host/op_api/aclnn_fused_matmul.h+4-4
@@ -30,9 +30,9 @@ ACLNN_API aclnnStatus aclnnFusedMatmulGetWorkspaceSize(const aclTensor* x1, cons
30 30 
31/**31/**
32 * @brief aclnnFusedMatmul的第二段接口，用于执行计算。32 * @brief aclnnFusedMatmul的第二段接口，用于执行计算。
33- * @param [in] workspace: 在npu device侧申请的workspace内存起址。33+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
34- * @param [in] workspace_size: 在npu device侧申请的workspace大小，由第一段接口aclnnFusedMatmulGetWorkspaceSize获取。34+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnFusedMatmulGetWorkspaceSize获取。
35- * @param [in] exector: op执行器，包含了算子计算流程。35+ * @param [in] executor: op执行器，包含了算子计算流程。
36 * @param [in] stream: acl stream流。36 * @param [in] stream: acl stream流。
37 * @return aclnnStatus: 返回状态码37 * @return aclnnStatus: 返回状态码
38 */38 */
@@ -43,4 +43,4 @@ ACLNN_API aclnnStatus aclnnFusedMatmul(void* workspace, uint64_t workspaceSize,
43}43}
44#endif44#endif
45 45 
46-#endif // OP_API_INC_FUSED_MATMUL_H46+#endif // OP_API_INC_FUSED_MATMUL_H
Mmatmul/fused_mat_mul/op_host/op_tiling/arch35/fused_matmul_builtin_tiling.cpp+1-1
@@ -145,7 +145,7 @@ void FusedMatMulBuiltInTiling::ExtractAttrFlags()
145{145{
146 args_.isHf32 = *context_->GetAttrs()->GetAttrPointer<bool>(ATTR_ENABLE_HF32_IDX);146 args_.isHf32 = *context_->GetAttrs()->GetAttrPointer<bool>(ATTR_ENABLE_HF32_IDX);
147 if (args_.isHf32 && arch_ != NpuArch::DAV_3510) {147 if (args_.isHf32 && arch_ != NpuArch::DAV_3510) {
148- OP_LOGW(args_.opName, "Hf32 flag is: %d, which is not support yet", args_.isHf32);148+ OP_LOGW(args_.opName, "Hf32 flag is: %d, which is not supported yet", args_.isHf32);
149 }149 }
150 innerPrecise_ = *context_->GetAttrs()->GetAttrPointer<int64_t>(ATTR_INNER_PRECISE_IDX);150 innerPrecise_ = *context_->GetAttrs()->GetAttrPointer<int64_t>(ATTR_INNER_PRECISE_IDX);
151 OP_LOGI(args_.opName, "FusedMatMul built-in tiling inner_precise is %ld", innerPrecise_);151 OP_LOGI(args_.opName, "FusedMatMul built-in tiling inner_precise is %ld", innerPrecise_);
Mmatmul/fused_mat_mul/op_host/op_tiling/arch35/fused_matmul_tiling.cpp+2-2
@@ -57,7 +57,7 @@ ge::graphStatus FusedMatMulTilingFunc(gert::TilingContext* context)
57 OP_TILING_CHECK(context == nullptr, CUBE_INNER_ERR_REPORT("FusedMatMul", "context is null"),57 OP_TILING_CHECK(context == nullptr, CUBE_INNER_ERR_REPORT("FusedMatMul", "context is null"),
58 return ge::GRAPH_FAILED);58 return ge::GRAPH_FAILED);
59 if (!IsAdvancedSocVersion(context)) {59 if (!IsAdvancedSocVersion(context)) {
60- OP_LOGE("FusedMatMul", "not support npu arch");60+ OP_LOGE("FusedMatMul", "does not support npu arch");
61 return ge::GRAPH_FAILED;61 return ge::GRAPH_FAILED;
62 }62 }
63 63 
@@ -77,7 +77,7 @@ ge::graphStatus FusedMatMulTilingFunc(gert::TilingContext* context)
77 const auto& supportedOps = it->second;77 const auto& supportedOps = it->second;
78 bool useBuiltInTiling = std::find(supportedOps.begin(), supportedOps.end(), fusedOpType) != supportedOps.end();78 bool useBuiltInTiling = std::find(supportedOps.begin(), supportedOps.end(), fusedOpType) != supportedOps.end();
79 OP_TILING_CHECK(!useBuiltInTiling,79 OP_TILING_CHECK(!useBuiltInTiling,
80- CUBE_INNER_ERR_REPORT(context->GetNodeName(), "unsupported fused op type: %s, supported: %s",80+ CUBE_INNER_ERR_REPORT(context->GetNodeName(), "unsupported fusedOpType: %s, supported: %s",
81 fusedOpType.c_str(), JoinOpTypes(supportedOps).c_str()),81 fusedOpType.c_str(), JoinOpTypes(supportedOps).c_str()),
82 return ge::GRAPH_FAILED);82 return ge::GRAPH_FAILED);
83 83 
Mmatmul/fused_mat_mul/op_kernel/fused_mat_mul.cpp+4-4
@@ -84,7 +84,7 @@ using L0TileShape = AscendC::Shape<_256, _256, _64>;
84enum class FusionOpType : uint8_t { ADD, MUL, GELU, GELU_ERF };84enum class FusionOpType : uint8_t { ADD, MUL, GELU, GELU_ERF };
85 85 
86#if __FIXED_POINT_ONLY_CUBE_TO_L0C__86#if __FIXED_POINT_ONLY_CUBE_TO_L0C__
87-#define BMMV3_IMPL_CLASS_COMMON_TRNAS(transA, transB, templateClass, ...) \87+#define BMMV3_IMPL_CLASS_COMMON_TRANS(transA, transB, templateClass, ...) \
88 do { \88 do { \
89 GET_TILING_DATA(tilingData, tilingGM); \89 GET_TILING_DATA(tilingData, tilingGM); \
90 using cType = MatmulType<AscendC::TPosition::GM, format_y, DTYPE_Y>; \90 using cType = MatmulType<AscendC::TPosition::GM, format_y, DTYPE_Y>; \
@@ -788,13 +788,13 @@ __global__ __aicore__ void fused_mat_mul(GM_ADDR x1GM, GM_ADDR x2GM, GM_ADDR bia
788 API_LEVEL == MAT_MUL_HIGH_LEVEL && FULL_LOAD == MAT_MUL_NO_FULL_LOAD &&788 API_LEVEL == MAT_MUL_HIGH_LEVEL && FULL_LOAD == MAT_MUL_NO_FULL_LOAD &&
789 BATCH_ITER_MODEL == MAT_MUL_FOR_FUSED_BATCH) {789 BATCH_ITER_MODEL == MAT_MUL_FOR_FUSED_BATCH) {
790 if constexpr (OPTYPE == F_OPTYPE_QUANT) {790 if constexpr (OPTYPE == F_OPTYPE_QUANT) {
791- BMMV3_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,791+ BMMV3_IMPL_CLASS_COMMON_TRANS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,
792 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD, true);792 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD, true);
793 } else if constexpr (OPTYPE == F_OPTYPE_RELU_QUANT) {793 } else if constexpr (OPTYPE == F_OPTYPE_RELU_QUANT) {
794- BMMV3_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,794+ BMMV3_IMPL_CLASS_COMMON_TRANS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,
795 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD_RELU, true);795 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD_RELU, true);
796 } else {796 } else {
797- BMMV3_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,797+ BMMV3_IMPL_CLASS_COMMON_TRANS(aTran, bTran, BatchMatMulV3Advanced::BatchMatMulAswKernel,
798 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD_RELU);798 BatchMatMulV3Advanced::BatchMatMulAswBlock, MM_CFG_NO_PRELOAD_RELU);
799 }799 }
800 } else if constexpr ( // basic, aswt, from bmmv3800 } else if constexpr ( // basic, aswt, from bmmv3
Mmatmul/gemm_v2/examples/test_geir_gemm_v2.cpp+40-40
@@ -40,49 +40,49 @@ using namespace ge;
40using std::map;40using std::map;
41using std::string;41using std::string;
42using std::vector;42using std::vector;
43-#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \43+#define ADD_INPUT(inputIndex, inputName, inputDtype, inputShape) \
44- vector<int64_t> placeholder##intputIndex##_shape = inputShape; \44+ vector<int64_t> placeholder##inputIndex##_shape = inputShape; \
45- auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \45+ auto placeholder##inputIndex = op::Data("placeholder" + inputIndex).set_attr_index(0); \
46- TensorDesc placeholder##intputIndex##_desc = TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, \46+ TensorDesc placeholder##inputIndex##_desc = TensorDesc(ge::Shape(placeholder##inputIndex##_shape), FORMAT_ND, \
47- intputDtype); \47+ inputDtype); \
48- placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \48+ placeholder##inputIndex##_desc.SetPlacement(ge::kPlacementHost); \
49- placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \49+ placeholder##inputIndex##_desc.SetFormat(FORMAT_ND); \
50- Tensor tensor_placeholder##intputIndex; \50+ Tensor tensor_placeholder##inputIndex; \
51- ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \51+ ret = GenOnesDataFloat32(placeholder##inputIndex##_shape, tensor_placeholder##inputIndex, \
52- placeholder##intputIndex##_desc, 2); \52+ placeholder##inputIndex##_desc, 2); \
53- if (ret != SUCCESS) { \53+ if (ret != SUCCESS) { \
54- printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \54+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
55- return FAILED; \55+ return FAILED; \
56- } \56+ } \
57- placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \57+ placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \
58- placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \58+ placeholder##inputIndex.update_output_desc_y(placeholder##inputIndex##_desc); \
59- input.push_back(tensor_placeholder##intputIndex); \59+ input.push_back(tensor_placeholder##inputIndex); \
60- graph.AddOp(placeholder##intputIndex); \60+ graph.AddOp(placeholder##inputIndex); \
61- add1.set_input_##intputName(placeholder##intputIndex); \61+ add1.set_input_##inputName(placeholder##inputIndex); \
62- inputs.push_back(placeholder##intputIndex);62+ inputs.push_back(placeholder##inputIndex);
63 63 
64#define ADD_INPUT_ATTR(attrName, attrValue) add1.set_attr_##attrName(attrValue);64#define ADD_INPUT_ATTR(attrName, attrValue) add1.set_attr_##attrName(attrValue);
65 65 
66-#define ADD_CONST_INPUT(intputIndex, intputName, intputDtype, inputShape) \66+#define ADD_CONST_INPUT(inputIndex, inputName, inputDtype, inputShape) \
67- vector<int64_t> placeholder##intputIndex##_shape = inputShape; \67+ vector<int64_t> placeholder##inputIndex##_shape = inputShape; \
68- auto placeholder##intputIndex = op::Const("placeholder" + intputIndex); \68+ auto placeholder##inputIndex = op::Const("placeholder" + inputIndex); \
69- TensorDesc placeholder##intputIndex##_desc = TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, \69+ TensorDesc placeholder##inputIndex##_desc = TensorDesc(ge::Shape(placeholder##inputIndex##_shape), FORMAT_ND, \
70- intputDtype); \70+ inputDtype); \
71- placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \71+ placeholder##inputIndex##_desc.SetPlacement(ge::kPlacementHost); \
72- placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \72+ placeholder##inputIndex##_desc.SetFormat(FORMAT_ND); \
73- Tensor tensor_placeholder##intputIndex; \73+ Tensor tensor_placeholder##inputIndex; \
74- ret = GenOnesData(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \74+ ret = GenOnesData(placeholder##inputIndex##_shape, tensor_placeholder##inputIndex, placeholder##inputIndex##_desc, \
75- placeholder##intputIndex##_desc, intputDtype, 2); \75+ inputDtype, 2); \
76- if (ret != SUCCESS) { \76+ if (ret != SUCCESS) { \
77- printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \77+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
78- return FAILED; \78+ return FAILED; \
79- } \79+ } \
80- placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \80+ placeholder##inputIndex.SetAttr("value", tensor_placeholder##inputIndex); \
81- placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \81+ placeholder##inputIndex.update_output_desc_y(placeholder##inputIndex##_desc); \
82- graph.AddOp(placeholder##intputIndex); \82+ graph.AddOp(placeholder##inputIndex); \
83- add1.set_input_##intputName(placeholder##intputIndex); \83+ add1.set_input_##inputName(placeholder##inputIndex); \
84- add1.update_input_desc_##intputName(placeholder##intputIndex##_desc); \84+ add1.update_input_desc_##inputName(placeholder##inputIndex##_desc); \
85- inputs.push_back(placeholder##intputIndex);85+ inputs.push_back(placeholder##inputIndex);
86 86 
87#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \87#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
88 TensorDesc outputName##outputIndex##_desc_ = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \88 TensorDesc outputName##outputIndex##_desc_ = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
Mmatmul/gemm_v2/op_host/gemm_v2_infershape.cpp+6-6
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file gemm_v2_infer.cpp12+ * \file gemm_v2_infershape.cpp
13 * \brief13 * \brief
14 */14 */
15#include "common/op_host/matmul_common_infershape.h"15#include "common/op_host/matmul_common_infershape.h"
@@ -49,7 +49,7 @@ static ge::graphStatus InferShapeForGemmV2(InferShapeContext* context)
49 OP_LOGD(op_name, "check the input shape length.");49 OP_LOGD(op_name, "check the input shape length.");
50 CHECK((shape_a->GetDimNum() != kMatmulV2MinShapeSize || shape_b->GetDimNum() != kMatmulV2MinShapeSize ||50 CHECK((shape_a->GetDimNum() != kMatmulV2MinShapeSize || shape_b->GetDimNum() != kMatmulV2MinShapeSize ||
51 shape_c->GetDimNum() != kMatmulV2MinShapeSize),51 shape_c->GetDimNum() != kMatmulV2MinShapeSize),
52- CUBE_INNER_ERR_REPORT(op_name, "input dim num[%zu] [%zu] [%zu]is not 2!", shape_a->GetDimNum(),52+ CUBE_INNER_ERR_REPORT(op_name, "input dim num[%zu] [%zu] [%zu] is not 2!", shape_a->GetDimNum(),
53 shape_b->GetDimNum(), shape_c->GetDimNum()),53 shape_b->GetDimNum(), shape_c->GetDimNum()),
54 return ge::GRAPH_FAILED);54 return ge::GRAPH_FAILED);
55 55 
@@ -64,8 +64,8 @@ static ge::graphStatus InferShapeForGemmV2(InferShapeContext* context)
64 return ge::GRAPH_FAILED);64 return ge::GRAPH_FAILED);
65 65 
66 CHECK(shape_a->GetDim(idx_m) != shape_c->GetDim(0) || shape_b->GetDim(idx_n) != shape_c->GetDim(1),66 CHECK(shape_a->GetDim(idx_m) != shape_c->GetDim(0) || shape_b->GetDim(idx_n) != shape_c->GetDim(1),
67- CUBE_INNER_ERR_REPORT(op_name, "The m(%ld), n(%ld) tensors must be the same c(%ld, %ld)",67+ CUBE_INNER_ERR_REPORT(op_name, "The m(%ld), n(%ld) tensors must match c(%ld, %ld)", shape_a->GetDim(idx_m),
68- shape_a->GetDim(idx_m), shape_b->GetDim(idx_n), shape_c->GetDim(0), shape_c->GetDim(1)),68+ shape_b->GetDim(idx_n), shape_c->GetDim(0), shape_c->GetDim(1)),
69 return ge::GRAPH_FAILED);69 return ge::GRAPH_FAILED);
70 70 
71 shape_out->SetDimNum(kMatmulV2MinShapeSize);71 shape_out->SetDimNum(kMatmulV2MinShapeSize);
@@ -84,7 +84,7 @@ static ge::graphStatus InferDataTypeForGemmV2(gert::InferDataTypeContext* contex
84 const ge::DataType c_data_type = context->GetInputDataType(4);84 const ge::DataType c_data_type = context->GetInputDataType(4);
85 CHECK(((a_data_type != ge::DT_FLOAT16 && a_data_type != ge::DT_BF16) ||85 CHECK(((a_data_type != ge::DT_FLOAT16 && a_data_type != ge::DT_BF16) ||
86 (b_data_type != ge::DT_FLOAT16 && b_data_type != ge::DT_BF16) || c_data_type != ge::DT_FLOAT),86 (b_data_type != ge::DT_FLOAT16 && b_data_type != ge::DT_BF16) || c_data_type != ge::DT_FLOAT),
87- CUBE_INNER_ERR_REPORT(op_name, "input dtype not support"), return ge::GRAPH_FAILED);87+ CUBE_INNER_ERR_REPORT(op_name, "input dtype is not supported"), return ge::GRAPH_FAILED);
88 ge::graphStatus ret = context->SetOutputDataType(0, ge::DT_FLOAT);88 ge::graphStatus ret = context->SetOutputDataType(0, ge::DT_FLOAT);
89 return ret;89 return ret;
90}90}
@@ -92,4 +92,4 @@ static ge::graphStatus InferDataTypeForGemmV2(gert::InferDataTypeContext* contex
92 92 
93namespace Ops::NN::MatMul {93namespace Ops::NN::MatMul {
94IMPL_OP_INFERSHAPE(GemmV2).InferShape(InferShapeForGemmV2).InferDataType(InferDataTypeForGemmV2);94IMPL_OP_INFERSHAPE(GemmV2).InferShape(InferShapeForGemmV2).InferDataType(InferDataTypeForGemmV2);
95-}95+}
Mmatmul/gemm_v2/op_host/op_tiling/gemm_v2_tiling.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file gemm_v2_tiling.cc12+ * \file gemm_v2_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15#include "gemm_v2_tiling.h"15#include "gemm_v2_tiling.h"
@@ -49,7 +49,7 @@ static ge::graphStatus TilingPrepareForGemmV2(gert::TilingParseContext* context)
49 platformInfo->GetPlatformRes("AICoreintrinsicDtypeMap", "Intrinsic_fix_pipe_l0c2out", val);49 platformInfo->GetPlatformRes("AICoreintrinsicDtypeMap", "Intrinsic_fix_pipe_l0c2out", val);
50 platformInfo->GetPlatformRes("AICoreintrinsicDtypeMap", "Intrinsic_data_move_l12bt", dataMoveL12Bt);50 platformInfo->GetPlatformRes("AICoreintrinsicDtypeMap", "Intrinsic_data_move_l12bt", dataMoveL12Bt);
51 compileInfoPtr->supportL0c2out = !val.empty();51 compileInfoPtr->supportL0c2out = !val.empty();
52- compileInfoPtr->supportL12BtBf16 = (dataMoveL12Bt.find("bf16") != string::npos);52+ compileInfoPtr->supportL12BtBf16 = (dataMoveL12Bt.find("bf16") != std::string::npos);
53 compileInfoPtr->aicNum = ascendcPlatform.GetCoreNumAic();53 compileInfoPtr->aicNum = ascendcPlatform.GetCoreNumAic();
54 compileInfoPtr->socVersion = ascendcPlatform.GetSocVersion();54 compileInfoPtr->socVersion = ascendcPlatform.GetSocVersion();
55 compileInfoPtr->npuArch = ascendcPlatform.GetCurNpuArch();55 compileInfoPtr->npuArch = ascendcPlatform.GetCurNpuArch();
Mmatmul/gemm_v2/tests/ut/op_host/test_gemm_v2_tiling.cpp+6-6
@@ -24,7 +24,7 @@ using namespace std;
24using namespace ge;24using namespace ge;
25 25 
26namespace {26namespace {
27-bool IsDisplayTilingdata(const string& case_name, size_t index)27+bool IsDisplayTilingData(const string& case_name, size_t index)
28{28{
29 if (index <= 18 || (index >= 22 && index <= 27) || (index >= 30 && index <= 32) || index >= 48) {29 if (index <= 18 || (index >= 22 && index <= 27) || (index >= 30 && index <= 32) || index >= 48) {
30 return true;30 return true;
@@ -41,7 +41,7 @@ static string TilingData2Str(const gert::TilingData* tiling_data, const string&
41 auto data = tiling_data->GetData();41 auto data = tiling_data->GetData();
42 string result;42 string result;
43 for (size_t i = 0; i < tiling_data->GetDataSize(); i += sizeof(int32_t)) {43 for (size_t i = 0; i < tiling_data->GetDataSize(); i += sizeof(int32_t)) {
44- if (IsDisplayTilingdata(case_name, i / sizeof(int32_t))) {44+ if (IsDisplayTilingData(case_name, i / sizeof(int32_t))) {
45 result += std::to_string((reinterpret_cast<const int32_t*>(tiling_data->GetData())[i / sizeof(int32_t)]));45 result += std::to_string((reinterpret_cast<const int32_t*>(tiling_data->GetData())[i / sizeof(int32_t)]));
46 result += " ";46 result += " ";
47 }47 }
@@ -59,7 +59,7 @@ static string GenGoldenTilingData(const string& tiling_data, const string& case_
59 }59 }
60 string golden_tiling_data;60 string golden_tiling_data;
61 for (size_t i = 0; i < data_list.size(); i++) {61 for (size_t i = 0; i < data_list.size(); i++) {
62- if (IsDisplayTilingdata(case_name, i)) {62+ if (IsDisplayTilingData(case_name, i)) {
63 golden_tiling_data += data_list[i];63 golden_tiling_data += data_list[i];
64 golden_tiling_data += " ";64 golden_tiling_data += " ";
65 }65 }
@@ -161,8 +161,8 @@ TEST_P(GemmV2TilingRuntime, general_cases)
161 ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);161 ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
162 162 
163 auto tiling_data = gert::TilingData::CreateCap(2048);163 auto tiling_data = gert::TilingData::CreateCap(2048);
164- auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);164+ auto workspace_size_holder = gert::ContinuousVector::Create<size_t>(4096);
165- auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());165+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holder.get());
166 166 
167 gert::KernelRunContextHolder holder;167 gert::KernelRunContextHolder holder;
168 holder = gert::TilingContextFaker()168 holder = gert::TilingContextFaker()
@@ -227,4 +227,4 @@ static TilingTestParam general_cases_params[] = {
227};227};
228 228 
229INSTANTIATE_TEST_CASE_P(GemmV2, GemmV2TilingRuntime, testing::ValuesIn(general_cases_params));229INSTANTIATE_TEST_CASE_P(GemmV2, GemmV2TilingRuntime, testing::ValuesIn(general_cases_params));
230-} // namespace230+} // namespace
Mmatmul/gemm_v3/op_host/gemm_v3_def.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file gemm_v3.cpp12+ * \file gemm_v3_def.cpp
13 * \brief13 * \brief
14 */14 */
15#include "register/op_def_registry.h"15#include "register/op_def_registry.h"
Mmatmul/gemm_v3/op_host/op_tiling/arch35/gemmv3_tiling_key.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file gemmv3_tiling_key.cc12+ * \file gemmv3_tiling_key.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/gemm_v3/op_host/op_tiling/gemm_v3_base_tiling.cpp+3-3
@@ -130,8 +130,8 @@ bool GemmV3BaseTiling::IsCapable()
130 bool is16In32Out = dtypeY == ge::DT_FLOAT && isLowPrecInput;130 bool is16In32Out = dtypeY == ge::DT_FLOAT && isLowPrecInput;
131 if ((!is16In32Out && dtypeC != dtypeY) || (is16In32Out && dtypeC != dtypeA && dtypeC != dtypeY)) {131 if ((!is16In32Out && dtypeC != dtypeY) || (is16In32Out && dtypeC != dtypeA && dtypeC != dtypeY)) {
132 OP_LOGW(params_.opName,132 OP_LOGW(params_.opName,
133- "Expected self dtype(%s) to be equal to mat dtype or out dtype(%s) in 16in32out scenario. or"133+ "Expected self dtype(%s) to be equal to mat dtype or out dtype(%s) in 16in32out scenario; "
134- "Expected self dtype(%s) to be equal to out dtype(%s).",134+ "or expected self dtype to be equal to out dtype in other scenarios.",
135 Ops::Base::ToString(dtypeC).c_str(), Ops::Base::ToString(dtypeY).c_str());135 Ops::Base::ToString(dtypeC).c_str(), Ops::Base::ToString(dtypeY).c_str());
136 return false;136 return false;
137 }137 }
@@ -139,7 +139,7 @@ bool GemmV3BaseTiling::IsCapable()
139 (dtypeY == ge::DT_FLOAT && isLowPrecInput));139 (dtypeY == ge::DT_FLOAT && isLowPrecInput));
140 if (!isValidOutput) {140 if (!isValidOutput) {
141 OP_LOGW(params_.opName,141 OP_LOGW(params_.opName,
142- "invalid out dtype (%s), only support half, bfloat16 or float(with low-precision input) output.",142+ "invalid out dtype (%s), only supports half, bfloat16 or float (with low-precision input) output.",
143 Ops::Base::ToString(dtypeY).c_str());143 Ops::Base::ToString(dtypeY).c_str());
144 return false;144 return false;
145 }145 }
Mmatmul/gemm_v3/op_host/op_tiling/gemmv3_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file gemmv3_tiling.cc12+ * \file gemmv3_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15#include "../../../mat_mul_v3/op_host/op_tiling/matmul_v3_tiling.h"15#include "../../../mat_mul_v3/op_host/op_tiling/matmul_v3_tiling.h"
Mmatmul/gemm_v3/op_kernel/arch35/gemm_v3.cpp+4-4
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file gemm_v3_apt.cpp12+ * \file gemm_v3.cpp
13 * \brief13 * \brief
14 */14 */
15#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)15#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)
@@ -40,7 +40,7 @@ using namespace matmul;
40#define FORMAT_FRACTAL_NZ40#define FORMAT_FRACTAL_NZ
41#endif41#endif
42 42 
43-#define MMV3_IMPL_CLASS_TRNAS(transA, transB, templateClass, ...) \43+#define MMV3_IMPL_CLASS_TRANS(transA, transB, templateClass, ...) \
44 do { \44 do { \
45 using cType = MatmulType<AscendC::TPosition::GM, CubeFormat::ND, DTYPE_C>; \45 using cType = MatmulType<AscendC::TPosition::GM, CubeFormat::ND, DTYPE_C>; \
46 using biasType = MatmulType<AscendC::TPosition::GM, CubeFormat::ND, DTYPE_BIAS>; \46 using biasType = MatmulType<AscendC::TPosition::GM, CubeFormat::ND, DTYPE_BIAS>; \
@@ -73,8 +73,8 @@ __global__ __aicore__ void gemm_v3(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR cGM, GM_ADD
73 // GemmV3复用matmulV3 kernel和tilingKey，暂只支持aswt模板73 // GemmV3复用matmulV3 kernel和tilingKey，暂只支持aswt模板
74 if constexpr (API_LEVEL == MAT_MUL_HIGH_LEVEL && FULL_LOAD == MAT_MUL_NO_FULL_LOAD && MODEL == MAT_MUL_BASIC &&74 if constexpr (API_LEVEL == MAT_MUL_HIGH_LEVEL && FULL_LOAD == MAT_MUL_NO_FULL_LOAD && MODEL == MAT_MUL_BASIC &&
75 L0C2OUT_MODEL == MAT_MUL_ON_THE_FLY) {75 L0C2OUT_MODEL == MAT_MUL_ON_THE_FLY) {
76- MMV3_IMPL_CLASS_TRNAS(aTran, bTran, MatmulV3Advanced::MatmulAswKernel, MatmulV3Advanced::MatmulAswBlock,76+ MMV3_IMPL_CLASS_TRANS(aTran, bTran, MatmulV3Advanced::MatmulAswKernel, MatmulV3Advanced::MatmulAswBlock,
77 MM_CFG_NO_PRELOAD);77 MM_CFG_NO_PRELOAD);
78 }78 }
79#endif79#endif
80-}80+}
Mmatmul/gemm_v3/tests/ut/op_kernel/gemm_v3_tiling_def.h+1-1
@@ -39,4 +39,4 @@ inline void InitTilingData(uint8_t* tiling, T* const_data)
39#define GET_TILING_DATA_WITH_STRUCT(tiling_struct, tiling_data, tiling_arg) \39#define GET_TILING_DATA_WITH_STRUCT(tiling_struct, tiling_data, tiling_arg) \
40 tiling_struct tiling_data; \40 tiling_struct tiling_data; \
41 InitTilingData<tiling_struct>(tiling_arg, &tiling_data);41 InitTilingData<tiling_struct>(tiling_arg, &tiling_data);
42-#endif // FOREACH_MINIMUM_SCALAR_TILING_DEF_H42+#endif // _TEST_GEMM_V3_TILING_DEF_H_
Mmatmul/mat_mul_v3/op_host/mat_mul_v3_infershape.cpp+5-5
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file mat_mul_v3_infer.cpp12+ * \file mat_mul_v3_infershape.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
@@ -66,7 +66,7 @@ bool CheckIsUnknownDimNum(const gert::Shape& shape)
66 return shape.GetDimNum() == 1 && shape.GetDim(0) == UNKNOWN_DIM_NUM;66 return shape.GetDimNum() == 1 && shape.GetDim(0) == UNKNOWN_DIM_NUM;
67}67}
68 68 
69-void UpdateUnknowDimNumToUnkownRank(gert::Shape& shape)69+void UpdateUnknownDimNumToUnknownRank(gert::Shape& shape)
70{70{
71 if (CheckIsUnknownDimNum(shape)) {71 if (CheckIsUnknownDimNum(shape)) {
72 shape.SetDimNum(MATMUL_MIN_SHAPE_SIZE);72 shape.SetDimNum(MATMUL_MIN_SHAPE_SIZE);
@@ -79,7 +79,7 @@ bool UpdateOutputShapeByBias(const std::string& op_name, gert::Shape* shape_out,
79{79{
80 if (shape_bias != nullptr && shape_bias->GetDimNum() > 0) {80 if (shape_bias != nullptr && shape_bias->GetDimNum() > 0) {
81 Shape shape_bias_new(*shape_bias);81 Shape shape_bias_new(*shape_bias);
82- UpdateUnknowDimNumToUnkownRank(shape_bias_new);82+ UpdateUnknownDimNumToUnknownRank(shape_bias_new);
83 int64_t bias_dim = shape_bias_new.GetDimNum();83 int64_t bias_dim = shape_bias_new.GetDimNum();
84 if (shape_bias_new.GetDim(bias_dim - 1) != UNKNOWN_DIM && shape_out->GetDim(1) != UNKNOWN_DIM) {84 if (shape_bias_new.GetDim(bias_dim - 1) != UNKNOWN_DIM && shape_out->GetDim(1) != UNKNOWN_DIM) {
85 OP_CHECK_IF(shape_bias_new.GetDim(bias_dim - 1) != shape_out->GetDim(1),85 OP_CHECK_IF(shape_bias_new.GetDim(bias_dim - 1) != shape_out->GetDim(1),
@@ -128,8 +128,8 @@ static ge::graphStatus InferShapeForMatMulV3(InferShapeContext* context)
128 128 
129 Shape shape_x1_new(*shape_x1);129 Shape shape_x1_new(*shape_x1);
130 Shape shape_x2_new(*shape_x2);130 Shape shape_x2_new(*shape_x2);
131- UpdateUnknowDimNumToUnkownRank(shape_x1_new);131+ UpdateUnknownDimNumToUnknownRank(shape_x1_new);
132- UpdateUnknowDimNumToUnkownRank(shape_x2_new);132+ UpdateUnknownDimNumToUnknownRank(shape_x2_new);
133 bool shape_x1_reshape_flag = false;133 bool shape_x1_reshape_flag = false;
134 bool shape_x2_reshape_flag = false;134 bool shape_x2_reshape_flag = false;
135 135 
Mmatmul/mat_mul_v3/op_host/op_api/aclnn_addmm.cpp+4-4
@@ -130,7 +130,7 @@ static inline bool CheckBroadcast(const aclTensor* self, const aclTensor* mat1,
130 return false;130 return false;
131 }131 }
132 if (self->GetStorageShape().GetDimNum() < DIM_SIZE_ONE) {132 if (self->GetStorageShape().GetDimNum() < DIM_SIZE_ONE) {
133- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self can not be empty.");133+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self cannot be empty.");
134 return false;134 return false;
135 }135 }
136 op::Shape matmulShape = {(mat1->GetViewShape())[0], (mat2->GetViewShape())[1]};136 op::Shape matmulShape = {(mat1->GetViewShape())[0], (mat2->GetViewShape())[1]};
@@ -139,7 +139,7 @@ static inline bool CheckBroadcast(const aclTensor* self, const aclTensor* mat1,
139 return true;139 return true;
140}140}
141 141 
142-// 假设mat1是 n x m，mat2是 m x p，out必须是 n x p 如果n / p为0，那么out为empty即可142+// 假设mat1是 n x k，mat2是 k x p，out必须是 n x p 如果n / p为0，那么out为empty即可
143static inline bool CheckOutShape(const aclTensor* mat1, const aclTensor* mat2, const aclTensor* out)143static inline bool CheckOutShape(const aclTensor* mat1, const aclTensor* mat2, const aclTensor* out)
144{144{
145 int64_t n = mat1->GetViewShape().GetDim(0);145 int64_t n = mat1->GetViewShape().GetDim(0);
@@ -391,7 +391,7 @@ static inline bool CheckMatmulWeightNz(const aclTensor* mat1, const aclTensor* m
391{391{
392 if (mat1->GetDataType() == op::DataType::DT_FLOAT || mat2->GetDataType() == op::DataType::DT_FLOAT ||392 if (mat1->GetDataType() == op::DataType::DT_FLOAT || mat2->GetDataType() == op::DataType::DT_FLOAT ||
393 mat1->GetDataType() != mat2->GetDataType()) {393 mat1->GetDataType() != mat2->GetDataType()) {
394- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "invalid mat1 dtype [%s] or mat2 dtype [%s] ",394+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "invalid mat1 dtype [%s] or mat2 dtype [%s]",
395 op::ToString(mat1->GetDataType()).GetString(), op::ToString(mat2->GetDataType()).GetString());395 op::ToString(mat1->GetDataType()).GetString(), op::ToString(mat2->GetDataType()).GetString());
396 return false;396 return false;
397 }397 }
@@ -421,7 +421,7 @@ static inline bool CheckMatmulWeightNz(const aclTensor* mat1, const aclTensor* m
421 }421 }
422 422 
423 if ((mat2->GetViewShape())[0] == 1 || (mat2->GetViewShape())[1] == 1) {423 if ((mat2->GetViewShape())[0] == 1 || (mat2->GetViewShape())[1] == 1) {
424- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The k-axis or n-axis can not be 1.");424+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The k-axis or n-axis cannot be 1.");
425 return false;425 return false;
426 }426 }
427 427 
Mmatmul/mat_mul_v3/op_host/op_api/aclnn_addmm.h+2-2
@@ -68,7 +68,7 @@ ACLNN_API aclnnStatus aclnnAddmmWeightNzGetWorkspaceSize(const aclTensor* self,
68 68 
69/**69/**
70 * @brief aclnnAddmmWeightNz的第二段接口，用于执行计算。70 * @brief aclnnAddmmWeightNz的第二段接口，用于执行计算。
71- * @param [in] workspace: 在npu device侧申请的workspace内存起址。71+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
72 * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnAddmmWeightNzGetWorkspaceSize获取。72 * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnAddmmWeightNzGetWorkspaceSize获取。
73 * @param [in] executor: op执行器，包含了算子计算流程。73 * @param [in] executor: op执行器，包含了算子计算流程。
74 * @param [in] stream: acl stream流。74 * @param [in] stream: acl stream流。
@@ -80,4 +80,4 @@ ACLNN_API aclnnStatus aclnnAddmmWeightNz(void* workspace, uint64_t workspaceSize
80}80}
81#endif81#endif
82 82 
83-#endif // OP_API_INC_ADD_H_83+#endif // OP_API_INC_ADDMM_H_
Mmatmul/mat_mul_v3/op_host/op_api/aclnn_matmul.cpp+7-7
@@ -98,7 +98,7 @@ inline static bool CheckWeightNzDtypeValid(const aclTensor* self, const aclTenso
98 if (curArch != NpuArch::DAV_2201 && mat2->GetStorageFormat() == Format::FORMAT_FRACTAL_NZ &&98 if (curArch != NpuArch::DAV_2201 && mat2->GetStorageFormat() == Format::FORMAT_FRACTAL_NZ &&
99 (self->GetDataType() == DataType::DT_FLOAT || mat2->GetDataType() == DataType::DT_FLOAT)) {99 (self->GetDataType() == DataType::DT_FLOAT || mat2->GetDataType() == DataType::DT_FLOAT)) {
100 OP_LOGE(ACLNN_ERR_PARAM_INVALID,100 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
101- "Float32 weight NZ is unsupported by the current SOC version [%s], now self is %s, mat2 is %s .",101+ "Float32 weight NZ is unsupported by the current SOC version [%s], now self is %s, mat2 is %s.",
102 op::ToString(socVersion).GetString(), op::ToString(self->GetDataType()).GetString(),102 op::ToString(socVersion).GetString(), op::ToString(self->GetDataType()).GetString(),
103 op::ToString(mat2->GetDataType()).GetString());103 op::ToString(mat2->GetDataType()).GetString());
104 return false;104 return false;
@@ -116,14 +116,14 @@ inline static bool CheckWeightNzDtypeValid(const aclTensor* self, const aclTenso
116 // keeptype模式支持类型检查116 // keeptype模式支持类型检查
117 if (cubeMathType == KEEP_DTYPE && !IsInputSupportFp32() &&117 if (cubeMathType == KEEP_DTYPE && !IsInputSupportFp32() &&
118 (self->GetDataType() == DataType::DT_FLOAT || mat2->GetDataType() == DataType::DT_FLOAT)) {118 (self->GetDataType() == DataType::DT_FLOAT || mat2->GetDataType() == DataType::DT_FLOAT)) {
119- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self dtype %s or mat2 dtype %s not support under keep type mode.",119+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self dtype %s or mat2 dtype %s are not supported under keep type mode.",
120 op::ToString(self->GetDataType()).GetString(), op::ToString(mat2->GetDataType()).GetString());120 op::ToString(self->GetDataType()).GetString(), op::ToString(mat2->GetDataType()).GetString());
121 return false;121 return false;
122 }122 }
123 if (cubeMathType == KEEP_DTYPE && out->GetDataType() == op::DataType::DT_FLOAT16 &&123 if (cubeMathType == KEEP_DTYPE && out->GetDataType() == op::DataType::DT_FLOAT16 &&
124 self->GetDataType() == op::DataType::DT_FLOAT) {124 self->GetDataType() == op::DataType::DT_FLOAT) {
125 OP_LOGE(ACLNN_ERR_PARAM_INVALID,125 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
126- "Input tensor's dtype[DT_FLOAT] should be same with output's dtype[DT_FLOAT16].");126+ "Input tensor's dtype[DT_FLOAT] should be the same as output's dtype[DT_FLOAT16].");
127 return false;127 return false;
128 }128 }
129 return CheckWeightNzDtype(self, mat2);129 return CheckWeightNzDtype(self, mat2);
@@ -160,7 +160,7 @@ static bool CheckShapeValid(const aclTensor* self, const aclTensor* mat2)
160 160 
161 // Tensor1 dims number is 0 OR error dims number is 0161 // Tensor1 dims number is 0 OR error dims number is 0
162 if (dimTensor1 == 0 || dimTensor2 == 0) {162 if (dimTensor1 == 0 || dimTensor2 == 0) {
163- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Matmul not support %s, %s", op::ToString(selfShape).GetString(),163+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Matmul does not support %s, %s", op::ToString(selfShape).GetString(),
164 op::ToString(mat2Shape).GetString());164 op::ToString(mat2Shape).GetString());
165 return false;165 return false;
166 } else if (dimTensor2 == 1 || dimTensor2 == 2) { // tensor1 dims number is 1 OR tensor2 dims number is 2166 } else if (dimTensor2 == 1 || dimTensor2 == 2) { // tensor1 dims number is 1 OR tensor2 dims number is 2
@@ -203,7 +203,7 @@ static bool CheckFormat(const aclTensor* selfTensor, [[maybe_unused]] const aclT
203 bool noSupportFormat = ((selfFormat == Format::FORMAT_FRACTAL_NZ) ||203 bool noSupportFormat = ((selfFormat == Format::FORMAT_FRACTAL_NZ) ||
204 (outTensorFormat == Format::FORMAT_FRACTAL_NZ));204 (outTensorFormat == Format::FORMAT_FRACTAL_NZ));
205 if (noSupportFormat) {205 if (noSupportFormat) {
206- OP_LOGE(ACLNN_ERR_PARAM_INVALID, " The 'self' or 'out' tensor currently does not support NZ format");206+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The 'self' or 'out' tensor currently does not support NZ format");
207 return false;207 return false;
208 }208 }
209 return true;209 return true;
@@ -422,7 +422,7 @@ static const aclTensor* BuildMatMulWeightNzGraph(const aclTensor* self, const ac
422 422 
423 const aclTensor* matmulOut = nullptr;423 const aclTensor* matmulOut = nullptr;
424 424 
425- // adpat for weightNz transpose scene425+ // adapt for weightNz transpose scene
426 bool transposeX2 = GetTransposeAttrValue(mat2);426 bool transposeX2 = GetTransposeAttrValue(mat2);
427 // swap last two dims value427 // swap last two dims value
428 if (transposeX2) {428 if (transposeX2) {
@@ -693,7 +693,7 @@ public:
693 if (dimTensor1 == 1) {693 if (dimTensor1 == 1) {
694 dimData = FVector<int64_t>{0}; // unsquee dim 0694 dimData = FVector<int64_t>{0}; // unsquee dim 0
695 } else {695 } else {
696- dimData = FVector<int64_t>{0, 1}; // unsquee dim 0,1696+ dimData = FVector<int64_t>{0, 1}; // unsquee dim 0,1
697 }697 }
698 auto selfUnsqueeze = ContiguousUnsqueezeNd(matA, dimData, executor);698 auto selfUnsqueeze = ContiguousUnsqueezeNd(matA, dimData, executor);
699 CHECK_RET(selfUnsqueeze != nullptr, ACLNN_ERR_INNER_NULLPTR);699 CHECK_RET(selfUnsqueeze != nullptr, ACLNN_ERR_INNER_NULLPTR);
Mmatmul/mat_mul_v3/op_host/op_api/aclnn_matmul.h+1-1
@@ -49,7 +49,7 @@ ACLNN_API aclnnStatus aclnnMatmulWeightNzGetWorkspaceSize(const aclTensor* self,
49 49 
50/**50/**
51 * @brief aclnnMatmulWeightNz的第二段接口，用于执行计算。51 * @brief aclnnMatmulWeightNz的第二段接口，用于执行计算。
52- * @param [in] workspace: 在npu device侧申请的workspace内存起址。52+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
53 * @param [in] workspace_size: 在npu device侧申请的workspace大小，由第一段接口aclnnMatmulWeightNzGetWorkspaceSize获取。53 * @param [in] workspace_size: 在npu device侧申请的workspace大小，由第一段接口aclnnMatmulWeightNzGetWorkspaceSize获取。
54 * @param [in] executor: op执行器，包含了算子计算流程。54 * @param [in] executor: op执行器，包含了算子计算流程。
55 * @param [in] stream: acl stream流。55 * @param [in] stream: acl stream流。
Mmatmul/mat_mul_v3/op_host/op_tiling/arch35/matmul_v3_asw_tiling.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file matmul_v3_asw_tiling.cc12+ * \file matmul_v3_asw_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15#include "matmul_v3_asw_tiling.h"15#include "matmul_v3_asw_tiling.h"
@@ -71,4 +71,4 @@ ge::graphStatus MatMulV3AswTiling::GetTilingData(TilingResult& tiling) const
71 return GetTilingDataImpl<MatMulV3TilingData>(tiling);71 return GetTilingDataImpl<MatMulV3TilingData>(tiling);
72}72}
73} // namespace matmul_v3_advanced73} // namespace matmul_v3_advanced
74-} // namespace optiling74+} // namespace optiling
Mmatmul/mat_mul_v3/op_host/op_tiling/arch35/matmul_v3_k_equal_zero_tiling.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file matmul_v3_k_equal_zero_tiling.cc12+ * \file matmul_v3_k_equal_zero_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
@@ -63,4 +63,4 @@ ge::graphStatus MatMulV3KEqZeroTiling::GetTilingData(TilingResult& tiling) const
63 return GetTilingDataImpl<MatMulV3KEqZeroBasicTilingData>(tiling);63 return GetTilingDataImpl<MatMulV3KEqZeroBasicTilingData>(tiling);
64}64}
65} // namespace matmul_v3_advanced65} // namespace matmul_v3_advanced
66-} // namespace optiling66+} // namespace optiling
Mmatmul/mat_mul_v3/op_host/op_tiling/arch35/matmul_v3_tiling_helper.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file matmul_v3_tiling_helper.cc12+ * \file matmul_v3_tiling_helper.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/mat_mul_v3/op_host/op_tiling/arch35/matmul_v3_tiling_key.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file matmul_v3_tiling_key.cc12+ * \file matmul_v3_tiling_key.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/mat_mul_v3/op_host/op_tiling/arch35/matmul_v3_to_mul_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file matmul_v3_to_mul_tiling.cc12+ * \file matmul_v3_to_mul_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/mat_mul_v3/op_host/op_tiling/arch35/matmul_v3_to_multi_mul_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file matmul_v3_to_multi_mul_tiling.cc12+ * \file matmul_v3_to_multi_mul_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/mat_mul_v3/op_host/op_tiling/matmul_v3_base_tiling.cpp+3-3
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file matmul_v3_base_tiling.cc12+ * \file matmul_v3_base_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
@@ -765,7 +765,7 @@ bool MatmulV3BaseTiling::IsPowerOfTwo(uint64_t x) const { return x > 0UL && (x &
765void MatmulV3BaseTiling::OptimizeLoadBalanceBasicKernel()765void MatmulV3BaseTiling::OptimizeLoadBalanceBasicKernel()
766{766{
767 OP_LOGI(args_.opName, "Optimize LoadBalance for BasicKernel");767 OP_LOGI(args_.opName, "Optimize LoadBalance for BasicKernel");
768- // 判决门限768+ // 判断门限
769 // 1. 需要tiling_key==10000000000000000001UL或10000000000000000000UL，表示时BasicKernel场景769 // 1. 需要tiling_key==10000000000000000001UL或10000000000000000000UL，表示时BasicKernel场景
770 // 2. baseM==128或256，baseN==256或128, baseK==128/dtypesize，刚好将L0 cache的利用最大化770 // 2. baseM==128或256，baseN==256或128, baseK==128/dtypesize，刚好将L0 cache的利用最大化
771 // 3. 要求m,n方向分合小于4轮或调小m可以不增加轮次771 // 3. 要求m,n方向分合小于4轮或调小m可以不增加轮次
@@ -807,7 +807,7 @@ void MatmulV3BaseTiling::OptimizeLoadBalanceBasicKernel()
807 807 
808void MatmulV3BaseTiling::OptimizeBasicKernelStepK()808void MatmulV3BaseTiling::OptimizeBasicKernelStepK()
809{809{
810- // 判决门限810+ // 判断门限
811 // 1. 需要tiling_key==10000000000000000001UL，表示时BasicKernel场景811 // 1. 需要tiling_key==10000000000000000001UL，表示时BasicKernel场景
812 // 2. baseM==128或256，baseN==256或128,baseK==64，刚好将L0 cache的利用最大化812 // 2. baseM==128或256，baseN==256或128,baseK==64，刚好将L0 cache的利用最大化
813 // 3. 要求m,n是256的倍数且大于等于768, k是256的倍数但不能是2的幂次方 或者m=[10368, 18000] && n,k=[1280, 5120]813 // 3. 要求m,n是256的倍数且大于等于768, k是256的倍数但不能是2的幂次方 或者m=[10368, 18000] && n,k=[1280, 5120]
Mmatmul/mat_mul_v3/op_host/op_tiling/matmul_v3_l2_cache.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_v3_l2_cache.cc12+ * \file matmul_v3_l2_cache.cpp
13 * \brief13 * \brief
14 */14 */
15#include "matmul_v3_l2_cache.h"15#include "matmul_v3_l2_cache.h"
@@ -167,4 +167,4 @@ void L2Cache::SetL2CacheFlag(TilingEnable tilingEnable, uint64_t l2Size, uint32_
167}167}
168 168 
169} // namespace matmul_v3169} // namespace matmul_v3
170-} // namespace optiling170+} // namespace optiling
Mmatmul/mat_mul_v3/op_host/op_tiling/matmul_v3_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file matmul_v3_tiling.cc12+ * \file matmul_v3_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15#include "matmul_v3_tiling.h"15#include "matmul_v3_tiling.h"
Mmatmul/mat_mul_v3/op_kernel/arch35/mat_mul_fixpipe_opti_basic_cmct.h+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file mat_mul_fixpipe_basic_cmct.h12+ * \file mat_mul_fixpipe_opti_basic_cmct.h
13 * \brief13 * \brief
14 */14 */
15#pragma once15#pragma once
Mmatmul/mat_mul_v3/op_kernel/mat_mul_base_block.h+1-1
@@ -346,4 +346,4 @@ __aicore__ inline void MatmulBaseBlock::UpdateBlockIndex()
346 }346 }
347}347}
348 348 
349-#endif // MMV3_MATMUL_BLOCK_H349+#endif // __OP_KERNEL_MATMUL_V3_BASE_BLOCK_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_base_kernel.h+1-1
@@ -207,4 +207,4 @@ __aicore__ inline void MatmulBaseKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BLOCK
207 mm_.SetHF32(false, 0);207 mm_.SetHF32(false, 0);
208 return;208 return;
209}209}
210-#endif // MMV3_MATMUL_KERNEL_H210+#endif // __OP_KERNEL_MATMUL_V3_BASE_KERNEL_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_base_kernel_vec_nz2nd.h+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file mat_mul_base_vector_nz2nd_kernel.h12+ * \file mat_mul_base_kernel_vec_nz2nd.h
13 * \brief Matmul kernel with AIV-based NZ2ND conversion for half/bfloat16 output13 * \brief Matmul kernel with AIV-based NZ2ND conversion for half/bfloat16 output
14 */14 */
15#ifndef OP_KERNEL_MATMUL_V3_BASE_VECTOR_NZ2ND_KERNEL_H15#ifndef OP_KERNEL_MATMUL_V3_BASE_VECTOR_NZ2ND_KERNEL_H
@@ -125,7 +125,7 @@ __aicore__ inline void MatmulBaseVectorNz2NdKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_
125 }125 }
126 126 
127 if (ubProcessMNum == 0UL) {127 if (ubProcessMNum == 0UL) {
128- //补充同步信号128+ // 补充同步信号
129 WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));129 WaitFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
130 CrossCoreSetFlag<0x2, PIPE_MTE2>(AIC_SYNC_AIV_FLAG + pingPongId);130 CrossCoreSetFlag<0x2, PIPE_MTE2>(AIC_SYNC_AIV_FLAG + pingPongId);
131 SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));131 SetFlag<HardEvent::MTE3_MTE2>(static_cast<event_t>(AIV_DB_SYNC_FLAG + pingPongId));
Mmatmul/mat_mul_v3/op_kernel/mat_mul_cvp_base_kernel.h+1-1
@@ -361,4 +361,4 @@ __aicore__ inline void MatmulCvpBaseKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, BL
361 return;361 return;
362}362}
363} // namespace MatmulV3363} // namespace MatmulV3
364-#endif // MMV3_MATMUL_KERNEL_H364+#endif // __OP_KERNEL_MATMUL_V3_CVP_BASE_KERNEL__H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_deterministic_splitk_kernel.h+1-1
@@ -1672,4 +1672,4 @@ __aicore__ inline void MatMulKernelDeterministicSplitK(GM_ADDR aGM, GM_ADDR bGM,
1672 }1672 }
1673}1673}
1674 1674 
1675-#endif // __OP_KERNEL_MATMUL_V3_H__1675+#endif // __OP_KERNEL_MATMUL_V3_DETERMINISTIC_SPLITK_KERNEL_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_l1_full_load.h+1-1
@@ -500,4 +500,4 @@ __aicore__ inline void MatmulBaseUnAlignedKernelBL1FullLoad<A_TYPE, B_TYPE, C_TY
500}500}
501 501 
502} // namespace MatmulV3502} // namespace MatmulV3
503-#endif // MMV3_MATMUL_BL1_FULL_LOAD_H503+#endif // __OP_KERNEL_MATMUL_V3_BL1_FULL_LOAD_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_sc_splitk_block.h+1-1
@@ -311,4 +311,4 @@ __aicore__ inline void MatmulSingleCoreSplitKBaseBlock::CalcGMOffset(uint64_t in
311 311 
312__aicore__ inline void MatmulSingleCoreSplitKBaseBlock::UpdateBlockIndex() { params_.index += 1; }312__aicore__ inline void MatmulSingleCoreSplitKBaseBlock::UpdateBlockIndex() { params_.index += 1; }
313 313 
314-#endif // MMV3_MATMUL_BLOCK_H314+#endif // __OP_KERNEL_MATMUL_V3_SC_SPLITK_BLOCK_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_sc_splitk_kernel.h+4-4
@@ -46,7 +46,7 @@ public:
46 46 
47 __aicore__ inline void SetParamAndExec(int kIndex, uint8_t enAtomic = 0);47 __aicore__ inline void SetParamAndExec(int kIndex, uint8_t enAtomic = 0);
48 48 
49- __aicore__ inline void Exector(uint8_t enAtomic = 0);49+ __aicore__ inline void Executor(uint8_t enAtomic = 0);
50 50 
51 __aicore__ inline void End() { mm_.End(); }51 __aicore__ inline void End() { mm_.End(); }
52 52 
@@ -319,7 +319,7 @@ __aicore__ inline void MatMulBaseKernelSingleCoreSplitK<A_TYPE, B_TYPE, L0C_TYPE
319template <class A_TYPE, class B_TYPE, class L0C_TYPE, class OUTPUT_TYPE, class BIAS_TYPE, class BLOCK_TYPE,319template <class A_TYPE, class B_TYPE, class L0C_TYPE, class OUTPUT_TYPE, class BIAS_TYPE, class BLOCK_TYPE,
320 const MatmulConfig& MM_CFG, const bool IS_NKM>320 const MatmulConfig& MM_CFG, const bool IS_NKM>
321__aicore__ inline void MatMulBaseKernelSingleCoreSplitK<A_TYPE, B_TYPE, L0C_TYPE, OUTPUT_TYPE, BIAS_TYPE, BLOCK_TYPE,321__aicore__ inline void MatMulBaseKernelSingleCoreSplitK<A_TYPE, B_TYPE, L0C_TYPE, OUTPUT_TYPE, BIAS_TYPE, BLOCK_TYPE,
322- MM_CFG, IS_NKM>::Exector(uint8_t enAtomic)322+ MM_CFG, IS_NKM>::Executor(uint8_t enAtomic)
323{323{
324 if constexpr (!IS_NKM) {324 if constexpr (!IS_NKM) {
325 for (uint64_t innerMIndex = 0; innerMIndex < block_.params_.innerLoopM; ++innerMIndex) {325 for (uint64_t innerMIndex = 0; innerMIndex < block_.params_.innerLoopM; ++innerMIndex) {
@@ -402,7 +402,7 @@ __aicore__ inline void MatMulBaseKernelSingleCoreSplitK<A_TYPE, B_TYPE, L0C_TYPE
402 block_.InitBlockIndex();402 block_.InitBlockIndex();
403 for (uint64_t j = 0; j < block_.params_.realRound; ++j) {403 for (uint64_t j = 0; j < block_.params_.realRound; ++j) {
404 block_.UpdateBlockCnt();404 block_.UpdateBlockCnt();
405- Exector(enAtomic);405+ Executor(enAtomic);
406 block_.UpdateBlockIndex();406 block_.UpdateBlockIndex();
407 }407 }
408 PipeBarrier<PIPE_ALL>();408 PipeBarrier<PIPE_ALL>();
@@ -666,4 +666,4 @@ __aicore__ inline void MatMulSingleCoreSplitKKernel<A_TYPE, B_TYPE, C_TYPE, BIAS
666 }666 }
667}667}
668 668 
669-#endif // MMV3_MATMUL_KERNEL_H669+#endif // __OP_KERNEL_MATMUL_V3_SC_SPLITK_KERNEL_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_sc_splitk_kernel_gm_to_l1.h+1-1
@@ -1156,4 +1156,4 @@ __aicore__ inline void MatMulSingleCoreSplitKKernelGmToL1<A_TYPE, B_TYPE, C_TYPE
1156 }1156 }
1157}1157}
1158 1158 
1159-#endif // MMV3_MATMUL_KERNEL_H1159+#endif // __OP_KERNEL_MATMUL_V3_SC_SPLITK_KERNEL_GM_TO_L1_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_unaligned_base_kernel.h+1-1
@@ -202,4 +202,4 @@ __aicore__ inline void MatmulBaseUnAlignedKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TY
202 }202 }
203}203}
204 204 
205-#endif // MMV3_MATMUL_KERNEL_H205+#endif // __OP_KERNEL_MATMUL_V3_UNALIGNED_BASE_KERNEL_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_unaligned_deterministic_splitk_kernel.h+1-1
@@ -237,4 +237,4 @@ __aicore__ inline void MatMulUnAlignedKernelDeterministicSplitK(GM_ADDR aGM, GM_
237 return;237 return;
238 }238 }
239}239}
240-#endif // __OP_KERNEL_MATMUL_V3_H__240+#endif // __OP_KERNEL_MATMUL_V3_UNALIGNED_DETERMINISTIC_SPLITK_KERNEL_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_unaligned_sc_splitk_kernel.h+1-1
@@ -27,4 +27,4 @@ class MatMulUnAlignedSingleCoreSplitKKernel
27public:27public:
28 __aicore__ inline MatMulUnAlignedSingleCoreSplitKKernel() = default;28 __aicore__ inline MatMulUnAlignedSingleCoreSplitKKernel() = default;
29};29};
30-#endif // MMV3_MATMUL_KERNEL_H30+#endif // __OP_KERNEL_MATMUL_V3_UNALIGNED_SC_SPLITK_KERNEL_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_unaligned_sc_splitk_kernel_gm_to_l1.h+1-1
@@ -29,4 +29,4 @@ class MatMulUnAlignedSingleCoreSplitKKernelGmToL1
29public:29public:
30 __aicore__ inline MatMulUnAlignedSingleCoreSplitKKernelGmToL1() = default;30 __aicore__ inline MatMulUnAlignedSingleCoreSplitKKernelGmToL1() = default;
31};31};
32-#endif // MMV3_MATMUL_KERNEL_H32+#endif // __OP_KERNEL_MATMUL_V3_UNALIGNED_SC_SPLITK_KERNEL_GM_TO_L1_H__
Mmatmul/mat_mul_v3/op_kernel/mat_mul_v3_common.h+1-1
@@ -440,4 +440,4 @@ __aicore__ inline void RemovePaddingImpl(GlobalTensor<T2> outputGlobal, GlobalTe
440}440}
441 441 
442#endif442#endif
443-#endif // __OP_KERNEL_MATMUL_V3_H__443+#endif // __OP_KERNEL_MATMUL_V3_COMMON_H__
Mmatmul/matmul_v2_compress_dequant/op_host/matmul_v2_compress_dequant_tiling.cpp+9-8
@@ -71,10 +71,11 @@ inline uint32_t CeilDivU32(uint32_t dividend, uint32_t divisor)
71 return (dividend + divisor - 1) / divisor;71 return (dividend + divisor - 1) / divisor;
72}72}
73 73 
74-// // ============================================================74+// ============================================================
75-// // Compute k0 and compressOverlapN75+// Compute k0 and compressOverlapN
76-// // Ported from PpTilingData310P::End76+// Ported from PpTilingData310P::End
77-// // ============================================================77+// ============================================================
78+ 
78uint32_t ComputeK0(uint32_t n0, uint32_t n, bool isCompress, uint32_t tilingNVal, uint32_t& compressOverlapN)79uint32_t ComputeK0(uint32_t n0, uint32_t n, bool isCompress, uint32_t tilingNVal, uint32_t& compressOverlapN)
79{80{
80 compressOverlapN = 0;81 compressOverlapN = 0;
@@ -99,7 +100,7 @@ ge::graphStatus ExtractMatmulDims(const gert::TilingContext* context, uint32_t&
99{100{
100 const auto* x1ShapePtr = context->GetInputShape(INPUT_IDX_X1);101 const auto* x1ShapePtr = context->GetInputShape(INPUT_IDX_X1);
101 if (x1ShapePtr == nullptr) {102 if (x1ShapePtr == nullptr) {
102- OP_LOGE("MatMulV2CompressDequant", "x1 shape is nullptr.");103+ OP_LOGE("MatMulV2CompressDequant", "x1 shape is null.");
103 return ge::GRAPH_FAILED;104 return ge::GRAPH_FAILED;
104 }105 }
105 const auto& x1Shape = x1ShapePtr->GetOriginShape();106 const auto& x1Shape = x1ShapePtr->GetOriginShape();
@@ -185,7 +186,7 @@ static ge::graphStatus TbmmEinsumTilingFunc(gert::TilingContext* context, uint32
185 tiling.set_kLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.kLoop);186 tiling.set_kLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.kLoop);
186 tiling.set_nLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.nLoop);187 tiling.set_nLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.nLoop);
187 tiling.set_coreLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.coreLoop);188 tiling.set_coreLoop(tbmmEinsumTiling.ppMatmulDefaultTilingData_.coreLoop);
188- tiling.set_swizzlCount(tbmmEinsumTiling.ppMatmulDefaultTilingData_.swizzlCount);189+ tiling.set_swizzleCount(tbmmEinsumTiling.ppMatmulDefaultTilingData_.swizzleCount);
189 tiling.set_tilingK(tilingK);190 tiling.set_tilingK(tilingK);
190 tiling.set_tilingN(tilingN);191 tiling.set_tilingN(tilingN);
191 tiling.set_compressOverlapN(compressOverlapN);192 tiling.set_compressOverlapN(compressOverlapN);
@@ -241,14 +242,14 @@ ge::graphStatus TilingForMatmulV2CompressDequant(gert::TilingContext* context)
241ge::graphStatus TilingPrepareForMatmulV2CompressDequant(gert::TilingParseContext* context)242ge::graphStatus TilingPrepareForMatmulV2CompressDequant(gert::TilingParseContext* context)
242{243{
243 if (context == nullptr) {244 if (context == nullptr) {
244- OP_LOGE("MatMulV2CompressDequant", "TilingParse context is nullptr.");245+ OP_LOGE("MatMulV2CompressDequant", "TilingParse context is null.");
245 return ge::GRAPH_FAILED;246 return ge::GRAPH_FAILED;
246 }247 }
247 OP_LOGD(context, "TilingPrepareForMatmulV2CompressDequant start.");248 OP_LOGD(context, "TilingPrepareForMatmulV2CompressDequant start.");
248 249 
249 fe::PlatFormInfos* platformInfo = context->GetPlatformInfo();250 fe::PlatFormInfos* platformInfo = context->GetPlatformInfo();
250 if (platformInfo == nullptr) {251 if (platformInfo == nullptr) {
251- OP_LOGE("MatMulV2CompressDequant", "platformInfoPtr is null");252+ OP_LOGE("MatMulV2CompressDequant", "platformInfo is null");
252 return ge::GRAPH_FAILED;253 return ge::GRAPH_FAILED;
253 }254 }
254 255 
Mmatmul/matmul_v2_compress_dequant/op_host/matmul_v2_compress_dequant_tiling.h+1-1
@@ -36,7 +36,7 @@ TILING_DATA_FIELD_DEF(uint32_t, mLoop);
36TILING_DATA_FIELD_DEF(uint32_t, kLoop);36TILING_DATA_FIELD_DEF(uint32_t, kLoop);
37TILING_DATA_FIELD_DEF(uint32_t, nLoop);37TILING_DATA_FIELD_DEF(uint32_t, nLoop);
38TILING_DATA_FIELD_DEF(uint32_t, coreLoop);38TILING_DATA_FIELD_DEF(uint32_t, coreLoop);
39-TILING_DATA_FIELD_DEF(uint32_t, swizzlCount);39+TILING_DATA_FIELD_DEF(uint32_t, swizzleCount);
40TILING_DATA_FIELD_DEF(uint32_t, tilingK);40TILING_DATA_FIELD_DEF(uint32_t, tilingK);
41TILING_DATA_FIELD_DEF(uint32_t, tilingN);41TILING_DATA_FIELD_DEF(uint32_t, tilingN);
42TILING_DATA_FIELD_DEF(uint32_t, compressOverlapN);42TILING_DATA_FIELD_DEF(uint32_t, compressOverlapN);
Mmatmul/matmul_v2_compress_dequant/op_host/op_api/aclnn_matmul_compress_dequant.cpp+4-4
@@ -51,7 +51,7 @@ static op::FVector<int64_t> GetShape(const aclTensor* tensor)
51 op::FVector<int64_t> shape;51 op::FVector<int64_t> shape;
52 if (tensor == nullptr) {52 if (tensor == nullptr) {
53 shape.push_back(1);53 shape.push_back(1);
54- OP_LOGW("The input tensor of Func GetShape is nullptr");54+ OP_LOGW("The input tensor of GetShape is nullptr.");
55 return shape;55 return shape;
56 }56 }
57 if (tensor->GetViewShape().GetDimNum() == 0U) {57 if (tensor->GetViewShape().GetDimNum() == 0U) {
@@ -118,8 +118,8 @@ static bool CheckShapeValid(const aclTensor* x1, const aclTensor* x2, const aclI
118 int64_t x1KDim = 0;118 int64_t x1KDim = 0;
119 int64_t x2KDim = 0;119 int64_t x2KDim = 0;
120 120 
121- if (dimTensor1 != 2 || dimTensor2 != 1) { // ND format dims > 2 for x1121+ if (dimTensor1 != 2 || dimTensor2 != 1) { // x1 must be 2D in ND format
122- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "MatmulUnzip not support x1 shape %s, x2 shape %s",122+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "MatmulUnzip does not support x1 shape %s, x2 shape %s",
123 op::ToString(x1Shape).GetString(), op::ToString(x2Shape).GetString());123 op::ToString(x1Shape).GetString(), op::ToString(x2Shape).GetString());
124 return false;124 return false;
125 } else {125 } else {
@@ -231,7 +231,7 @@ static const aclTensor* BuildMatMulUnzipGraph(MatmulUnzipInput matmulUnzipInput,
231 if (matmulUnzipInput.deqScale->Numel() % DEQUANT_SCALE_ALIGN_SIZE == 0) {231 if (matmulUnzipInput.deqScale->Numel() % DEQUANT_SCALE_ALIGN_SIZE == 0) {
232 deqScale5HD = TensorReformat(matmulUnzipInput.deqScale, op::Format::FORMAT_NC1HWC0, executor);232 deqScale5HD = TensorReformat(matmulUnzipInput.deqScale, op::Format::FORMAT_NC1HWC0, executor);
233 } else {233 } else {
234- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Dequant Scale is invalid Data.");234+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Dequant scale data is invalid.");
235 return ProcessEmptyTensor(matmulUnzipInput.x1, out, executor);235 return ProcessEmptyTensor(matmulUnzipInput.x1, out, executor);
236 }236 }
237 const aclTensor* x2ReFormatFractalZ = TensorReformat(matmulUnzipInput.x2, op::Format::FORMAT_FRACTAL_Z, executor);237 const aclTensor* x2ReFormatFractalZ = TensorReformat(matmulUnzipInput.x2, op::Format::FORMAT_FRACTAL_Z, executor);
Mmatmul/matmul_v2_compress_dequant/op_host/op_api/matmul_compress_dequant.cpp+2-2
@@ -21,7 +21,7 @@ using namespace op;
21namespace l0op {21namespace l0op {
22 22 
23OP_TYPE_REGISTER(MatMulV2CompressDequant);23OP_TYPE_REGISTER(MatMulV2CompressDequant);
24-// 用大小写区分是否可以区分两个不同的kernel24+// 用大小写区分两个不同的 kernel
25const aclTensor* MatMulCompressDequant(const aclTensor* x1, const aclTensor* x2, const aclTensor* compressIndex,25const aclTensor* MatMulCompressDequant(const aclTensor* x1, const aclTensor* x2, const aclTensor* compressIndex,
26 const aclTensor* deqScale, const aclTensor* bias, const aclTensor* offsetW,26 const aclTensor* deqScale, const aclTensor* bias, const aclTensor* offsetW,
27 const bool transposeX1, const bool transposeX2, const aclIntArray* compressInfo,27 const bool transposeX1, const bool transposeX2, const aclIntArray* compressInfo,
@@ -31,7 +31,7 @@ const aclTensor* MatMulCompressDequant(const aclTensor* x1, const aclTensor* x2,
31 compressInfo, offsetX, algStr);31 compressInfo, offsetX, algStr);
32 auto mmCompressDequantOut = executor->AllocTensor(DataType::DT_FLOAT16, Format::FORMAT_FRACTAL_NZ,32 auto mmCompressDequantOut = executor->AllocTensor(DataType::DT_FLOAT16, Format::FORMAT_FRACTAL_NZ,
33 Format::FORMAT_ND);33 Format::FORMAT_ND);
34- // 是否可以复用tbe的infershape34+ // 复用 tbe 的 infershape
35 auto ret = INFER_SHAPE(MatMulV2CompressDequant, OP_INPUT(x1, x2, compressIndex, deqScale, bias, offsetW),35 auto ret = INFER_SHAPE(MatMulV2CompressDequant, OP_INPUT(x1, x2, compressIndex, deqScale, bias, offsetW),
36 OP_OUTPUT(mmCompressDequantOut),36 OP_OUTPUT(mmCompressDequantOut),
37 OP_ATTR(transposeX1, transposeX2, compressInfo, offsetX, algStr));37 OP_ATTR(transposeX1, transposeX2, compressInfo, offsetX, algStr));
Mmatmul/matmul_v2_compress_dequant/op_host/pp_matmul_common_tiling.h+5-5
@@ -217,7 +217,7 @@ template <typename PpTilingDataType>
217uint64_t Swizzl(PpTilingDataType& tilingData)217uint64_t Swizzl(PpTilingDataType& tilingData)
218{218{
219 uint64_t swizzlDirect = 0UL;219 uint64_t swizzlDirect = 0UL;
220- uint64_t swizzlCount = 1UL;220+ uint64_t swizzleCount = 1UL;
221 float m0 = tilingData.opShape.m0;221 float m0 = tilingData.opShape.m0;
222 float n0 = tilingData.opShape.n0;222 float n0 = tilingData.opShape.n0;
223 float m = tilingData.opShape.m;223 float m = tilingData.opShape.m;
@@ -234,22 +234,22 @@ uint64_t Swizzl(PpTilingDataType& tilingData)
234 cost = n0 * i + m0 * c;234 cost = n0 * i + m0 * c;
235 if (cost <= mincost) {235 if (cost <= mincost) {
236 mincost = cost;236 mincost = cost;
237- swizzlCount = i;237+ swizzleCount = i;
238 }238 }
239 } else {239 } else {
240 swizzlDirect = 0UL; // Zn240 swizzlDirect = 0UL; // Zn
241 cost = m0 * i + n0 * c;241 cost = m0 * i + n0 * c;
242 if (cost < mincost) {242 if (cost < mincost) {
243 mincost = cost;243 mincost = cost;
244- swizzlCount = i;244+ swizzleCount = i;
245 }245 }
246 }246 }
247 }247 }
248 tilingData.swizzlDirect = swizzlDirect;248 tilingData.swizzlDirect = swizzlDirect;
249- tilingData.swizzlCount = swizzlCount;249+ tilingData.swizzleCount = swizzleCount;
250 return swizzlDirect;250 return swizzlDirect;
251}251}
252} // namespace matmulCompressDequant252} // namespace matmulCompressDequant
253} // namespace pp_matmul253} // namespace pp_matmul
254} // namespace optiling254} // namespace optiling
255-#endif255+#endif
Mmatmul/matmul_v2_compress_dequant/op_host/pp_matmul_default.h+1-1
@@ -174,7 +174,7 @@ void PpMatMulDefault::PrintTiling()
174 OP_LOGD(context_->GetNodeName(), "PpMatMul kLoop: %ld.", ppMatmulDefaultTilingData_.kLoop);174 OP_LOGD(context_->GetNodeName(), "PpMatMul kLoop: %ld.", ppMatmulDefaultTilingData_.kLoop);
175 OP_LOGD(context_->GetNodeName(), "PpMatMul nLoop: %ld.", ppMatmulDefaultTilingData_.nLoop);175 OP_LOGD(context_->GetNodeName(), "PpMatMul nLoop: %ld.", ppMatmulDefaultTilingData_.nLoop);
176 OP_LOGD(context_->GetNodeName(), "PpMatMul coreLoop: %ld.", ppMatmulDefaultTilingData_.coreLoop);176 OP_LOGD(context_->GetNodeName(), "PpMatMul coreLoop: %ld.", ppMatmulDefaultTilingData_.coreLoop);
177- OP_LOGD(context_->GetNodeName(), "PpMatMul swizzlCount: %ld.", ppMatmulDefaultTilingData_.swizzlCount);177+ OP_LOGD(context_->GetNodeName(), "PpMatMul swizzleCount: %ld.", ppMatmulDefaultTilingData_.swizzleCount);
178 OP_LOGD(context_->GetNodeName(), "PpMatMul tilingKey: %d.", ppMatmulDefaultTilingData_.tilingKey);178 OP_LOGD(context_->GetNodeName(), "PpMatMul tilingKey: %d.", ppMatmulDefaultTilingData_.tilingKey);
179 OP_LOGD(context_->GetNodeName(), "PpMatMul blockDim: %ld.", ppMatmulDefaultTilingData_.blockDim);179 OP_LOGD(context_->GetNodeName(), "PpMatMul blockDim: %ld.", ppMatmulDefaultTilingData_.blockDim);
180 OP_LOGD(context_->GetNodeName(), "PpMatMul swizzlDirect: %ld.", ppMatmulDefaultTilingData_.swizzlDirect);180 OP_LOGD(context_->GetNodeName(), "PpMatMul swizzlDirect: %ld.", ppMatmulDefaultTilingData_.swizzlDirect);
Mmatmul/matmul_v2_compress_dequant/op_host/pp_matmul_info.h+2-2
@@ -72,7 +72,7 @@ struct PpMatmulDefaultTilingData {
72 uint64_t kLoop{1};72 uint64_t kLoop{1};
73 uint64_t nLoop{1};73 uint64_t nLoop{1};
74 uint64_t coreLoop{1};74 uint64_t coreLoop{1};
75- uint64_t swizzlCount{1};75+ uint64_t swizzleCount{1};
76 uint32_t tilingKey{0};76 uint32_t tilingKey{0};
77 uint64_t blockDim{1};77 uint64_t blockDim{1};
78 uint64_t swizzlDirect{0};78 uint64_t swizzlDirect{0};
@@ -87,4 +87,4 @@ struct PpMatmulDefaultTilingData {
87} // namespace matmulCompressDequant87} // namespace matmulCompressDequant
88} // namespace pp_matmul88} // namespace pp_matmul
89} // namespace optiling89} // namespace optiling
90-#endif90+#endif
Mmatmul/matmul_v2_compress_dequant/op_kernel/matmul_v2_compress_dequant.cpp+2-2
@@ -23,7 +23,7 @@ extern "C" __global__ __aicore__ void mat_mul_v2_compress_dequant(GM_ADDR x1, GM
23{23{
24 SetPadding<uint64_t>((uint64_t)0x0);24 SetPadding<uint64_t>((uint64_t)0x0);
25 SetVectorMask<int8_t>((uint64_t)-1, (uint64_t)-1);25 SetVectorMask<int8_t>((uint64_t)-1, (uint64_t)-1);
26- SetAtomicnone();26+ AscendC::SetAtomicNone();
27 27 
28 PpMatmulI8NzCompress<0, false, true, false, int8_t, uint64_t, int32_t> kernel;28 PpMatmulI8NzCompress<0, false, true, false, int8_t, uint64_t, int32_t> kernel;
29 29 
@@ -48,7 +48,7 @@ extern "C" __global__ __aicore__ void mat_mul_v2_compress_dequant(GM_ADDR x1, GM
48 ub_tiling.GetValue(8), // kLoop48 ub_tiling.GetValue(8), // kLoop
49 ub_tiling.GetValue(9), // nLoop49 ub_tiling.GetValue(9), // nLoop
50 ub_tiling.GetValue(10), // coreLoop50 ub_tiling.GetValue(10), // coreLoop
51- ub_tiling.GetValue(11), // swizzlCount51+ ub_tiling.GetValue(11), // swizzleCount
52 ub_tiling.GetValue(12), // tilingK52 ub_tiling.GetValue(12), // tilingK
53 ub_tiling.GetValue(13), // tilingN53 ub_tiling.GetValue(13), // tilingN
54 ub_tiling.GetValue(14)); // compressOverlapN54 ub_tiling.GetValue(14)); // compressOverlapN
Mmatmul/matmul_v2_compress_dequant/op_kernel/pp_matmul_i8_nz_compress.h+1-1
@@ -56,7 +56,7 @@ public:
56 __aicore__ explicit PpMatmulI8NzCompress()56 __aicore__ explicit PpMatmulI8NzCompress()
57 {57 {
58 SetPadding<uint64_t>((uint64_t)0x0);58 SetPadding<uint64_t>((uint64_t)0x0);
59- SetAtomicnone();59+ AscendC::SetAtomicNone();
60 set_ctrl(sbitset1(get_ctrl(), kSparseCtrlBitPos));60 set_ctrl(sbitset1(get_ctrl(), kSparseCtrlBitPos));
61 SetMasknorm();61 SetMasknorm();
62 }62 }
Mmatmul/mv/docs/aclnnMv.md+1-1
@@ -27,7 +27,7 @@
27- 计算公式：27- 计算公式：
28 28 
29 $$29 $$
30- out = input * vec30+ out = input @ vec
31 $$31 $$
32 32 
33 其中$input$为n*m的二维张量，$vec$为长度为m的一维张量，$out$为长度为n的一维张量。33 其中$input$为n*m的二维张量，$vec$为长度为m的一维张量，$out$为长度为n的一维张量。
Mmatmul/mv/op_host/op_api/aclnn_mv.cpp+1-1
@@ -85,7 +85,7 @@ static aclnnStatus CheckInputParams(const aclTensor* self, const aclTensor* vec,
85 // 2. self的dtype支持 + vec、out的数据类型要与self一致85 // 2. self的dtype支持 + vec、out的数据类型要与self一致
86 CHECK_RET(CheckDtypeSame(self, vec, out), ACLNN_ERR_PARAM_INVALID);86 CHECK_RET(CheckDtypeSame(self, vec, out), ACLNN_ERR_PARAM_INVALID);
87 87 
88- // self dtype 按soc校验。88+ // self dtype 按soc校验。
89 auto archRule = BuildRule();89 auto archRule = BuildRule();
90 CHECK_RET(archRule != nullptr, ACLNN_ERR_PARAM_INVALID);90 CHECK_RET(archRule != nullptr, ACLNN_ERR_PARAM_INVALID);
91 CHECK_RET(archRule->CheckInput(self, vec, nullptr, out, cubeMathType), ACLNN_ERR_PARAM_INVALID);91 CHECK_RET(archRule->CheckInput(self, vec, nullptr, out, cubeMathType), ACLNN_ERR_PARAM_INVALID);
Mmatmul/mv/op_host/op_api/aclnn_mv.h+5-5
@@ -23,12 +23,12 @@ extern "C" {
23 * @domain aclnn_ops_infer23 * @domain aclnn_ops_infer
24 *24 *
25 * 算子功能： 计算矩阵input与向量vec的乘积25 * 算子功能： 计算矩阵input与向量vec的乘积
26- * @param [in] self: npu device侧的aclTensor，数据类型支持FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128类型。支持26+ * @param [in] self: npu device侧的aclTensor，数据类型支持FLOAT16、BFLOAT16、FLOAT类型。支持
27 * [非连续的Tensor](#)，shape为n*m的二维张量，数据格式支持ND([参考](#))。27 * [非连续的Tensor](#)，shape为n*m的二维张量，数据格式支持ND([参考](#))。
28- * @param [in] vec: npu device侧的aclTensor。数据类型支持FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128类型，且数据类型与28+ * @param [in] vec: npu device侧的aclTensor。数据类型支持FLOAT16、BFLOAT16、FLOAT类型，且数据类型与
29 * self保持一致。支持[非连续的Tensor](#)，shape为长度为m的一维张量，数据格式支持ND([参考](#))。29 * self保持一致。支持[非连续的Tensor](#)，shape为长度为m的一维张量，数据格式支持ND([参考](#))。
30 * @param [in] out: npu30 * @param [in] out: npu
31- * device侧的aclTensor，数据类型支持FLOAT16、FLOAT、DOUBLE、COMPLEX64、COMPLEX128类型，且数据类型与self31+ * device侧的aclTensor，数据类型支持FLOAT16、BFLOAT16、FLOAT类型，且数据类型与self
32 * 保持一致。支持[非连续的Tensor](#)，shape为长度为n的一维张量，数据格式支持ND([参考](#))。32 * 保持一致。支持[非连续的Tensor](#)，shape为长度为n的一维张量，数据格式支持ND([参考](#))。
33 * @param [in] cubeMathType(INT8,33 * @param [in] cubeMathType(INT8,
34 * 计算输入)：INT8类型的枚举值，用于判断Cube单元应该使用哪种计算逻辑进行运算，0：KEEP_DTYPE, 保34 * 计算输入)：INT8类型的枚举值，用于判断Cube单元应该使用哪种计算逻辑进行运算，0：KEEP_DTYPE, 保
@@ -45,7 +45,7 @@ ACLNN_API aclnnStatus aclnnMvGetWorkspaceSize(const aclTensor* self, const aclTe
45 * @brief: aclnnMv的第二段接口，用于执行计算45 * @brief: aclnnMv的第二段接口，用于执行计算
46 *46 *
47 * 算子功能： 计算矩阵input与向量vec的乘积47 * 算子功能： 计算矩阵input与向量vec的乘积
48- * @param [in] workspace: 在npu device侧申请的workspace内存起址。48+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
49 * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnMvGetWorkspaceSize获取。49 * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnMvGetWorkspaceSize获取。
50 * @param [in] executor: op执行器，包含了算子计算流程。50 * @param [in] executor: op执行器，包含了算子计算流程。
51 * @param [in] stream: acl stream流。51 * @param [in] stream: acl stream流。
@@ -57,4 +57,4 @@ ACLNN_API aclnnStatus aclnnMv(void* workspace, uint64_t workspaceSize, aclOpExec
57}57}
58#endif58#endif
59 59 
60-#endif // OP_API_INC_LEVEL2_ACLNN_MV_H_60+#endif // OP_API_INC_LEVEL2_ACLNN_MV_H_
Mmatmul/transpose_batch_mat_mul/op_api/aclnn_transpose_batch_mat_mul.cpp+38-38
@@ -37,16 +37,16 @@
37using namespace std;37using namespace std;
38using namespace op;38using namespace op;
39using namespace Ops::NN;39using namespace Ops::NN;
40-static const std::initializer_list<op::DataType> x1_SUPPORT_LIST = {DataType::DT_FLOAT, DataType::DT_FLOAT16,40+static const std::initializer_list<op::DataType> X1_SUPPORT_LIST = {DataType::DT_FLOAT, DataType::DT_FLOAT16,
41 DataType::DT_BF16};41 DataType::DT_BF16};
42-static const std::initializer_list<op::DataType> x2_SUPPORT_LIST = {DataType::DT_FLOAT, DataType::DT_FLOAT16,42+static const std::initializer_list<op::DataType> X2_SUPPORT_LIST = {DataType::DT_FLOAT, DataType::DT_FLOAT16,
43 DataType::DT_BF16};43 DataType::DT_BF16};
44static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST_WEIGHTNZ = {op::DataType::DT_FLOAT16,44static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST_WEIGHTNZ = {op::DataType::DT_FLOAT16,
45 op::DataType::DT_BF16};45 op::DataType::DT_BF16};
46static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST_WEIGHTNZ = {46static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST_WEIGHTNZ = {
47 op::DataType::DT_FLOAT16, op::DataType::DT_BF16, DataType::DT_INT8};47 op::DataType::DT_FLOAT16, op::DataType::DT_BF16, DataType::DT_INT8};
48-static const std::initializer_list<op::DataType> x1_SCALE_SUPPORT_LIST = {DataType::DT_FLOAT16};48+static const std::initializer_list<op::DataType> X1_SCALE_SUPPORT_LIST = {DataType::DT_FLOAT16};
49-static const std::initializer_list<op::DataType> x2_SCALE_SUPPORT_LIST = {DataType::DT_FLOAT16};49+static const std::initializer_list<op::DataType> X2_SCALE_SUPPORT_LIST = {DataType::DT_FLOAT16};
50static const std::initializer_list<op::DataType> SCALE_DTYPE_SUPPORT_LIST = {DataType::DT_INT64, DataType::DT_UINT64};50static const std::initializer_list<op::DataType> SCALE_DTYPE_SUPPORT_LIST = {DataType::DT_INT64, DataType::DT_UINT64};
51static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST = {DataType::DT_FLOAT, DataType::DT_FLOAT16,51static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST = {DataType::DT_FLOAT, DataType::DT_FLOAT16,
52 DataType::DT_BF16, DataType::DT_INT8};52 DataType::DT_BF16, DataType::DT_INT8};
@@ -81,7 +81,7 @@ inline static bool CheckDtypeValid(const aclTensor* x1, const aclTensor* x2, con
81 auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();81 auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
82 if ((npuArch != NpuArch::DAV_2201) && !IsNpuArch3510Series()) {82 if ((npuArch != NpuArch::DAV_2201) && !IsNpuArch3510Series()) {
83 OP_LOGE(ACLNN_ERR_PARAM_INVALID,83 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
84- "transposebatchmatmulweightnz is unsupported by the current SOC version [%s].",84+ "TransposeBatchMatMulWeightNz is unsupported by the current SOC version [%s].",
85 op::ToString(socVersion).GetString());85 op::ToString(socVersion).GetString());
86 return false;86 return false;
87 }87 }
@@ -90,8 +90,8 @@ inline static bool CheckDtypeValid(const aclTensor* x1, const aclTensor* x2, con
90 OP_CHECK_DTYPE_NOT_SUPPORT(out, OUT_DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);90 OP_CHECK_DTYPE_NOT_SUPPORT(out, OUT_DTYPE_SUPPORT_LIST_WEIGHTNZ, return false);
91 if (scale != nullptr) {91 if (scale != nullptr) {
92 OP_CHECK_DTYPE_NOT_SUPPORT(scale, SCALE_DTYPE_SUPPORT_LIST, return false);92 OP_CHECK_DTYPE_NOT_SUPPORT(scale, SCALE_DTYPE_SUPPORT_LIST, return false);
93- OP_CHECK_DTYPE_NOT_SUPPORT(x1, x1_SCALE_SUPPORT_LIST, return false);93+ OP_CHECK_DTYPE_NOT_SUPPORT(x1, X1_SCALE_SUPPORT_LIST, return false);
94- OP_CHECK_DTYPE_NOT_SUPPORT(x2, x2_SCALE_SUPPORT_LIST, return false);94+ OP_CHECK_DTYPE_NOT_SUPPORT(x2, X2_SCALE_SUPPORT_LIST, return false);
95 }95 }
96 if ((x1->GetDataType() != out->GetDataType()) && (npuArch == NpuArch::DAV_2201)) {96 if ((x1->GetDataType() != out->GetDataType()) && (npuArch == NpuArch::DAV_2201)) {
97 OP_LOGE(ACLNN_ERR_PARAM_INVALID, "x1's dtype [%s] and out's dtype [%s] are not equal in DAV_2201.",97 OP_LOGE(ACLNN_ERR_PARAM_INVALID, "x1's dtype [%s] and out's dtype [%s] are not equal in DAV_2201.",
@@ -102,12 +102,12 @@ inline static bool CheckDtypeValid(const aclTensor* x1, const aclTensor* x2, con
102 }102 }
103 103 
104 // Regular ND format checks104 // Regular ND format checks
105- OP_CHECK_DTYPE_NOT_SUPPORT(x1, x1_SUPPORT_LIST, return false);105+ OP_CHECK_DTYPE_NOT_SUPPORT(x1, X1_SUPPORT_LIST, return false);
106- OP_CHECK_DTYPE_NOT_SUPPORT(x2, x2_SUPPORT_LIST, return false);106+ OP_CHECK_DTYPE_NOT_SUPPORT(x2, X2_SUPPORT_LIST, return false);
107 if (scale != nullptr) {107 if (scale != nullptr) {
108 OP_CHECK_DTYPE_NOT_SUPPORT(scale, SCALE_DTYPE_SUPPORT_LIST, return false);108 OP_CHECK_DTYPE_NOT_SUPPORT(scale, SCALE_DTYPE_SUPPORT_LIST, return false);
109- OP_CHECK_DTYPE_NOT_SUPPORT(x1, x1_SCALE_SUPPORT_LIST, return false);109+ OP_CHECK_DTYPE_NOT_SUPPORT(x1, X1_SCALE_SUPPORT_LIST, return false);
110- OP_CHECK_DTYPE_NOT_SUPPORT(x2, x2_SCALE_SUPPORT_LIST, return false);110+ OP_CHECK_DTYPE_NOT_SUPPORT(x2, X2_SCALE_SUPPORT_LIST, return false);
111 }111 }
112 OP_CHECK_DTYPE_NOT_SUPPORT(out, OUT_DTYPE_SUPPORT_LIST, return false);112 OP_CHECK_DTYPE_NOT_SUPPORT(out, OUT_DTYPE_SUPPORT_LIST, return false);
113 return true;113 return true;
@@ -120,16 +120,16 @@ inline static bool CheckScaleValid(const aclTensor* scale, int64_t batch, int64_
120 auto dimTensorScale = scale->GetViewShape().GetDimNum();120 auto dimTensorScale = scale->GetViewShape().GetDimNum();
121 int64_t scaleDim = scale->GetViewShape().GetDim(0);121 int64_t scaleDim = scale->GetViewShape().GetDim(0);
122 if (ops::FloorDiv(INT64_MAX, batch) < n) {122 if (ops::FloorDiv(INT64_MAX, batch) < n) {
123- OP_LOGE(ACLNN_ERR_PARAM_INVALID, " batch mul N > INT64_MAX");123+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "batch * N > INT64_MAX");
124 return false;124 return false;
125 }125 }
126 if ((dimTensorScale != 1) || (scaleDim != batch * n)) {126 if ((dimTensorScale != 1) || (scaleDim != batch * n)) {
127 OP_LOGE(ACLNN_ERR_PARAM_INVALID,127 OP_LOGE(ACLNN_ERR_PARAM_INVALID,
128- "dimTensorScale[%zu] != 1 or the length of the first dim of scale != batch mul N", dimTensorScale);128+ "dimTensorScale[%zu] != 1 or the length of the first dim of scale != batch * N", dimTensorScale);
129 return false;129 return false;
130 }130 }
131 if (!IsNpuArch3510Series() && scaleDim >= SUPPORTED_INNER_AXIS) {131 if (!IsNpuArch3510Series() && scaleDim >= SUPPORTED_INNER_AXIS) {
132- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "batch mul N should be less than 65536.");132+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "batch * N should be less than 65536.");
133 return false;133 return false;
134 }134 }
135 }135 }
@@ -222,12 +222,12 @@ static inline bool CheckNzStorageShape(const aclTensor* x2)
222 auto storageShape = x2->GetStorageShape();222 auto storageShape = x2->GetStorageShape();
223 auto storageShapeDim = storageShape.GetDimNum();223 auto storageShapeDim = storageShape.GetDimNum();
224 if (x2Shape[1] == 1 || x2Shape[2] == 1) {224 if (x2Shape[1] == 1 || x2Shape[2] == 1) {
225- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The k-axis or n-axis can not be 1 when the format is nz.");225+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The k-axis or n-axis cannot be 1 when the format is nz.");
226 return false;226 return false;
227 }227 }
228 OP_CHECK(228 OP_CHECK(
229 storageShapeDim == EXPECTED_NZ_DIM,229 storageShapeDim == EXPECTED_NZ_DIM,
230- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Only supports x2 storageShapeDim is 5, which are [%zu].", storageShapeDim),230+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Only supports x2 storageShapeDim is 5, which is [%zu].", storageShapeDim),
231 return false);231 return false);
232 return true;232 return true;
233}233}
@@ -245,7 +245,7 @@ inline static aclnnStatus CheckParams(const aclTensor* x1, const aclTensor* x2,
245 }245 }
246 // perm必须为3维246 // perm必须为3维
247 if (perm_x1->Size() != EXPECTED_DIM || perm_x2->Size() != EXPECTED_DIM || perm_y->Size() != EXPECTED_DIM) {247 if (perm_x1->Size() != EXPECTED_DIM || perm_x2->Size() != EXPECTED_DIM || perm_y->Size() != EXPECTED_DIM) {
248- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The perm parameter must be three-dimensional!");248+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The perm parameter must be three-dimensional.");
249 return ACLNN_ERR_PARAM_INVALID;249 return ACLNN_ERR_PARAM_INVALID;
250 }250 }
251 if (IsNpuArch3510Series() && cubeMathType == -1) {251 if (IsNpuArch3510Series() && cubeMathType == -1) {
@@ -264,7 +264,7 @@ inline static aclnnStatus CheckParams(const aclTensor* x1, const aclTensor* x2,
264 // 不支持x1Format、 outFormat为NZ264 // 不支持x1Format、 outFormat为NZ
265 if (ge::GetPrimaryFormat(x1->GetStorageFormat()) == Format::FORMAT_FRACTAL_NZ ||265 if (ge::GetPrimaryFormat(x1->GetStorageFormat()) == Format::FORMAT_FRACTAL_NZ ||
266 ge::GetPrimaryFormat(out->GetStorageFormat()) == Format::FORMAT_FRACTAL_NZ) {266 ge::GetPrimaryFormat(out->GetStorageFormat()) == Format::FORMAT_FRACTAL_NZ) {
267- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Format of x1 or out can not be FORMAT_FRACTAL_NZ.");267+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Format of x1 or out cannot be FORMAT_FRACTAL_NZ.");
268 return ACLNN_ERR_PARAM_INVALID;268 return ACLNN_ERR_PARAM_INVALID;
269 }269 }
270 CHECK_RET(CheckDtypeValid(x1, x2, scale, out), ACLNN_ERR_PARAM_INVALID);270 CHECK_RET(CheckDtypeValid(x1, x2, scale, out), ACLNN_ERR_PARAM_INVALID);
@@ -375,8 +375,8 @@ aclnnStatus aclnnTransposeBatchMatMulGetWorkspaceSize(const aclTensor* x1, const
375 DFX_IN(x1, x2, bias, scale, permX1, permX2, permY, cubeMathType, batchSplitFactor), DFX_OUT(out));375 DFX_IN(x1, x2, bias, scale, permX1, permX2, permY, cubeMathType, batchSplitFactor), DFX_OUT(out));
376 376 
377 // 固定写法, 创建OpExecutor377 // 固定写法, 创建OpExecutor
378- auto unique_executor = CREATE_EXECUTOR();378+ auto uniqueExecutor = CREATE_EXECUTOR();
379- CHECK_RET(unique_executor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);379+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
380 380 
381 // 路由cubeMathType4到cubeMathType0, 该接口不支持cubeMathType=4的场景381 // 路由cubeMathType4到cubeMathType0, 该接口不支持cubeMathType=4的场景
382 cubeMathType = routeCubeMathType4ToCubeMathType0DAV_2201(cubeMathType);382 cubeMathType = routeCubeMathType4ToCubeMathType0DAV_2201(cubeMathType);
@@ -387,34 +387,34 @@ aclnnStatus aclnnTransposeBatchMatMulGetWorkspaceSize(const aclTensor* x1, const
387 387 
388 // 空tensor 处理388 // 空tensor 处理
389 if (x1->IsEmpty() || x2->IsEmpty()) {389 if (x1->IsEmpty() || x2->IsEmpty()) {
390- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "aclnnTransposeBatchMatMul do not support empty tensor!");390+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "aclnnTransposeBatchMatMul does not support empty tensor.");
391 return ACLNN_ERR_PARAM_INVALID;391 return ACLNN_ERR_PARAM_INVALID;
392 }392 }
393 393 
394 if (bias != nullptr) {394 if (bias != nullptr) {
395- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The bias is not support in TBMM.");395+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The bias is not supported in TBMM.");
396 return ACLNN_ERR_PARAM_INVALID;396 return ACLNN_ERR_PARAM_INVALID;
397 }397 }
398 398 
399 // 构建matmul计算图399 // 构建matmul计算图
400 const aclTensor* tbmmOut = nullptr;400 const aclTensor* tbmmOut = nullptr;
401 tbmmOut = BuildTransposeBatchMatMulGraph(x1, x2, scale, permX1, permX2, permY, cubeMathType, batchSplitFactor,401 tbmmOut = BuildTransposeBatchMatMulGraph(x1, x2, scale, permX1, permX2, permY, cubeMathType, batchSplitFactor,
402- unique_executor.get());402+ uniqueExecutor.get());
403 CHECK_RET(tbmmOut != nullptr, ACLNN_ERR_PARAM_INVALID);403 CHECK_RET(tbmmOut != nullptr, ACLNN_ERR_PARAM_INVALID);
404 404 
405 if (tbmmOut->IsEmpty()) {405 if (tbmmOut->IsEmpty()) {
406 *workspaceSize = 0;406 *workspaceSize = 0;
407- unique_executor.ReleaseTo(executor);407+ uniqueExecutor.ReleaseTo(executor);
408 return ACLNN_SUCCESS;408 return ACLNN_SUCCESS;
409 }409 }
410 410 
411- tbmmOut = l0op::Cast(tbmmOut, out->GetDataType(), unique_executor.get());411+ tbmmOut = l0op::Cast(tbmmOut, out->GetDataType(), uniqueExecutor.get());
412 CHECK_RET(tbmmOut != nullptr, ACLNN_ERR_PARAM_INVALID);412 CHECK_RET(tbmmOut != nullptr, ACLNN_ERR_PARAM_INVALID);
413- auto viewCopyResult = l0op::ViewCopy(tbmmOut, out, unique_executor.get());413+ auto viewCopyResult = l0op::ViewCopy(tbmmOut, out, uniqueExecutor.get());
414 CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_PARAM_INVALID);414 CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_PARAM_INVALID);
415 415 
416- *workspaceSize = unique_executor->GetWorkspaceSize();416+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
417- unique_executor.ReleaseTo(executor);417+ uniqueExecutor.ReleaseTo(executor);
418 return ACLNN_SUCCESS;418 return ACLNN_SUCCESS;
419}419}
420 420 
@@ -436,8 +436,8 @@ aclnnStatus aclnnTransposeBatchMatMulWeightNzGetWorkspaceSize(const aclTensor* x
436 DFX_IN(x1, x2, bias, scale, permX1, permX2, permY, cubeMathType, batchSplitFactor), DFX_OUT(out));436 DFX_IN(x1, x2, bias, scale, permX1, permX2, permY, cubeMathType, batchSplitFactor), DFX_OUT(out));
437 437 
438 // 固定写法, 创建OpExecutor438 // 固定写法, 创建OpExecutor
439- auto unique_executor = CREATE_EXECUTOR();439+ auto uniqueExecutor = CREATE_EXECUTOR();
440- CHECK_RET(unique_executor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);440+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
441 441 
442 // 路由cubeMathType4到cubeMathType0, 该接口不支持cubeMathType=4的场景442 // 路由cubeMathType4到cubeMathType0, 该接口不支持cubeMathType=4的场景
443 cubeMathType = routeCubeMathType4ToCubeMathType0DAV_2201(cubeMathType);443 cubeMathType = routeCubeMathType4ToCubeMathType0DAV_2201(cubeMathType);
@@ -455,34 +455,34 @@ aclnnStatus aclnnTransposeBatchMatMulWeightNzGetWorkspaceSize(const aclTensor* x
455 455 
456 // 空tensor 处理456 // 空tensor 处理
457 if (x1->IsEmpty() || x2->IsEmpty()) {457 if (x1->IsEmpty() || x2->IsEmpty()) {
458- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "aclnnTransposeBatchMatMulWeightNz do not support empty tensor!");458+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "aclnnTransposeBatchMatMulWeightNz does not support empty tensor.");
459 return ACLNN_ERR_PARAM_INVALID;459 return ACLNN_ERR_PARAM_INVALID;
460 }460 }
461 461 
462 if (bias != nullptr) {462 if (bias != nullptr) {
463- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The bias is not support in TBMM.");463+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The bias is not supported in TBMM.");
464 return ACLNN_ERR_PARAM_INVALID;464 return ACLNN_ERR_PARAM_INVALID;
465 }465 }
466 466 
467 // 构建matmul计算图467 // 构建matmul计算图
468 const aclTensor* tbmmOut = nullptr;468 const aclTensor* tbmmOut = nullptr;
469 tbmmOut = BuildTransposeBatchMatMulWeightNzGraph(x1, x2, scale, permX1, permX2, permY, cubeMathType,469 tbmmOut = BuildTransposeBatchMatMulWeightNzGraph(x1, x2, scale, permX1, permX2, permY, cubeMathType,
470- batchSplitFactor, unique_executor.get());470+ batchSplitFactor, uniqueExecutor.get());
471 CHECK_RET(tbmmOut != nullptr, ACLNN_ERR_PARAM_INVALID);471 CHECK_RET(tbmmOut != nullptr, ACLNN_ERR_PARAM_INVALID);
472 472 
473 if (tbmmOut->IsEmpty()) {473 if (tbmmOut->IsEmpty()) {
474 *workspaceSize = 0;474 *workspaceSize = 0;
475- unique_executor.ReleaseTo(executor);475+ uniqueExecutor.ReleaseTo(executor);
476 return ACLNN_SUCCESS;476 return ACLNN_SUCCESS;
477 }477 }
478 478 
479- tbmmOut = l0op::Cast(tbmmOut, out->GetDataType(), unique_executor.get());479+ tbmmOut = l0op::Cast(tbmmOut, out->GetDataType(), uniqueExecutor.get());
480 CHECK_RET(tbmmOut != nullptr, ACLNN_ERR_PARAM_INVALID);480 CHECK_RET(tbmmOut != nullptr, ACLNN_ERR_PARAM_INVALID);
481- auto viewCopyResult = l0op::ViewCopy(tbmmOut, out, unique_executor.get());481+ auto viewCopyResult = l0op::ViewCopy(tbmmOut, out, uniqueExecutor.get());
482 CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_PARAM_INVALID);482 CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_PARAM_INVALID);
483 483 
484- *workspaceSize = unique_executor->GetWorkspaceSize();484+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
485- unique_executor.ReleaseTo(executor);485+ uniqueExecutor.ReleaseTo(executor);
486 return ACLNN_SUCCESS;486 return ACLNN_SUCCESS;
487}487}
488 488 
Mmatmul/transpose_batch_mat_mul/op_api/aclnn_transpose_batch_mat_mul.h+6-6
@@ -19,7 +19,7 @@ extern "C" {
19#endif19#endif
20 20 
21/**21/**
22- * @brief aclnnTransposeBatchMatmul的第一段接口，根据具体的计算流程，计算workspace大小。22+ * @brief aclnnTransposeBatchMatMul的第一段接口，根据具体的计算流程，计算workspace大小。
23 * @domain aclnn_ops_infer23 * @domain aclnn_ops_infer
24 * 算子功能：实现TransposeBatchMatMul算子和Quant的融合算子，TransposeBatchMatMul可支持参数为int8。24 * 算子功能：实现TransposeBatchMatMul算子和Quant的融合算子，TransposeBatchMatMul可支持参数为int8。
25 * @param [in] x1: matmul左矩阵，数据类型支持：float16、float32、bfloat16。25 * @param [in] x1: matmul左矩阵，数据类型支持：float16、float32、bfloat16。
@@ -44,7 +44,7 @@ ACLNN_API aclnnStatus aclnnTransposeBatchMatMulGetWorkspaceSize(const aclTensor*
44 uint64_t* workspaceSize, aclOpExecutor** executor);44 uint64_t* workspaceSize, aclOpExecutor** executor);
45 45 
46/**46/**
47- * @brief aclnnTransposeBatchMatmul的第二段接口，用于执行计算。47+ * @brief aclnnTransposeBatchMatMul的第二段接口，用于执行计算。
48 */48 */
49ACLNN_API aclnnStatus aclnnTransposeBatchMatMul(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,49ACLNN_API aclnnStatus aclnnTransposeBatchMatMul(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
50 const aclrtStream stream);50 const aclrtStream stream);
@@ -52,7 +52,7 @@ ACLNN_API aclnnStatus aclnnTransposeBatchMatMul(void* workspace, uint64_t worksp
52/**52/**
53 * @brief aclnnTransposeBatchMatMulWeightNz的第一段接口，根据具体的计算流程，计算workspace大小。53 * @brief aclnnTransposeBatchMatMulWeightNz的第一段接口，根据具体的计算流程，计算workspace大小。
54 * @domain aclnn_ops_infer54 * @domain aclnn_ops_infer
55- * 算子功能：相对于aclnnTransposeBatchMatmul, mat2为NZ格式。55+ * 算子功能：相对于aclnnTransposeBatchMatMul, mat2为NZ格式。
56 * @param [in] x1: matmul左矩阵，数据类型支持：float16、bfloat16。数据格式支持ND。56 * @param [in] x1: matmul左矩阵，数据类型支持：float16、bfloat16。数据格式支持ND。
57 * @param [in] x2: matmul右矩阵，数据类型支持：float16、bfloat16。数据格式支持NZ。57 * @param [in] x2: matmul右矩阵，数据类型支持：float16、bfloat16。数据格式支持NZ。
58 * @param [in] bias: 偏置，当前不支持。58 * @param [in] bias: 偏置，当前不支持。
@@ -74,9 +74,9 @@ ACLNN_API aclnnStatus aclnnTransposeBatchMatMulWeightNzGetWorkspaceSize(
74 74 
75/**75/**
76 * @brief aclnnTransposeBatchMatMulWeightNz的第二段接口，用于执行计算。76 * @brief aclnnTransposeBatchMatMulWeightNz的第二段接口，用于执行计算。
77- * @param [in] workspace: 在npu device侧申请的workspace内存起址。77+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
78- * @param [in] workspace_size: 在npu78+ * @param [in] workspaceSize: 在npu
79- * device侧申请的workspace大小，由第一段接口aclnnBatchMatMulWeightNzGetWorkspaceSize获取。79+ * device侧申请的workspace大小，由第一段接口aclnnTransposeBatchMatMulWeightNzGetWorkspaceSize获取。
80 * @param [in] executor: op执行器，包含了算子计算流程。80 * @param [in] executor: op执行器，包含了算子计算流程。
81 * @param [in] stream: acl stream流。81 * @param [in] stream: acl stream流。
82 * @return aclnnStatus: 返回状态码。82 * @return aclnnStatus: 返回状态码。
Mmatmul/transpose_batch_mat_mul/op_host/op_tiling/arch35/transpose_batch_mat_mul_asw_tiling.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file transpose_batch_mat_mul_asw_tiling.cc12+ * \file transpose_batch_mat_mul_asw_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/transpose_batch_mat_mul/op_host/op_tiling/arch35/transpose_batch_mat_mul_tiling_advanced.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file transpose_batch_mat_mul_tiling_advanced.cc12+ * \file transpose_batch_mat_mul_tiling_advanced.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
@@ -408,4 +408,4 @@ ge::graphStatus TransposeBatchMatMulTiling::GetBatchInfo(const gert::TilingConte
408 return ge::GRAPH_SUCCESS;408 return ge::GRAPH_SUCCESS;
409}409}
410} // namespace transpose_batch_mat_mul_advanced410} // namespace transpose_batch_mat_mul_advanced
411-} // namespace optiling411+} // namespace optiling
Mmatmul/transpose_batch_mat_mul/op_host/op_tiling/arch35/transpose_batch_mat_mul_tiling_key.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file transpose_batch_mat_mul_tiling_key.cc12+ * \file transpose_batch_mat_mul_tiling_key.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
Mmatmul/transpose_batch_mat_mul/op_host/op_tiling/pp_matmul_default.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file pp_matmul_default.cc12+ * \file pp_matmul_default.cpp
13 * \brief13 * \brief
14 */14 */
15#include "pp_matmul_default.h"15#include "pp_matmul_default.h"
@@ -199,4 +199,4 @@ void PpMatMulDefault::PrintTiling()
199}199}
200 200 
201} // namespace pp_matmul201} // namespace pp_matmul
202-} // namespace optiling202+} // namespace optiling
Mmatmul/transpose_batch_mat_mul/op_host/op_tiling/transpose_batch_mat_mul_base_tiling.cpp+5-5
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file transpose_batch_mat_mul_base_tiling.cc12+ * \file transpose_batch_mat_mul_base_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15#include "transpose_batch_mat_mul_base_tiling.h"15#include "transpose_batch_mat_mul_base_tiling.h"
@@ -237,7 +237,7 @@ static void TuneBaseMKN(matmul_v3::MatmulV3RunInfo& runInfo, const matmul_v3::Ma
237void TransposeBatchMatMulBaseTiling::ResetBasicBlock(uint64_t tempBaseM, uint64_t tempBaseN)237void TransposeBatchMatMulBaseTiling::ResetBasicBlock(uint64_t tempBaseM, uint64_t tempBaseN)
238{238{
239 OP_TILING_CHECK(tempBaseM == 0 && tempBaseN == 0,239 OP_TILING_CHECK(tempBaseM == 0 && tempBaseN == 0,
240- OP_LOGW(args_.opName, "tempBaseM == 0 && tempBaseN == 0 is invalid"), return );240+ OP_LOGW(args_.opName, "tempBaseM == 0 && tempBaseN == 0 is invalid"), return);
241 uint64_t baseKAlignNum = (!args_.isATrans && args_.isBTrans) ?241 uint64_t baseKAlignNum = (!args_.isATrans && args_.isBTrans) ?
242 GetAlignNumWithDataType(BASIC_BLOCK_SIZE_256, args_.aType) :242 GetAlignNumWithDataType(BASIC_BLOCK_SIZE_256, args_.aType) :
243 BLOCK_CUBE;243 BLOCK_CUBE;
@@ -421,7 +421,7 @@ ge::graphStatus TransposeBatchMatMulBaseTiling::CheckArgs()
421 auto* shape_scale = context_->GetOptionalInputShape(SCALE_IDX);421 auto* shape_scale = context_->GetOptionalInputShape(SCALE_IDX);
422 if (shape_scale != nullptr) {422 if (shape_scale != nullptr) {
423 OP_TILING_CHECK(shape_scale->GetShape().GetDim(0) >= kSupportedInnerAxis,423 OP_TILING_CHECK(shape_scale->GetShape().GetDim(0) >= kSupportedInnerAxis,
424- CUBE_INNER_ERR_REPORT(args_.opName, "batch mul n should be less than 65536."),424+ CUBE_INNER_ERR_REPORT(args_.opName, "batch * n should be less than 65536."),
425 return ge::GRAPH_FAILED);425 return ge::GRAPH_FAILED);
426 }426 }
427 if (attrs->GetAttrNum() >= ATTR_NUM) {427 if (attrs->GetAttrNum() >= ATTR_NUM) {
@@ -643,9 +643,9 @@ ge::graphStatus TransposeBatchMatMulBaseTiling::GetArgs()
643 CUBE_INNER_ERR_REPORT(args_.opName, "format and dtype check failed"), return ge::GRAPH_FAILED);643 CUBE_INNER_ERR_REPORT(args_.opName, "format and dtype check failed"), return ge::GRAPH_FAILED);
644 OP_TILING_CHECK(644 OP_TILING_CHECK(
645 (args_.bFormat == ge::FORMAT_FRACTAL_NZ) && ((transA_ != 213UL) || (transB_ != 123UL) || batchSplitFactor_ > 1),645 (args_.bFormat == ge::FORMAT_FRACTAL_NZ) && ((transA_ != 213UL) || (transB_ != 123UL) || batchSplitFactor_ > 1),
646- CUBE_INNER_ERR_REPORT(args_.opName, "The current attrs is not support weightNZ."), return ge::GRAPH_FAILED);646+ CUBE_INNER_ERR_REPORT(args_.opName, "The current attrs does not support weightNZ."), return ge::GRAPH_FAILED);
647 return ge::GRAPH_SUCCESS;647 return ge::GRAPH_SUCCESS;
648}648}
649 649 
650} // namespace transpose_batch_mat_mul650} // namespace transpose_batch_mat_mul
651-} // namespace optiling651+} // namespace optiling
Mmatmul/transpose_batch_mat_mul/op_host/op_tiling/transpose_batch_mat_mul_einsum_tiling.cpp+3-3
@@ -78,7 +78,7 @@ ge::graphStatus IsPpMatmulEinsumMode(gert::TilingContext* context)
78 constexpr size_t ALLOW_DIM = 3;78 constexpr size_t ALLOW_DIM = 3;
79 constexpr int64_t SUPPORTED_INNER_AXIS = 65536;79 constexpr int64_t SUPPORTED_INNER_AXIS = 65536;
80 if (CheckInputArgs(context) != ge::GRAPH_SUCCESS) {80 if (CheckInputArgs(context) != ge::GRAPH_SUCCESS) {
81- OP_LOGI(context->GetNodeName(), "Current scenario is not support PpMatmulEinsum.");81+ OP_LOGI(context->GetNodeName(), "Current scenario does not support PpMatmulEinsum.");
82 return ge::GRAPH_FAILED;82 return ge::GRAPH_FAILED;
83 }83 }
84 auto attrs = context->GetAttrs();84 auto attrs = context->GetAttrs();
@@ -111,7 +111,7 @@ ge::graphStatus IsPpMatmulEinsumMode(gert::TilingContext* context)
111 return ge::GRAPH_SUCCESS;111 return ge::GRAPH_SUCCESS;
112 }112 }
113 }113 }
114- OP_LOGI(context->GetNodeName(), "Current scenario is not support PpMatmulEinsum.");114+ OP_LOGI(context->GetNodeName(), "Current scenario does not support PpMatmulEinsum.");
115 return ge::GRAPH_FAILED;115 return ge::GRAPH_FAILED;
116}116}
117 117 
@@ -259,4 +259,4 @@ ge::graphStatus TransposeBatchMatMulEinsumTiling::PostTiling()
259 return ge::GRAPH_SUCCESS;259 return ge::GRAPH_SUCCESS;
260}260}
261} // namespace transpose_batch_mat_mul261} // namespace transpose_batch_mat_mul
262-} // namespace optiling262+} // namespace optiling
Mmatmul/transpose_batch_mat_mul/op_host/op_tiling/transpose_batch_mat_mul_tiling.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file transpose_batch_mat_mul_tiling.cc12+ * \file transpose_batch_mat_mul_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15#include "transpose_batch_mat_mul_tiling.h"15#include "transpose_batch_mat_mul_tiling.h"
@@ -96,4 +96,4 @@ IMPL_OP_OPTILING(TransposeBatchMatMul)
96 .Tiling(TransposeBatchMatMulTilingFunc)96 .Tiling(TransposeBatchMatMulTilingFunc)
97 .TilingParse<MatmulV3CompileInfo>(TilingPrepareForTransposeBatchMatMul)97 .TilingParse<MatmulV3CompileInfo>(TilingPrepareForTransposeBatchMatMul)
98 .GenSimplifiedKey(transpose_batch_matmul::GenSimplifiedKey);98 .GenSimplifiedKey(transpose_batch_matmul::GenSimplifiedKey);
99-} // namespace optiling99+} // namespace optiling
Mmatmul/transpose_batch_mat_mul/op_host/transpose_batch_mat_mul_infershape.cpp+11-13
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file transpose_batch_mat_mul_infer.cpp12+ * \file transpose_batch_mat_mul_infershape.cpp
13 * \brief13 * \brief
14 */14 */
15#include <string>15#include <string>
@@ -121,7 +121,7 @@ static ge::graphStatus CheckPermForTransposeBatchMatMul(const TypedContinuousVec
121 121 
122 const auto perm_y_attr = perm_y.GetData();122 const auto perm_y_attr = perm_y.GetData();
123 auto check_perm_y = *perm_y_attr == 1 && *(perm_y_attr + 1) == 0 && *(perm_y_attr + 2) == 2;123 auto check_perm_y = *perm_y_attr == 1 && *(perm_y_attr + 1) == 0 && *(perm_y_attr + 2) == 2;
124- CHECK(!check_perm_y, CUBE_INNER_ERR_REPORT("TBMM", "[InferShape] perm_y should {1, 0, 2}"),124+ CHECK(!check_perm_y, CUBE_INNER_ERR_REPORT("TBMM", "[InferShape] perm_y should be {1, 0, 2}"),
125 return ge::GRAPH_FAILED);125 return ge::GRAPH_FAILED);
126 return ge::GRAPH_SUCCESS;126 return ge::GRAPH_SUCCESS;
127}127}
@@ -137,14 +137,14 @@ static ge::graphStatus InferShapeForTransposeBatchMatMul(InferShapeContext* cont
137 auto attrs = context->GetAttrs();137 auto attrs = context->GetAttrs();
138 auto name_op = context->GetNodeName();138 auto name_op = context->GetNodeName();
139 CHECK(shape_x1 == nullptr || shape_x2 == nullptr || shape_y == nullptr || attrs == nullptr,139 CHECK(shape_x1 == nullptr || shape_x2 == nullptr || shape_y == nullptr || attrs == nullptr,
140- CUBE_INNER_ERR_REPORT(name_op, "[Infershape]shape or attrs is null"), return ge::GRAPH_FAILED);140+ CUBE_INNER_ERR_REPORT(name_op, "[InferShape] shape or attrs is null"), return ge::GRAPH_FAILED);
141 141 
142 const auto perm_x1 = attrs->GetListInt(0);142 const auto perm_x1 = attrs->GetListInt(0);
143 const auto perm_x2 = attrs->GetListInt(1);143 const auto perm_x2 = attrs->GetListInt(1);
144 const auto perm_y = attrs->GetListInt(2);144 const auto perm_y = attrs->GetListInt(2);
145 const auto batch_split_factor = attrs->GetAttrPointer<int32_t>(4); // batch_split_factor index is 4145 const auto batch_split_factor = attrs->GetAttrPointer<int32_t>(4); // batch_split_factor index is 4
146 CHECK(perm_x1 == nullptr || perm_x2 == nullptr || perm_y == nullptr || batch_split_factor == nullptr,146 CHECK(perm_x1 == nullptr || perm_x2 == nullptr || perm_y == nullptr || batch_split_factor == nullptr,
147- CUBE_INNER_ERR_REPORT(name_op, "[Infershape] null"), return ge::GRAPH_FAILED);147+ CUBE_INNER_ERR_REPORT(name_op, "[InferShape] null"), return ge::GRAPH_FAILED);
148 148 
149 CHECK(CheckPermForTransposeBatchMatMul(*perm_x1, *perm_x2, *perm_y) != ge::GRAPH_SUCCESS,149 CHECK(CheckPermForTransposeBatchMatMul(*perm_x1, *perm_x2, *perm_y) != ge::GRAPH_SUCCESS,
150 CUBE_INNER_ERR_REPORT(name_op, "[InferShape] Failed to check perm"), return ge::GRAPH_FAILED);150 CUBE_INNER_ERR_REPORT(name_op, "[InferShape] Failed to check perm"), return ge::GRAPH_FAILED);
@@ -175,7 +175,7 @@ static ge::graphStatus InferShapeForTransposeBatchMatMul(InferShapeContext* cont
175 return ge::GRAPH_FAILED);175 return ge::GRAPH_FAILED);
176 CHECK(!CheckIsUnknownDimNum(*shape_scale) &&176 CHECK(!CheckIsUnknownDimNum(*shape_scale) &&
177 shape_scale->GetDim(0) != shape_x1_transposed.GetDim(0) * shape_x2_transposed.GetDim(2),177 shape_scale->GetDim(0) != shape_x1_transposed.GetDim(0) * shape_x2_transposed.GetDim(2),
178- CUBE_INNER_ERR_REPORT(name_op, "The dimension of n mul b [%ld] and scale [%ld] tensors must be the same",178+ CUBE_INNER_ERR_REPORT(name_op, "The dimension of n * b [%ld] and scale [%ld] tensors must be the same",
179 shape_x1_transposed.GetDim(0) * shape_x2_transposed.GetDim(2),179 shape_x1_transposed.GetDim(0) * shape_x2_transposed.GetDim(2),
180 shape_scale->GetDim(0)),180 shape_scale->GetDim(0)),
181 return ge::GRAPH_FAILED);181 return ge::GRAPH_FAILED);
@@ -183,14 +183,12 @@ static ge::graphStatus InferShapeForTransposeBatchMatMul(InferShapeContext* cont
183 auto tensor_x2 = context->GetInputDesc(1);183 auto tensor_x2 = context->GetInputDesc(1);
184 ge::DataType dtype_x1 = tensor_x1->GetDataType();184 ge::DataType dtype_x1 = tensor_x1->GetDataType();
185 ge::DataType dtype_x2 = tensor_x2->GetDataType();185 ge::DataType dtype_x2 = tensor_x2->GetDataType();
186- CHECK(186+ CHECK(dtype_x1 != ge::DT_FLOAT16,
187- dtype_x1 != ge::DT_FLOAT16,187+ CUBE_INNER_ERR_REPORT(name_op, "the dtype of input only supports FLOAT16 when the scale takes effect."),
188- CUBE_INNER_ERR_REPORT(name_op, "the dtype of input is only supported FLOAT16 when the scale takes effect."),188+ return ge::GRAPH_FAILED);
189- return ge::GRAPH_FAILED);189+ CHECK(dtype_x2 != ge::DT_FLOAT16,
190- CHECK(190+ CUBE_INNER_ERR_REPORT(name_op, "the dtype of input only supports FLOAT16 when the scale takes effect."),
191- dtype_x2 != ge::DT_FLOAT16,191+ return ge::GRAPH_FAILED);
192- CUBE_INNER_ERR_REPORT(name_op, "the dtype of input is only supported FLOAT16 when the scale takes effect."),
193- return ge::GRAPH_FAILED);
194 }192 }
195 ge::graphStatus ret = SetShapeY(*shape_y, shape_x1_transposed, shape_x2_transposed, *perm_y, *batch_split_factor,193 ge::graphStatus ret = SetShapeY(*shape_y, shape_x1_transposed, shape_x2_transposed, *perm_y, *batch_split_factor,
196 shape_scale != nullptr);194 shape_scale != nullptr);
Mmatmul/transpose_batch_mat_mul/op_kernel/arch35/transpose_batch_mat_mul.cpp+5-5
@@ -40,7 +40,7 @@ constexpr CubeFormat format_x2 = CubeFormat::NZ;
40constexpr CubeFormat format_x2 = CubeFormat::ND;40constexpr CubeFormat format_x2 = CubeFormat::ND;
41#endif41#endif
42 42 
43-#define TBMM_IMPL_CLASS_COMMON_TRNAS(transA, transB, Mode, templateClass, ...) \43+#define TBMM_IMPL_CLASS_COMMON_TRANS(transA, transB, Mode, templateClass, ...) \
44 do { \44 do { \
45 using cType = MatmulType<AscendC::TPosition::GM, format_y, DTYPE_Y>; \45 using cType = MatmulType<AscendC::TPosition::GM, format_y, DTYPE_Y>; \
46 using biasType = MatmulType<AscendC::TPosition::GM, CubeFormat::ND, DTYPE_BIAS>; \46 using biasType = MatmulType<AscendC::TPosition::GM, CubeFormat::ND, DTYPE_BIAS>; \
@@ -82,25 +82,25 @@ __global__ __aicore__ void transpose_batch_mat_mul(GM_ADDR aGM, GM_ADDR bGM, GM_
82 if constexpr (API_LEVEL == TRANSPOSE_BATCH_MAT_MUL_HIGH_LEVEL &&82 if constexpr (API_LEVEL == TRANSPOSE_BATCH_MAT_MUL_HIGH_LEVEL &&
83 BATCH_SPLIT == TRANSPOSE_BATCH_MAT_MUL_BATCH_SPLIT_FALSE &&83 BATCH_SPLIT == TRANSPOSE_BATCH_MAT_MUL_BATCH_SPLIT_FALSE &&
84 PERM_X1 == TRANSPOSE_BATCH_MAT_MUL_PERM_X1_0_1_2) {84 PERM_X1 == TRANSPOSE_BATCH_MAT_MUL_PERM_X1_0_1_2) {
85- TBMM_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, TBMM_MODE::BMM_TRANS,85+ TBMM_IMPL_CLASS_COMMON_TRANS(aTran, bTran, TBMM_MODE::BMM_TRANS,
86 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswKernel,86 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswKernel,
87 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswBlock, MM_CFG_NO_PRELOAD);87 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswBlock, MM_CFG_NO_PRELOAD);
88 } else if constexpr (API_LEVEL == TRANSPOSE_BATCH_MAT_MUL_HIGH_LEVEL &&88 } else if constexpr (API_LEVEL == TRANSPOSE_BATCH_MAT_MUL_HIGH_LEVEL &&
89 BATCH_SPLIT == TRANSPOSE_BATCH_MAT_MUL_BATCH_SPLIT_FALSE &&89 BATCH_SPLIT == TRANSPOSE_BATCH_MAT_MUL_BATCH_SPLIT_FALSE &&
90 PERM_X1 == TRANSPOSE_BATCH_MAT_MUL_PERM_X1_1_0_2) {90 PERM_X1 == TRANSPOSE_BATCH_MAT_MUL_PERM_X1_1_0_2) {
91- TBMM_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, TBMM_MODE::TRANS_BMM_TRANS,91+ TBMM_IMPL_CLASS_COMMON_TRANS(aTran, bTran, TBMM_MODE::TRANS_BMM_TRANS,
92 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswKernel,92 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswKernel,
93 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswBlock, MM_CFG_NO_PRELOAD);93 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswBlock, MM_CFG_NO_PRELOAD);
94 } else if constexpr (API_LEVEL == TRANSPOSE_BATCH_MAT_MUL_HIGH_LEVEL &&94 } else if constexpr (API_LEVEL == TRANSPOSE_BATCH_MAT_MUL_HIGH_LEVEL &&
95 BATCH_SPLIT == TRANSPOSE_BATCH_MAT_MUL_BATCH_SPLIT_TRUE &&95 BATCH_SPLIT == TRANSPOSE_BATCH_MAT_MUL_BATCH_SPLIT_TRUE &&
96 PERM_X1 == TRANSPOSE_BATCH_MAT_MUL_PERM_X1_0_1_2) {96 PERM_X1 == TRANSPOSE_BATCH_MAT_MUL_PERM_X1_0_1_2) {
97- TBMM_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, TBMM_MODE::BMM_TRANS_TRANS,97+ TBMM_IMPL_CLASS_COMMON_TRANS(aTran, bTran, TBMM_MODE::BMM_TRANS_TRANS,
98 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswKernel,98 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswKernel,
99 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswBlock, MM_CFG_NO_PRELOAD);99 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswBlock, MM_CFG_NO_PRELOAD);
100 } else if constexpr (API_LEVEL == TRANSPOSE_BATCH_MAT_MUL_HIGH_LEVEL &&100 } else if constexpr (API_LEVEL == TRANSPOSE_BATCH_MAT_MUL_HIGH_LEVEL &&
101 BATCH_SPLIT == TRANSPOSE_BATCH_MAT_MUL_BATCH_SPLIT_TRUE &&101 BATCH_SPLIT == TRANSPOSE_BATCH_MAT_MUL_BATCH_SPLIT_TRUE &&
102 PERM_X1 == TRANSPOSE_BATCH_MAT_MUL_PERM_X1_1_0_2) {102 PERM_X1 == TRANSPOSE_BATCH_MAT_MUL_PERM_X1_1_0_2) {
103- TBMM_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, TBMM_MODE::TRANS_BMM_TRANS_TRANS,103+ TBMM_IMPL_CLASS_COMMON_TRANS(aTran, bTran, TBMM_MODE::TRANS_BMM_TRANS_TRANS,
104 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswKernel,104 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswKernel,
105 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswBlock, MM_CFG_NO_PRELOAD);105 TransposeBatchMatMulAdvanced::TransposeBatchMatMulAswBlock, MM_CFG_NO_PRELOAD);
106#if !__FIXED_POINT_ONLY_CUBE_TO_L0C__106#if !__FIXED_POINT_ONLY_CUBE_TO_L0C__
Mmatmul/transpose_batch_mat_mul/op_kernel/arch35/transpose_batch_mat_mul_asw_kernel.h+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/* !11/* !
12- * \file mat_mul_bl1_full_load.h12+ * \file transpose_batch_mat_mul_asw_kernel.h
13 * \brief13 * \brief
14 */14 */
15#ifndef TRANSPOSE_BATCH_MAT_MUL_ASW_KERNEL_H15#ifndef TRANSPOSE_BATCH_MAT_MUL_ASW_KERNEL_H
Mmatmul/transpose_batch_mat_mul/op_kernel/transpose_batch_mat_mul.cpp+2-2
@@ -53,7 +53,7 @@ constexpr CubeFormat format_y = CubeFormat::ND;
53 do { \53 do { \
54 PpMatMulNS::SetPadding<uint64_t>((uint64_t)0); \54 PpMatMulNS::SetPadding<uint64_t>((uint64_t)0); \
55 PpMatMulNS::SetNdpara(1, 0, 0); \55 PpMatMulNS::SetNdpara(1, 0, 0); \
56- PpMatMulNS::SetAtomicnone(); \56+ PpMatMulNS::SetAtomicNone(); \
57 if (tilingData.swizzleDirect == 0) { \57 if (tilingData.swizzleDirect == 0) { \
58 templateClass<0, transA, transB, DTYPE_X1, DTYPE_Y, PpMatMulNS::DataFormat::ND> op; \58 templateClass<0, transA, transB, DTYPE_X1, DTYPE_Y, PpMatMulNS::DataFormat::ND> op; \
59 op.Init(aGM, bGM, cGM, &tilingData); \59 op.Init(aGM, bGM, cGM, &tilingData); \
@@ -130,4 +130,4 @@ __global__ __aicore__ void transpose_batch_mat_mul(GM_ADDR aGM, GM_ADDR bGM, GM_
130 PPMATMUL_EINSUM_CLASS(PpMatMulNS::PpMatmulEinSum, false, true);130 PPMATMUL_EINSUM_CLASS(PpMatMulNS::PpMatmulEinSum, false, true);
131 }131 }
132#endif132#endif
133-}133+}
Mmatmul/transpose_batch_mat_mul/op_kernel/transpose_batch_mat_mul.h+3-3
@@ -12,8 +12,8 @@
12 * \file transpose_batch_mat_mul.h12 * \file transpose_batch_mat_mul.h
13 * \brief13 * \brief
14 */14 */
15-#ifndef TRNASPOSE_BATCH_MAT_MUL_H15+#ifndef TRANSPOSE_BATCH_MAT_MUL_H
16-#define TRNASPOSE_BATCH_MAT_MUL_H16+#define TRANSPOSE_BATCH_MAT_MUL_H
17 17 
18#include "kernel_operator.h"18#include "kernel_operator.h"
19#include "lib/matmul_intf.h"19#include "lib/matmul_intf.h"
@@ -234,4 +234,4 @@ TransposeBatchMatMulKernel<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, MODE, BLOCK_TYPE,
234 mm_.SetHF32(false, 0);234 mm_.SetHF32(false, 0);
235 return;235 return;
236}236}
237-#endif // TRNASPOSE_BATCH_MAT_MUL_H237+#endif // TRANSPOSE_BATCH_MAT_MUL_H
Mmatmul/transpose_batch_mat_mul/op_kernel/utils/utils.h+2-2
@@ -30,7 +30,7 @@ __aicore__ inline void SetPadding(IN_DTYPE padValue)
30{30{
31 AscendC::SetLoadDataPaddingValue<IN_DTYPE>(padValue);31 AscendC::SetLoadDataPaddingValue<IN_DTYPE>(padValue);
32}32}
33-__aicore__ inline void SetAtomicnone() { AscendC::SetAtomicNone(); }33+__aicore__ inline void SetAtomicNone() { AscendC::SetAtomicNone(); }
34__aicore__ inline void SetMasknorm()34__aicore__ inline void SetMasknorm()
35{35{
36#if __CCE_AICORE__ == 10036#if __CCE_AICORE__ == 100
@@ -61,4 +61,4 @@ __aicore__ inline void SetFpc(const AscendC::LocalTensor<IN_DTYPE>& preTensor, b
61}61}
62 62 
63} // namespace PpMatMulNS63} // namespace PpMatMulNS
64-#endif64+#endif
Mmatmul/transpose_quant_batch_mat_mul/op_api/aclnn_transpose_quant_batch_mat_mul.cpp+12-12
@@ -284,7 +284,7 @@ inline static bool CheckScalex1Valid(const aclTensor* x1Scale, int64_t batch, in
284 if (dimTensorScale != EXPECTED_MX_SCALE_DIM) {284 if (dimTensorScale != EXPECTED_MX_SCALE_DIM) {
285 OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON(285 OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON(
286 OP_NAME, "x1Scale", Ops::NN::FormatString("%zu", dimTensorScale).c_str(),286 OP_NAME, "x1Scale", Ops::NN::FormatString("%zu", dimTensorScale).c_str(),
287- Ops::NN::FormatString("In %s scene, the shape dim of %s must be %d", "MXFp8", "x1Scale",287+ Ops::NN::FormatString("In %s scene, the shape dim of %s must be %d", "MXFP8", "x1Scale",
288 static_cast<int>(EXPECTED_MX_SCALE_DIM))288 static_cast<int>(EXPECTED_MX_SCALE_DIM))
289 .c_str());289 .c_str());
290 return false;290 return false;
@@ -294,7 +294,7 @@ inline static bool CheckScalex1Valid(const aclTensor* x1Scale, int64_t batch, in
294 x1Scale->GetViewShape().GetDim(NUM_THREE) != NUM_TWO) {294 x1Scale->GetViewShape().GetDim(NUM_THREE) != NUM_TWO) {
295 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(295 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
296 OP_NAME, "x1Scale", op::ToString(x1Scale->GetViewShape()).GetString(),296 OP_NAME, "x1Scale", op::ToString(x1Scale->GetViewShape()).GetString(),
297- Ops::NN::FormatString("In %s scene, the shape of %s must be %s", "MXFp8", "x1Scale",297+ Ops::NN::FormatString("In %s scene, the shape of %s must be %s", "MXFP8", "x1Scale",
298 Ops::NN::FormatString("[%ld, %ld, %ld, 2]", m, batch, numGroup).c_str())298 Ops::NN::FormatString("[%ld, %ld, %ld, 2]", m, batch, numGroup).c_str())
299 .c_str());299 .c_str());
300 return false;300 return false;
@@ -329,7 +329,7 @@ inline static bool CheckScalex2Valid(const aclTensor* x2Scale, int64_t batch, in
329 if (dimTensorScale != EXPECTED_MX_SCALE_DIM) {329 if (dimTensorScale != EXPECTED_MX_SCALE_DIM) {
330 OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON(330 OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON(
331 OP_NAME, "x2Scale", Ops::NN::FormatString("%zu", dimTensorScale).c_str(),331 OP_NAME, "x2Scale", Ops::NN::FormatString("%zu", dimTensorScale).c_str(),
332- Ops::NN::FormatString("In %s scene, the shape dim of %s must be %d", "MXFp8", "x2Scale",332+ Ops::NN::FormatString("In %s scene, the shape dim of %s must be %d", "MXFP8", "x2Scale",
333 static_cast<int>(EXPECTED_MX_SCALE_DIM))333 static_cast<int>(EXPECTED_MX_SCALE_DIM))
334 .c_str());334 .c_str());
335 return false;335 return false;
@@ -340,7 +340,7 @@ inline static bool CheckScalex2Valid(const aclTensor* x2Scale, int64_t batch, in
340 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(340 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
341 OP_NAME, "x2Scale", op::ToString(x2Scale->GetViewShape()).GetString(),341 OP_NAME, "x2Scale", op::ToString(x2Scale->GetViewShape()).GetString(),
342 Ops::NN::FormatString(342 Ops::NN::FormatString(
343- "In %s scene, the shape of %s must be %s", "MXFp8", "x2Scale",343+ "In %s scene, the shape of %s must be %s", "MXFP8", "x2Scale",
344 Ops::NN::FormatString("[%ld, %ld, %ld, 2]", batch, dims[(*permX2)[1]], dims[(*permX2)[NUM_TWO]])344 Ops::NN::FormatString("[%ld, %ld, %ld, 2]", batch, dims[(*permX2)[1]], dims[(*permX2)[NUM_TWO]])
345 .c_str())345 .c_str())
346 .c_str());346 .c_str());
@@ -611,22 +611,22 @@ static const aclTensor* BuildTransposeQuantBatchMatMulGraph(const aclTensor* x1,
611 // 连续性转换611 // 连续性转换
612 auto contiguousX1 = l0op::Contiguous(x1, executor);612 auto contiguousX1 = l0op::Contiguous(x1, executor);
613 OP_CHECK(contiguousX1 != nullptr,613 OP_CHECK(contiguousX1 != nullptr,
614- OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x1 perprocess failed, contiguous return nullptr."),614+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x1 preprocess failed, contiguous return nullptr."),
615 return nullptr);615 return nullptr);
616 auto reformX1 = l0op::ReFormat(contiguousX1, op::Format::FORMAT_ND);616 auto reformX1 = l0op::ReFormat(contiguousX1, op::Format::FORMAT_ND);
617 OP_CHECK(reformX1 != nullptr,617 OP_CHECK(reformX1 != nullptr,
618- OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x1 perprocess failed, reformat return nullptr."),618+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x1 preprocess failed, reformat return nullptr."),
619 return nullptr);619 return nullptr);
620 620 
621 auto contiguousX2 = l0op::Contiguous(x2, executor);621 auto contiguousX2 = l0op::Contiguous(x2, executor);
622 OP_CHECK(contiguousX2 != nullptr,622 OP_CHECK(contiguousX2 != nullptr,
623- OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x2 perprocess failed, contiguous return nullptr."),623+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x2 preprocess failed, contiguous return nullptr."),
624 return nullptr);624 return nullptr);
625 auto reformX2 = contiguousX2;625 auto reformX2 = contiguousX2;
626 if (ge::GetPrimaryFormat(x2->GetStorageFormat()) != op::Format::FORMAT_FRACTAL_NZ) {626 if (ge::GetPrimaryFormat(x2->GetStorageFormat()) != op::Format::FORMAT_FRACTAL_NZ) {
627 reformX2 = l0op::ReFormat(contiguousX2, op::Format::FORMAT_ND);627 reformX2 = l0op::ReFormat(contiguousX2, op::Format::FORMAT_ND);
628 OP_CHECK(reformX2 != nullptr,628 OP_CHECK(reformX2 != nullptr,
629- OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x2 perprocess failed, reformat return nullptr."),629+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x2 preprocess failed, reformat return nullptr."),
630 return nullptr);630 return nullptr);
631 } else {631 } else {
632 reformX2->SetStorageShape(x2->GetStorageShape());632 reformX2->SetStorageShape(x2->GetStorageShape());
@@ -636,18 +636,18 @@ static const aclTensor* BuildTransposeQuantBatchMatMulGraph(const aclTensor* x1,
636 if (contiguousX1Scale != nullptr) {636 if (contiguousX1Scale != nullptr) {
637 contiguousX1Scale = l0op::Contiguous(x1Scale, executor);637 contiguousX1Scale = l0op::Contiguous(x1Scale, executor);
638 OP_CHECK(contiguousX1Scale != nullptr,638 OP_CHECK(contiguousX1Scale != nullptr,
639- OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x1Scale perprocess failed, contiguous return nullptr."),639+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x1Scale preprocess failed, contiguous return nullptr."),
640 return nullptr);640 return nullptr);
641 }641 }
642 auto contiguousX2Scale = x2Scale;642 auto contiguousX2Scale = x2Scale;
643 if (contiguousX2Scale != nullptr) {643 if (contiguousX2Scale != nullptr) {
644 contiguousX2Scale = l0op::Contiguous(x2Scale, executor);644 contiguousX2Scale = l0op::Contiguous(x2Scale, executor);
645 OP_CHECK(contiguousX2Scale != nullptr,645 OP_CHECK(contiguousX2Scale != nullptr,
646- OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x2Scale perprocess failed, contiguous return nullptr."),646+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "The input x2Scale preprocess failed, contiguous return nullptr."),
647 return nullptr);647 return nullptr);
648 }648 }
649 649 
650- // Invoke tqbmmm l0 api650+ // Invoke tqbmm l0 api
651 return l0op::TransposeQuantBatchMatMul(reformX1, reformX2, nullptr, contiguousX1Scale, contiguousX2Scale, dtype,651 return l0op::TransposeQuantBatchMatMul(reformX1, reformX2, nullptr, contiguousX1Scale, contiguousX2Scale, dtype,
652 groupSize, permX1, permX2, permY, batchSplitFactor, executor);652 groupSize, permX1, permX2, permY, batchSplitFactor, executor);
653}653}
@@ -769,4 +769,4 @@ aclnnStatus aclnnTransposeQuantBatchMatMulWeightNz(void* workspace, uint64_t wor
769{769{
770 L2_DFX_PHASE_2(aclnnTransposeQuantBatchMatMulWeightNz);770 L2_DFX_PHASE_2(aclnnTransposeQuantBatchMatMulWeightNz);
771 return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);771 return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
772-}772+}
Mmatmul/transpose_quant_batch_mat_mul/op_api/aclnn_transpose_quant_batch_mat_mul.h+5-5
@@ -17,9 +17,9 @@
17extern "C" {17extern "C" {
18#endif18#endif
19/**19/**
20- * @brief aclnnTransposeQuantBatchMatmul的第一段接口，根据具体的计算流程，计算workspace大小。20+ * @brief aclnnTransposeQuantBatchMatMul的第一段接口，根据具体的计算流程，计算workspace大小。
21 * @domain aclnn_ops_infer21 * @domain aclnn_ops_infer
22- * 算子功能：实现TransposeQuantBatchMatmul计算。22+ * 算子功能：实现TransposeQuantBatchMatMul计算。
23 * @param [in] x1: matmul左矩阵，数据类型支持：float8_e4m3fn, float8_e5m2。23 * @param [in] x1: matmul左矩阵，数据类型支持：float8_e4m3fn, float8_e5m2。
24 * @param [in] x2: matmul右矩阵，数据类型支持：float8_e4m3fn, float8_e5m2。24 * @param [in] x2: matmul右矩阵，数据类型支持：float8_e4m3fn, float8_e5m2。
25 * @param [in] bias: 偏置，当前不支持。25 * @param [in] bias: 偏置，当前不支持。
@@ -42,7 +42,7 @@ ACLNN_API aclnnStatus aclnnTransposeQuantBatchMatMulGetWorkspaceSize(
42 int32_t batchSplitFactor, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);42 int32_t batchSplitFactor, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);
43 43 
44/**44/**
45- * @brief aclnnTransposeBatchMatmul的第二段接口，用于执行计算。45+ * @brief aclnnTransposeQuantBatchMatMul的第二段接口，用于执行计算。
46 */46 */
47ACLNN_API aclnnStatus aclnnTransposeQuantBatchMatMul(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,47ACLNN_API aclnnStatus aclnnTransposeQuantBatchMatMul(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
48 const aclrtStream stream);48 const aclrtStream stream);
@@ -74,8 +74,8 @@ ACLNN_API aclnnStatus aclnnTransposeQuantBatchMatMulWeightNzGetWorkspaceSize(
74 74 
75/**75/**
76 * @brief aclnnTransposeQuantBatchMatMulWeightNz的第二段接口，用于执行计算。76 * @brief aclnnTransposeQuantBatchMatMulWeightNz的第二段接口，用于执行计算。
77- * @param [in] workspace: 在npu device侧申请的workspace内存起址。77+ * @param [in] workspace: 在npu device侧申请的workspace内存地址。
78- * @param [in] workspace_size: 在npu78+ * @param [in] workspaceSize: 在npu
79 * device侧申请的workspace大小，由第一段接口aclnnTransposeQuantBatchMatMulWeightNzGetWorkspaceSize获取。79 * device侧申请的workspace大小，由第一段接口aclnnTransposeQuantBatchMatMulWeightNzGetWorkspaceSize获取。
80 * @param [in] executor: op执行器，包含了算子计算流程。80 * @param [in] executor: op执行器，包含了算子计算流程。
81 * @param [in] stream: acl stream流。81 * @param [in] stream: acl stream流。
Mmatmul/transpose_quant_batch_mat_mul/op_host/op_tiling/arch35/transpose_quant_batch_mat_mul_tiling_advanced.cpp+3-3
@@ -321,7 +321,7 @@ ge::graphStatus TransposeQuantBatchMatMulTiling::CheckScale(const int64_t b, con
321 if (invalidScaleX1) {321 if (invalidScaleX1) {
322 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(322 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
323 args_.opName, "x1Scale", Ops::Base::ToString(scaleX1ShapePtr->GetStorageShape()).c_str(),323 args_.opName, "x1Scale", Ops::Base::ToString(scaleX1ShapePtr->GetStorageShape()).c_str(),
324- Ops::NN::FormatString("In %s case, the shape of %s must be %s", "MXFp8", "x1Scale",324+ Ops::NN::FormatString("In %s case, the shape of %s must be %s", "MXFP8", "x1Scale",
325 Ops::NN::FormatString("[%ld, %ld, %ld, 2]", m, b, numGroup).c_str())325 Ops::NN::FormatString("[%ld, %ld, %ld, 2]", m, b, numGroup).c_str())
326 .c_str());326 .c_str());
327 return ge::GRAPH_FAILED;327 return ge::GRAPH_FAILED;
@@ -335,7 +335,7 @@ ge::graphStatus TransposeQuantBatchMatMulTiling::CheckScale(const int64_t b, con
335 if (invalidScaleX2) {335 if (invalidScaleX2) {
336 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(336 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
337 args_.opName, "x2Scale", Ops::Base::ToString(scaleX2ShapePtr->GetStorageShape()).c_str(),337 args_.opName, "x2Scale", Ops::Base::ToString(scaleX2ShapePtr->GetStorageShape()).c_str(),
338- Ops::NN::FormatString("In %s case, the shape of %s must be %s", "MXFp8", "x2Scale",338+ Ops::NN::FormatString("In %s case, the shape of %s must be %s", "MXFP8", "x2Scale",
339 "[b, numGroup, n, 2] after permX2")339 "[b, numGroup, n, 2] after permX2")
340 .c_str());340 .c_str());
341 return ge::GRAPH_FAILED;341 return ge::GRAPH_FAILED;
@@ -513,4 +513,4 @@ ge::graphStatus TransposeQuantBatchMatMulTiling::GetBatchInfo(const gert::Tiling
513}513}
514 514 
515} // namespace transpose_quant_batch_mat_mul_advanced515} // namespace transpose_quant_batch_mat_mul_advanced
516-} // namespace optiling516+} // namespace optiling
Mmatmul/transpose_quant_batch_mat_mul/op_host/op_tiling/transpose_quant_batch_mat_mul_tiling.cpp+2-2
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file transpose_quant_batch_mat_mul_tiling.cc12+ * \file transpose_quant_batch_mat_mul_tiling.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
@@ -74,4 +74,4 @@ IMPL_OP_OPTILING(TransposeQuantBatchMatMul)
74 .Tiling(TransposeQuantBatchMatMulTilingFunc)74 .Tiling(TransposeQuantBatchMatMulTilingFunc)
75 .TilingParse<MatmulV3CompileInfo>(TilingPrepareForTransposeQuantBatchMatMul)75 .TilingParse<MatmulV3CompileInfo>(TilingPrepareForTransposeQuantBatchMatMul)
76 .GenSimplifiedKey(transpose_quant_batch_matmul::GenSimplifiedKey);76 .GenSimplifiedKey(transpose_quant_batch_matmul::GenSimplifiedKey);
77-} // namespace optiling77+} // namespace optiling
Mmatmul/transpose_quant_batch_mat_mul/op_host/transpose_quant_batch_mat_mul_infershape.cpp+15-16
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file transpose_quant_batch_mat_mul_infer.cpp12+ * \file transpose_quant_batch_mat_mul_infershape.cpp
13 * \brief13 * \brief
14 */14 */
15#include <string>15#include <string>
@@ -59,23 +59,22 @@ static bool CheckDtypeValid(const ge::DataType& dtypeX1, const ge::DataType& dty
59 // MXFP859 // MXFP8
60 if (IsMicroScaling(dtypeX1Scale, dtypeX2Scale)) {60 if (IsMicroScaling(dtypeX1Scale, dtypeX2Scale)) {
61 CHECK(dtypeX1 != ge::DT_FLOAT8_E4M3FN || dtypeX2 != ge::DT_FLOAT8_E4M3FN,61 CHECK(dtypeX1 != ge::DT_FLOAT8_E4M3FN || dtypeX2 != ge::DT_FLOAT8_E4M3FN,
62- CUBE_INNER_ERR_REPORT("TQBMM", "the dtype of input is only supported FLOAT8_E4M3FN."), return false);62+ CUBE_INNER_ERR_REPORT("TQBMM", "the dtype of input only supports FLOAT8_E4M3FN."), return false);
63 // FP863 // FP8
64 } else if (!IsHIFP8(dtypeX1, dtypeX2)) {64 } else if (!IsHIFP8(dtypeX1, dtypeX2)) {
65 CHECK((dtypeX1 != ge::DT_FLOAT8_E4M3FN && dtypeX1 != ge::DT_FLOAT8_E5M2) ||65 CHECK((dtypeX1 != ge::DT_FLOAT8_E4M3FN && dtypeX1 != ge::DT_FLOAT8_E5M2) ||
66 (dtypeX2 != ge::DT_FLOAT8_E4M3FN && dtypeX2 != ge::DT_FLOAT8_E5M2),66 (dtypeX2 != ge::DT_FLOAT8_E4M3FN && dtypeX2 != ge::DT_FLOAT8_E5M2),
67- CUBE_INNER_ERR_REPORT("TQBMM", "the dtype of input is only supported FLOAT8_E4M3FN or FLOAT8_E5M2."),67+ CUBE_INNER_ERR_REPORT("TQBMM", "the dtype of input only supports FLOAT8_E4M3FN or FLOAT8_E5M2."),
68 return false);68 return false);
69 }69 }
70 // MXFP8/FP8 scale dtype check70 // MXFP8/FP8 scale dtype check
71 if (!IsHIFP8(dtypeX1, dtypeX2)) {71 if (!IsHIFP8(dtypeX1, dtypeX2)) {
72 CHECK((dtypeX1Scale != ge::DT_FLOAT && dtypeX1Scale != ge::DT_FLOAT8_E8M0) ||72 CHECK((dtypeX1Scale != ge::DT_FLOAT && dtypeX1Scale != ge::DT_FLOAT8_E8M0) ||
73 (dtypeX2Scale != ge::DT_FLOAT && dtypeX2Scale != ge::DT_FLOAT8_E8M0),73 (dtypeX2Scale != ge::DT_FLOAT && dtypeX2Scale != ge::DT_FLOAT8_E8M0),
74- CUBE_INNER_ERR_REPORT("TQBMM", "the dtype of scale is only supported FLOAT or FLOAT8_E8M0."),74+ CUBE_INNER_ERR_REPORT("TQBMM", "the dtype of scale only supports FLOAT or FLOAT8_E8M0."), return false);
75- return false);
76 } else {75 } else {
77- CHECK(dtypeX2Scale != ge::DT_UINT64,76+ CHECK(dtypeX2Scale != ge::DT_UINT64, CUBE_INNER_ERR_REPORT("TQBMM", "the dtype of scale only supports UINT64."),
78- CUBE_INNER_ERR_REPORT("TQBMM", "the dtype of scale is only supported UINT64."), return false);77+ return false);
79 }78 }
80 return true;79 return true;
81}80}
@@ -84,7 +83,7 @@ static bool CheckPerm(const TypedContinuousVector<int64_t>* permX1, const TypedC
84 const TypedContinuousVector<int64_t>* permY)83 const TypedContinuousVector<int64_t>* permY)
85{84{
86 CHECK(permX1 == nullptr || permX2 == nullptr || permY == nullptr,85 CHECK(permX1 == nullptr || permX2 == nullptr || permY == nullptr,
87- CUBE_INNER_ERR_REPORT("TQBMM", "[Infershape] attr is nullptr."), return false);86+ CUBE_INNER_ERR_REPORT("TQBMM", "[InferShape] attr is null."), return false);
88 87 
89 CHECK(permX1->GetSize() != PERM_DIM_NUM || permX2->GetSize() != PERM_DIM_NUM || permY->GetSize() != PERM_DIM_NUM,88 CHECK(permX1->GetSize() != PERM_DIM_NUM || permX2->GetSize() != PERM_DIM_NUM || permY->GetSize() != PERM_DIM_NUM,
90 CUBE_INNER_ERR_REPORT("TQBMM", "[InferShape] The dims of the perm intArray should be 3"), return false);89 CUBE_INNER_ERR_REPORT("TQBMM", "[InferShape] The dims of the perm intArray should be 3"), return false);
@@ -100,7 +99,7 @@ static bool CheckPerm(const TypedContinuousVector<int64_t>* permX1, const TypedC
100 99 
101 const auto permYAttr = permY->GetData();100 const auto permYAttr = permY->GetData();
102 auto checkPermY = *permYAttr == 1 && *(permYAttr + 1) == 0 && *(permYAttr + 2) == 2;101 auto checkPermY = *permYAttr == 1 && *(permYAttr + 1) == 0 && *(permYAttr + 2) == 2;
103- CHECK(!checkPermY, CUBE_INNER_ERR_REPORT("TQBMM", "[InferShape] perm_y should {1, 0, 2}"), return false);102+ CHECK(!checkPermY, CUBE_INNER_ERR_REPORT("TQBMM", "[InferShape] perm_y should be {1, 0, 2}"), return false);
104 return true;103 return true;
105}104}
106 105 
@@ -162,10 +161,10 @@ static ge::graphStatus InferShapeForTransposeQuantBatchMatMul(InferShapeContext*
162 auto attrs = context->GetAttrs();161 auto attrs = context->GetAttrs();
163 auto nameOp = context->GetNodeName();162 auto nameOp = context->GetNodeName();
164 CHECK(shapeX1 == nullptr || shapeX2 == nullptr || shapeY == nullptr || attrs == nullptr,163 CHECK(shapeX1 == nullptr || shapeX2 == nullptr || shapeY == nullptr || attrs == nullptr,
165- CUBE_INNER_ERR_REPORT(nameOp, "[Infershape]shape or attrs is null."), return ge::GRAPH_FAILED);164+ CUBE_INNER_ERR_REPORT(nameOp, "[InferShape] shape or attrs is null."), return ge::GRAPH_FAILED);
166 165 
167 const auto dtype = attrs->GetAttrPointer<int64_t>(0); // dtype index is 0166 const auto dtype = attrs->GetAttrPointer<int64_t>(0); // dtype index is 0
168- CHECK(dtype == nullptr, CUBE_INNER_ERR_REPORT(nameOp, "[Infershape] attr dtype is null."), return ge::GRAPH_FAILED);167+ CHECK(dtype == nullptr, CUBE_INNER_ERR_REPORT(nameOp, "[InferShape] attr dtype is null."), return ge::GRAPH_FAILED);
169 auto tensorX1 = context->GetInputDesc(0);168 auto tensorX1 = context->GetInputDesc(0);
170 auto tensorX2 = context->GetInputDesc(1);169 auto tensorX2 = context->GetInputDesc(1);
171 CHECK(tensorX1 == nullptr || tensorX2 == nullptr, CUBE_INNER_ERR_REPORT(nameOp, "x1 or x2 is null."),170 CHECK(tensorX1 == nullptr || tensorX2 == nullptr, CUBE_INNER_ERR_REPORT(nameOp, "x1 or x2 is null."),
@@ -177,7 +176,7 @@ static ge::graphStatus InferShapeForTransposeQuantBatchMatMul(InferShapeContext*
177 auto tensorX1Scale = context->GetOptionalInputDesc(kX1ScaleIdx);176 auto tensorX1Scale = context->GetOptionalInputDesc(kX1ScaleIdx);
178 auto tensorX2Scale = context->GetOptionalInputDesc(kX2ScaleIdx);177 auto tensorX2Scale = context->GetOptionalInputDesc(kX2ScaleIdx);
179 CHECK((!isHIFP8 && tensorX1Scale == nullptr) || tensorX2Scale == nullptr,178 CHECK((!isHIFP8 && tensorX1Scale == nullptr) || tensorX2Scale == nullptr,
180- CUBE_INNER_ERR_REPORT(nameOp, "X1Scale or x2Scale is null."), return ge::GRAPH_FAILED);179+ CUBE_INNER_ERR_REPORT(nameOp, "x1Scale or x2Scale is null."), return ge::GRAPH_FAILED);
181 180 
182 ge::DataType dtypeX1Scale = isHIFP8 ? ge::DT_UINT64 : tensorX1Scale->GetDataType();181 ge::DataType dtypeX1Scale = isHIFP8 ? ge::DT_UINT64 : tensorX1Scale->GetDataType();
183 ge::DataType dtypeX2Scale = tensorX2Scale->GetDataType();182 ge::DataType dtypeX2Scale = tensorX2Scale->GetDataType();
@@ -214,7 +213,7 @@ static ge::graphStatus InferShapeForTransposeQuantBatchMatMul(InferShapeContext*
214 213 
215 // batchSplitFactor only support 1214 // batchSplitFactor only support 1
216 CHECK(batchSplitFactor != nullptr && *batchSplitFactor != VALID_BATCH_SPLIT_FACTOR,215 CHECK(batchSplitFactor != nullptr && *batchSplitFactor != VALID_BATCH_SPLIT_FACTOR,
217- CUBE_INNER_ERR_REPORT(nameOp, "batchSplitFactor should be 1 ."), return ge::GRAPH_FAILED);216+ CUBE_INNER_ERR_REPORT(nameOp, "batchSplitFactor should be 1."), return ge::GRAPH_FAILED);
218 217 
219 // Set shapeY218 // Set shapeY
220 ge::graphStatus ret = SetShapeY(*shapeY, shapeX1Transposed, shapeX2Transposed, *permY, *batchSplitFactor);219 ge::graphStatus ret = SetShapeY(*shapeY, shapeX1Transposed, shapeX2Transposed, *permY, *batchSplitFactor);
@@ -236,14 +235,14 @@ static ge::graphStatus TransposeQuantBatchMatMulInferDataType(gert::InferDataTyp
236 OP_LOGD(context, "TransposeQuantBatchMatMulInferDataType begin");235 OP_LOGD(context, "TransposeQuantBatchMatMulInferDataType begin");
237 auto nameOp = context->GetNodeName();236 auto nameOp = context->GetNodeName();
238 auto* attrs = context->GetAttrs();237 auto* attrs = context->GetAttrs();
239- CHECK(attrs == nullptr, CUBE_INNER_ERR_REPORT(nameOp, "[Infershape] attr is nullptr."), return ge::GRAPH_FAILED);238+ CHECK(attrs == nullptr, CUBE_INNER_ERR_REPORT(nameOp, "[InferShape] attr is nullptr."), return ge::GRAPH_FAILED);
240 const int32_t* dtype = attrs->GetAttrPointer<int32_t>(ATTR_INDEX_DST_TYPE);239 const int32_t* dtype = attrs->GetAttrPointer<int32_t>(ATTR_INDEX_DST_TYPE);
241- CHECK(dtype == nullptr, CUBE_INNER_ERR_REPORT(nameOp, "[Infershape] dtype is nullptr."), return ge::GRAPH_FAILED);240+ CHECK(dtype == nullptr, CUBE_INNER_ERR_REPORT(nameOp, "[InferShape] dtype is nullptr."), return ge::GRAPH_FAILED);
242 int32_t dstDtype = *dtype;241 int32_t dstDtype = *dtype;
243 ge::DataType yDtype = static_cast<ge::DataType>(dstDtype);242 ge::DataType yDtype = static_cast<ge::DataType>(dstDtype);
244 243 
245 OP_CHECK_IF(std::find(OUT_TYPE_LIST.begin(), OUT_TYPE_LIST.end(), yDtype) == OUT_TYPE_LIST.end(),244 OP_CHECK_IF(std::find(OUT_TYPE_LIST.begin(), OUT_TYPE_LIST.end(), yDtype) == OUT_TYPE_LIST.end(),
246- OP_LOGE(context, "attr dtype only support float16, bfloat16"), return ge::GRAPH_FAILED);245+ OP_LOGE(nameOp, "attr dtype only supports float16, bfloat16"), return ge::GRAPH_FAILED);
247 246 
248 context->SetOutputDataType(Y_INDEX, yDtype);247 context->SetOutputDataType(Y_INDEX, yDtype);
249 248 
Mmatmul/transpose_quant_batch_mat_mul/op_kernel/arch35/transpose_quant_batch_mat_mul.cpp+5-5
@@ -48,7 +48,7 @@ constexpr CubeFormat format_x2 = CubeFormat::ND;
48constexpr MatmulConfig MM_CFG_NO_PRELOAD_OPEN_UNIT_FLAG = GetMDLConfig(false, false, 0, false, false, false, true, true,48constexpr MatmulConfig MM_CFG_NO_PRELOAD_OPEN_UNIT_FLAG = GetMDLConfig(false, false, 0, false, false, false, true, true,
49 false, false, false);49 false, false, false);
50 50 
51-#define TQBMM_IMPL_CLASS_COMMON_TRNAS(transposeX1, transposeX2, precisionMode, templateClass, ...) \51+#define TQBMM_IMPL_CLASS_COMMON_TRANS(transposeX1, transposeX2, precisionMode, templateClass, ...) \
52 do { \52 do { \
53 templateClass<DTYPE_X1, DTYPE_X2, DTYPE_X2_SCALE, DTYPE_BIAS, DTYPE_X1_SCALE, DTYPE_Y, precisionMode, \53 templateClass<DTYPE_X1, DTYPE_X2, DTYPE_X2_SCALE, DTYPE_BIAS, DTYPE_X1_SCALE, DTYPE_Y, precisionMode, \
54 transposeX1, transposeX2, format_x2, DTYPE_LOC_LOCAL, __VA_ARGS__> \54 transposeX1, transposeX2, format_x2, DTYPE_LOC_LOCAL, __VA_ARGS__> \
@@ -69,19 +69,19 @@ __global__ __aicore__ void transpose_quant_batch_mat_mul(GM_ADDR aGM, GM_ADDR bG
69 REGISTER_TILING_DEFAULT(BatchMatMulV3TilingData);69 REGISTER_TILING_DEFAULT(BatchMatMulV3TilingData);
70 GET_TILING_DATA(tilingData, tilingGM);70 GET_TILING_DATA(tilingData, tilingGM);
71 if constexpr (sizeof(DTYPE_X2_SCALE) == sizeof(uint64_t)) {71 if constexpr (sizeof(DTYPE_X2_SCALE) == sizeof(uint64_t)) {
72- TQBMM_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, static_cast<int8_t>(TQBMMPrecisionMode::PRECISION_MODE_HIFP8),72+ TQBMM_IMPL_CLASS_COMMON_TRANS(aTran, bTran, static_cast<int8_t>(TQBMMPrecisionMode::PRECISION_MODE_HIFP8),
73 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswKernel,73 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswKernel,
74 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswBlock,74 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswBlock,
75 MM_CFG_NO_PRELOAD_OPEN_UNIT_FLAG);75 MM_CFG_NO_PRELOAD_OPEN_UNIT_FLAG);
76 } else if constexpr (sizeof(DTYPE_X2_SCALE) == sizeof(float)) {76 } else if constexpr (sizeof(DTYPE_X2_SCALE) == sizeof(float)) {
77- TQBMM_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, static_cast<int8_t>(TQBMMPrecisionMode::PRECISION_MODE_FP8),77+ TQBMM_IMPL_CLASS_COMMON_TRANS(aTran, bTran, static_cast<int8_t>(TQBMMPrecisionMode::PRECISION_MODE_FP8),
78 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswKernel,78 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswKernel,
79 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswBlock,79 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswBlock,
80 MM_CFG_NO_PRELOAD_OPEN_UNIT_FLAG);80 MM_CFG_NO_PRELOAD_OPEN_UNIT_FLAG);
81 } else if constexpr (sizeof(DTYPE_X2_SCALE) == sizeof(uint8_t)) {81 } else if constexpr (sizeof(DTYPE_X2_SCALE) == sizeof(uint8_t)) {
82- TQBMM_IMPL_CLASS_COMMON_TRNAS(aTran, bTran, static_cast<int8_t>(TQBMMPrecisionMode::PRECISION_MODE_MXFP8),82+ TQBMM_IMPL_CLASS_COMMON_TRANS(aTran, bTran, static_cast<int8_t>(TQBMMPrecisionMode::PRECISION_MODE_MXFP8),
83 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswKernel,83 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswKernel,
84 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswBlock,84 TransposeQuantBatchMatMulAdvanced::TransposeQuantBatchMatMulAswBlock,
85 MM_CFG_NO_PRELOAD_OPEN_UNIT_FLAG);85 MM_CFG_NO_PRELOAD_OPEN_UNIT_FLAG);
86 }86 }
87-}87+}
Mmatmul/transpose_quant_batch_mat_mul/tests/ut/op_kernel/transpose_quant_batch_mat_mul_tiling_def.h+2-2
@@ -24,11 +24,11 @@ inline void InitTqbmmTilingData(void* tiling, void* const_data)
24 InitTqbmmTilingData(tiling_arg, &tiling_data);24 InitTqbmmTilingData(tiling_arg, &tiling_data);
25#endif25#endif
26 26 
27-#define TQBMM_IMPL_CLASS_COMMON_TRNAS(transposeX1, transposeX2, templateClass, ...) \27+#define TQBMM_IMPL_CLASS_COMMON_TRANS(transposeX1, transposeX2, templateClass, ...) \
28 do { \28 do { \
29 templateClass<DTYPE_X1, DTYPE_X2, DTYPE_X2_SCALE, DTYPE_BIAS, DTYPE_X1_SCALE, DTYPE_Y, transposeX1, \29 templateClass<DTYPE_X1, DTYPE_X2, DTYPE_X2_SCALE, DTYPE_BIAS, DTYPE_X1_SCALE, DTYPE_Y, transposeX1, \
30 transposeX2, DTYPE_LOC_LOCAL, __VA_ARGS__> \30 transposeX2, DTYPE_LOC_LOCAL, __VA_ARGS__> \
31 op; \31 op; \
32 op.Init(aGM, bGM, x2_scaleGM, x1_scaleGM, cGM, user, &tilingData, &pipe); \32 op.Init(aGM, bGM, x2_scaleGM, x1_scaleGM, cGM, user, &tilingData, &pipe); \
33 op.Process(); \33 op.Process(); \
34- } while (0)34+ } while (0)