Merged
resize类算子tiling侧增加单轴是否超INT32_MAX的判断 #1133
wkx12138Add on Jul 16
resize类算子tiling侧增加单轴是否超INT32_MAX的判断 #1133
Merged
Merged fromwkx12138/ops-cvresizeintocann/ops-cv9.1.0
Total: 12 files changed+129-7
| @@ -69,6 +69,20 @@ | |||
| 69 | <td>FLOAT16、FLOAT32、BFLOAT16</td> | 69 | <td>FLOAT16、FLOAT32、BFLOAT16</td> |
| 70 | <td>NCHW、NHWC</td> | 70 | <td>NCHW、NHWC</td> |
| 71 | </tr> | 71 | </tr> |
| 72 | + <tr> | ||
| 73 | + <td>align_corners</td> | ||
| 74 | + <td>属性</td> | ||
| 75 | + <td>是否对齐输入与输出张量四个角像素的中心以保留角点像素值。默认值：false。</td> | ||
| 76 | + <td>Bool</td> | ||
| 77 | + <td>-</td> | ||
| 78 | + </tr> | ||
| 79 | + <tr> | ||
| 80 | + <td>scales</td> | ||
| 81 | + <td>属性</td> | ||
| 82 | + <td>空间维度（高、宽）的缩放系数。默认值：{0.0, 0.0}。</td> | ||
| 83 | + <td>ListFloat</td> | ||
| 84 | + <td>-</td> | ||
| 85 | + </tr> | ||
| 72 | </tbody> | 86 | </tbody> |
| 73 | </table> | 87 | </table> |
| 74 | 88 | ||
| @@ -299,7 +299,7 @@ ge::graphStatus ResizeBicubicV2Tiling::CheckParams() | |||
| 299 | return ge::GRAPH_FAILED); | 299 | return ge::GRAPH_FAILED); |
| 300 | 300 | ||
| 301 | isNchw_ = DIM_1; | 301 | isNchw_ = DIM_1; |
| 302 | - if ((dataFormat == ge::FORMAT_NHWC) && (lenDesH_ != lenSrcH_ && lenDesW_ != lenSrcW_)) { | 302 | + if ((dataFormat == ge::FORMAT_NHWC) && (lenDesH_ != lenSrcH_ || lenDesW_ != lenSrcW_)) { |
| 303 | isNchw_ = DIM_0; | 303 | isNchw_ = DIM_0; |
| 304 | } | 304 | } |
| 305 | 305 | ||
| @@ -483,6 +483,9 @@ void ResizeBicubicV2Tiling::DoTilingKeyPostProcess() | |||
| 483 | OP_LOGI(context_->GetNodeName(), "ySize or xSize is too large"); | 483 | OP_LOGI(context_->GetNodeName(), "ySize or xSize is too large"); |
| 484 | isInt32_ = DIM_0; | 484 | isInt32_ = DIM_0; |
| 485 | } | 485 | } |
| 486 | + if (lenSrcH_ > INT32_MAX || lenSrcW_ > INT32_MAX || lenDesH_ > INT32_MAX || lenDesW_ > INT32_MAX) { | ||
| 487 | + isInt32_ = DIM_0; | ||
| 488 | + } | ||
| 486 | if (schId_ == DIM_5 || schId_ == DIM_6) { // simd scenario ignores isInt32_ and isNchw_ params | 489 | if (schId_ == DIM_5 || schId_ == DIM_6) { // simd scenario ignores isInt32_ and isNchw_ params |
| 487 | isInt32_ = DIM_0; | 490 | isInt32_ = DIM_0; |
| 488 | isNchw_ = DIM_0; | 491 | isNchw_ = DIM_0; |
| @@ -34,10 +34,12 @@ __global__ __aicore__ void resize_bicubic_v2(GM_ADDR x, GM_ADDR size, GM_ADDR y, | |||
| 34 | ResizeBicubicV2::ResizeBicubicV2AllCopy<DTYPE_X, isHalfPixel, schId, uint32_t, int32_t> op; | 34 | ResizeBicubicV2::ResizeBicubicV2AllCopy<DTYPE_X, isHalfPixel, schId, uint32_t, int32_t> op; |
| 35 | op.Init(x, size, y, &pipe, &tilingData); | 35 | op.Init(x, size, y, &pipe, &tilingData); |
| 36 | op.Process(); | 36 | op.Process(); |
| 37 | + return; | ||
| 37 | } else if (schId == RESIZE_BICUBIC_V2_TPL_SCH_MODE_6) { | 38 | } else if (schId == RESIZE_BICUBIC_V2_TPL_SCH_MODE_6) { |
| 38 | ResizeBicubicV2::ResizeBicubicV2PointCopy<DTYPE_X, isHalfPixel, schId, uint32_t, int32_t> op; | 39 | ResizeBicubicV2::ResizeBicubicV2PointCopy<DTYPE_X, isHalfPixel, schId, uint32_t, int32_t> op; |
| 39 | op.Init(x, size, y, &pipe, &tilingData); | 40 | op.Init(x, size, y, &pipe, &tilingData); |
| 40 | op.Process(); | 41 | op.Process(); |
| 42 | + return; | ||
| 41 | } | 43 | } |
| 42 | if constexpr (isNchw == 1) { | 44 | if constexpr (isNchw == 1) { |
| 43 | if constexpr (isInt32 == 1) { | 45 | if constexpr (isInt32 == 1) { |
| @@ -49,6 +51,7 @@ __global__ __aicore__ void resize_bicubic_v2(GM_ADDR x, GM_ADDR size, GM_ADDR y, | |||
| 49 | op.Init(x, size, y, &tilingData); | 51 | op.Init(x, size, y, &tilingData); |
| 50 | op.Process(); | 52 | op.Process(); |
| 51 | } | 53 | } |
| 54 | + return; | ||
| 52 | } else { | 55 | } else { |
| 53 | if constexpr (isInt32 == 1) { | 56 | if constexpr (isInt32 == 1) { |
| 54 | ResizeBicubicV2::ResizeBicubicV2SimtNHWC<DTYPE_X, isHalfPixel, schId, uint32_t, int32_t> op; | 57 | ResizeBicubicV2::ResizeBicubicV2SimtNHWC<DTYPE_X, isHalfPixel, schId, uint32_t, int32_t> op; |
| @@ -59,5 +62,6 @@ __global__ __aicore__ void resize_bicubic_v2(GM_ADDR x, GM_ADDR size, GM_ADDR y, | |||
| 59 | op.Init(x, size, y, &tilingData); | 62 | op.Init(x, size, y, &tilingData); |
| 60 | op.Process(); | 63 | op.Process(); |
| 61 | } | 64 | } |
| 65 | + return; | ||
| 62 | } | 66 | } |
| 63 | } | 67 | } |
| @@ -77,4 +77,88 @@ TEST_F(ResizeBicubicV2TilingTest, resize_bicubic_v2_tiling_02) | |||
| 77 | std::vector<size_t> expectWorkspaces = {16777216}; | 77 | std::vector<size_t> expectWorkspaces = {16777216}; |
| 78 | 78 | ||
| 79 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | 79 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); |
| 80 | +} | ||
| 81 | + | ||
| 82 | +// tilingKey 位布局(见 resize_bicubic_v2_tiling_key.h 的 ASCENDC_TPL_UINT_DECL 顺序,各 8bit): | ||
| 83 | +// byte0=schId, byte1=isInt32, byte2=isHalfPixel, byte3=isNchw | ||
| 84 | +static inline int64_t GetIsNchwFromTilingKey(int64_t tilingKey) | ||
| 85 | +{ | ||
| 86 | + return (tilingKey >> 24) & 0xFF; | ||
| 87 | +} | ||
| 88 | +static inline int64_t GetIsInt32FromTilingKey(int64_t tilingKey) | ||
| 89 | +{ | ||
| 90 | + return (tilingKey >> 8) & 0xFF; | ||
| 91 | +} | ||
| 92 | + | ||
| 93 | +// Test 07 [回归]: NHWC 单空间维变化(仅 H 变、W 不变)必须走 NHWC 布局核(isNchw=0)。 | ||
| 94 | +// 缺陷(修复前 tiling arch35 L297 用 &&):NHWC 仅当 H、W 都变才置 isNchw_=0,单维变时 isNchw_ 停在 1, | ||
| 95 | +// 用 NCHW 布局核处理 channel-last 数据 → 真机 ~99.8% 元素错(见 op-quality-review DEFECT 报告)。 | ||
| 96 | +// 修复(&& -> ||):NHWC 任一空间维变化即置 isNchw_=0。此用例严格断言 isNchw 字节=0,修复前必失败、修复后通过。 | ||
| 97 | +TEST_F(ResizeBicubicV2TilingTest, resize_bicubic_v2_tiling_07_nhwc_single_dim_h) | ||
| 98 | +{ | ||
| 99 | + gert::StorageShape inputXShape = {{1, 16, 16, 64}, {1, 16, 16, 64}}; | ||
| 100 | + gert::StorageShape inputSizeShape = {{1, 2}, {1, 2}}; | ||
| 101 | + gert::StorageShape outputShape = {{1, 32, 16, 64}, {1, 32, 16, 64}}; // H:16->32, W:16 不变 | ||
| 102 | + int size_value[2] = {32, 16}; | ||
| 103 | + | ||
| 104 | + ResizeBicubicV2CompileInfo compileInfo = {64, 200704}; | ||
| 105 | + | ||
| 106 | + gert::TilingContextPara tilingContextPara( | ||
| 107 | + "ResizeBicubicV2", | ||
| 108 | + {{inputXShape, ge::DT_FLOAT, ge::FORMAT_NHWC}, {inputSizeShape, ge::DT_INT32, ge::FORMAT_ND, true, size_value}}, | ||
| 109 | + {{outputShape, ge::DT_FLOAT, ge::FORMAT_NHWC}}, | ||
| 110 | + {gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 111 | + gert::TilingContextPara::OpAttr("scales", Ops::Cv::AnyValue::CreateFrom<vector<float>>({0.0f, 0.0f}))}, | ||
| 112 | + &compileInfo); | ||
| 113 | + TilingInfo tilingInfo; | ||
| 114 | + EXPECT_TRUE(ExecuteTiling(tilingContextPara, tilingInfo)); | ||
| 115 | + EXPECT_EQ(GetIsNchwFromTilingKey(tilingInfo.tilingKey), 0); // NHWC 必走 NHWC 布局核 | ||
| 116 | +} | ||
| 117 | + | ||
| 118 | +// Test 08 [回归]: NHWC 单空间维变化(仅 W 变、H 不变)同样必须 isNchw=0。 | ||
| 119 | +TEST_F(ResizeBicubicV2TilingTest, resize_bicubic_v2_tiling_08_nhwc_single_dim_w) | ||
| 120 | +{ | ||
| 121 | + gert::StorageShape inputXShape = {{1, 16, 16, 64}, {1, 16, 16, 64}}; | ||
| 122 | + gert::StorageShape inputSizeShape = {{1, 2}, {1, 2}}; | ||
| 123 | + gert::StorageShape outputShape = {{1, 16, 32, 64}, {1, 16, 32, 64}}; // W:16->32, H:16 不变 | ||
| 124 | + int size_value[2] = {16, 32}; | ||
| 125 | + | ||
| 126 | + ResizeBicubicV2CompileInfo compileInfo = {64, 200704}; | ||
| 127 | + | ||
| 128 | + gert::TilingContextPara tilingContextPara( | ||
| 129 | + "ResizeBicubicV2", | ||
| 130 | + {{inputXShape, ge::DT_FLOAT, ge::FORMAT_NHWC}, {inputSizeShape, ge::DT_INT32, ge::FORMAT_ND, true, size_value}}, | ||
| 131 | + {{outputShape, ge::DT_FLOAT, ge::FORMAT_NHWC}}, | ||
| 132 | + {gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 133 | + gert::TilingContextPara::OpAttr("scales", Ops::Cv::AnyValue::CreateFrom<vector<float>>({0.0f, 0.0f}))}, | ||
| 134 | + &compileInfo); | ||
| 135 | + TilingInfo tilingInfo; | ||
| 136 | + EXPECT_TRUE(ExecuteTiling(tilingContextPara, tilingInfo)); | ||
| 137 | + EXPECT_EQ(GetIsNchwFromTilingKey(tilingInfo.tilingKey), 0); | ||
| 138 | +} | ||
| 139 | + | ||
| 140 | +// Test 09 [回归]: 输入输出的H和W轴超 INT32_MAX 必须走 int64 索引核(isInt32=0)。 | ||
| 141 | +// 缺陷(修复前用 `> UINT32_MAX`):isInt32_==1 时 SIMT 核绑定 <T_IDX=uint32_t, T_IDX2=int32_t>(apt.cpp), | ||
| 142 | +// T_IDX2 是**有符号** int32,承接空间坐标(topY/leftX、lenSrcH1=lenSrcH-1),上限 INT32_MAX(2^31-1)。 | ||
| 143 | +// 旧 UINT32_MAX 阈值只护无符号 T_IDX,输入输出的H和W轴(INT32_MAX, UINT32_MAX] 区间被误判 isInt32=1 → int32 有符号 | ||
| 144 | +// 溢出为负 → GetSrc 钳位错乱 → GM 偏移错乱静默数据错。本用例lenSrcH = 2147483648 > INT32_MAX | ||
| 145 | +TEST_F(ResizeBicubicV2TilingTest, resize_bicubic_v2_tiling_09_int32_overflow_coarse) | ||
| 146 | +{ | ||
| 147 | + gert::StorageShape inputXShape = {{1, 1, 2147483648, 1}, {1, 1, 2147483648, 1}}; // lenSrcH = 2147483648 > INT32_MAX | ||
| 148 | + gert::StorageShape inputSizeShape = {{1, 2}, {1, 2}}; | ||
| 149 | + gert::StorageShape outputShape = {{1, 1, 70, 83}, {1, 1, 70, 83}}; | ||
| 150 | + int size_value[2] = {70, 83}; | ||
| 151 | + | ||
| 152 | + ResizeBicubicV2CompileInfo compileInfo = {64, 200704}; | ||
| 153 | + | ||
| 154 | + gert::TilingContextPara tilingContextPara( | ||
| 155 | + "ResizeBicubicV2", | ||
| 156 | + {{inputXShape, ge::DT_FLOAT, ge::FORMAT_NCHW}, {inputSizeShape, ge::DT_INT32, ge::FORMAT_ND, true, size_value}}, | ||
| 157 | + {{outputShape, ge::DT_FLOAT, ge::FORMAT_NCHW}}, | ||
| 158 | + {gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(true)), | ||
| 159 | + gert::TilingContextPara::OpAttr("scales", Ops::Cv::AnyValue::CreateFrom<vector<float>>({0.0f, 0.0f}))}, | ||
| 160 | + &compileInfo); | ||
| 161 | + TilingInfo tilingInfo; | ||
| 162 | + EXPECT_TRUE(ExecuteTiling(tilingContextPara, tilingInfo)); | ||
| 163 | + EXPECT_EQ(GetIsInt32FromTilingKey(tilingInfo.tilingKey), 0); // lenSrcH超 INT32_MAX 必走 int64 | ||
| 80 | } | 164 | } |
| @@ -371,7 +371,8 @@ void ResizeBicubicV2GradBaseTiling::SetScales() | |||
| 371 | bool ResizeBicubicV2GradBaseTiling::IsUseIdx32() const | 371 | bool ResizeBicubicV2GradBaseTiling::IsUseIdx32() const |
| 372 | { | 372 | { |
| 373 | return calcInfo_.gradsShapeSize <= UINT32_MAX && calcInfo_.yShapeSize <= UINT32_MAX && | 373 | return calcInfo_.gradsShapeSize <= UINT32_MAX && calcInfo_.yShapeSize <= UINT32_MAX && |
| 374 | - inputInfo_.lenSrcH <= INT32_MAX && inputInfo_.lenSrcW <= INT32_MAX; | 374 | + inputInfo_.lenSrcH <= INT32_MAX && inputInfo_.lenSrcW <= INT32_MAX && inputInfo_.lenDstH <= INT32_MAX && |
| 375 | + inputInfo_.lenDstW <= INT32_MAX; | ||
 | |||
| 375 | } | 376 | } |
| 376 | 377 | ||
| 377 | ge::graphStatus Tiling4ResizeBicubicV2Grad(gert::TilingContext* context) | 378 | ge::graphStatus Tiling4ResizeBicubicV2Grad(gert::TilingContext* context) |
| @@ -596,7 +596,8 @@ void ResizeBilinearV2AscendCTilingImpl::DoTilingSIMT_HW() | |||
| 596 | 596 | ||
| 597 | void ResizeBilinearV2AscendCTilingImpl::MatchTilingStrategyAndSetTilingKey() | 597 | void ResizeBilinearV2AscendCTilingImpl::MatchTilingStrategyAndSetTilingKey() |
| 598 | { | 598 | { |
| 599 | - bool useIdx32 = yShape_.GetShapeSize() < UINT32_MAX && xShape_.GetShapeSize() < UINT32_MAX; | 599 | + bool useIdx32 = yShape_.GetShapeSize() < UINT32_MAX && xShape_.GetShapeSize() < UINT32_MAX && |
| 600 | + lenSrcH_ <= INT32_MAX && lenSrcW_ <= INT32_MAX && lenDesH_ <= INT32_MAX && lenDesW_ <= INT32_MAX; | ||
| 600 | if (IsMatchAllCopy()) { | 601 | if (IsMatchAllCopy()) { |
| 601 | tilingKey_ = TILING_KEY_ALL_COPY; | 602 | tilingKey_ = TILING_KEY_ALL_COPY; |
| 602 | } else if (IsMatchPointCopy()) { | 603 | } else if (IsMatchPointCopy()) { |
| @@ -418,7 +418,10 @@ void ResizeBilinearV2GradTilingAscendC::SetFactors() | |||
| 418 | 418 | ||
| 419 | void ResizeBilinearV2GradTilingAscendC::SetSimtTilingKey(bool isDetermine) | 419 | void ResizeBilinearV2GradTilingAscendC::SetSimtTilingKey(bool isDetermine) |
| 420 | { | 420 | { |
| 421 | - bool isIdx32 = lenN_ * lenC_ * lenDesH_ * lenDesW_ < UINT32_MAX && lenN_ * lenC_ * lenSrcH_ * lenSrcW_ < UINT32_MAX; | 421 | + int64_t ySize_ = lenN_ * lenC_ * lenDesH_ * lenDesW_; |
| 422 | + int64_t xSize_ = lenN_ * lenC_ * lenSrcH_ * lenSrcW_; | ||
| 423 | + bool isIdx32 = ySize_ < UINT32_MAX && xSize_ < UINT32_MAX && lenSrcH_ <= INT32_MAX && lenSrcW_ <= INT32_MAX && | ||
| 424 | + lenDesH_ <= INT32_MAX && lenDesW_ <= INT32_MAX; | ||
| 422 | if ((isDetermine && alignCorners_) || | 425 | if ((isDetermine && alignCorners_) || |
| 423 | (isDetermine && !alignCorners_ && | 426 | (isDetermine && !alignCorners_ && |
| 424 | (std::fabs(static_cast<float>(lenSrcH_) / lenDesH_ - scaleH_) < FLT_EPSILON) && | 427 | (std::fabs(static_cast<float>(lenSrcH_) / lenDesH_ - scaleH_) < FLT_EPSILON) && |
| @@ -214,6 +214,9 @@ ge::graphStatus ResizeLinearTiling::LinearCompute() | |||
| 214 | OP_LOGI(context_->GetNodeName(), "ySize or xSize is too large"); | 214 | OP_LOGI(context_->GetNodeName(), "ySize or xSize is too large"); |
| 215 | isInt32_ = DIM_0; | 215 | isInt32_ = DIM_0; |
| 216 | } | 216 | } |
| 217 | + if (lenSrcL_ > INT32_MAX || lenDesL_ > INT32_MAX) { | ||
| 218 | + isInt32_ = DIM_0; | ||
| 219 | + } | ||
| 217 | ComputeKey(); | 220 | ComputeKey(); |
| 218 | return ge::GRAPH_SUCCESS; | 221 | return ge::GRAPH_SUCCESS; |
| 219 | } | 222 | } |
| @@ -374,6 +374,9 @@ ge::graphStatus ResizeLinearGradTiling::LinearComputeGrad() | |||
| 374 | OP_LOGI(context_->GetNodeName(), "ySize or xSize is large than UINT32_MAX"); | 374 | OP_LOGI(context_->GetNodeName(), "ySize or xSize is large than UINT32_MAX"); |
| 375 | isInt32_ = DIM_0; | 375 | isInt32_ = DIM_0; |
| 376 | } | 376 | } |
| 377 | + if (lenSrcLOrUb_ > INT32_MAX || lenDesL_ > INT32_MAX) { | ||
| 378 | + isInt32_ = DIM_0; | ||
| 379 | + } | ||
| 377 | ComputeKey(); | 380 | ComputeKey(); |
| 378 | return ge::GRAPH_SUCCESS; | 381 | return ge::GRAPH_SUCCESS; |
| 379 | } | 382 | } |
| @@ -698,7 +698,8 @@ void ResizeNearestNeighborV2AscendCTilingImpl::MatchTilingStrategyAndSetTilingKe | |||
| 698 | idxUseInt32_ = static_cast<uint64_t>(1); | 698 | idxUseInt32_ = static_cast<uint64_t>(1); |
| 699 | format_ = ge::FORMAT_NCHW; | 699 | format_ = ge::FORMAT_NCHW; |
| 700 | } else { | 700 | } else { |
| 701 | - idxUseInt32_ = xShape_.GetShapeSize() < UINT32_MAX && yShape_.GetShapeSize() < UINT32_MAX; | 701 | + idxUseInt32_ = xShape_.GetShapeSize() < UINT32_MAX && yShape_.GetShapeSize() < UINT32_MAX && |
| 702 | + lenSrcH_ <= INT32_MAX && lenSrcW_ <= INT32_MAX && lenDesH_ <= INT32_MAX && lenDesW_ <= INT32_MAX; | ||
| 702 | if (lenSrcH_ == lenDesH_ && lenSrcW_ == lenDesW_) { | 703 | if (lenSrcH_ == lenDesH_ && lenSrcW_ == lenDesW_) { |
| 703 | schId_ = SCHEDULE_ID_SIMT_INPUT_EQ_OUTPUT; | 704 | schId_ = SCHEDULE_ID_SIMT_INPUT_EQ_OUTPUT; |
| 704 | } else if (lenSrcH_ == DIM_1 && lenSrcW_ == DIM_1) { | 705 | } else if (lenSrcH_ == DIM_1 && lenSrcW_ == DIM_1) { |
Mimage/resize_nearest_neighbor_v2_grad/op_host/arch35/resize_nearest_neighbor_v2_grad_tiling_base.cpp+2-1
| @@ -427,7 +427,8 @@ bool ResizeNearestNeighborV2GradTiling::IsSimtNotDetermineHW() | |||
| 427 | 427 | ||
| 428 | void ResizeNearestNeighborV2GradTiling::SetIndexType() | 428 | void ResizeNearestNeighborV2GradTiling::SetIndexType() |
| 429 | { | 429 | { |
| 430 | - bool isIdx32 = yShape_.GetShapeSize() < UINT32_MAX && gradsShape_.GetShapeSize() < UINT32_MAX; | 430 | + bool isIdx32 = yShape_.GetShapeSize() < UINT32_MAX && gradsShape_.GetShapeSize() < UINT32_MAX && |
| 431 | + lenSrcH_ <= INT32_MAX && lenSrcW_ <= INT32_MAX && lenDstH_ <= INT32_MAX && lenDstW_ <= INT32_MAX; | ||
| 431 | idxType_ = static_cast<uint64_t>(isIdx32 ? TPL_IDX_INT32 : TPL_IDX_INT64); | 432 | idxType_ = static_cast<uint64_t>(isIdx32 ? TPL_IDX_INT32 : TPL_IDX_INT64); |
| 432 | } | 433 | } |
| 433 | 434 | ||
| @@ -320,7 +320,11 @@ ge::graphStatus UpsampleNearest3dRegbaseTiling::GetAndCheckShapes() | |||
| 320 | return ge::GRAPH_FAILED; | 320 | return ge::GRAPH_FAILED; |
| 321 | } | 321 | } |
| 322 | int64_t uint32Max = static_cast<int64_t>(std::numeric_limits<uint32_t>::max()); | 322 | int64_t uint32Max = static_cast<int64_t>(std::numeric_limits<uint32_t>::max()); |
| 323 | - int32_t isUint32 = static_cast<int32_t>((inputSize <= uint32Max) && (outputSize <= uint32Max)); | 323 | + int64_t int32Max = static_cast<int64_t>(std::numeric_limits<int32_t>::max()); |
| 324 | + int32_t isUint32 = static_cast<int32_t>( | ||
| 325 | + (inputSize <= uint32Max) && (outputSize <= uint32Max) && baseTiling_.inD <= int32Max && | ||
| 326 | + baseTiling_.inH <= int32Max && baseTiling_.inW <= int32Max && baseTiling_.outD <= int32Max && | ||
| 327 | + baseTiling_.outH <= int32Max && baseTiling_.outW <= int32Max); | ||
| 324 | baseTiling_.isUint32 = isUint32; | 328 | baseTiling_.isUint32 = isUint32; |
| 325 | return ge::GRAPH_SUCCESS; | 329 | return ge::GRAPH_SUCCESS; |
| 326 | } | 330 | } |
🔴 Critical
diff 中
IsUseIdx32()函数原有的闭合花括号}（旧行375）被删除，新代码行390return ... <= INT32_MAX;之后没有恢复}。导致后续的Tiling4ResizeBicubicV2Grad、TilingPrepare4ResizeBicubicV2Grad两个函数和IMPL_OP_OPTILING宏被错误地嵌入到IsUseIdx32()成员函数体内。后果：
return语句之后，形成不可达死代码；IMPL_OP_OPTILING宏展开需要命名空间级别的上下文，在函数体内会编译失败；Tiling4ResizeBicubicV2Grad和TilingPrepare4ResizeBicubicV2Grad被解析为类成员函数定义，但无法匹配任何声明。这是 diff 格式化重构过程中的误删。在行390
inputInfo_.lenDstW <= INT32_MAX;之后、行392ge::graphStatus Tiling4ResizeBicubicV2Grad(...)之前，必须补回}。建议：在行390
inputInfo_.lenDstW <= INT32_MAX;之后、行391空行之后、行392之前，插入}闭合IsUseIdx32()函数。