已合并
fix: 整改责任田算子日志质量(修复日志规范问题) #9154
StoneChan_创建于 12 天前
fix: 整改责任田算子日志质量(修复日志规范问题) #9154
已合并
共 23 个文件变更+82-111
| @@ -41,13 +41,13 @@ static ge::graphStatus TilingPrepare4Glu(gert::TilingParseContext* context) | |||
| 41 | auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo); | 41 | auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo); |
| 42 | compileInfo->totalCoreNum = ascendcPlatform.GetCoreNumAiv(); | 42 | compileInfo->totalCoreNum = ascendcPlatform.GetCoreNumAiv(); |
| 43 | OP_LOGD(context, "Tiling totalCoreNum: %d", compileInfo->totalCoreNum); | 43 | OP_LOGD(context, "Tiling totalCoreNum: %d", compileInfo->totalCoreNum); |
| 44 | - OP_CHECK_IF((compileInfo->totalCoreNum <= 0), OP_LOGE(context, "TilingPrepare4Glu fail to get core num."), | 44 | + OP_CHECK_IF((compileInfo->totalCoreNum <= 0), OP_LOGE(context, "TilingPrepare4Glu Failed to get core num."), |
| 45 | return ge::GRAPH_FAILED); | 45 | return ge::GRAPH_FAILED); |
| 46 | 46 | ||
| 47 | uint64_t ubSizePlatForm; | 47 | uint64_t ubSizePlatForm; |
| 48 | ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm); | 48 | ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm); |
| 49 | compileInfo->ubSizePlatForm = static_cast<int64_t>(ubSizePlatForm); | 49 | compileInfo->ubSizePlatForm = static_cast<int64_t>(ubSizePlatForm); |
| 50 | - OP_CHECK_IF((compileInfo->ubSizePlatForm <= 0), OP_LOGE(context, "TilingPrepare4Glu fail to get ub size."), | 50 | + OP_CHECK_IF((compileInfo->ubSizePlatForm <= 0), OP_LOGE(context, "TilingPrepare4Glu Failed to get ub size."), |
| 51 | return ge::GRAPH_FAILED); | 51 | return ge::GRAPH_FAILED); |
| 52 | 52 | ||
| 53 | OP_LOGD(context, "TilingPrepare4Glu exit. coreNum: %d ubSize: %lu", compileInfo->totalCoreNum, | 53 | OP_LOGD(context, "TilingPrepare4Glu exit. coreNum: %d ubSize: %lu", compileInfo->totalCoreNum, |
| @@ -76,4 +76,4 @@ static ge::graphStatus Tiling4Glu(gert::TilingContext* context) | |||
| 76 | } | 76 | } |
| 77 | 77 | ||
| 78 | IMPL_OP_OPTILING(GLU).Tiling(Tiling4Glu).TilingParse<GluCompileInfo>(TilingPrepare4Glu); | 78 | IMPL_OP_OPTILING(GLU).Tiling(Tiling4Glu).TilingParse<GluCompileInfo>(TilingPrepare4Glu); |
| 79 | -} // namespace optiling | 79 | +} // namespace optiling |
| @@ -33,7 +33,7 @@ ge::graphStatus GluRegbaseTiling::RunFusionKernelTiling(gert::TilingContext* con | |||
| 33 | OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo); | 33 | OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo); |
| 34 | 34 | ||
| 35 | int64_t ubSizePlatForm = compileInfo->ubSizePlatForm - RESERVED_SIZE_8K; | 35 | int64_t ubSizePlatForm = compileInfo->ubSizePlatForm - RESERVED_SIZE_8K; |
| 36 | - OP_CHECK_IF((ubSizePlatForm <= 0), OP_LOGE(context, "RunFusionKernelTiling fail to get ub size."), | 36 | + OP_CHECK_IF((ubSizePlatForm <= 0), OP_LOGE(context, "RunFusionKernelTiling Failed to get ub size."), |
| 37 | return ge::GRAPH_FAILED); | 37 | return ge::GRAPH_FAILED); |
| 38 | 38 | ||
| 39 | int64_t commonBufferSize = ubSizePlatForm / (BUFFER_SIZE_FACTOR * sizeof(float)); | 39 | int64_t commonBufferSize = ubSizePlatForm / (BUFFER_SIZE_FACTOR * sizeof(float)); |
| @@ -46,4 +46,4 @@ ge::graphStatus GluRegbaseTiling::RunFusionKernelTiling(gert::TilingContext* con | |||
| 46 | return RunCommonTiling(context, commonBufferSize, commonBufferSize, 0); | 46 | return RunCommonTiling(context, commonBufferSize, commonBufferSize, 0); |
| 47 | } | 47 | } |
| 48 | 48 | ||
| 49 | -} // namespace optiling | 49 | +} // namespace optiling |
| @@ -138,7 +138,7 @@ inline ge::graphStatus CheckInputParams(gert::TilingContext* context) | |||
| 138 | int32_t typeSize = ge::GetSizeByDataType(dtype); | 138 | int32_t typeSize = ge::GetSizeByDataType(dtype); |
| 139 | 139 | ||
| 140 | OP_CHECK_IF(dtype != ge::DT_FLOAT16 && dtype != ge::DT_BF16 && dtype != ge::DT_FLOAT, | 140 | OP_CHECK_IF(dtype != ge::DT_FLOAT16 && dtype != ge::DT_BF16 && dtype != ge::DT_FLOAT, |
| 141 | - OP_LOGE(context, "input dtype only support fp16, fp32, bf16 currently, please check."), | 141 | + OP_LOGE(context, "input dtype only supports fp16, fp32, bf16 currently, please check."), |
| 142 | return ge::GRAPH_FAILED); | 142 | return ge::GRAPH_FAILED); |
| 143 | 143 | ||
| 144 | OP_CHECK_IF((typeSize <= 0), OP_LOGE(context, "typeSize is invalid %d, please check.", typeSize), | 144 | OP_CHECK_IF((typeSize <= 0), OP_LOGE(context, "typeSize is invalid %d, please check.", typeSize), |
| @@ -227,7 +227,7 @@ inline ge::graphStatus RunCommonTiling(gert::TilingContext* context, int64_t com | |||
| 227 | tilingData.set_ny(tilingParam.ny); | 227 | tilingData.set_ny(tilingParam.ny); |
| 228 | 228 | ||
| 229 | OP_CHECK_IF(SetTilingDataForGlu(context, tilingData) != ge::GRAPH_SUCCESS, | 229 | OP_CHECK_IF(SetTilingDataForGlu(context, tilingData) != ge::GRAPH_SUCCESS, |
| 230 | - OP_LOGE(context, "GluSetTilingData set tiling data fail."), return ge::GRAPH_FAILED); | 230 | + OP_LOGE(context, "GluSetTilingData failed to set tiling data."), return ge::GRAPH_FAILED); |
| 231 | 231 | ||
| 232 | context->SetBlockDim(tilingData.get_realCoreNum()); | 232 | context->SetBlockDim(tilingData.get_realCoreNum()); |
| 233 | context->SetTilingKey(tilingData.get_tilingKey()); | 233 | context->SetTilingKey(tilingData.get_tilingKey()); |
| @@ -251,4 +251,4 @@ inline ge::graphStatus RunCommonTiling(gert::TilingContext* context, int64_t com | |||
| 251 | } // namespace glu_common | 251 | } // namespace glu_common |
| 252 | } // namespace optiling | 252 | } // namespace optiling |
| 253 | 253 | ||
| 254 | -#endif // GLU_TILING_COMMON_H | 254 | +#endif // GLU_TILING_COMMON_H |
| @@ -20,9 +20,9 @@ using namespace std; | |||
| 20 | 20 | ||
| 21 | class l2_glu_test : public testing::Test { | 21 | class l2_glu_test : public testing::Test { |
| 22 | protected: | 22 | protected: |
| 23 | - static void SetUpTestCase() { std::cout << "glu_tensor_test SetUp" << std::endl; } | 23 | + static void SetUpTestCase() {} |
| 24 | 24 | ||
| 25 | - static void TearDownTestCase() { std::cout << "glu_tensor_test TearDown" << std::endl; } | 25 | + static void TearDownTestCase() {} |
| 26 | }; | 26 | }; |
| 27 | 27 | ||
| 28 | // 正常场景 float dim正数 | 28 | // 正常场景 float dim正数 |
| @@ -387,4 +387,4 @@ TEST_F(l2_glu_test, test_glu_date_range_f1_1) | |||
| 387 | aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize); | 387 | aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize); |
| 388 | EXPECT_EQ(aclRet, ACLNN_SUCCESS); | 388 | EXPECT_EQ(aclRet, ACLNN_SUCCESS); |
| 389 | ut.TestPrecision(); | 389 | ut.TestPrecision(); |
| 390 | -} | 390 | +} |
| @@ -23,9 +23,9 @@ | |||
| 23 | 23 | ||
| 24 | class GLU : public testing::Test { | 24 | class GLU : public testing::Test { |
| 25 | protected: | 25 | protected: |
| 26 | - static void SetUpTestCase() { std::cout << "GLU Proto Test SetUp" << std::endl; } | 26 | + static void SetUpTestCase() {} |
| 27 | 27 | ||
| 28 | - static void TearDownTestCase() { std::cout << "GLU Proto Test TearDown" << std::endl; } | 28 | + static void TearDownTestCase() {} |
| 29 | }; | 29 | }; |
| 30 | 30 | ||
| 31 | TEST_F(GLU, GLU_infershape_diff_test_legal_input) | 31 | TEST_F(GLU, GLU_infershape_diff_test_legal_input) |
| @@ -28,9 +28,9 @@ using namespace ge; | |||
| 28 | 28 | ||
| 29 | class GluTiling : public testing::Test { | 29 | class GluTiling : public testing::Test { |
| 30 | protected: | 30 | protected: |
| 31 | - static void SetUpTestCase() { std::cout << "GluTiling SetUp" << std::endl; } | 31 | + static void SetUpTestCase() {} |
| 32 | 32 | ||
| 33 | - static void TearDownTestCase() { std::cout << "GluTiling TearDown" << std::endl; } | 33 | + static void TearDownTestCase() {} |
| 34 | }; | 34 | }; |
| 35 | 35 | ||
| 36 | TEST_F(GluTiling, glu_tiling_bigshape) | 36 | TEST_F(GluTiling, glu_tiling_bigshape) |
| @@ -268,4 +268,4 @@ TEST_F(GluTiling, glu_tiling_singleshape) | |||
| 268 | // todo check tiling result | 268 | // todo check tiling result |
| 269 | auto tiling_key = tiling_context->GetTilingKey(); | 269 | auto tiling_key = tiling_context->GetTilingKey(); |
| 270 | ASSERT_EQ(tiling_key, 1); | 270 | ASSERT_EQ(tiling_key, 1); |
| 271 | -} | 271 | +} |
| @@ -16,8 +16,8 @@ using namespace ge; | |||
| 16 | 16 | ||
| 17 | class SleepInferShapeTest : public testing::Test { | 17 | class SleepInferShapeTest : public testing::Test { |
| 18 | protected: | 18 | protected: |
| 19 | - static void SetUpTestCase() { std::cout << "SleepInferShapeTest SetUp" << std::endl; } | 19 | + static void SetUpTestCase() {} |
| 20 | - static void TearDownTestCase() { std::cout << "SleepInferShapeTest TearDown" << std::endl; } | 20 | + static void TearDownTestCase() {} |
| 21 | }; | 21 | }; |
| 22 | 22 | ||
| 23 | TEST_F(SleepInferShapeTest, InferShape_success) | 23 | TEST_F(SleepInferShapeTest, InferShape_success) |
| @@ -29,8 +29,8 @@ using namespace ge; | |||
| 29 | 29 | ||
| 30 | class SleepTilingTest : public testing::Test { | 30 | class SleepTilingTest : public testing::Test { |
| 31 | protected: | 31 | protected: |
| 32 | - static void SetUpTestCase() { std::cout << "SleepTilingTest SetUp" << std::endl; } | 32 | + static void SetUpTestCase() {} |
| 33 | - static void TearDownTestCase() { std::cout << "SleepTilingTest TearDown" << std::endl; } | 33 | + static void TearDownTestCase() {} |
| 34 | }; | 34 | }; |
| 35 | 35 | ||
| 36 | template <typename T> | 36 | template <typename T> |
| @@ -35,8 +35,8 @@ using namespace std; | |||
| 35 | 35 | ||
| 36 | class sleep_kernel_test : public testing::Test { | 36 | class sleep_kernel_test : public testing::Test { |
| 37 | protected: | 37 | protected: |
| 38 | - static void SetUpTestCase() { cout << "sleep_kernel_test SetUp" << endl; } | 38 | + static void SetUpTestCase() {} |
| 39 | - static void TearDownTestCase() { cout << "sleep_kernel_test TearDown" << endl; } | 39 | + static void TearDownTestCase() {} |
| 40 | }; | 40 | }; |
| 41 | 41 | ||
| 42 | // 正常 cycles:kernel 应正常执行(CPU stub 下 SIMT 路径不编译,仅验证不崩溃) | 42 | // 正常 cycles:kernel 应正常执行(CPU stub 下 SIMT 路径不编译,仅验证不崩溃) |
| @@ -33,9 +33,9 @@ using namespace ge; | |||
| 33 | 33 | ||
| 34 | class ConcatOffsetTiling : public testing::Test { | 34 | class ConcatOffsetTiling : public testing::Test { |
| 35 | protected: | 35 | protected: |
| 36 | - static void SetUpTestCase() { std::cout << "ConcatOffsetTiling SetUp" << std::endl; } | 36 | + static void SetUpTestCase() {} |
| 37 | 37 | ||
| 38 | - static void TearDownTestCase() { std::cout << "ConcatOffsetTiling TearDown" << std::endl; } | 38 | + static void TearDownTestCase() {} |
| 39 | }; | 39 | }; |
| 40 | 40 | ||
| 41 | template <typename T> | 41 | template <typename T> |
| @@ -173,9 +173,5 @@ TEST_F(ConcatOffsetTiling, concat_offset_simt_tiling_1) | |||
| 173 | 173 | ||
| 174 | // todo check tiling result | 174 | // todo check tiling result |
| 175 | auto tiling_key = tiling_context->GetTilingKey(); | 175 | auto tiling_key = tiling_context->GetTilingKey(); |
| 176 | - auto block_dim = tiling_context->GetBlockDim(); | ||
| 177 | - auto raw_tiling_data = tiling_context->GetRawTilingData(); | ||
| 178 | - auto tiling_data_result = to_string<int64_t>(raw_tiling_data->GetData(), raw_tiling_data->GetDataSize()); | ||
| 179 | EXPECT_EQ(tiling_key, 1000); | 176 | EXPECT_EQ(tiling_key, 1000); |
| 180 | - std::cout << tiling_data_result << std::endl; | 177 | +} |
| 181 | -} | ||
| @@ -54,7 +54,7 @@ using std::vector; | |||
| 54 | ret = GenOnesDataInt64(placeholder##inputIndex##_shape, tensor_placeholder##inputIndex, \ | 54 | ret = GenOnesDataInt64(placeholder##inputIndex##_shape, tensor_placeholder##inputIndex, \ |
| 55 | placeholder##inputIndex##_desc, val); \ | 55 | placeholder##inputIndex##_desc, val); \ |
| 56 | if (ret != SUCCESS) { \ | 56 | if (ret != SUCCESS) { \ |
| 57 | - printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \ | 57 | + LOG_PRINT("%s - ERROR - [REVERSE_SEQUENCE_GE_IR]: Generate input data failed\n", GetTime().c_str()); \ |
| 58 | return FAILED; \ | 58 | return FAILED; \ |
| 59 | } \ | 59 | } \ |
| 60 | placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \ | 60 | placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \ |
| @@ -74,7 +74,7 @@ using std::vector; | |||
| 74 | ret = GenOnesDataDouble(placeholder##inputIndex##_shape, tensor_placeholder##inputIndex, \ | 74 | ret = GenOnesDataDouble(placeholder##inputIndex##_shape, tensor_placeholder##inputIndex, \ |
| 75 | placeholder##inputIndex##_desc, val); \ | 75 | placeholder##inputIndex##_desc, val); \ |
| 76 | if (ret != SUCCESS) { \ | 76 | if (ret != SUCCESS) { \ |
| 77 | - printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \ | 77 | + LOG_PRINT("%s - ERROR - [REVERSE_SEQUENCE_GE_IR]: Generate input data failed\n", GetTime().c_str()); \ |
| 78 | return FAILED; \ | 78 | return FAILED; \ |
| 79 | } \ | 79 | } \ |
| 80 | placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \ | 80 | placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \ |
| @@ -179,7 +179,8 @@ bool InitEnv() | |||
| 179 | std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}}; | 179 | std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}}; |
| 180 | Status ret = ge::GEInitialize(global_options); | 180 | Status ret = ge::GEInitialize(global_options); |
| 181 | if (ret != SUCCESS) { | 181 | if (ret != SUCCESS) { |
| 182 | - LOG_PRINT("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str()); | 182 | + LOG_PRINT("%s - ERROR - [REVERSE_SEQUENCE_GE_IR]: Initialize ge using ge global options failed\n", |
| 183 | + GetTime().c_str()); | ||
| 183 | return false; | 184 | return false; |
| 184 | } | 185 | } |
| 185 | return true; | 186 | return true; |
| @@ -192,7 +193,8 @@ bool CreateAndConfigGraph(Graph& graph, std::vector<ge::Tensor>& input) | |||
| 192 | 193 | ||
| 193 | Status ret = CreateOppInGraph(DT_DOUBLE, DT_INT64, input, inputs, outputs, graph); | 194 | Status ret = CreateOppInGraph(DT_DOUBLE, DT_INT64, input, inputs, outputs, graph); |
| 194 | if (ret != SUCCESS) { | 195 | if (ret != SUCCESS) { |
| 195 | - LOG_PRINT("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str()); | 196 | + LOG_PRINT("%s - ERROR - [REVERSE_SEQUENCE_GE_IR]: Create ir session using build options failed\n", |
| 197 | + GetTime().c_str()); | ||
| 196 | return false; | 198 | return false; |
| 197 | } | 199 | } |
| 198 | 200 | ||
| @@ -210,11 +212,11 @@ bool AddGraphToSession(ge::Session* session, Graph& graph, uint32_t graph_id) | |||
| 210 | if (ret != SUCCESS) { | 212 | if (ret != SUCCESS) { |
| 211 | ge::AscendString error_msg = ge::GEGetErrorMsgV2(); | 213 | ge::AscendString error_msg = ge::GEGetErrorMsgV2(); |
| 212 | std::string error_str(error_msg.GetString()); | 214 | std::string error_str(error_msg.GetString()); |
| 213 | - std::cout << "Error message: " << error_str << std::endl; | 215 | + LOG_PRINT("Error message: %s\n", error_str.c_str()); |
| 214 | ge::AscendString warning_msg = ge::GEGetWarningMsgV2(); | 216 | ge::AscendString warning_msg = ge::GEGetWarningMsgV2(); |
| 215 | std::string warning_str(warning_msg.GetString()); | 217 | std::string warning_str(warning_msg.GetString()); |
| 216 | - std::cout << "Warning message: " << warning_str << std::endl; | 218 | + LOG_PRINT("Warning message: %s\n", warning_str.c_str()); |
| 217 | - LOG_PRINT("%s - INFO - [XIR]: Add graph failed\n", GetTime().c_str()); | 219 | + LOG_PRINT("%s - ERROR - [REVERSE_SEQUENCE_GE_IR]: Add graph failed\n", GetTime().c_str()); |
| 218 | delete session; | 220 | delete session; |
| 219 | ge::GEFinalize(); | 221 | ge::GEFinalize(); |
| 220 | return false; | 222 | return false; |
| @@ -230,7 +232,7 @@ bool DumpAndRunGraph(ge::Session* session, Graph& graph, std::vector<ge::Tensor> | |||
| 230 | 232 | ||
| 231 | Status ret = session->RunGraph(graph_id, input, output); | 233 | Status ret = session->RunGraph(graph_id, input, output); |
| 232 | if (ret != SUCCESS) { | 234 | if (ret != SUCCESS) { |
| 233 | - LOG_PRINT("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str()); | 235 | + LOG_PRINT("%s - ERROR - [REVERSE_SEQUENCE_GE_IR]: Run graph failed\n", GetTime().c_str()); |
| 234 | delete session; | 236 | delete session; |
| 235 | ge::GEFinalize(); | 237 | ge::GEFinalize(); |
| 236 | return false; | 238 | return false; |
| @@ -243,17 +245,17 @@ void ProcessOutputData(std::vector<ge::Tensor>& output) | |||
| 243 | int output_num = output.size(); | 245 | int output_num = output.size(); |
| 244 | double epsilon = 1e-9; | 246 | double epsilon = 1e-9; |
| 245 | for (int i = 0; i < output_num; i++) { | 247 | for (int i = 0; i < output_num; i++) { |
| 246 | - std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl; | 248 | + LOG_PRINT("output %d dtype: %d\n", i, static_cast<int>(output[i].GetTensorDesc().GetDataType())); |
| 247 | double* output_data_i = (double*)output[i].GetData(); | 249 | double* output_data_i = (double*)output[i].GetData(); |
| 248 | int64_t output_size = output[i].GetTensorDesc().GetShape().GetShapeSize(); | 250 | int64_t output_size = output[i].GetTensorDesc().GetShape().GetShapeSize(); |
| 249 | double expect_out[9] = {1.0, 5.0, 9.0, 4.0, 2.0, 6.0, 7.0, 8.0, 3.0}; | 251 | double expect_out[9] = {1.0, 5.0, 9.0, 4.0, 2.0, 6.0, 7.0, 8.0, 3.0}; |
| 250 | for (int64_t j = 0; j < output_size; j++) { | 252 | for (int64_t j = 0; j < output_size; j++) { |
| 251 | if (std::abs(expect_out[j] - output_data_i[j]) > epsilon) { | 253 | if (std::abs(expect_out[j] - output_data_i[j]) > epsilon) { |
| 252 | - LOG_PRINT("ERROR - [XIR]: Precision is fail, please check. \n"); | 254 | + LOG_PRINT("ERROR - [REVERSE_SEQUENCE_GE_IR]: Precision check failed, please check. \n"); |
| 253 | return; | 255 | return; |
| 254 | } | 256 | } |
| 255 | } | 257 | } |
| 256 | - LOG_PRINT("INFO - [XIR]: Precison is ok. \n"); | 258 | + LOG_PRINT("INFO - [REVERSE_SEQUENCE_GE_IR]: Precision is ok. \n"); |
| 257 | } | 259 | } |
| 258 | } | 260 | } |
| 259 | 261 | ||
| @@ -261,7 +263,7 @@ int FinalizeRes() | |||
| 261 | { | 263 | { |
| 262 | Status ret = ge::GEFinalize(); | 264 | Status ret = ge::GEFinalize(); |
| 263 | if (ret != SUCCESS) { | 265 | if (ret != SUCCESS) { |
| 264 | - LOG_PRINT("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str()); | 266 | + LOG_PRINT("%s - ERROR - [REVERSE_SEQUENCE_GE_IR]: Finalize ir graph session failed\n", GetTime().c_str()); |
| 265 | return FAILED; | 267 | return FAILED; |
| 266 | } | 268 | } |
| 267 | 269 | ||
| @@ -123,7 +123,7 @@ void ReverseSequenceBASTiling::CalcSplitDimA() | |||
| 123 | ubFactorA_ = inDimALower; | 123 | ubFactorA_ = inDimALower; |
| 124 | 124 | ||
| 125 | if (ubFactorA_ <= 0) { | 125 | if (ubFactorA_ <= 0) { |
| 126 | - OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorA_ is %ld.", ubFactorA_); | 126 | + OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorA_ must be greater than 0, but got %ld.", ubFactorA_); |
| 127 | return; | 127 | return; |
| 128 | } | 128 | } |
| 129 | ubFactorB_ = 1; | 129 | ubFactorB_ = 1; |
| @@ -139,7 +139,7 @@ void ReverseSequenceBASTiling::CalcSplitDimS() | |||
| 139 | ubFactorS_ = inDimS; | 139 | ubFactorS_ = inDimS; |
| 140 | 140 | ||
| 141 | if (ubFactorS_ <= 0) { | 141 | if (ubFactorS_ <= 0) { |
| 142 | - OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorS_ is %ld.", ubFactorS_); | 142 | + OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorS_ must be greater than 0, but got %ld.", ubFactorS_); |
| 143 | return; | 143 | return; |
| 144 | } | 144 | } |
| 145 | ubFactorB_ = 1; | 145 | ubFactorB_ = 1; |
| @@ -313,4 +313,4 @@ ge::graphStatus ReverseSequenceBASTiling::GetShapeAttrsInfo() | |||
| 313 | 313 | ||
| 314 | REGISTER_TILING_TEMPLATE("ReverseSequence", ReverseSequenceBASTiling, 2); | 314 | REGISTER_TILING_TEMPLATE("ReverseSequence", ReverseSequenceBASTiling, 2); |
| 315 | 315 | ||
| 316 | -} // namespace optiling | 316 | +} // namespace optiling |
| @@ -85,7 +85,7 @@ ge::graphStatus ReverseSequenceBSTiling::InitializationVars() | |||
| 85 | { | 85 | { |
| 86 | oneBlockNum_ = Ops::Base::GetUbBlockSize(context_) / inputData_.xDtypeSize; | 86 | oneBlockNum_ = Ops::Base::GetUbBlockSize(context_) / inputData_.xDtypeSize; |
| 87 | OP_CHECK_IF((ubSize_ <= DCACHE_SIZE), | 87 | OP_CHECK_IF((ubSize_ <= DCACHE_SIZE), |
| 88 | - OP_LOGE(context_->GetNodeName(), "ub size:%lu less than Dcache Size:128k", ubSize_), | 88 | + OP_LOGE(context_->GetNodeName(), "ub size:%lu is less than Dcache size:128k", ubSize_), |
| 89 | return ge::GRAPH_FAILED); | 89 | return ge::GRAPH_FAILED); |
| 90 | ubSize_ = ubSize_ - DCACHE_SIZE; | 90 | ubSize_ = ubSize_ - DCACHE_SIZE; |
| 91 | availableUb_ = static_cast<int64_t>(ubSize_) / inputData_.xDtypeSize; | 91 | availableUb_ = static_cast<int64_t>(ubSize_) / inputData_.xDtypeSize; |
| @@ -129,7 +129,7 @@ void ReverseSequenceBSTiling::CalcSplitDimB() | |||
| 129 | ubFactorB_ = inDimBLower; | 129 | ubFactorB_ = inDimBLower; |
| 130 | ubFactorS_ = inputData_.inputDim[dimS_]; | 130 | ubFactorS_ = inputData_.inputDim[dimS_]; |
| 131 | if (ubFactorB_ <= 0) { | 131 | if (ubFactorB_ <= 0) { |
| 132 | - OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorB_ is %ld.", ubFactorB_); | 132 | + OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorB_ must be greater than 0, but got %ld.", ubFactorB_); |
| 133 | return; | 133 | return; |
| 134 | } | 134 | } |
| 135 | 135 | ||
| @@ -143,7 +143,7 @@ void ReverseSequenceBSTiling::CalcSplitDimS() | |||
| 143 | { | 143 | { |
| 144 | ubFactorS_ = std::min(inputData_.inputDim[dimS_], (availableUb_ / DIGIT_FOUR)); | 144 | ubFactorS_ = std::min(inputData_.inputDim[dimS_], (availableUb_ / DIGIT_FOUR)); |
| 145 | if (ubFactorS_ <= 0) { | 145 | if (ubFactorS_ <= 0) { |
| 146 | - OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorS_ is %ld.", ubFactorS_); | 146 | + OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorS_ must be greater than 0, but got %ld.", ubFactorS_); |
| 147 | return; | 147 | return; |
| 148 | } | 148 | } |
| 149 | 149 | ||
| @@ -220,7 +220,7 @@ void ReverseSequenceBSTiling::DoUBTiling() | |||
| 220 | ge::graphStatus ReverseSequenceBSTiling::DoOpTiling() | 220 | ge::graphStatus ReverseSequenceBSTiling::DoOpTiling() |
| 221 | { | 221 | { |
| 222 | OP_CHECK_IF(InitializationVars() != ge::GRAPH_SUCCESS, | 222 | OP_CHECK_IF(InitializationVars() != ge::GRAPH_SUCCESS, |
| 223 | - OP_LOGE(context_->GetNodeName(), "ub size:%lu less than Dcache Size:128k", ubSize_), | 223 | + OP_LOGE(context_->GetNodeName(), "ub size:%lu is less than Dcache size:128k", ubSize_), |
| 224 | return ge::GRAPH_FAILED); | 224 | return ge::GRAPH_FAILED); |
| 225 | DoUBTiling(); | 225 | DoUBTiling(); |
| 226 | DoBlockTiling(); | 226 | DoBlockTiling(); |
| @@ -318,4 +318,4 @@ ge::graphStatus ReverseSequenceBSTiling::GetShapeAttrsInfo() | |||
| 318 | 318 | ||
| 319 | REGISTER_TILING_TEMPLATE("ReverseSequence", ReverseSequenceBSTiling, 3); | 319 | REGISTER_TILING_TEMPLATE("ReverseSequence", ReverseSequenceBSTiling, 3); |
| 320 | 320 | ||
| 321 | -} // namespace optiling | 321 | +} // namespace optiling |
| @@ -143,7 +143,7 @@ void ReverseSequenceBSATiling::CalcSplitDimS() | |||
| 143 | ubFactorS_ = MAX_INPUT_ELEMENTS / inputData_.inputDim[DIM_A]; | 143 | ubFactorS_ = MAX_INPUT_ELEMENTS / inputData_.inputDim[DIM_A]; |
| 144 | } | 144 | } |
| 145 | if (ubFactorS_ <= 0) { | 145 | if (ubFactorS_ <= 0) { |
| 146 | - OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorS_ is %ld.", ubFactorS_); | 146 | + OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorS_ must be greater than 0, but got %ld.", ubFactorS_); |
| 147 | return; | 147 | return; |
| 148 | } | 148 | } |
| 149 | ubFactorB_ = 1; | 149 | ubFactorB_ = 1; |
| @@ -162,7 +162,7 @@ void ReverseSequenceBSATiling::CalcSplitDimA() | |||
| 162 | ubFactorA_ = MAX_INPUT_ELEMENTS; | 162 | ubFactorA_ = MAX_INPUT_ELEMENTS; |
| 163 | } | 163 | } |
| 164 | if (ubFactorA_ <= 0) { | 164 | if (ubFactorA_ <= 0) { |
| 165 | - OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorA_ is %ld.", ubFactorA_); | 165 | + OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorA_ must be greater than 0, but got %ld.", ubFactorA_); |
| 166 | return; | 166 | return; |
| 167 | } | 167 | } |
| 168 | ubFactorB_ = 1; | 168 | ubFactorB_ = 1; |
| @@ -348,4 +348,4 @@ ge::graphStatus ReverseSequenceBSATiling::GetShapeAttrsInfo() | |||
| 348 | 348 | ||
| 349 | REGISTER_TILING_TEMPLATE("ReverseSequence", ReverseSequenceBSATiling, 1); | 349 | REGISTER_TILING_TEMPLATE("ReverseSequence", ReverseSequenceBSATiling, 1); |
| 350 | 350 | ||
| 351 | -} // namespace optiling | 351 | +} // namespace optiling |
| @@ -139,7 +139,7 @@ void ReverseSequenceSBACommonTiling::CalcSplitDimA() | |||
| 139 | ubFactorA_ = inDimA; | 139 | ubFactorA_ = inDimA; |
| 140 | 140 | ||
| 141 | if (ubFactorA_ <= 0) { | 141 | if (ubFactorA_ <= 0) { |
| 142 | - OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorA_ is %ld.", ubFactorA_); | 142 | + OP_LOGE(context_->GetNodeName(), "ReverseSequence ubFactorA_ must be greater than 0, but got %ld.", ubFactorA_); |
| 143 | return; | 143 | return; |
| 144 | } | 144 | } |
| 145 | ubFactorB_ = 1; | 145 | ubFactorB_ = 1; |
| @@ -357,4 +357,4 @@ ge::graphStatus ReverseSequenceSBACommonTiling::GetShapeAttrsInfo() | |||
| 357 | OP_LOGD("ReverseSequenceSBACommonTiling::GetShapeAttrsInfo begin"); | 357 | OP_LOGD("ReverseSequenceSBACommonTiling::GetShapeAttrsInfo begin"); |
| 358 | return GetReverseSequenceShapeAttrsInfo(context_, inputData_); | 358 | return GetReverseSequenceShapeAttrsInfo(context_, inputData_); |
| 359 | } | 359 | } |
| 360 | -} // namespace optiling | 360 | +} // namespace optiling |
| @@ -60,7 +60,7 @@ ge::graphStatus ReverseSequenceSimtTiling::DoReverseSequenceSimtTiling() | |||
| 60 | 60 | ||
| 61 | int64_t oneBlockNum = Ops::Base::GetUbBlockSize(context_) / inputData.xDtypeSize; | 61 | int64_t oneBlockNum = Ops::Base::GetUbBlockSize(context_) / inputData.xDtypeSize; |
| 62 | OP_CHECK_IF((ubSize <= DCACHE_SIZE), | 62 | OP_CHECK_IF((ubSize <= DCACHE_SIZE), |
| 63 | - OP_LOGE(context_->GetNodeName(), "ub size:%lu less than Dcache Size:128k", ubSize), | 63 | + OP_LOGE(context_->GetNodeName(), "ub size:%lu is less than Dcache size:128k", ubSize), |
| 64 | return ge::GRAPH_FAILED); | 64 | return ge::GRAPH_FAILED); |
| 65 | ubSize = ubSize - DCACHE_SIZE; | 65 | ubSize = ubSize - DCACHE_SIZE; |
| 66 | ubSize = ubSize - RESERVE_BUF_SIZE; | 66 | ubSize = ubSize - RESERVE_BUF_SIZE; |
| @@ -177,4 +177,4 @@ ge::graphStatus ReverseSequenceSimtTiling::GetShapeAttrsInfo() | |||
| 177 | 177 | ||
| 178 | REGISTER_TILING_TEMPLATE("ReverseSequence", ReverseSequenceSimtTiling, 9); | 178 | REGISTER_TILING_TEMPLATE("ReverseSequence", ReverseSequenceSimtTiling, 9); |
| 179 | 179 | ||
| 180 | -} // namespace optiling | 180 | +} // namespace optiling |
| @@ -1,12 +1,11 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * This program is free software, you can redistribute it and/or modify. | ||
| 3 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | * Copyright (c) 2025 Huawei Technologies Co., Ltd. |
| 4 | - * This file is a part of the CANN Open Software. | 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 5 | - * Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 6 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| 7 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, |
| 8 | - * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of | 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. |
| 9 | - * the software repository for the full text of the License. | 8 | + * See LICENSE in the root of the software repository for the full text of the License. |
| 10 | */ | 9 | */ |
| 11 | 10 | ||
| 12 | /*! | 11 | /*! |
| @@ -36,9 +35,9 @@ using namespace ge; | |||
| 36 | 35 | ||
| 37 | class ReverseSequenceTiling : public testing::Test { | 36 | class ReverseSequenceTiling : public testing::Test { |
| 38 | protected: | 37 | protected: |
| 39 | - static void SetUpTestCase() { std::cout << "ReverseSequenceTiling SetUp" << std::endl; } | 38 | + static void SetUpTestCase() {} |
| 40 | 39 | ||
| 41 | - static void TearDownTestCase() { std::cout << "ReverseSequenceTiling TearDown" << std::endl; } | 40 | + static void TearDownTestCase() {} |
| 42 | }; | 41 | }; |
| 43 | 42 | ||
| 44 | template <typename T> | 43 | template <typename T> |
| @@ -237,4 +236,4 @@ TEST_F(ReverseSequenceTiling, test_tiling_ascendc_int8) // SAB | |||
| 237 | int64_t seqDim = 0; | 236 | int64_t seqDim = 0; |
| 238 | 237 | ||
| 239 | ExecuteTestCase(ge::DT_INT8, ge::DT_INT64, shape1, shape2, batchDim, seqDim, tilingKeyValue, expectTilingData, 0); | 238 | ExecuteTestCase(ge::DT_INT8, ge::DT_INT64, shape1, shape2, batchDim, seqDim, tilingKeyValue, expectTilingData, 0); |
| 240 | -} | 239 | +} |
| @@ -22,9 +22,9 @@ | |||
| 22 | 22 | ||
| 23 | class ReverseSequenceInfershapeTest : public testing::Test { | 23 | class ReverseSequenceInfershapeTest : public testing::Test { |
| 24 | protected: | 24 | protected: |
| 25 | - static void SetUpTestCase() { std::cout << "ReverseSequenceInfershapeTest SetUp" << std::endl; } | 25 | + static void SetUpTestCase() {} |
| 26 | 26 | ||
| 27 | - static void TearDownTestCase() { std::cout << "ReverseSequenceInfershapeTest TearDown" << std::endl; } | 27 | + static void TearDownTestCase() {} |
| 28 | }; | 28 | }; |
| 29 | 29 | ||
| 30 | TEST_F(ReverseSequenceInfershapeTest, reverse_sequence_infershape_float16_2d) | 30 | TEST_F(ReverseSequenceInfershapeTest, reverse_sequence_infershape_float16_2d) |
| @@ -60,7 +60,7 @@ static bool IsShapeEquals(const aclTensor* tensor1, const aclTensor* tensor2) | |||
| 60 | { | 60 | { |
| 61 | auto shape1 = tensor1->GetViewShape(); | 61 | auto shape1 = tensor1->GetViewShape(); |
| 62 | auto shape2 = tensor2->GetViewShape(); | 62 | auto shape2 = tensor2->GetViewShape(); |
| 63 | - OP_LOGE(ACLNN_ERR_PARAM_INVALID, "x shape %ld, y shape %ld", shape1.GetShapeSize(), shape2.GetShapeSize()); | 63 | + OP_LOGD("QuantMax", "x shape %ld, y shape %ld", shape1.GetShapeSize(), shape2.GetShapeSize()); |
| 64 | if (shape1.GetDimNum() != shape2.GetDimNum()) { | 64 | if (shape1.GetDimNum() != shape2.GetDimNum()) { |
| 65 | return false; | 65 | return false; |
| 66 | } | 66 | } |
| @@ -2,7 +2,7 @@ | |||
| 2 | * Copyright (c) 2026 Huawei Technologies Co., Ltd. | 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 | 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"). | 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 5 | - * Please refer to the License for details. You may not use 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, | 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. | 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. | 8 | * See LICENSE in the root of the software repository for the full text of the License. |
| @@ -15,9 +15,9 @@ | |||
| 15 | 15 | ||
| 16 | class l2QuantMaxTest : public testing::Test { | 16 | class l2QuantMaxTest : public testing::Test { |
| 17 | protected: | 17 | protected: |
| 18 | - static void SetUpTestCase() { std::cout << "l2QuantMaxTest SetUp" << std::endl; } | 18 | + static void SetUpTestCase() {} |
| 19 | 19 | ||
| 20 | - static void TearDownTestCase() { std::cout << "l2QuantMaxTest TearDown" << std::endl; } | 20 | + static void TearDownTestCase() {} |
| 21 | }; | 21 | }; |
| 22 | 22 | ||
| 23 | // Test null input x | 23 | // Test null input x |
| @@ -702,4 +702,4 @@ TEST_F(l2QuantMaxTest, l2_quant_max_0dim_tensor) | |||
| 702 | 702 | ||
| 703 | aclnnStatus ret = aclnnQuantMaxGetWorkspaceSize(x, scale, roundMode, dstType, y, amax, &workspaceSize, &executor); | 703 | aclnnStatus ret = aclnnQuantMaxGetWorkspaceSize(x, scale, roundMode, dstType, y, amax, &workspaceSize, &executor); |
| 704 | EXPECT_EQ(ret, ACLNN_ERR_PARAM_INVALID); | 704 | EXPECT_EQ(ret, ACLNN_ERR_PARAM_INVALID); |
| 705 | -} | 705 | +} |
| @@ -2,7 +2,7 @@ | |||
| 2 | * Copyright (c) 2026 Huawei Technologies Co., Ltd. | 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 | 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"). | 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 5 | - * Please refer to the License for details. You may not use 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, | 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. | 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. | 8 | * See LICENSE in the root of the software repository for the full text of the License. |
| @@ -30,9 +30,9 @@ using namespace ge; | |||
| 30 | 30 | ||
| 31 | class QuantMaxTilingTest : public testing::Test { | 31 | class QuantMaxTilingTest : public testing::Test { |
| 32 | protected: | 32 | protected: |
| 33 | - static void SetUpTestCase() { std::cout << "QuantMaxTilingTest SetUp" << std::endl; } | 33 | + static void SetUpTestCase() {} |
| 34 | 34 | ||
| 35 | - static void TearDownTestCase() { std::cout << "QuantMaxTilingTest TearDown" << std::endl; } | 35 | + static void TearDownTestCase() {} |
| 36 | }; | 36 | }; |
| 37 | 37 | ||
| 38 | // Test that QuantMax OpImpl is registered correctly | 38 | // Test that QuantMax OpImpl is registered correctly |
| @@ -51,8 +51,6 @@ TEST_F(QuantMaxTilingTest, quant_max_op_impl_registered) | |||
| 51 | // Check if tiling_parse function is registered | 51 | // Check if tiling_parse function is registered |
| 52 | auto tiling_parse_func = op_impl->tiling_parse; | 52 | auto tiling_parse_func = op_impl->tiling_parse; |
| 53 | ASSERT_NE(tiling_parse_func, nullptr); | 53 | ASSERT_NE(tiling_parse_func, nullptr); |
| 54 | - | ||
| 55 | - std::cout << "QuantMax OpImpl registered successfully" << std::endl; | ||
| 56 | } | 54 | } |
| 57 | 55 | ||
| 58 | // DISABLED: These tests require real hardware platform info (coreNum, ubSize) | 56 | // DISABLED: These tests require real hardware platform info (coreNum, ubSize) |
| @@ -160,8 +158,6 @@ TEST_F(QuantMaxTilingTest, DISABLED_quant_max_tiling_fp32_success) | |||
| 160 | 158 | ||
| 161 | // Actually call tiling function and verify result | 159 | // Actually call tiling function and verify result |
| 162 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); | 160 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); |
| 163 | - | ||
| 164 | - std::cout << "QuantMax tiling executed successfully for FP32 input" << std::endl; | ||
| 165 | } | 161 | } |
| 166 | 162 | ||
| 167 | // DISABLED: Requires real hardware platform info | 163 | // DISABLED: Requires real hardware platform info |
| @@ -225,8 +221,6 @@ TEST_F(QuantMaxTilingTest, DISABLED_quant_max_tiling_fp16_hifloat8_success) | |||
| 225 | 221 | ||
| 226 | // Actually call tiling function and verify result | 222 | // Actually call tiling function and verify result |
| 227 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); | 223 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); |
| 228 | - | ||
| 229 | - std::cout << "QuantMax tiling executed successfully for FP16 to HIFLOAT8" << std::endl; | ||
| 230 | } | 224 | } |
| 231 | 225 | ||
| 232 | // DISABLED: Requires real hardware platform info | 226 | // DISABLED: Requires real hardware platform info |
| @@ -290,8 +284,6 @@ TEST_F(QuantMaxTilingTest, DISABLED_quant_max_tiling_bf16_success) | |||
| 290 | 284 | ||
| 291 | // Actually call tiling function and verify result | 285 | // Actually call tiling function and verify result |
| 292 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); | 286 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); |
| 293 | - | ||
| 294 | - std::cout << "QuantMax tiling executed successfully for BF16 input" << std::endl; | ||
| 295 | } | 287 | } |
| 296 | 288 | ||
| 297 | // DISABLED: Requires real hardware platform info | 289 | // DISABLED: Requires real hardware platform info |
| @@ -355,8 +347,6 @@ TEST_F(QuantMaxTilingTest, DISABLED_quant_max_tiling_fp8_e4m3fn_success) | |||
| 355 | 347 | ||
| 356 | // Actually call tiling function and verify result | 348 | // Actually call tiling function and verify result |
| 357 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); | 349 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); |
| 358 | - | ||
| 359 | - std::cout << "QuantMax tiling executed successfully for FLOAT8_E4M3FN output" << std::endl; | ||
| 360 | } | 350 | } |
| 361 | 351 | ||
| 362 | // DISABLED: Requires real hardware platform info | 352 | // DISABLED: Requires real hardware platform info |
| @@ -420,8 +410,6 @@ TEST_F(QuantMaxTilingTest, DISABLED_quant_max_tiling_hifloat8_hybrid_success) | |||
| 420 | 410 | ||
| 421 | // Actually call tiling function and verify result | 411 | // Actually call tiling function and verify result |
| 422 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); | 412 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); |
| 423 | - | ||
| 424 | - std::cout << "QuantMax tiling executed successfully for HIFLOAT8 with Hybrid mode" << std::endl; | ||
| 425 | } | 413 | } |
| 426 | 414 | ||
| 427 | // DISABLED: Requires real hardware platform info | 415 | // DISABLED: Requires real hardware platform info |
| @@ -485,8 +473,6 @@ TEST_F(QuantMaxTilingTest, DISABLED_quant_max_tiling_large_shape_success) | |||
| 485 | 473 | ||
| 486 | // Actually call tiling function and verify result | 474 | // Actually call tiling function and verify result |
| 487 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); | 475 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS); |
| 488 | - | ||
| 489 | - std::cout << "QuantMax tiling executed successfully for large shape" << std::endl; | ||
| 490 | } | 476 | } |
| 491 | 477 | ||
| 492 | // Test CheckDtype failure: x dtype not in FLOAT/FLOAT16/BF16 | 478 | // Test CheckDtype failure: x dtype not in FLOAT/FLOAT16/BF16 |
| @@ -541,8 +527,6 @@ TEST_F(QuantMaxTilingTest, quant_max_tiling_invalid_x_dtype) | |||
| 541 | ASSERT_NE(tiling_func, nullptr); | 527 | ASSERT_NE(tiling_func, nullptr); |
| 542 | 528 | ||
| 543 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); | 529 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); |
| 544 | - | ||
| 545 | - std::cout << "QuantMax tiling check dtype failure for INT32 input" << std::endl; | ||
| 546 | } | 530 | } |
| 547 | 531 | ||
| 548 | // Test CheckShape failure: x dim > 8 | 532 | // Test CheckShape failure: x dim > 8 |
| @@ -598,8 +582,6 @@ TEST_F(QuantMaxTilingTest, quant_max_tiling_x_dim_exceed_max) | |||
| 598 | ASSERT_NE(tiling_func, nullptr); | 582 | ASSERT_NE(tiling_func, nullptr); |
| 599 | 583 | ||
| 600 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); | 584 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); |
| 601 | - | ||
| 602 | - std::cout << "QuantMax tiling check shape failure for 9-dim input" << std::endl; | ||
| 603 | } | 585 | } |
| 604 | 586 | ||
| 605 | // Test CheckShape failure: scale dim != 1 | 587 | // Test CheckShape failure: scale dim != 1 |
| @@ -654,8 +636,6 @@ TEST_F(QuantMaxTilingTest, quant_max_tiling_invalid_scale_dim) | |||
| 654 | ASSERT_NE(tiling_func, nullptr); | 636 | ASSERT_NE(tiling_func, nullptr); |
| 655 | 637 | ||
| 656 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); | 638 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); |
| 657 | - | ||
| 658 | - std::cout << "QuantMax tiling check shape failure for invalid scale dim" << std::endl; | ||
| 659 | } | 639 | } |
| 660 | 640 | ||
| 661 | // Test CheckShape failure: amax dim != 1 | 641 | // Test CheckShape failure: amax dim != 1 |
| @@ -710,8 +690,6 @@ TEST_F(QuantMaxTilingTest, quant_max_tiling_invalid_amax_dim) | |||
| 710 | ASSERT_NE(tiling_func, nullptr); | 690 | ASSERT_NE(tiling_func, nullptr); |
| 711 | 691 | ||
| 712 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); | 692 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); |
| 713 | - | ||
| 714 | - std::cout << "QuantMax tiling check shape failure for invalid amax dim" << std::endl; | ||
| 715 | } | 693 | } |
| 716 | 694 | ||
| 717 | // Test CheckAttrs failure: dstType not in 34/35/36 | 695 | // Test CheckAttrs failure: dstType not in 34/35/36 |
| @@ -768,8 +746,6 @@ TEST_F(QuantMaxTilingTest, quant_max_tiling_invalid_dst_type) | |||
| 768 | ASSERT_NE(tiling_func, nullptr); | 746 | ASSERT_NE(tiling_func, nullptr); |
| 769 | 747 | ||
| 770 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); | 748 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); |
| 771 | - | ||
| 772 | - std::cout << "QuantMax tiling check attrs failure for invalid dstType" << std::endl; | ||
| 773 | } | 749 | } |
| 774 | 750 | ||
| 775 | // Test CheckAttrs failure: roundMode mismatch with dstType | 751 | // Test CheckAttrs failure: roundMode mismatch with dstType |
| @@ -825,6 +801,4 @@ TEST_F(QuantMaxTilingTest, quant_max_tiling_roundmode_mismatch) | |||
| 825 | ASSERT_NE(tiling_func, nullptr); | 801 | ASSERT_NE(tiling_func, nullptr); |
| 826 | 802 | ||
| 827 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); | 803 | EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED); |
| 828 | - | 804 | +} |
| 829 | - std::cout << "QuantMax tiling check attrs failure for roundMode mismatch" << std::endl; | ||
| 830 | -} | ||
| @@ -1,8 +1,8 @@ | |||
| 1 | /** | 1 | /** |
| 2 | * Copyright (c) 2026 Huawei Technologies Co., Ltd. | 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 | 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 Version 2.0 (the "License"). | 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 5 | - * Please refer to the License for details. You may not use 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, | 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. | 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. | 8 | * See LICENSE in the root of the software repository for the full text of the License. |
| @@ -28,9 +28,9 @@ using namespace ge; | |||
| 28 | 28 | ||
| 29 | class QuantMaxInferShapeTest : public testing::Test { | 29 | class QuantMaxInferShapeTest : public testing::Test { |
| 30 | protected: | 30 | protected: |
| 31 | - static void SetUpTestCase() { std::cout << "QuantMaxInferShapeTest SetUp" << std::endl; } | 31 | + static void SetUpTestCase() {} |
| 32 | 32 | ||
| 33 | - static void TearDownTestCase() { std::cout << "QuantMaxInferShapeTest TearDown" << std::endl; } | 33 | + static void TearDownTestCase() {} |
| 34 | }; | 34 | }; |
| 35 | 35 | ||
| 36 | TEST_F(QuantMaxInferShapeTest, quant_max_infershape_basic) | 36 | TEST_F(QuantMaxInferShapeTest, quant_max_infershape_basic) |
| @@ -489,4 +489,4 @@ TEST_F(QuantMaxInferShapeTest, quant_max_infershape_empty_tensor) | |||
| 489 | ASSERT_NE(output_y_shape, nullptr); | 489 | ASSERT_NE(output_y_shape, nullptr); |
| 490 | EXPECT_EQ(output_y_shape->GetDimNum(), 1); | 490 | EXPECT_EQ(output_y_shape->GetDimNum(), 1); |
| 491 | EXPECT_EQ(output_y_shape->GetDim(0), 0); | 491 | EXPECT_EQ(output_y_shape->GetDim(0), 0); |
| 492 | -} | 492 | +} |
| @@ -2,7 +2,7 @@ | |||
| 2 | * Copyright (c) 2026 Huawei Technologies Co., Ltd. | 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 | 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"). | 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 5 | - * Please refer to the License for details. You may not use 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, | 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. | 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. | 8 | * See LICENSE in the root of the software repository for the full text of the License. |
| @@ -39,8 +39,8 @@ extern "C" __global__ __aicore__ void quant_max(GM_ADDR x, GM_ADDR scale, GM_ADD | |||
| 39 | 39 | ||
| 40 | class quant_max_test : public testing::Test { | 40 | class quant_max_test : public testing::Test { |
| 41 | protected: | 41 | protected: |
| 42 | - static void SetUpTestCase() { std::cout << "quant_max_test SetUp" << std::endl; } | 42 | + static void SetUpTestCase() {} |
| 43 | - static void TearDownTestCase() { std::cout << "quant_max_test TearDown" << std::endl; } | 43 | + static void TearDownTestCase() {} |
| 44 | }; | 44 | }; |
| 45 | 45 | ||
| 46 | // Helper function to generate test data | 46 | // Helper function to generate test data |
| @@ -426,4 +426,4 @@ TEST_F(quant_max_test, test_bf16_to_fp8_e5m2) | |||
| 426 | AscendC::GmFree((void*)amax); | 426 | AscendC::GmFree((void*)amax); |
| 427 | AscendC::GmFree((void*)workspace); | 427 | AscendC::GmFree((void*)workspace); |
| 428 | AscendC::GmFree((void*)tiling); | 428 | AscendC::GmFree((void*)tiling); |
| 429 | -} | 429 | +} |