已合并
【qbmm】公共校验逻辑拆分与安全性增强 #6889
zhaoyingchao2创建于 7月2日
【qbmm】公共校验逻辑拆分与安全性增强 #6889
已合并
共 5 个文件变更+261-213
| @@ -30,136 +30,17 @@ | |||
| 30 | 30 | ||
| 31 | 31 | ||
| 32 | 32 | ||
| 33 | + | ||
| 33 | 34 | ||
| 34 | using namespace op; | 35 | using namespace op; |
| 36 | +using namespace quant_matmul_v4; | ||
| 35 | using Ops::NN::SwapLastTwoDimValue; | 37 | using Ops::NN::SwapLastTwoDimValue; |
| 36 | using Ops::NN::BoolToString; | 38 | using Ops::NN::BoolToString; |
| 37 | using Ops::NN::FormatString; | 39 | using Ops::NN::FormatString; |
| 38 | using Ops::NN::IsTransposeLastTwoDims; | 40 | using Ops::NN::IsTransposeLastTwoDims; |
| 39 | using Ops::Base::CeilDiv; | 41 | using Ops::Base::CeilDiv; |
| 40 | -using TupleTensor = std::tuple<const aclTensor *, const aclTensor *, const aclTensor *>; | ||
| 41 | -using TupleOptional = std::tuple<const aclTensor *, const aclTensor *, const aclTensor *, const aclTensor *, | ||
| 42 | - const aclTensor *, const int64_t &>; | ||
| 43 | -using TupleInput = std::tuple<const aclTensor *, const aclTensor *>; | ||
| 44 | -using TupleQuant = std::tuple<const aclTensor *, const aclTensor *, const aclTensor *, const aclTensor *, | ||
| 45 | - const aclTensor *, const aclTensor *, const aclTensor *, const int64_t &, const int64_t &>; | ||
| 46 | -using TupleAttr = std::tuple<bool, bool>; | ||
| 47 | 42 | ||
| 48 | namespace { | 43 | namespace { |
| 49 | -static constexpr int MX_SCALE_LAST_DIM_INDEX = 2; | ||
| 50 | -static constexpr int MX_SCALE_LAST_DIM = 2; | ||
| 51 | - | ||
| 52 | -static constexpr int INDEX_X1_IN_MANDTORY_TUPLE = 0; | ||
| 53 | -static constexpr int INDEX_X2_IN_MANDTORY_TUPLE = 1; | ||
| 54 | -static constexpr int INDEX_SCALE_IN_MANDTORY_TUPLE = 2; | ||
| 55 | -static constexpr int INDEX_OFFSET_IN_OPTIONAL_TUPLE = 0; | ||
| 56 | -static constexpr int INDEX_PERTOKEN_IN_OPTIONAL_TUPLE = 1; | ||
| 57 | -static constexpr int INDEX_BIAS_IN_OPTIONAL_TUPLE = 2; | ||
| 58 | -static constexpr int INDEX_Y_SCALE_IN_OPTIONAL_TUPLE = 3; | ||
| 59 | -static constexpr int INDEX_Y_OFFSET_IN_OPTIONAL_TUPLE = 4; | ||
| 60 | -static constexpr int INDEX_GROUP_SIZE_IN_OPTIONAL_TUPLE = 5; | ||
| 61 | -static constexpr int INDEX_OUT_IN_TUPLE = 2; | ||
| 62 | -static constexpr int INDEX_ISA4W4_IN_BOOL_TUPLE = 2; | ||
| 63 | -static constexpr size_t LAST_SECOND_DIM_INDEX = 2; | ||
| 64 | - | ||
| 65 | -static constexpr int MIN_DIM_NUM_ND = 2; | ||
| 66 | -static constexpr int MAX_DIM_NUM_ND = 6; | ||
| 67 | -static constexpr int MIN_DIM_NUM_NZ = 4; | ||
| 68 | -static constexpr int MAX_DIM_NUM_NZ = 8; | ||
| 69 | -static constexpr int PENULTIMATE_DIM = 2; | ||
| 70 | -static constexpr int NZ_K1_INDEX = 3; | ||
| 71 | -static constexpr int NZ_K1_INDEX_TRANS = 4; | ||
| 72 | -static constexpr int NZ_STORAGE_PENULTIMATE_DIM = 16; | ||
| 73 | -static constexpr int NZ_STORAGE_LAST_DIM = 32; | ||
| 74 | -static constexpr int64_t NZ_K0_VALUE_BMM_BLOCK_NUM = 16; | ||
| 75 | -static constexpr int64_t NZ_K0_VALUE_INT32_TRANS = 8; | ||
| 76 | -static constexpr int64_t NZ_K0_VALUE_INT8_TRANS = 32; | ||
| 77 | -static constexpr int64_t NZ_K0_VALUE_INT4_TRANS = 64; | ||
| 78 | -static constexpr int64_t OUTPUT_INFER_FAIL = -1L; | ||
| 79 | -static constexpr int64_t LAST_AXIS_LIMIT = 65535; | ||
| 80 | -static constexpr int X2_FIXED_DIM_NUM_A4W4 = 2; | ||
| 81 | -static constexpr int64_t INT4_NUMS_IN_INT8 = 2; | ||
| 82 | -static constexpr int64_t INT4_NUMS_IN_INT32 = 8; | ||
| 83 | -static constexpr int64_t INNER_SIZE_MULTIPLE = 64; | ||
| 84 | -static constexpr int64_t K_VALUE = 3696; | ||
| 85 | -static constexpr int64_t N_VALUE = 8192; | ||
| 86 | -static constexpr int64_t M_RANGE1_LEFT = 128; | ||
| 87 | -static constexpr int64_t M_RANGE1_RIGHT = 512; | ||
| 88 | -static constexpr int32_t CORE_NUM_20 = 20; | ||
| 89 | -static constexpr int64_t SUPPORTED_GROUP_SIZE = 32; | ||
| 90 | -static constexpr uint64_t B4_PER_B32 = 8UL; | ||
| 91 | -static constexpr int64_t SUPPORTED_TCG_A8W4_K_ALIGN_NUM = 32; | ||
| 92 | -static constexpr int64_t SUPPORTED_MX_A8W4_K_ALIGN_NUM = 8; | ||
| 93 | -static constexpr int64_t SUPPORTED_N_ALIGN_NUM = 8; | ||
| 94 | -static constexpr size_t MAX_DIM_VALUE = 2; | ||
| 95 | -static constexpr size_t MX_SCALE_DIM_VALUE = 3; | ||
| 96 | -static constexpr uint64_t GROUP_M_OFFSET = 32; | ||
| 97 | -static constexpr uint64_t GROUP_N_OFFSET = 16; | ||
| 98 | -static constexpr uint64_t GROUP_MNK_BIT_SIZE = 0xFFFF; | ||
| 99 | -static constexpr size_t MX_SCALE_MAX_DIM = 3; | ||
| 100 | -static constexpr size_t MX_SCALE_DIM_NUM = 3; | ||
| 101 | -static constexpr int64_t MAX_SHAPE_SIZE_A8W4_INT = 29576; | ||
| 102 | -static constexpr int64_t PPMATMUL_PRIORITY_M = 1024; | ||
| 103 | -static constexpr int64_t NO_BATCH_DIM_SUM = 2; | ||
| 104 | - | ||
| 105 | -static const std::initializer_list<op::DataType> IN_TYPE_SUPPORT_LIST = {op::DataType::DT_INT4, | ||
| 106 | - op::DataType::DT_INT8}; | ||
| 107 | -static const std::initializer_list<op::DataType> INT4_TYPE_SUPPORT_LIST = {op::DataType::DT_INT4, | ||
| 108 | - op::DataType::DT_INT32}; | ||
| 109 | -static const std::initializer_list<op::DataType> OUT_TYPE_SUPPORT_LIST = {op::DataType::DT_INT8, | ||
| 110 | - op::DataType::DT_FLOAT16, | ||
| 111 | - op::DataType::DT_BF16, | ||
| 112 | - op::DataType::DT_INT32}; | ||
| 113 | -static const std::initializer_list<op::DataType> SCALE_TYPE_SUPPORT_LIST = {op::DataType::DT_UINT64, | ||
| 114 | - op::DataType::DT_BF16, | ||
| 115 | - op::DataType::DT_INT64, | ||
| 116 | - op::DataType::DT_FLOAT}; | ||
| 117 | -static const std::initializer_list<op::DataType> BIAS_TYPE_SUPPORT_LIST = {op::DataType::DT_INT32, | ||
| 118 | - op::DataType::DT_BF16, | ||
| 119 | - op::DataType::DT_FLOAT16, | ||
| 120 | - op::DataType::DT_FLOAT}; | ||
| 121 | -static const std::initializer_list<op::DataType> Y_SCALE_SUPPORT_LIST = {op::DataType::DT_UINT64}; | ||
| 122 | - | ||
| 123 | -static inline bool isA8W4Float(const aclTensor* x1, const aclTensor* x2) | ||
| 124 | -{ | ||
| 125 | - return x1->GetDataType() == op::DataType::DT_FLOAT8_E4M3FN && | ||
| 126 | - (x2->GetDataType() == op::DataType::DT_FLOAT || x2->GetDataType() == op::DataType::DT_FLOAT4_E2M1); | ||
| 127 | -} | ||
| 128 | - | ||
| 129 | -static inline bool isA8W4Int(const aclTensor* x1, const aclTensor* x2) | ||
| 130 | -{ | ||
| 131 | - return x1->GetDataType() == op::DataType::DT_INT8 && | ||
| 132 | - (x2->GetDataType() == op::DataType::DT_INT4 || x2->GetDataType() == op::DataType::DT_INT32); | ||
| 133 | -} | ||
| 134 | - | ||
| 135 | -static inline bool isMx(const aclTensor* scale) | ||
| 136 | -{ | ||
| 137 | - return scale->GetDataType() == op::DataType::DT_FLOAT8_E8M0; | ||
| 138 | -} | ||
| 139 | - | ||
| 140 | -static inline bool isA8W4Msd(const aclTensor* x1, const aclTensor* x2, const aclTensor* scale, | ||
| 141 | - const aclTensor* pertokenScale) | ||
| 142 | -{ | ||
| 143 | - if (x1->GetDataType() != op::DataType::DT_INT8) { | ||
| 144 | - return false; | ||
| 145 | - } | ||
| 146 | - | ||
| 147 | - if (std::find(INT4_TYPE_SUPPORT_LIST.begin(), INT4_TYPE_SUPPORT_LIST.end(), | ||
| 148 | - x2->GetDataType()) == INT4_TYPE_SUPPORT_LIST.end()) { | ||
| 149 | - return false; | ||
| 150 | - } | ||
| 151 | - | ||
| 152 | - if (scale->GetDataType() != op::DataType::DT_UINT64) { | ||
| 153 | - return false; | ||
| 154 | - } | ||
| 155 | - | ||
| 156 | - if (pertokenScale == nullptr || pertokenScale->GetDataType() != op::DataType::DT_FLOAT) { | ||
| 157 | - return false; | ||
| 158 | - } | ||
| 159 | - | ||
| 160 | - return true; | ||
| 161 | -} | ||
| 162 | - | ||
| 163 | static inline bool CheckNotNull(TupleTensor mandatoryTensors, const aclTensor *out) { | 44 | static inline bool CheckNotNull(TupleTensor mandatoryTensors, const aclTensor *out) { |
| 164 | auto x1 = std::get<INDEX_X1_IN_MANDTORY_TUPLE>(mandatoryTensors); | 45 | auto x1 = std::get<INDEX_X1_IN_MANDTORY_TUPLE>(mandatoryTensors); |
| 165 | auto x2 = std::get<INDEX_X2_IN_MANDTORY_TUPLE>(mandatoryTensors); | 46 | auto x2 = std::get<INDEX_X2_IN_MANDTORY_TUPLE>(mandatoryTensors); |
| @@ -670,23 +551,6 @@ static inline bool MaxDimCheck(int64_t x1DimNum, int64_t x2DimNum, const op::Sha | |||
| 670 | return true; | 551 | return true; |
| 671 | } | 552 | } |
| 672 | 553 | ||
| 673 | -static inline int64_t SelectNzK0Value(op::DataType dataType, const bool isA8W4Float) { | ||
| 674 | - switch (dataType) { | ||
| 675 | - case op::DataType::DT_INT4: | ||
| 676 | - return NZ_K0_VALUE_INT4_TRANS; | ||
| 677 | - case op::DataType::DT_FLOAT4_E2M1: | ||
| 678 | - if (isA8W4Float) { | ||
| 679 | - return NZ_K0_VALUE_INT8_TRANS; | ||
| 680 | - } else { | ||
| 681 | - return NZ_K0_VALUE_INT4_TRANS; | ||
| 682 | - } | ||
| 683 | - case op::DataType::DT_INT32: | ||
| 684 | - return NZ_K0_VALUE_INT32_TRANS; | ||
| 685 | - default: | ||
| 686 | - return NZ_K0_VALUE_INT8_TRANS; | ||
| 687 | - } | ||
| 688 | -} | ||
| 689 | - | ||
| 690 | static inline bool CheckShapeForWeightNz(const aclTensor *x1, const aclTensor *x2, bool transposeX1, bool transposeX2) { | 554 | static inline bool CheckShapeForWeightNz(const aclTensor *x1, const aclTensor *x2, bool transposeX1, bool transposeX2) { |
| 691 | const op::Shape x1Shape = x1->GetViewShape(); | 555 | const op::Shape x1Shape = x1->GetViewShape(); |
| 692 | const op::Shape x2Shape = x2->GetStorageShape(); | 556 | const op::Shape x2Shape = x2->GetStorageShape(); |
| @@ -706,17 +570,6 @@ static inline bool CheckShapeForWeightNz(const aclTensor *x1, const aclTensor *x | |||
| 706 | return true; | 570 | return true; |
| 707 | } | 571 | } |
| 708 | 572 | ||
| 709 | -template <typename T> | ||
| 710 | -static inline bool IsAligned(T num, T factor) | ||
| 711 | -{ | ||
| 712 | - if (factor == 0) { | ||
| 713 | - OP_LOGE(ACLNN_ERR_PARAM_INVALID, | ||
| 714 | - "The divisor cannot be zero."); | ||
| 715 | - return false; | ||
| 716 | - } | ||
| 717 | - return num > 0 && num % factor == 0; | ||
| 718 | -} | ||
| 719 | - | ||
| 720 | static inline bool CheckShapeInt4(const aclTensor *x1, const aclTensor *x2, bool transposeX1, bool transposeX2, | 573 | static inline bool CheckShapeInt4(const aclTensor *x1, const aclTensor *x2, bool transposeX1, bool transposeX2, |
| 721 | const aclTensor *bias) | 574 | const aclTensor *bias) |
| 722 | { | 575 | { |
| @@ -821,19 +674,6 @@ static inline bool CheckEmptyTensor(TupleTensor mandatoryTensors) { | |||
| 821 | return true; | 674 | return true; |
| 822 | } | 675 | } |
| 823 | 676 | ||
| 824 | -static inline bool IsMicroScaling(const aclTensor *x1Scale, const aclTensor *x2Scale) { | ||
| 825 | - if (x1Scale == nullptr) { | ||
| 826 | - return false; | ||
| 827 | - } | ||
| 828 | - return x1Scale->GetDataType() == op::DataType::DT_FLOAT8_E8M0 && | ||
| 829 | - x2Scale->GetDataType() == op::DataType::DT_FLOAT8_E8M0; | ||
| 830 | -} | ||
| 831 | - | ||
| 832 | -static inline bool IsTCG(const aclTensor *x1Scale, const aclTensor *x2Scale) { | ||
| 833 | - return x1Scale == nullptr && x2Scale != nullptr && | ||
| 834 | - (x2Scale->GetDataType() == op::DataType::DT_BF16 || x2Scale->GetDataType() == op::DataType::DT_FLOAT16); | ||
| 835 | -} | ||
| 836 | - | ||
| 837 | static inline bool CheckA8W4FloatQuantType(const aclTensor *x1, const aclTensor *x2, | 677 | static inline bool CheckA8W4FloatQuantType(const aclTensor *x1, const aclTensor *x2, |
| 838 | const aclTensor *perTokenScale, const aclTensor *scale) { | 678 | const aclTensor *perTokenScale, const aclTensor *scale) { |
| 839 | if (isA8W4Float(x1, x2)) { | 679 | if (isA8W4Float(x1, x2)) { |
| @@ -996,7 +836,7 @@ static inline bool CheckA8W4ScaleX2Shape( | |||
| 996 | x2Scale->GetViewShape().GetDim(MX_SCALE_LAST_DIM_INDEX) != MX_SCALE_LAST_DIM) { | 836 | x2Scale->GetViewShape().GetDim(MX_SCALE_LAST_DIM_INDEX) != MX_SCALE_LAST_DIM) { |
| 997 | OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON( | 837 | OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON( |
| 998 | "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x2Scale", | 838 | "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x2Scale", |
| 999 | - FormatString("%ld, %ld, %ld", x2ScaleNDim, x2ScaleGroupDim, | 839 | + FormatString("%ld, %ld, %ld", x2ScaleNDim, x2ScaleGroupDim, |
| 1000 | x2Scale->GetViewShape().GetDim(MX_SCALE_LAST_DIM_INDEX)).c_str(), | 840 | x2Scale->GetViewShape().GetDim(MX_SCALE_LAST_DIM_INDEX)).c_str(), |
| 1001 | FormatString("the shape of x2Scale must be [%ld, %ld, 2]", groupDimN, | 841 | FormatString("the shape of x2Scale must be [%ld, %ld, 2]", groupDimN, |
| 1002 | CeilDiv(groupDimK, x2ScaleReshapeFactor)).c_str()); | 842 | CeilDiv(groupDimK, x2ScaleReshapeFactor)).c_str()); |
| @@ -1070,7 +910,7 @@ static inline bool CheckA8W4X1X2Shape(int64_t x1KDim, int64_t x2KDim, int64_t x2 | |||
| 1070 | if (isMx && (x1KDim % SUPPORTED_MX_A8W4_K_ALIGN_NUM != 0)) { // Mx量化k方向8对齐 | 910 | if (isMx && (x1KDim % SUPPORTED_MX_A8W4_K_ALIGN_NUM != 0)) { // Mx量化k方向8对齐 |
| 1071 | OP_LOGE_FOR_INVALID_VALUE_WITH_REASON( | 911 | OP_LOGE_FOR_INVALID_VALUE_WITH_REASON( |
| 1072 | "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1", std::to_string(x1KDim).c_str(), | 912 | "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1", std::to_string(x1KDim).c_str(), |
| 1073 | - FormatString("the k dimension of x1 must be aligned to %ld for MX quantization", | 913 | + FormatString("the k dimension of x1 must be aligned to %ld for MX quantization", |
| 1074 | SUPPORTED_MX_A8W4_K_ALIGN_NUM).c_str()); | 914 | SUPPORTED_MX_A8W4_K_ALIGN_NUM).c_str()); |
| 1075 | return false; | 915 | return false; |
| 1076 | } | 916 | } |
| @@ -1172,12 +1012,6 @@ static aclnnStatus CheckParamsDAV3510(TupleTensor mandatoryTensors, TupleOptiona | |||
| 1172 | return qmmV3Checker.CheckParams(); | 1012 | return qmmV3Checker.CheckParams(); |
| 1173 | } | 1013 | } |
| 1174 | 1014 | ||
| 1175 | -static bool IsFormatNZ(const aclTensor* tensor) { | ||
| 1176 | - return ge::GetPrimaryFormat(tensor->GetStorageFormat()) == op::Format::FORMAT_FRACTAL_NZ || | ||
| 1177 | - ge::GetPrimaryFormat(tensor->GetStorageFormat()) == op::Format::FORMAT_FRACTAL_NZ_C0_4 || | ||
| 1178 | - ge::GetPrimaryFormat(tensor->GetStorageFormat()) == op::Format::FORMAT_FRACTAL_NZ_C0_32; | ||
| 1179 | -} | ||
| 1180 | - | ||
| 1181 | static aclnnStatus CheckWeightNzParamsDAV3510(const aclTensor *x1, const aclTensor *x2, const aclTensor *out) | 1015 | static aclnnStatus CheckWeightNzParamsDAV3510(const aclTensor *x1, const aclTensor *x2, const aclTensor *out) |
| 1182 | { | 1016 | { |
| 1183 | if (op::GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_3510) { | 1017 | if (op::GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_3510) { |
| @@ -1528,7 +1362,7 @@ static inline bool CheckInputAttrExistence(const TupleAttr &boolsTrans, const Tu | |||
| 1528 | "aclnnQuantMatmulWeightNzGetWorkspaceSize", "transposeX2", transposeX2 ? "true" : "false", | 1362 | "aclnnQuantMatmulWeightNzGetWorkspaceSize", "transposeX2", transposeX2 ? "true" : "false", |
| 1529 | "in A8W4 scenario with NZ format, when the quantization mode is t-cg, transposeX2 must be false"); | 1363 | "in A8W4 scenario with NZ format, when the quantization mode is t-cg, transposeX2 must be false"); |
| 1530 | return false; | 1364 | return false; |
| 1531 | - } else if (IsMicroScaling(x1Scale, x2Scale) && !transposeX2) { | 1365 | + } else if (IsMicroScaling(x1Scale, x2Scale) && !transposeX2) { |
| 1532 | // A8W4 scenario with mx quant mode | 1366 | // A8W4 scenario with mx quant mode |
| 1533 | OP_LOGE_FOR_INVALID_VALUE_WITH_REASON( | 1367 | OP_LOGE_FOR_INVALID_VALUE_WITH_REASON( |
| 1534 | "aclnnQuantMatmulWeightNzGetWorkspaceSize", "transposeX2", transposeX2 ? "true" : "false", | 1368 | "aclnnQuantMatmulWeightNzGetWorkspaceSize", "transposeX2", transposeX2 ? "true" : "false", |
| @@ -1552,7 +1386,7 @@ static inline bool CheckDimRangeA8W4(const TupleTensor& mandatoryTensors, const | |||
| 1552 | auto bias = std::get<INDEX_BIAS_IN_OPTIONAL_TUPLE>(optionalTensors); | 1386 | auto bias = std::get<INDEX_BIAS_IN_OPTIONAL_TUPLE>(optionalTensors); |
| 1553 | 1387 | ||
| 1554 | if (x1->GetViewShape().GetDimNum() != MAX_DIM_VALUE) { | 1388 | if (x1->GetViewShape().GetDimNum() != MAX_DIM_VALUE) { |
| 1555 | - OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON("aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1", | 1389 | + OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON("aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1", |
| 1556 | FormatString("%zuD", x1->GetViewShape().GetDimNum()).c_str(), "the shape dim of x1 must be 2"); | 1390 | FormatString("%zuD", x1->GetViewShape().GetDimNum()).c_str(), "the shape dim of x1 must be 2"); |
| 1557 | return false; | 1391 | return false; |
| 1558 | } | 1392 | } |
| @@ -0,0 +1,214 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | + | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | +namespace quant_matmul_v4 { | ||
| 26 | +// Constants | ||
| 27 | +static constexpr int MX_SCALE_LAST_DIM_INDEX = 2; | ||
| 28 | +static constexpr int MX_SCALE_LAST_DIM = 2; | ||
| 29 | + | ||
| 30 | +static constexpr int INDEX_X1_IN_MANDTORY_TUPLE = 0; | ||
| 31 | +static constexpr int INDEX_X2_IN_MANDTORY_TUPLE = 1; | ||
| 32 | +static constexpr int INDEX_SCALE_IN_MANDTORY_TUPLE = 2; | ||
| 33 | +static constexpr int INDEX_OFFSET_IN_OPTIONAL_TUPLE = 0; | ||
| 34 | +static constexpr int INDEX_PERTOKEN_IN_OPTIONAL_TUPLE = 1; | ||
| 35 | +static constexpr int INDEX_BIAS_IN_OPTIONAL_TUPLE = 2; | ||
| 36 | +static constexpr int INDEX_Y_SCALE_IN_OPTIONAL_TUPLE = 3; | ||
| 37 | +static constexpr int INDEX_Y_OFFSET_IN_OPTIONAL_TUPLE = 4; | ||
| 38 | +static constexpr int INDEX_GROUP_SIZE_IN_OPTIONAL_TUPLE = 5; | ||
| 39 | +static constexpr int INDEX_OUT_IN_TUPLE = 2; | ||
| 40 | +static constexpr int INDEX_ISA4W4_IN_BOOL_TUPLE = 2; | ||
| 41 | +static constexpr size_t LAST_SECOND_DIM_INDEX = 2; | ||
| 42 | + | ||
| 43 | +static constexpr int MIN_DIM_NUM_ND = 2; | ||
| 44 | +static constexpr int MAX_DIM_NUM_ND = 6; | ||
| 45 | +static constexpr int MIN_DIM_NUM_NZ = 4; | ||
| 46 | +static constexpr int MAX_DIM_NUM_NZ = 8; | ||
| 47 | +static constexpr int PENULTIMATE_DIM = 2; | ||
| 48 | +static constexpr int NZ_K1_INDEX = 3; | ||
| 49 | +static constexpr int NZ_K1_INDEX_TRANS = 4; | ||
| 50 | +static constexpr int NZ_STORAGE_PENULTIMATE_DIM = 16; | ||
| 51 | +static constexpr int NZ_STORAGE_LAST_DIM = 32; | ||
| 52 | +static constexpr int64_t NZ_K0_VALUE_BMM_BLOCK_NUM = 16; | ||
| 53 | +static constexpr int64_t NZ_K0_VALUE_INT32_TRANS = 8; | ||
| 54 | +static constexpr int64_t NZ_K0_VALUE_INT8_TRANS = 32; | ||
| 55 | +static constexpr int64_t NZ_K0_VALUE_INT4_TRANS = 64; | ||
| 56 | +static constexpr int64_t OUTPUT_INFER_FAIL = -1L; | ||
| 57 | +static constexpr int64_t LAST_AXIS_LIMIT = 65535; | ||
| 58 | +static constexpr int X2_FIXED_DIM_NUM_A4W4 = 2; | ||
| 59 | +static constexpr int64_t INT4_NUMS_IN_INT8 = 2; | ||
| 60 | +static constexpr int64_t INT4_NUMS_IN_INT32 = 8; | ||
| 61 | +static constexpr int64_t INNER_SIZE_MULTIPLE = 64; | ||
| 62 | +static constexpr int64_t K_VALUE = 3696; | ||
| 63 | +static constexpr int64_t N_VALUE = 8192; | ||
| 64 | +static constexpr int64_t M_RANGE1_LEFT = 128; | ||
| 65 | +static constexpr int64_t M_RANGE1_RIGHT = 512; | ||
| 66 | +static constexpr int32_t CORE_NUM_20 = 20; | ||
| 67 | +static constexpr int64_t SUPPORTED_GROUP_SIZE = 32; | ||
| 68 | +static constexpr uint64_t B4_PER_B32 = 8UL; | ||
| 69 | +static constexpr int64_t SUPPORTED_TCG_A8W4_K_ALIGN_NUM = 32; | ||
| 70 | +static constexpr int64_t SUPPORTED_MX_A8W4_K_ALIGN_NUM = 8; | ||
| 71 | +static constexpr int64_t SUPPORTED_N_ALIGN_NUM = 8; | ||
| 72 | +static constexpr size_t MAX_DIM_VALUE = 2; | ||
| 73 | +static constexpr size_t MX_SCALE_DIM_VALUE = 3; | ||
| 74 | +static constexpr uint64_t GROUP_M_OFFSET = 32; | ||
| 75 | +static constexpr uint64_t GROUP_N_OFFSET = 16; | ||
| 76 | +static constexpr uint64_t GROUP_MNK_BIT_SIZE = 0xFFFF; | ||
| 77 | +static constexpr size_t MX_SCALE_MAX_DIM = 3; | ||
| 78 | +static constexpr size_t MX_SCALE_DIM_NUM = 3; | ||
| 79 | +static constexpr int64_t MAX_SHAPE_SIZE_A8W4_INT = 29576; | ||
| 80 | +static constexpr int64_t PPMATMUL_PRIORITY_M = 1024; | ||
| 81 | +static constexpr int64_t NO_BATCH_DIM_SUM = 2; | ||
| 82 | + | ||
| 83 | +static const std::initializer_list<op::DataType> IN_TYPE_SUPPORT_LIST = {op::DataType::DT_INT4, | ||
| 84 | + op::DataType::DT_INT8}; | ||
| 85 | +static const std::initializer_list<op::DataType> INT4_TYPE_SUPPORT_LIST = {op::DataType::DT_INT4, | ||
| 86 | + op::DataType::DT_INT32}; | ||
| 87 | +static const std::initializer_list<op::DataType> OUT_TYPE_SUPPORT_LIST = {op::DataType::DT_INT8, | ||
| 88 | + op::DataType::DT_FLOAT16, | ||
| 89 | + op::DataType::DT_BF16, | ||
| 90 | + op::DataType::DT_INT32}; | ||
| 91 | +static const std::initializer_list<op::DataType> SCALE_TYPE_SUPPORT_LIST = {op::DataType::DT_UINT64, | ||
| 92 | + op::DataType::DT_BF16, | ||
| 93 | + op::DataType::DT_INT64, | ||
| 94 | + op::DataType::DT_FLOAT}; | ||
| 95 | +static const std::initializer_list<op::DataType> BIAS_TYPE_SUPPORT_LIST = {op::DataType::DT_INT32, | ||
| 96 | + op::DataType::DT_BF16, | ||
| 97 | + op::DataType::DT_FLOAT16, | ||
| 98 | + op::DataType::DT_FLOAT}; | ||
| 99 | +static const std::initializer_list<op::DataType> Y_SCALE_SUPPORT_LIST = {op::DataType::DT_UINT64}; | ||
| 100 | + | ||
| 101 | +using TupleTensor = std::tuple<const aclTensor *, const aclTensor *, const aclTensor *>; | ||
| 102 | +using TupleOptional = std::tuple<const aclTensor *, const aclTensor *, const aclTensor *, const aclTensor *, | ||
| 103 | + const aclTensor *, const int64_t &>; | ||
| 104 | +using TupleInput = std::tuple<const aclTensor *, const aclTensor *>; | ||
| 105 | +using TupleQuant = std::tuple<const aclTensor *, const aclTensor *, const aclTensor *, const aclTensor *, | ||
| 106 | + const aclTensor *, const aclTensor *, const aclTensor *, const int64_t &, const int64_t &>; | ||
| 107 | +using TupleAttr = std::tuple<bool, bool>; | ||
| 108 | + | ||
| 109 | +static inline bool isA8W4Float(const aclTensor* x1, const aclTensor* x2) | ||
| 110 | +{ | ||
| 111 | + if (x1 == nullptr || x2 == nullptr) { | ||
| 112 | + return false; | ||
| 113 | + } | ||
| 114 | + return x1->GetDataType() == op::DataType::DT_FLOAT8_E4M3FN && | ||
| 115 | + (x2->GetDataType() == op::DataType::DT_FLOAT || x2->GetDataType() == op::DataType::DT_FLOAT4_E2M1); | ||
| 116 | +} | ||
| 117 | + | ||
| 118 | +static inline bool isA8W4Int(const aclTensor* x1, const aclTensor* x2) | ||
| 119 | +{ | ||
| 120 | + if (x1 == nullptr || x2 == nullptr) { | ||
| 121 | + return false; | ||
| 122 | + } | ||
| 123 | + return x1->GetDataType() == op::DataType::DT_INT8 && | ||
| 124 | + (x2->GetDataType() == op::DataType::DT_INT4 || x2->GetDataType() == op::DataType::DT_INT32); | ||
| 125 | +} | ||
| 126 | + | ||
| 127 | +static inline bool isMx(const aclTensor* scale) | ||
| 128 | +{ | ||
| 129 | + if (scale == nullptr) { | ||
| 130 | + return false; | ||
| 131 | + } | ||
| 132 | + return scale->GetDataType() == op::DataType::DT_FLOAT8_E8M0; | ||
| 133 | +} | ||
| 134 | + | ||
| 135 | +static inline bool isA8W4Msd(const aclTensor* x1, const aclTensor* x2, const aclTensor* scale, | ||
| 136 | + const aclTensor* pertokenScale) | ||
| 137 | +{ | ||
| 138 | + if (x1 == nullptr || x2 == nullptr || scale == nullptr) { | ||
| 139 | + return false; | ||
| 140 | + } | ||
| 141 | + if (x1->GetDataType() != op::DataType::DT_INT8) { | ||
| 142 | + return false; | ||
| 143 | + } | ||
| 144 | + | ||
| 145 | + if (std::find(INT4_TYPE_SUPPORT_LIST.begin(), INT4_TYPE_SUPPORT_LIST.end(), | ||
| 146 | + x2->GetDataType()) == INT4_TYPE_SUPPORT_LIST.end()) { | ||
| 147 | + return false; | ||
| 148 | + } | ||
| 149 | + | ||
| 150 | + if (scale->GetDataType() != op::DataType::DT_UINT64) { | ||
| 151 | + return false; | ||
| 152 | + } | ||
| 153 | + | ||
| 154 | + if (pertokenScale == nullptr || pertokenScale->GetDataType() != op::DataType::DT_FLOAT) { | ||
| 155 | + return false; | ||
| 156 | + } | ||
| 157 | + | ||
| 158 | + return true; | ||
| 159 | +} | ||
| 160 | + | ||
| 161 | +static inline bool IsMicroScaling(const aclTensor *x1Scale, const aclTensor *x2Scale) { | ||
| 162 | + if (x1Scale == nullptr) { | ||
| 163 | + return false; | ||
| 164 | + } | ||
| 165 | + if (x2Scale == nullptr) { | ||
| 166 | + return false; | ||
| 167 | + } | ||
| 168 | + return x1Scale->GetDataType() == op::DataType::DT_FLOAT8_E8M0 && | ||
| 169 | + x2Scale->GetDataType() == op::DataType::DT_FLOAT8_E8M0; | ||
| 170 | +} | ||
 | |||
| 171 | + | ||
| 172 | +static inline bool IsTCG(const aclTensor *x1Scale, const aclTensor *x2Scale) { | ||
| 173 | + return x1Scale == nullptr && x2Scale != nullptr && | ||
| 174 | + (x2Scale->GetDataType() == op::DataType::DT_BF16 || x2Scale->GetDataType() == op::DataType::DT_FLOAT16); | ||
| 175 | +} | ||
| 176 | + | ||
| 177 | +static inline int64_t SelectNzK0Value(op::DataType dataType, const bool isA8W4Float) { | ||
| 178 | + switch (dataType) { | ||
| 179 | + case op::DataType::DT_INT4: | ||
| 180 | + return NZ_K0_VALUE_INT4_TRANS; | ||
| 181 | + case op::DataType::DT_FLOAT4_E2M1: | ||
| 182 | + if (isA8W4Float) { | ||
| 183 | + return NZ_K0_VALUE_INT8_TRANS; | ||
| 184 | + } else { | ||
| 185 | + return NZ_K0_VALUE_INT4_TRANS; | ||
| 186 | + } | ||
| 187 | + case op::DataType::DT_INT32: | ||
| 188 | + return NZ_K0_VALUE_INT32_TRANS; | ||
| 189 | + default: | ||
| 190 | + return NZ_K0_VALUE_INT8_TRANS; | ||
| 191 | + } | ||
| 192 | +} | ||
| 193 | + | ||
| 194 | +template <typename T> | ||
| 195 | +static inline bool IsAligned(T num, T factor) | ||
| 196 | +{ | ||
| 197 | + if (factor == 0) { | ||
| 198 | + OP_LOGE(ACLNN_ERR_PARAM_INVALID, | ||
| 199 | + "The divisor cannot be zero."); | ||
| 200 | + return false; | ||
| 201 | + } | ||
| 202 | + return num > 0 && num % factor == 0; | ||
| 203 | +} | ||
| 204 | + | ||
| 205 | +static inline bool IsFormatNZ(const aclTensor* tensor) { | ||
| 206 | + return tensor != nullptr && | ||
| 207 | + (ge::GetPrimaryFormat(tensor->GetStorageFormat()) == op::Format::FORMAT_FRACTAL_NZ || | ||
| 208 | + ge::GetPrimaryFormat(tensor->GetStorageFormat()) == op::Format::FORMAT_FRACTAL_NZ_C0_4 || | ||
| 209 | + ge::GetPrimaryFormat(tensor->GetStorageFormat()) == op::Format::FORMAT_FRACTAL_NZ_C0_32); | ||
| 210 | +} | ||
| 211 | + | ||
| 212 | +} // namespace quant_matmul_v4 | ||
| 213 | + | ||
| 214 | + | ||
| @@ -94,7 +94,7 @@ const void* AdaptiveSlidingWindowMXBasicAPITiling::GetTilingData() const | |||
| 94 | static_cast<const void*>(&tilingData_); | 94 | static_cast<const void*>(&tilingData_); |
| 95 | } | 95 | } |
| 96 | 96 | ||
| 97 | -uint64_t AdaptiveSlidingWindowMXBasicAPITiling::GetApiLevel(NpuArch) const | 97 | +uint64_t AdaptiveSlidingWindowMXBasicAPITiling::GetApiLevel(NpuArch npuArch) const |
| 98 | { | 98 | { |
| 99 | return IsTensorapiCapable() ? static_cast<uint64_t>(QMMApiLevel::BLAZE_LEVEL) : | 99 | return IsTensorapiCapable() ? static_cast<uint64_t>(QMMApiLevel::BLAZE_LEVEL) : |
| 100 | static_cast<uint64_t>(QMMApiLevel::BASIC_LEVEL); | 100 | static_cast<uint64_t>(QMMApiLevel::BASIC_LEVEL); |
| @@ -15,14 +15,6 @@ | |||
| 15 | 15 | ||
| 16 | 16 | ||
| 17 | 17 | ||
| 18 | - | ||
| 19 | - | ||
| 20 | - | ||
| 21 | - | ||
| 22 | - | ||
| 23 | - | ||
| 24 | - | ||
| 25 | - | ||
| 26 | 18 | ||
| 27 | 19 | ||
| 28 | 20 | ||
| @@ -776,44 +768,44 @@ UT_STATIC __global__ __aicore__ void quant_batch_matmul_v3( | |||
| 776 | 768 | ||
| 777 | if constexpr (TPL_KERNELTYPE == TPL_NO_VEC_EPILOGUE_WITH_MMAPI && | 769 | if constexpr (TPL_KERNELTYPE == TPL_NO_VEC_EPILOGUE_WITH_MMAPI && |
| 778 | TPL_APILEVEL == TPL_API_LEVEL_HIGH) { // Kernel Type = 0; | 770 | TPL_APILEVEL == TPL_API_LEVEL_HIGH) { // Kernel Type = 0; |
| 779 | - GET_TILING_DATA_WITH_STRUCT(DequantBmm::QuantBatchMatmulV3TilingDataParams, tilingData, tiling); | 771 | + GET_TILING_DATA_WITH_STRUCT(DequantBmm::QuantBatchMatmulV3TilingDataParams, tilingData, tiling); |
| 780 | - MatMulASWKernel< | 772 | + MatMulASWKernel< |
| 781 | - DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_BIAS, DTYPE_Y, format_x1, format_x2, format_y, | 773 | + DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_BIAS, DTYPE_Y, format_x1, format_x2, format_y, |
| 782 | - static_cast<bool>(TPL_ATRANS), static_cast<bool>(TPL_BTRANS)> | 774 | + static_cast<bool>(TPL_ATRANS), static_cast<bool>(TPL_BTRANS)> |
| 783 | - op; | 775 | + op; |
| 784 | - op.Init(x1, x2, bias, scale, pertokenScale, y, user1, &tilingData, &tPipe); | 776 | + op.Init(x1, x2, bias, scale, pertokenScale, y, user1, &tilingData, &tPipe); |
| 785 | - op.Process(); | 777 | + op.Process(); |
| 786 | - } | 778 | + } |
| 787 | if constexpr (TPL_KERNELTYPE == TPL_NO_VEC_EPILOGUE_CUSTOM_GMTOAL1_WITH_MMAPI && | 779 | if constexpr (TPL_KERNELTYPE == TPL_NO_VEC_EPILOGUE_CUSTOM_GMTOAL1_WITH_MMAPI && |
| 788 | TPL_APILEVEL == TPL_API_LEVEL_HIGH) { // Kernel Type = 1; | 780 | TPL_APILEVEL == TPL_API_LEVEL_HIGH) { // Kernel Type = 1; |
| 789 | - GET_TILING_DATA_WITH_STRUCT(DequantBmm::QuantBatchMatmulV3TilingDataParams, tilingData, tiling); | 781 | + GET_TILING_DATA_WITH_STRUCT(DequantBmm::QuantBatchMatmulV3TilingDataParams, tilingData, tiling); |
| 790 | - QuantBatchMatmulV3::MatmulAswKernelAL1FullLoad< | 782 | + QuantBatchMatmulV3::MatmulAswKernelAL1FullLoad< |
| 791 | - DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_BIAS, DTYPE_Y, format_x1, format_x2, format_y, | 783 | + DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_BIAS, DTYPE_Y, format_x1, format_x2, format_y, |
| 792 | - static_cast<bool>(TPL_ATRANS), static_cast<bool>(TPL_BTRANS)> | 784 | + static_cast<bool>(TPL_ATRANS), static_cast<bool>(TPL_BTRANS)> |
| 793 | - op; | 785 | + op; |
| 794 | - op.Init(x1, x2, bias, scale, pertokenScale, y, user1, &tilingData, &tPipe); | 786 | + op.Init(x1, x2, bias, scale, pertokenScale, y, user1, &tilingData, &tPipe); |
| 795 | - op.Process(); | 787 | + op.Process(); |
| 796 | - } | 788 | + } |
| 797 | if constexpr (TPL_KERNELTYPE == TPL_NO_VEC_EPILOGUE_CUSTOM_GMTOABL1_WITH_MMAPI && | 789 | if constexpr (TPL_KERNELTYPE == TPL_NO_VEC_EPILOGUE_CUSTOM_GMTOABL1_WITH_MMAPI && |
| 798 | TPL_APILEVEL == TPL_API_LEVEL_HIGH) { | 790 | TPL_APILEVEL == TPL_API_LEVEL_HIGH) { |
| 799 | - GET_TILING_DATA_WITH_STRUCT(DequantBmm::QuantBatchMatmulV3TilingDataParams, tilingData, tiling); | 791 | + GET_TILING_DATA_WITH_STRUCT(DequantBmm::QuantBatchMatmulV3TilingDataParams, tilingData, tiling); |
| 800 | - QuantBatchMatmulV3::MatmulAswKernelABL1FullLoad< | 792 | + QuantBatchMatmulV3::MatmulAswKernelABL1FullLoad< |
| 801 | - DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_BIAS, DTYPE_Y, format_x1, format_x2, format_y, | 793 | + DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_BIAS, DTYPE_Y, format_x1, format_x2, format_y, |
| 802 | - static_cast<bool>(TPL_ATRANS), static_cast<bool>(TPL_BTRANS)> | 794 | + static_cast<bool>(TPL_ATRANS), static_cast<bool>(TPL_BTRANS)> |
| 803 | - op; | 795 | + op; |
| 804 | - op.Init(x1, x2, bias, scale, pertokenScale, y, user1, &tilingData, &tPipe); | 796 | + op.Init(x1, x2, bias, scale, pertokenScale, y, user1, &tilingData, &tPipe); |
| 805 | - op.Process(); | 797 | + op.Process(); |
| 806 | - } | 798 | + } |
| 807 | if constexpr (TPL_KERNELTYPE == TPL_NO_VEC_EPILOGUE_CUSTOM_GMTOBL1_WITH_MMAPI && | 799 | if constexpr (TPL_KERNELTYPE == TPL_NO_VEC_EPILOGUE_CUSTOM_GMTOBL1_WITH_MMAPI && |
| 808 | TPL_APILEVEL == TPL_API_LEVEL_HIGH) { | 800 | TPL_APILEVEL == TPL_API_LEVEL_HIGH) { |
| 809 | - GET_TILING_DATA_WITH_STRUCT(DequantBmm::QuantBatchMatmulV3TilingDataParams, tilingData, tiling); | 801 | + GET_TILING_DATA_WITH_STRUCT(DequantBmm::QuantBatchMatmulV3TilingDataParams, tilingData, tiling); |
| 810 | - QuantBatchMatmulV3::MatmulAswKernelBL1FullLoad< | 802 | + QuantBatchMatmulV3::MatmulAswKernelBL1FullLoad< |
| 811 | - DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_BIAS, DTYPE_Y, format_x1, format_x2, format_y, | 803 | + DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_BIAS, DTYPE_Y, format_x1, format_x2, format_y, |
| 812 | - static_cast<bool>(TPL_ATRANS), static_cast<bool>(TPL_BTRANS)> | 804 | + static_cast<bool>(TPL_ATRANS), static_cast<bool>(TPL_BTRANS)> |
| 813 | - op; | 805 | + op; |
| 814 | - op.Init(x1, x2, bias, scale, pertokenScale, y, user1, &tilingData, &tPipe); | 806 | + op.Init(x1, x2, bias, scale, pertokenScale, y, user1, &tilingData, &tPipe); |
| 815 | - op.Process(); | 807 | + op.Process(); |
| 816 | - } | 808 | + } |
| 817 | 809 | ||
| 818 | } | 810 | } |
| 819 | 811 | ||
| @@ -30,6 +30,14 @@ | |||
| 30 | 30 | ||
| 31 | 31 | ||
| 32 | 32 | ||
| 33 | + | ||
| 34 | + | ||
| 35 | + | ||
| 36 | + | ||
| 37 | + | ||
| 38 | + | ||
| 39 | + | ||
| 40 | + | ||
| 33 | 41 | ||
| 34 | defined(ORIG_DTYPE_X1) && defined(ORIG_DTYPE_X2) && defined(DT_INT4) && \ | 42 | defined(ORIG_DTYPE_X1) && defined(ORIG_DTYPE_X2) && defined(DT_INT4) && \ |
| 35 | ORIG_DTYPE_X1 == DT_INT4 && ORIG_DTYPE_X2 == DT_INT4 | 43 | ORIG_DTYPE_X1 == DT_INT4 && ORIG_DTYPE_X2 == DT_INT4 |
🟡 Medium Priority
函数
IsMicroScaling（quant_matmul_v4_common.h第 149-155 行）在检查x1Scale == nullptr提前返回后，无条件解引用了x2Scale->GetDataType()，但并未对x2Scale做空指针检查。变更影响：该函数原本是
aclnn_quant_matmul_v4.cpp匿名命名空间内的 file-local 函数，所有现有调用方传入的x2Scale均来自 mandatory tensors（已通过CheckNotNull校验为非空），因此旧代码中不会被触发。但本次改动将函数移入公共头文件quant_matmul_v4_common.h，使其可被任意翻译单元调用，调用方可能传入空x2Scale而仅依赖该函数内部防护——当前缺乏此防护。失败模式：当
x1Scale != nullptr而x2Scale == nullptr时，x2Scale->GetDataType()引发空指针解引用崩溃。建议：在解引用 x2Scale 前增加空指针检查，例如在第 153 行前插入
if (x2Scale == nullptr) { return false; }。