已合并
优化AdaptiveAvgPool2dGrad性能 #6899
小王！创建于 7月2日
优化AdaptiveAvgPool2dGrad性能 #6899
已合并
共 5 个文件变更+1031-238
| @@ -16,6 +16,7 @@ | |||
| 16 | 16 | ||
| 17 | 17 | ||
| 18 | 18 | ||
| 19 | + | ||
| 19 | 20 | ||
| 20 | namespace optiling { | 21 | namespace optiling { |
| 21 | using namespace AdaptiveAvgPool2dGradOp; | 22 | using namespace AdaptiveAvgPool2dGradOp; |
| @@ -28,6 +29,41 @@ constexpr int64_t WINSIZE_THRESHOLD = 16; | |||
| 28 | constexpr int64_t INPUTW_FLOAT_THRESHOLD = 8; | 29 | constexpr int64_t INPUTW_FLOAT_THRESHOLD = 8; |
| 29 | constexpr int64_t INPUTW_BFLOAT_THRESHOLD = 16; | 30 | constexpr int64_t INPUTW_BFLOAT_THRESHOLD = 16; |
| 30 | constexpr int64_t LIMIT = 1; | 31 | constexpr int64_t LIMIT = 1; |
| 32 | +constexpr int64_t NC_SEARCH_MAX = 256; | ||
| 33 | +constexpr int64_t SEARCH_HW_SIZE_LIMIT = 200000; | ||
| 34 | + | ||
| 35 | +constexpr int64_t PREFER_SIMT_W_THRESHOLD = 4; | ||
| 36 | +constexpr int64_t PREFER_SIMT_H_LOW_THRESHOLD = 64; | ||
| 37 | +constexpr int64_t STRONG_RESIZE_RATIO = 4; | ||
| 38 | +constexpr int64_t UPSAMPLE_AREA_EXPAND_RATIO = 3; | ||
| 39 | +constexpr int64_t STRONG_COLLAPSE_RATIO = 3; | ||
| 40 | +constexpr int64_t RESIZE_W_EXPAND_RATIO = 2; | ||
| 41 | +constexpr int64_t NC_SEARCH_INPUT_VL_MULTIPLIER = 2; | ||
| 42 | +constexpr int64_t HW_INNER_SAFE_MARGIN = 1; | ||
| 43 | +constexpr int64_t OUTPUT_FP32_FACTOR = 2; | ||
| 44 | +constexpr int64_t WORK_PER_BLOCK_UB_OVERHEAD = 64; | ||
| 45 | +constexpr long double COST_HIGH_AXIS_PADDING_FACTOR = 4.0L; | ||
| 46 | +constexpr int64_t HIGH_AXIS_TAIL_OPT_THRESHOLD = 8; | ||
| 47 | +constexpr long double COST_PARTIAL_VL_PENALTY_HW = 1024.0L; | ||
| 48 | +constexpr long double COST_PARTIAL_VL_PENALTY_BASELINE = 2048.0L; | ||
| 49 | +constexpr long double COST_TRANS_ALIGN_FACTOR = 0.25L; | ||
| 50 | +constexpr long double COST_IDLE_CORE_FACTOR = 0.15L; | ||
| 51 | + | ||
| 52 | +static inline int64_t ShrinkInnerStrict(int64_t total, int64_t curInner) | ||
| 53 | +{ | ||
| 54 | + if (curInner <= LIMIT) { | ||
| 55 | + return LIMIT; | ||
| 56 | + } | ||
| 57 | + | ||
| 58 | + const int64_t curOuter = Ops::Base::CeilDiv(total, curInner); | ||
| 59 | + int64_t nextInner = Ops::Base::CeilDiv(total, curOuter + 1); | ||
| 60 | + | ||
| 61 | + if (nextInner >= curInner) { | ||
| 62 | + nextInner = curInner - 1; | ||
| 63 | + } | ||
| 64 | + | ||
| 65 | + return std::max<int64_t>(static_cast<int64_t>(LIMIT), nextInner); | ||
| 66 | +} | ||
| 31 | 67 | ||
| 32 | void AdaptiveAvgPool2dGradTilingSmallKernel::InitializationVars() | 68 | void AdaptiveAvgPool2dGradTilingSmallKernel::InitializationVars() |
| 33 | { | 69 | { |
| @@ -50,6 +86,8 @@ void AdaptiveAvgPool2dGradTilingSmallKernel::InitializationVars() | |||
| 50 | baseData.maxDataNumInOneBlock = baseData.ubBlockSize / baseData.inputBytes; | 86 | baseData.maxDataNumInOneBlock = baseData.ubBlockSize / baseData.inputBytes; |
| 51 | baseData.proDataNumInOneBeatT2 = baseData.vRegSize / baseData.inputBytes; | 87 | baseData.proDataNumInOneBeatT2 = baseData.vRegSize / baseData.inputBytes; |
| 52 | baseData.inputNCSize = gradOutputN * gradOutputC; | 88 | baseData.inputNCSize = gradOutputN * gradOutputC; |
| 89 | + | ||
| 90 | + const int64_t computeVl = baseData.vRegSize / FLOAT32_SIZE; | ||
| 53 | } | 91 | } |
| 54 | 92 | ||
| 55 | void AdaptiveAvgPool2dGradTilingSmallKernel::DoBufferCalculate() | 93 | void AdaptiveAvgPool2dGradTilingSmallKernel::DoBufferCalculate() |
| @@ -96,12 +134,40 @@ bool AdaptiveAvgPool2dGradTilingSmallKernel::IsCapable() | |||
| 96 | { | 134 | { |
| 97 | InitializationVars(); | 135 | InitializationVars(); |
| 98 | 136 | ||
| 137 | + const bool preferSimtInputAsGradY = | ||
| 138 | + gradInputH > gradOutputH && | ||
| 139 | + gradInputW <= gradOutputW && | ||
| 140 | + gradInputW <= PREFER_SIMT_W_THRESHOLD && | ||
| 141 | + (gradInputH >= gradOutputH * STRONG_RESIZE_RATIO || gradOutputH <= PREFER_SIMT_H_LOW_THRESHOLD); | ||
| 142 | + | ||
| 143 | + const bool preferSimtOutputAsGradY = | ||
| 144 | + gradOutputH > gradInputH && | ||
| 145 | + gradOutputW <= gradInputW && | ||
| 146 | + gradOutputW <= PREFER_SIMT_W_THRESHOLD && | ||
| 147 | + (gradOutputH >= gradInputH * STRONG_RESIZE_RATIO || gradInputH <= PREFER_SIMT_H_LOW_THRESHOLD); | ||
| 148 | + | ||
| 149 | + if (preferSimtInputAsGradY || preferSimtOutputAsGradY) { | ||
| 150 | + return false; | ||
| 151 | + } | ||
| 152 | + | ||
| 99 | kernelH = Ops::Base::CeilDiv(gradOutputH, gradInputH); | 153 | kernelH = Ops::Base::CeilDiv(gradOutputH, gradInputH); |
| 100 | kernelW = Ops::Base::CeilDiv(gradOutputW, gradInputW); | 154 | kernelW = Ops::Base::CeilDiv(gradOutputW, gradInputW); |
| 101 | 155 | ||
| 102 | - if (kernelH * kernelW >= KERNEL_SIZE_MAX || | 156 | + const int64_t kernelSize = kernelH * kernelW; |
| 103 | - baseData.inputNCSize < HIGH_THRESHOLD || | 157 | + const int64_t inputWinSize = gradInputW * gradInputH; |
| 104 | - gradInputW * gradInputH < WINSIZE_THRESHOLD) { | 158 | + const int64_t highAxis = baseData.inputNCSize; |
| 159 | + const int64_t inputHW = gradInputH * gradInputW; | ||
| 160 | + const int64_t outputHW = gradOutputH * gradOutputW; | ||
| 161 | + | ||
| 162 | + if (kernelSize >= KERNEL_SIZE_MAX) { | ||
| 163 | + return false; | ||
| 164 | + } | ||
| 165 | + | ||
| 166 | + if (baseData.inputNCSize < HIGH_THRESHOLD) { | ||
| 167 | + return false; | ||
| 168 | + } | ||
| 169 | + | ||
| 170 | + if (inputWinSize < WINSIZE_THRESHOLD) { | ||
| 105 | return false; | 171 | return false; |
| 106 | } | 172 | } |
| 107 | 173 | ||
| @@ -115,12 +181,54 @@ bool AdaptiveAvgPool2dGradTilingSmallKernel::IsCapable() | |||
| 115 | } | 181 | } |
| 116 | } | 182 | } |
| 117 | 183 | ||
| 184 | + constexpr int64_t SIMT_RESIZE_KERNEL_SIZE_MAX = 7; | ||
| 185 | + constexpr int64_t SIMT_RESIZE_NC_MAX = 4096; | ||
| 186 | + constexpr int64_t SIMT_BOTH_UPSAMPLE_NC_MAX = 700; | ||
| 187 | + | ||
| 188 | + const bool hUpsample = gradOutputH > gradInputH; | ||
| 189 | + const bool wUpsample = gradOutputW > gradInputW; | ||
| 190 | + const bool hDownsample = gradOutputH < gradInputH; | ||
| 191 | + const bool wDownsample = gradOutputW < gradInputW; | ||
| 192 | + const bool hResize = gradOutputH != gradInputH; | ||
| 193 | + const bool wResize = gradOutputW != gradInputW; | ||
| 194 | + const bool twoAxisResize = hResize && wResize; | ||
| 195 | + | ||
| 196 | + const bool bothUpsampleAreaExpand = | ||
| 197 | + hUpsample && | ||
| 198 | + wUpsample && | ||
| 199 | + highAxis <= SIMT_BOTH_UPSAMPLE_NC_MAX && | ||
| 200 | + outputHW >= inputHW * UPSAMPLE_AREA_EXPAND_RATIO; | ||
| 201 | + | ||
| 202 | + const bool hExpandWCollapseStrong = | ||
| 203 | + hUpsample && | ||
| 204 | + wDownsample && | ||
| 205 | + gradOutputH >= gradInputH * STRONG_RESIZE_RATIO && | ||
| 206 | + gradInputW >= gradOutputW * STRONG_COLLAPSE_RATIO; | ||
| 207 | + | ||
| 208 | + const bool hCollapseWExpandStrong = | ||
| 209 | + hDownsample && | ||
| 210 | + wUpsample && | ||
| 211 | + gradInputH >= gradOutputH * STRONG_COLLAPSE_RATIO && | ||
| 212 | + gradOutputW * RESIZE_W_EXPAND_RATIO >= gradInputW * STRONG_COLLAPSE_RATIO; | ||
| 213 | + | ||
| 214 | + const bool preferSimtUnfriendlyResize = | ||
| 215 | + kernelSize <= SIMT_RESIZE_KERNEL_SIZE_MAX && | ||
| 216 | + highAxis >= HIGH_THRESHOLD && | ||
| 217 | + highAxis <= SIMT_RESIZE_NC_MAX && | ||
| 218 | + twoAxisResize && | ||
| 219 | + (bothUpsampleAreaExpand || hExpandWCollapseStrong || hCollapseWExpandStrong); | ||
| 220 | + | ||
| 221 | + if (preferSimtUnfriendlyResize) { | ||
| 222 | + return false; | ||
| 223 | + } | ||
| 224 | + | ||
| 118 | splitData.highAxisInner = baseData.proDataNumInOneBeatT2; | 225 | splitData.highAxisInner = baseData.proDataNumInOneBeatT2; |
| 119 | splitData.hOutputInner = LIMIT; | 226 | splitData.hOutputInner = LIMIT; |
| 120 | splitData.wOutputInner = LIMIT; | 227 | splitData.wOutputInner = LIMIT; |
| 121 | DoBufferCalculate(); | 228 | DoBufferCalculate(); |
| 122 | 229 | ||
| 123 | - return splitData.totalBufferSize <= baseData.availableUb; | 230 | + const bool capable = splitData.totalBufferSize <= baseData.availableUb; |
| 231 | + return capable; | ||
| 124 | } | 232 | } |
| 125 | 233 | ||
| 126 | bool AdaptiveAvgPool2dGradTilingSmallKernel::IsMeetTargetCoreNum() | 234 | bool AdaptiveAvgPool2dGradTilingSmallKernel::IsMeetTargetCoreNum() |
| @@ -150,15 +258,13 @@ bool AdaptiveAvgPool2dGradTilingSmallKernel::TrySplitNC() | |||
| 150 | 258 | ||
| 151 | void AdaptiveAvgPool2dGradTilingSmallKernel::DynamicAdjustmentHW() | 259 | void AdaptiveAvgPool2dGradTilingSmallKernel::DynamicAdjustmentHW() |
| 152 | { | 260 | { |
| 153 | - if (splitData.hOutputInner != LIMIT) { | 261 | + if (splitData.hOutputInner > LIMIT) { |
| 154 | - splitData.hOutputOuter++; | 262 | + splitData.hOutputInner = ShrinkInnerStrict(gradOutputH, splitData.hOutputInner); |
| 155 | - splitData.hOutputInner = Ops::Base::CeilDiv(gradOutputH, splitData.hOutputOuter); | ||
| 156 | return; | 263 | return; |
| 157 | } | 264 | } |
| 158 | 265 | ||
| 159 | - if (splitData.wOutputInner != LIMIT) { | 266 | + if (splitData.wOutputInner > LIMIT) { |
| 160 | - splitData.wOutputOuter++; | 267 | + splitData.wOutputInner = ShrinkInnerStrict(gradOutputW, splitData.wOutputInner); |
| 161 | - splitData.wOutputInner = Ops::Base::CeilDiv(gradOutputW, splitData.wOutputOuter); | ||
| 162 | return; | 268 | return; |
| 163 | } | 269 | } |
| 164 | } | 270 | } |
| @@ -166,18 +272,19 @@ void AdaptiveAvgPool2dGradTilingSmallKernel::DynamicAdjustmentHW() | |||
| 166 | void AdaptiveAvgPool2dGradTilingSmallKernel::SplitUnalignHW() | 272 | void AdaptiveAvgPool2dGradTilingSmallKernel::SplitUnalignHW() |
| 167 | { | 273 | { |
| 168 | splitData.highAxisInner = baseData.proDataNumInOneBeatT2; | 274 | splitData.highAxisInner = baseData.proDataNumInOneBeatT2; |
| 169 | - | ||
| 170 | splitData.hOutputInner = gradOutputH; | 275 | splitData.hOutputInner = gradOutputH; |
| 171 | splitData.wOutputInner = gradOutputW; | 276 | splitData.wOutputInner = gradOutputW; |
| 172 | 277 | ||
| 173 | - splitData.hOutputOuter = Ops::Base::CeilDiv(gradOutputH, splitData.hOutputInner); | 278 | + while (!IsMeetTargetCoreNum() || !IsMeetUBSize()) { |
| 174 | - splitData.wOutputOuter = Ops::Base::CeilDiv(gradOutputW, splitData.wOutputInner); | 279 | + const int64_t oldH = splitData.hOutputInner; |
| 280 | + const int64_t oldW = splitData.wOutputInner; | ||
| 175 | 281 | ||
| 176 | - while (splitData.hOutputInner != LIMIT || splitData.wOutputInner != LIMIT) { | ||
| 177 | - if (IsMeetTargetCoreNum() && IsMeetUBSize()) { | ||
| 178 | - return; | ||
| 179 | - } | ||
| 180 | DynamicAdjustmentHW(); | 282 | DynamicAdjustmentHW(); |
| 283 | + | ||
| 284 | + if (oldH == splitData.hOutputInner && | ||
| 285 | + oldW == splitData.wOutputInner) { | ||
| 286 | + break; | ||
| 287 | + } | ||
| 181 | } | 288 | } |
| 182 | 289 | ||
| 183 | DoBufferCalculate(); | 290 | DoBufferCalculate(); |
| @@ -189,6 +296,254 @@ void AdaptiveAvgPool2dGradTilingSmallKernel::SearchBestTiling() | |||
| 189 | return; | 296 | return; |
| 190 | } | 297 | } |
| 191 | 298 | ||
| 299 | + const int64_t computeVl = std::max<int64_t>(TRANS_ADDR_LEN, baseData.vRegSize / FLOAT32_SIZE); | ||
| 300 | + const int64_t inputVl = baseData.proDataNumInOneBeatT2; | ||
| 301 | + const int64_t searchHwSize = gradOutputH * gradOutputW; | ||
| 302 | + | ||
| 303 | + bool found = false; | ||
| 304 | + int64_t bestHighAxisInner = 0; | ||
| 305 | + int64_t bestHOutputInner = 0; | ||
| 306 | + int64_t bestWOutputInner = 0; | ||
| 307 | + int64_t bestBlockNum = 0; | ||
| 308 | + int64_t bestUsedCoreNum = 0; | ||
| 309 | + int64_t bestHighAxisPadding = 0; | ||
| 310 | + int64_t bestHighAxisTail = 0; | ||
| 311 | + int64_t bestBufferSize = 0; | ||
| 312 | + long double bestCost = std::numeric_limits<long double>::max(); | ||
| 313 | + | ||
| 314 | + if (searchHwSize <= SEARCH_HW_SIZE_LIMIT) { | ||
| 315 | + int64_t ncSearchMax = std::max<int64_t>( | ||
| 316 | + Ops::Base::CeilAlign(baseData.inputNCSize, TRANS_ADDR_LEN), | ||
| 317 | + inputVl * NC_SEARCH_INPUT_VL_MULTIPLIER); | ||
| 318 | + ncSearchMax = std::min<int64_t>(ncSearchMax, static_cast<int64_t>(NC_SEARCH_MAX)); | ||
| 319 | + ncSearchMax = std::max<int64_t>(ncSearchMax, computeVl); | ||
| 320 | + | ||
| 321 | + ExhaustiveSearchBestTiling( | ||
| 322 | + computeVl, ncSearchMax, | ||
| 323 | + bestHighAxisInner, bestHOutputInner, bestWOutputInner, | ||
| 324 | + bestBlockNum, bestUsedCoreNum, bestHighAxisPadding, | ||
| 325 | + bestHighAxisTail, bestBufferSize, bestCost, found); | ||
| 326 | + } | ||
| 327 | + | ||
| 328 | + if (found) { | ||
| 329 | + splitData.highAxisInner = bestHighAxisInner; | ||
| 330 | + splitData.hOutputInner = bestHOutputInner; | ||
| 331 | + splitData.wOutputInner = bestWOutputInner; | ||
| 332 | + DoBufferCalculate(); | ||
| 333 | + return; | ||
| 334 | + } | ||
| 335 | + | ||
| 336 | + ApplyCoarseFallback(); | ||
| 337 | +} | ||
| 338 | + | ||
| 339 | +bool AdaptiveAvgPool2dGradTilingSmallKernel::ExhaustiveSearchBestTiling( | ||
| 340 | + int64_t computeVl, int64_t ncSearchMax, | ||
| 341 | + int64_t& bestHighAxisInner, int64_t& bestHOutputInner, int64_t& bestWOutputInner, | ||
| 342 | + int64_t& bestBlockNum, int64_t& bestUsedCoreNum, | ||
| 343 | + int64_t& bestHighAxisPadding, int64_t& bestHighAxisTail, | ||
| 344 | + int64_t& bestBufferSize, long double& bestCost, bool& found) | ||
| 345 | +{ | ||
| 346 | + for (int64_t highAxisInner = computeVl; | ||
| 347 | + highAxisInner <= ncSearchMax; highAxisInner += TRANS_ADDR_LEN) { | ||
| 348 | + splitData.highAxisInner = highAxisInner; | ||
| 349 | + const int64_t highAxisOuter = | ||
| 350 | + Ops::Base::CeilDiv(baseData.inputNCSize, highAxisInner); | ||
| 351 | + const int64_t highAxisTail = | ||
| 352 | + (baseData.inputNCSize % highAxisInner == 0) ? highAxisInner : | ||
| 353 | + (baseData.inputNCSize % highAxisInner); | ||
| 354 | + const int64_t highAxisPadding = | ||
| 355 | + highAxisOuter * highAxisInner - baseData.inputNCSize; | ||
| 356 | + for (int64_t hOutputInner = LIMIT; hOutputInner <= gradOutputH; ++hOutputInner) { | ||
| 357 | + splitData.hOutputInner = hOutputInner; | ||
| 358 | + const int64_t hOutputOuter = | ||
| 359 | + Ops::Base::CeilDiv(gradOutputH, hOutputInner); | ||
| 360 | + for (int64_t wOutputInner = LIMIT; wOutputInner <= gradOutputW; ++wOutputInner) { | ||
| 361 | + splitData.wOutputInner = wOutputInner; | ||
| 362 | + DoBufferCalculate(); | ||
| 363 | + if (splitData.totalBufferSize > baseData.availableUb) { continue; } | ||
| 364 | + const int64_t wOutputOuter = | ||
| 365 | + Ops::Base::CeilDiv(gradOutputW, wOutputInner); | ||
| 366 | + const int64_t blockNum = highAxisOuter * hOutputOuter * wOutputOuter; | ||
| 367 | + if (blockNum < baseData.coreUsedForBestPerformance) { continue; } | ||
| 368 | + const int64_t normalCoreProcessNum = | ||
| 369 | + Ops::Base::CeilDiv(blockNum, baseData.totalCoreNum); | ||
| 370 | + long double cost = EvalTilingCandidate( | ||
| 371 | + highAxisInner, highAxisOuter, highAxisTail, highAxisPadding, | ||
| 372 | + hOutputInner, hOutputOuter, wOutputInner, wOutputOuter, | ||
| 373 | + blockNum, computeVl, normalCoreProcessNum); | ||
| 374 | + if (cost < 0.0L) { continue; } | ||
| 375 | + const int64_t usedCoreNum = | ||
| 376 | + Ops::Base::CeilDiv(blockNum, normalCoreProcessNum); | ||
| 377 | + TryRecordBetterTiling(cost, hOutputInner, wOutputInner, | ||
| 378 | + blockNum, usedCoreNum, highAxisInner, | ||
| 379 | + highAxisPadding, highAxisTail, | ||
| 380 | + bestHighAxisInner, bestHOutputInner, bestWOutputInner, | ||
| 381 | + bestBlockNum, bestUsedCoreNum, bestHighAxisPadding, | ||
| 382 | + bestHighAxisTail, bestBufferSize, bestCost, found); | ||
| 383 | + } | ||
| 384 | + } | ||
| 385 | + } | ||
| 386 | + return found; | ||
| 387 | +} | ||
| 388 | + | ||
| 389 | +long double AdaptiveAvgPool2dGradTilingSmallKernel::EvalTilingCandidate( | ||
| 390 | + int64_t highAxisInner, int64_t highAxisOuter, int64_t highAxisTail, | ||
| 391 | + int64_t highAxisPadding, | ||
| 392 | + int64_t hOutputInner, int64_t hOutputOuter, | ||
| 393 | + int64_t wOutputInner, int64_t wOutputOuter, | ||
| 394 | + int64_t blockNum, int64_t computeVl, | ||
| 395 | + int64_t normalCoreProcessNum) | ||
| 396 | +{ | ||
| 397 | + const int64_t oneBufferSize = | ||
| 398 | + splitData.inputQueBufferSize + splitData.transQueBufferSize + | ||
| 399 | + splitData.transOutQueBufferSize; | ||
| 400 | + const int64_t hInputInner = | ||
| 401 | + Ops::Base::CeilDiv(hOutputInner * gradInputH, gradOutputH) + | ||
| 402 | + HW_INNER_SAFE_MARGIN; | ||
| 403 | + const int64_t wInputInner = | ||
| 404 | + Ops::Base::CeilDiv(wOutputInner * gradInputW, gradOutputW) + | ||
| 405 | + HW_INNER_SAFE_MARGIN; | ||
| 406 | + const int64_t actualInputElem = highAxisInner * hInputInner * wInputInner; | ||
| 407 | + const int64_t actualOutputElem = highAxisInner * hOutputInner * wOutputInner; | ||
| 408 | + const int64_t oneBlockWork = | ||
| 409 | + oneBufferSize + | ||
| 410 | + actualInputElem * (baseData.inputBytes + FLOAT32_SIZE) + | ||
| 411 | + actualOutputElem * FLOAT32_SIZE * OUTPUT_FP32_FACTOR + | ||
| 412 | + baseData.ubBlockSize * WORK_PER_BLOCK_UB_OVERHEAD; | ||
| 413 | + | ||
| 414 | + long double cost = | ||
| 415 | + static_cast<long double>(normalCoreProcessNum) * | ||
| 416 | + static_cast<long double>(oneBlockWork); | ||
| 417 | + cost += static_cast<long double>(highAxisPadding) * | ||
| 418 | + static_cast<long double>(gradOutputH) * | ||
| 419 | + static_cast<long double>(gradOutputW) * COST_HIGH_AXIS_PADDING_FACTOR; | ||
| 420 | + | ||
| 421 | + cost = AddCostPenalties( | ||
| 422 | + cost, highAxisInner, highAxisOuter, highAxisTail, | ||
| 423 | + hOutputInner, wOutputInner, | ||
| 424 | + blockNum, computeVl, | ||
| 425 | + normalCoreProcessNum, oneBlockWork); | ||
| 426 | + return cost; | ||
| 427 | +} | ||
| 428 | + | ||
| 429 | +long double AdaptiveAvgPool2dGradTilingSmallKernel::AddCostPenalties( | ||
| 430 | + long double cost, int64_t highAxisInner, int64_t highAxisOuter, | ||
| 431 | + int64_t highAxisTail, | ||
| 432 | + int64_t hOutputInner, int64_t wOutputInner, | ||
| 433 | + int64_t blockNum, int64_t computeVl, | ||
| 434 | + int64_t normalCoreProcessNum, int64_t oneBlockWork) | ||
| 435 | +{ | ||
| 436 | + if (highAxisOuter > 1 && highAxisTail < computeVl) { | ||
| 437 | + if (highAxisOuter >= HIGH_AXIS_TAIL_OPT_THRESHOLD && | ||
| 438 | + gradInputW > gradOutputW) { | ||
| 439 | + cost += static_cast<long double>(normalCoreProcessNum) * | ||
| 440 | + static_cast<long double>(oneBlockWork) / | ||
| 441 | + static_cast<long double>(highAxisOuter); | ||
| 442 | + } else { | ||
| 443 | + cost += static_cast<long double>(normalCoreProcessNum) * | ||
| 444 | + static_cast<long double>(oneBlockWork); | ||
| 445 | + } | ||
| 446 | + } | ||
| 447 | + | ||
| 448 | + if (computeVl > 0 && highAxisInner % computeVl != 0) { | ||
| 449 | + const long double partialVlPenalty = gradInputW > gradOutputW ? | ||
| 450 | + COST_PARTIAL_VL_PENALTY_HW : COST_PARTIAL_VL_PENALTY_BASELINE; | ||
| 451 | + cost += static_cast<long double>(highAxisInner % computeVl) * | ||
| 452 | + static_cast<long double>(normalCoreProcessNum) * partialVlPenalty; | ||
| 453 | + } | ||
| 454 | + | ||
| 455 | + if (gradInputW > gradOutputW && | ||
| 456 | + gradInputW >= gradOutputW * STRONG_COLLAPSE_RATIO) { | ||
| 457 | + const int64_t alignedOutputRow = Ops::Base::CeilAlign( | ||
| 458 | + hOutputInner * | ||
| 459 | + Ops::Base::CeilAlign(wOutputInner, baseData.maxDataNumInOneBlock), | ||
| 460 | + TRANS_ADDR_LEN); | ||
| 461 | + cost += static_cast<long double>(blockNum) * | ||
| 462 | + static_cast<long double>(alignedOutputRow) * | ||
| 463 | + static_cast<long double>(highAxisInner) * COST_TRANS_ALIGN_FACTOR; | ||
| 464 | + } | ||
| 465 | + | ||
| 466 | + const int64_t idleCoreNum = | ||
| 467 | + baseData.totalCoreNum - | ||
| 468 | + Ops::Base::CeilDiv(blockNum, normalCoreProcessNum); | ||
| 469 | + cost += static_cast<long double>(std::max<int64_t>(0, idleCoreNum)) * | ||
| 470 | + static_cast<long double>(oneBlockWork) * COST_IDLE_CORE_FACTOR; | ||
| 471 | + | ||
| 472 | + if (hOutputInner == LIMIT && gradOutputH > LIMIT) { | ||
| 473 | + cost += static_cast<long double>(normalCoreProcessNum) * | ||
| 474 | + static_cast<long double>(oneBlockWork); | ||
| 475 | + } | ||
| 476 | + | ||
| 477 | + return cost; | ||
| 478 | +} | ||
| 479 | + | ||
| 480 | +bool AdaptiveAvgPool2dGradTilingSmallKernel::TryRecordBetterTiling( | ||
| 481 | + long double cost, int64_t hOutputInner, int64_t wOutputInner, | ||
| 482 | + int64_t blockNum, int64_t usedCoreNum, | ||
| 483 | + int64_t highAxisInner, int64_t highAxisPadding, int64_t highAxisTail, | ||
| 484 | + int64_t& bestHighAxisInner, int64_t& bestHOutputInner, int64_t& bestWOutputInner, | ||
| 485 | + int64_t& bestBlockNum, int64_t& bestUsedCoreNum, | ||
| 486 | + int64_t& bestHighAxisPadding, int64_t& bestHighAxisTail, | ||
| 487 | + int64_t& bestBufferSize, long double& bestCost, | ||
| 488 | + bool& found) | ||
| 489 | +{ | ||
| 490 | + bool better = false; | ||
| 491 | + if (!found || cost < bestCost) { | ||
| 492 | + better = true; | ||
| 493 | + } else if (cost == bestCost) { | ||
| 494 | + const int64_t curArea = hOutputInner * wOutputInner; | ||
| 495 | + const int64_t bestArea = bestHOutputInner * bestWOutputInner; | ||
| 496 | + if (blockNum < bestBlockNum || | ||
| 497 | + (blockNum == bestBlockNum && curArea > bestArea) || | ||
| 498 | + (blockNum == bestBlockNum && curArea == bestArea && | ||
| 499 | + highAxisPadding < bestHighAxisPadding)) { | ||
| 500 | + better = true; | ||
| 501 | + } | ||
| 502 | + } | ||
| 503 | + | ||
| 504 | + if (!better) { | ||
| 505 | + return false; | ||
| 506 | + } | ||
| 507 | + | ||
| 508 | + found = true; | ||
| 509 | + bestCost = cost; | ||
| 510 | + bestHighAxisInner = highAxisInner; | ||
| 511 | + bestHOutputInner = hOutputInner; | ||
| 512 | + bestWOutputInner = wOutputInner; | ||
| 513 | + bestBlockNum = blockNum; | ||
| 514 | + bestUsedCoreNum = usedCoreNum; | ||
| 515 | + bestHighAxisPadding = highAxisPadding; | ||
| 516 | + bestHighAxisTail = highAxisTail; | ||
| 517 | + bestBufferSize = splitData.totalBufferSize; | ||
| 518 | + return true; | ||
| 519 | +} | ||
| 520 | + | ||
| 521 | +void AdaptiveAvgPool2dGradTilingSmallKernel::ApplyCoarseFallback() | ||
| 522 | +{ | ||
| 523 | + splitData.highAxisInner = baseData.proDataNumInOneBeatT2; | ||
| 524 | + splitData.hOutputInner = gradOutputH; | ||
| 525 | + splitData.wOutputInner = gradOutputW; | ||
| 526 | + | ||
| 527 | + while (splitData.hOutputInner > kernelH || splitData.wOutputInner > kernelW) { | ||
| 528 | + if (IsMeetTargetCoreNum() && IsMeetUBSize()) { | ||
| 529 | + return; | ||
| 530 | + } | ||
| 531 | + | ||
| 532 | + if (splitData.hOutputInner > kernelH) { | ||
| 533 | + splitData.hOutputInner -= kernelH; | ||
| 534 | + continue; | ||
| 535 | + } | ||
| 536 | + | ||
| 537 | + if (splitData.wOutputInner > kernelW) { | ||
| 538 | + splitData.wOutputInner -= kernelW; | ||
| 539 | + continue; | ||
| 540 | + } | ||
| 541 | + } | ||
| 542 | + | ||
| 543 | + if (IsMeetUBSize() && IsMeetTargetCoreNum()) { | ||
| 544 | + return; | ||
| 545 | + } | ||
| 546 | + | ||
| 192 | SplitUnalignHW(); | 547 | SplitUnalignHW(); |
| 193 | } | 548 | } |
| 194 | 549 | ||
| @@ -255,34 +610,15 @@ void AdaptiveAvgPool2dGradTilingSmallKernel::SetTilingData() | |||
| 255 | 610 | ||
| 256 | void AdaptiveAvgPool2dGradTilingSmallKernel::PrintSplitData() const | 611 | void AdaptiveAvgPool2dGradTilingSmallKernel::PrintSplitData() const |
| 257 | { | 612 | { |
| 258 | - OP_LOGD("AdaptiveAvgPool2dGradNCHW", "[AdaptiveAvgPool2dGradNCHW] PrintSplitData start running"); | 613 | + const int64_t highAxisTotalCapacity = splitData.highAxisOuter * splitData.highAxisInner; |
| 614 | + const int64_t highAxisPadding = highAxisTotalCapacity - baseData.inputNCSize; | ||
| 615 | + const double highAxisValidRate = highAxisTotalCapacity == 0 ? 0.0 : | ||
| 616 | + static_cast<double>(baseData.inputNCSize) / static_cast<double>(highAxisTotalCapacity); | ||
| 617 | + const double ubUseRate = baseData.availableUb == 0 ? 0.0 : | ||
| 618 | + static_cast<double>(splitData.totalBufferSize) / static_cast<double>(baseData.availableUb); | ||
| 619 | + const double coreUseRate = baseData.totalCoreNum == 0 ? 0.0 : | ||
| 620 | + static_cast<double>(splitData.usedCoreNum) / static_cast<double>(baseData.totalCoreNum); | ||
| 259 | 621 | ||
| 260 | - std::ostringstream info; | ||
| 261 | - info << "baseData.availableUb: " << baseData.availableUb << std::endl; | ||
| 262 | - | ||
| 263 | - info << "splitData.highAxisInner: " << splitData.highAxisInner << std::endl; | ||
| 264 | - info << "splitData.highAxisTail: " << splitData.highAxisTail << std::endl; | ||
| 265 | - info << "splitData.highAxisOuter: " << splitData.highAxisOuter << std::endl; | ||
| 266 | - | ||
| 267 | - info << "splitData.hOutputInner: " << splitData.hOutputInner << std::endl; | ||
| 268 | - info << "splitData.hOutputTail: " << splitData.hOutputTail << std::endl; | ||
| 269 | - info << "splitData.hOutputOuter: " << splitData.hOutputOuter << std::endl; | ||
| 270 | - | ||
| 271 | - info << "splitData.wOutputInner: " << splitData.wOutputInner << std::endl; | ||
| 272 | - info << "splitData.wOutputTail: " << splitData.wOutputTail << std::endl; | ||
| 273 | - info << "splitData.wOutputOuter: " << splitData.wOutputOuter << std::endl; | ||
| 274 | - | ||
| 275 | - info << "splitData.normalCoreProcessNum: " << splitData.normalCoreProcessNum << std::endl; | ||
| 276 | - info << "splitData.tailCoreProcessNum: " << splitData.tailCoreProcessNum << std::endl; | ||
| 277 | - info << "splitData.usedCoreNum: " << splitData.usedCoreNum << std::endl; | ||
| 278 | - info << "splitData.totalBaseBlockNum: " << splitData.totalBaseBlockNum << std::endl; | ||
| 279 | - | ||
| 280 | - info << "splitData.inputQueBufferSize: " << splitData.inputQueBufferSize << std::endl; | ||
| 281 | - info << "splitData.transQueBufferSize: " << splitData.transQueBufferSize << std::endl; | ||
| 282 | - info << "splitData.transOutQueBufferSize: " << splitData.transOutQueBufferSize << std::endl; | ||
| 283 | - info << "splitData.totalBufferSize: " << splitData.totalBufferSize << std::endl; | ||
| 284 | - | ||
| 285 | - OP_LOGI("AdaptiveAvgPool2dGradNCHW", "%s", info.str().c_str()); | ||
| 286 | } | 622 | } |
 | |||
| 287 | 623 | ||
| 288 | ge::graphStatus AdaptiveAvgPool2dGradTilingSmallKernel::DoOpTiling() | 624 | ge::graphStatus AdaptiveAvgPool2dGradTilingSmallKernel::DoOpTiling() |
| @@ -312,7 +648,8 @@ uint64_t AdaptiveAvgPool2dGradTilingSmallKernel::GetTilingKey() const | |||
| 312 | 648 | ||
| 313 | ge::graphStatus AdaptiveAvgPool2dGradTilingSmallKernel::PostTiling() | 649 | ge::graphStatus AdaptiveAvgPool2dGradTilingSmallKernel::PostTiling() |
| 314 | { | 650 | { |
| 315 | - context_->SetTilingKey(GetTilingKey()); | 651 | + const uint64_t tilingKey = GetTilingKey(); |
| 652 | + context_->SetTilingKey(tilingKey); | ||
| 316 | context_->SetBlockDim(tilingData->usedCoreNum); | 653 | context_->SetBlockDim(tilingData->usedCoreNum); |
| 317 | return ge::GRAPH_SUCCESS; | 654 | return ge::GRAPH_SUCCESS; |
| 318 | } | 655 | } |
| @@ -323,4 +660,4 @@ ge::graphStatus AdaptiveAvgPool2dGradTilingSmallKernel::DoLibApiTiling() | |||
| 323 | } | 660 | } |
| 324 | 661 | ||
| 325 | REGISTER_OPS_TILING_TEMPLATE(AdaptiveAvgPool2dGrad, AdaptiveAvgPool2dGradTilingSmallKernel, 20); | 662 | REGISTER_OPS_TILING_TEMPLATE(AdaptiveAvgPool2dGrad, AdaptiveAvgPool2dGradTilingSmallKernel, 20); |
| 326 | -} // namespace optiling | 663 | +} // namespace optiling |
Mpooling/adaptive_avg_pool2d_grad/op_host/arch35/adaptive_avg_pool2d_grad_nchw_small_kernel_tiling.h+35-0
| @@ -115,6 +115,41 @@ protected: | |||
| 115 | void DynamicAdjustmentHW(); | 115 | void DynamicAdjustmentHW(); |
| 116 | 116 | ||
| 117 | void SearchBestTiling(); | 117 | void SearchBestTiling(); |
| 118 | + | ||
| 119 | + bool ExhaustiveSearchBestTiling( | ||
| 120 | + int64_t computeVl, int64_t ncSearchMax, | ||
| 121 | + int64_t& bestHighAxisInner, int64_t& bestHOutputInner, int64_t& bestWOutputInner, | ||
| 122 | + int64_t& bestBlockNum, int64_t& bestUsedCoreNum, | ||
| 123 | + int64_t& bestHighAxisPadding, int64_t& bestHighAxisTail, | ||
| 124 | + int64_t& bestBufferSize, long double& bestCost, bool& found); | ||
| 125 | + | ||
| 126 | + long double EvalTilingCandidate( | ||
| 127 | + int64_t highAxisInner, int64_t highAxisOuter, int64_t highAxisTail, | ||
| 128 | + int64_t highAxisPadding, | ||
| 129 | + int64_t hOutputInner, int64_t hOutputOuter, | ||
| 130 | + int64_t wOutputInner, int64_t wOutputOuter, | ||
| 131 | + int64_t blockNum, int64_t computeVl, | ||
| 132 | + int64_t normalCoreProcessNum); | ||
| 133 | + | ||
| 134 | + long double AddCostPenalties( | ||
| 135 | + long double cost, int64_t highAxisInner, int64_t highAxisOuter, | ||
| 136 | + int64_t highAxisTail, | ||
| 137 | + int64_t hOutputInner, int64_t wOutputInner, | ||
| 138 | + int64_t blockNum, int64_t computeVl, | ||
| 139 | + int64_t normalCoreProcessNum, int64_t oneBlockWork); | ||
| 140 | + | ||
| 141 | + bool TryRecordBetterTiling( | ||
| 142 | + long double cost, int64_t hOutputInner, int64_t wOutputInner, | ||
| 143 | + int64_t blockNum, int64_t usedCoreNum, | ||
| 144 | + int64_t highAxisInner, int64_t highAxisPadding, int64_t highAxisTail, | ||
| 145 | + int64_t& bestHighAxisInner, int64_t& bestHOutputInner, int64_t& bestWOutputInner, | ||
| 146 | + int64_t& bestBlockNum, int64_t& bestUsedCoreNum, | ||
| 147 | + int64_t& bestHighAxisPadding, int64_t& bestHighAxisTail, | ||
| 148 | + int64_t& bestBufferSize, long double& bestCost, | ||
| 149 | + bool& found); | ||
| 150 | + | ||
| 151 | + void ApplyCoarseFallback(); | ||
| 152 | + | ||
| 118 | void DoUBTiling(); | 153 | void DoUBTiling(); |
| 119 | void DoBlockTiling(); | 154 | void DoBlockTiling(); |
| 120 | void SetTilingData(); | 155 | void SetTilingData(); |
Mpooling/adaptive_avg_pool2d_grad/op_kernel/arch35/adaptive_avg_pool2d_grad_nchw_small_kernel.h+141-87
| @@ -3,8 +3,8 @@ | |||
| 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 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. |
| 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. |
| 9 | */ | 9 | */ |
| 10 | 10 | ||
| @@ -56,10 +56,6 @@ public: | |||
| 56 | __aicore__ inline void TransposeB32(LocalTensor<I> dst, LocalTensor<I> src, uint32_t rowNum, uint32_t colNum); | 56 | __aicore__ inline void TransposeB32(LocalTensor<I> dst, LocalTensor<I> src, uint32_t rowNum, uint32_t colNum); |
| 57 | 57 | ||
| 58 | private: | 58 | private: |
| 59 | - __aicore__ inline void CalcOutputRangeFromInputIndex( | ||
| 60 | - int64_t inputIdxGlobal, int64_t outputSize, int64_t inputSize, int64_t axisTileIndex, int64_t axisInner, | ||
| 61 | - int64_t axisOutputActual, int64_t& stLocal, int64_t& edLocal, int64_t& coverCount) const; | ||
| 62 | - | ||
| 63 | __aicore__ inline void AccumulateOutputRowsForInputPointRegFp32( | 59 | __aicore__ inline void AccumulateOutputRowsForInputPointRegFp32( |
| 64 | LocalTensor<COMPUTE_TYPE> srcLocal, LocalTensor<COMPUTE_TYPE> dstLocal, int64_t inBase, COMPUTE_TYPE scale, | 60 | LocalTensor<COMPUTE_TYPE> srcLocal, LocalTensor<COMPUTE_TYPE> dstLocal, int64_t inBase, COMPUTE_TYPE scale, |
| 65 | int64_t stH, int64_t edH, int64_t stW, int64_t edW); | 61 | int64_t stH, int64_t edH, int64_t stW, int64_t edW); |
| @@ -76,6 +72,10 @@ private: | |||
| 76 | TBuf<QuePosition::VECCALC> stWRegBuf_; | 72 | TBuf<QuePosition::VECCALC> stWRegBuf_; |
| 77 | TBuf<QuePosition::VECCALC> edWRegBuf_; | 73 | TBuf<QuePosition::VECCALC> edWRegBuf_; |
| 78 | TBuf<QuePosition::VECCALC> coverWRegBuf_; | 74 | TBuf<QuePosition::VECCALC> coverWRegBuf_; |
| 75 | + TBuf<QuePosition::VECCALC> stHRegBuf_; | ||
| 76 | + TBuf<QuePosition::VECCALC> edHRegBuf_; | ||
| 77 | + TBuf<QuePosition::VECCALC> coverHRegBuf_; | ||
| 78 | + TBuf<QuePosition::VECCALC> invCoverWRegBuf_; | ||
| 79 | 79 | ||
| 80 | GlobalTensor<T> gradInputGm_; | 80 | GlobalTensor<T> gradInputGm_; |
| 81 | GlobalTensor<T> yGm_; | 81 | GlobalTensor<T> yGm_; |
| @@ -85,6 +85,7 @@ private: | |||
| 85 | uint32_t blockIdx_ = 0; | 85 | uint32_t blockIdx_ = 0; |
| 86 | 86 | ||
| 87 | int64_t highAxisActual_ = 1; | 87 | int64_t highAxisActual_ = 1; |
| 88 | + int64_t highAxisLocalStride_ = 1; | ||
| 88 | int64_t hOutputActual_ = 1; | 89 | int64_t hOutputActual_ = 1; |
| 89 | int64_t wOutputActual_ = 1; | 90 | int64_t wOutputActual_ = 1; |
| 90 | int64_t curCoreProcessNum_ = 1; | 91 | int64_t curCoreProcessNum_ = 1; |
| @@ -148,6 +149,10 @@ __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::Init( | |||
| 148 | pipe_.InitBuffer(stWRegBuf_, platform::GetVRegSize()); | 149 | pipe_.InitBuffer(stWRegBuf_, platform::GetVRegSize()); |
| 149 | pipe_.InitBuffer(edWRegBuf_, platform::GetVRegSize()); | 150 | pipe_.InitBuffer(edWRegBuf_, platform::GetVRegSize()); |
| 150 | pipe_.InitBuffer(coverWRegBuf_, platform::GetVRegSize()); | 151 | pipe_.InitBuffer(coverWRegBuf_, platform::GetVRegSize()); |
| 152 | + pipe_.InitBuffer(stHRegBuf_, platform::GetVRegSize()); | ||
| 153 | + pipe_.InitBuffer(edHRegBuf_, platform::GetVRegSize()); | ||
| 154 | + pipe_.InitBuffer(coverHRegBuf_, platform::GetVRegSize()); | ||
| 155 | + pipe_.InitBuffer(invCoverWRegBuf_, platform::GetVRegSize()); | ||
| 151 | } | 156 | } |
| 152 | 157 | ||
| 153 | template <typename T, typename INDEX> | 158 | template <typename T, typename INDEX> |
| @@ -157,6 +162,7 @@ __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::ScalarCom | |||
| 157 | 162 | ||
| 158 | highAxisIndex_ = baseBlockIdx / (tiling_->hOutputOuter * tiling_->wOutputOuter); | 163 | highAxisIndex_ = baseBlockIdx / (tiling_->hOutputOuter * tiling_->wOutputOuter); |
| 159 | highAxisActual_ = (highAxisIndex_ == (tiling_->highAxisOuter - 1)) ? tiling_->highAxisTail : tiling_->highAxisInner; | 164 | highAxisActual_ = (highAxisIndex_ == (tiling_->highAxisOuter - 1)) ? tiling_->highAxisTail : tiling_->highAxisInner; |
| 165 | + highAxisLocalStride_ = CeilAlign(highAxisActual_, static_cast<int64_t>(TRANS_ADDR_LEN)); | ||
| 160 | 166 | ||
| 161 | int64_t tempTail = baseBlockIdx % (tiling_->hOutputOuter * tiling_->wOutputOuter); | 167 | int64_t tempTail = baseBlockIdx % (tiling_->hOutputOuter * tiling_->wOutputOuter); |
| 162 | hAxisIndex_ = tempTail / tiling_->wOutputOuter; | 168 | hAxisIndex_ = tempTail / tiling_->wOutputOuter; |
| @@ -335,25 +341,6 @@ __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::TransInpu | |||
| 335 | transQue_.EnQue(dstLocal); | 341 | transQue_.EnQue(dstLocal); |
| 336 | } | 342 | } |
| 337 | 343 | ||
| 338 | -template <typename T, typename INDEX> | ||
| 339 | -__aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::CalcOutputRangeFromInputIndex( | ||
| 340 | - int64_t inputIdxGlobal, int64_t outputSize, int64_t inputSize, int64_t axisTileIndex, int64_t axisInner, | ||
| 341 | - int64_t axisOutputActual, int64_t& stLocal, int64_t& edLocal, int64_t& coverCount) const | ||
| 342 | -{ | ||
| 343 | - const int64_t stGlobal = GetStartFromOutputInputSize(inputIdxGlobal, outputSize, inputSize); | ||
| 344 | - const int64_t edGlobal = GetEndFromOutputInputSize(inputIdxGlobal, outputSize, inputSize); | ||
| 345 | - | ||
| 346 | - const int64_t tileStart = axisTileIndex * axisInner; | ||
| 347 | - const int64_t tileEnd = tileStart + axisOutputActual; | ||
| 348 | - | ||
| 349 | - const int64_t stClamped = stGlobal > tileStart ? stGlobal : tileStart; | ||
| 350 | - const int64_t edClamped = edGlobal < tileEnd ? edGlobal : tileEnd; | ||
| 351 | - | ||
| 352 | - stLocal = stClamped - tileStart; | ||
| 353 | - edLocal = edClamped - tileStart; | ||
| 354 | - coverCount = edGlobal - stGlobal; | ||
| 355 | -} | ||
| 356 | - | ||
| 357 | template <typename T, typename INDEX> | 344 | template <typename T, typename INDEX> |
| 358 | __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::AccumulateOutputRowsForInputPointRegFp32( | 345 | __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::AccumulateOutputRowsForInputPointRegFp32( |
| 359 | LocalTensor<COMPUTE_TYPE> srcLocal, LocalTensor<COMPUTE_TYPE> dstLocal, int64_t inBase, COMPUTE_TYPE scale, | 346 | LocalTensor<COMPUTE_TYPE> srcLocal, LocalTensor<COMPUTE_TYPE> dstLocal, int64_t inBase, COMPUTE_TYPE scale, |
| @@ -386,7 +373,7 @@ __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::Accumulat | |||
| 386 | const int64_t hRowBase = hRowBase0 + static_cast<int64_t>(oh) * wOutputAligned_; | 373 | const int64_t hRowBase = hRowBase0 + static_cast<int64_t>(oh) * wOutputAligned_; |
| 387 | for (uint16_t ow = 0; ow < wLoopCount; ++ow) { | 374 | for (uint16_t ow = 0; ow < wLoopCount; ++ow) { |
| 388 | const int64_t outRow = hRowBase + static_cast<int64_t>(stW + ow); | 375 | const int64_t outRow = hRowBase + static_cast<int64_t>(stW + ow); |
| 389 | - const int64_t outBase = outRow * tiling_->highAxisInner; | 376 | + const int64_t outBase = outRow * highAxisLocalStride_; |
| 390 | __local_mem__ COMPUTE_TYPE* dstAddr = | 377 | __local_mem__ COMPUTE_TYPE* dstAddr = |
| 391 | (__local_mem__ COMPUTE_TYPE*)dstLocal[outBase + processed].GetPhyAddr(); | 378 | (__local_mem__ COMPUTE_TYPE*)dstLocal[outBase + processed].GetPhyAddr(); |
| 392 | 379 | ||
| @@ -417,7 +404,7 @@ __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::Accumulat | |||
| 417 | const int64_t hRowBase = hRowBase0 + static_cast<int64_t>(oh) * wOutputAligned_; | 404 | const int64_t hRowBase = hRowBase0 + static_cast<int64_t>(oh) * wOutputAligned_; |
| 418 | for (uint16_t ow = 0; ow < wLoopCount; ++ow) { | 405 | for (uint16_t ow = 0; ow < wLoopCount; ++ow) { |
| 419 | const int64_t outRow = hRowBase + static_cast<int64_t>(stW + ow); | 406 | const int64_t outRow = hRowBase + static_cast<int64_t>(stW + ow); |
| 420 | - const int64_t outBase = outRow * tiling_->highAxisInner; | 407 | + const int64_t outBase = outRow * highAxisLocalStride_; |
| 421 | __local_mem__ COMPUTE_TYPE* dstAddr = | 408 | __local_mem__ COMPUTE_TYPE* dstAddr = |
| 422 | (__local_mem__ COMPUTE_TYPE*)dstLocal[outBase + processed].GetPhyAddr(); | 409 | (__local_mem__ COMPUTE_TYPE*)dstLocal[outBase + processed].GetPhyAddr(); |
| 423 | 410 | ||
| @@ -445,92 +432,159 @@ __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::ComputeFp | |||
| 445 | __local_mem__ INDEX* edWAddr = reinterpret_cast<__local_mem__ INDEX*>(edWLocal.GetPhyAddr()); | 432 | __local_mem__ INDEX* edWAddr = reinterpret_cast<__local_mem__ INDEX*>(edWLocal.GetPhyAddr()); |
| 446 | __local_mem__ INDEX* coverWAddr = reinterpret_cast<__local_mem__ INDEX*>(coverWLocal.GetPhyAddr()); | 433 | __local_mem__ INDEX* coverWAddr = reinterpret_cast<__local_mem__ INDEX*>(coverWLocal.GetPhyAddr()); |
| 447 | 434 | ||
| 435 | + LocalTensor<INDEX> stHLocal = stHRegBuf_.Get<INDEX>(); | ||
| 436 | + LocalTensor<INDEX> edHLocal = edHRegBuf_.Get<INDEX>(); | ||
| 437 | + LocalTensor<INDEX> coverHLocal = coverHRegBuf_.Get<INDEX>(); | ||
| 438 | + | ||
| 439 | + __local_mem__ INDEX* stHAddr = reinterpret_cast<__local_mem__ INDEX*>(stHLocal.GetPhyAddr()); | ||
| 440 | + __local_mem__ INDEX* edHAddr = reinterpret_cast<__local_mem__ INDEX*>(edHLocal.GetPhyAddr()); | ||
| 441 | + __local_mem__ INDEX* coverHAddr = reinterpret_cast<__local_mem__ INDEX*>(coverHLocal.GetPhyAddr()); | ||
| 442 | + | ||
| 443 | + LocalTensor<COMPUTE_TYPE> invCoverWLocal = invCoverWRegBuf_.Get<COMPUTE_TYPE>(); | ||
| 444 | + __local_mem__ COMPUTE_TYPE* invCoverWAddr = | ||
| 445 | + reinterpret_cast<__local_mem__ COMPUTE_TYPE*>(invCoverWLocal.GetPhyAddr()); | ||
| 446 | + | ||
| 448 | const INDEX wTileStart = static_cast<INDEX>(wAxisIndex_ * tiling_->wOutputInner); | 447 | const INDEX wTileStart = static_cast<INDEX>(wAxisIndex_ * tiling_->wOutputInner); |
| 449 | const INDEX wTileEnd = static_cast<INDEX>(wTileStart + wOutputActual_); | 448 | const INDEX wTileEnd = static_cast<INDEX>(wTileStart + wOutputActual_); |
| 450 | const INDEX wOutput = static_cast<INDEX>(tiling_->wOutput); | 449 | const INDEX wOutput = static_cast<INDEX>(tiling_->wOutput); |
| 451 | const INDEX wGradInput = static_cast<INDEX>(tiling_->wInput); | 450 | const INDEX wGradInput = static_cast<INDEX>(tiling_->wInput); |
| 452 | 451 | ||
| 453 | - for (int64_t swLocalBatch = 0; swLocalBatch < wGradInputActual_; swLocalBatch += INDEX_VF_LEN) { | 452 | + const INDEX hTileStart = static_cast<INDEX>(hAxisIndex_ * tiling_->hOutputInner); |
| 454 | - int64_t curBatchCount = wGradInputActual_ - swLocalBatch; | 453 | + const INDEX hTileEnd = static_cast<INDEX>(hTileStart + hOutputActual_); |
| 455 | - curBatchCount = curBatchCount > INDEX_VF_LEN ? INDEX_VF_LEN : curBatchCount; | 454 | + const INDEX hOutput = static_cast<INDEX>(tiling_->hOutput); |
| 455 | + const INDEX hGradInput = static_cast<INDEX>(tiling_->hInput); | ||
| 456 | 456 | ||
| 457 | - const INDEX wBaseIdx = static_cast<INDEX>(wStLeftCornerIdx_ + swLocalBatch); | 457 | + for (int64_t shLocalBatch = 0; shLocalBatch < hGradInputActual_; shLocalBatch += INDEX_VF_LEN) { |
| 458 | - uint32_t batchCountMask = static_cast<uint32_t>(curBatchCount); | 458 | + int64_t curHBatchCount = hGradInputActual_ - shLocalBatch; |
| 459 | + curHBatchCount = curHBatchCount > INDEX_VF_LEN ? INDEX_VF_LEN : curHBatchCount; | ||
| 460 | + | ||
| 461 | + const INDEX hBaseIdx = static_cast<INDEX>(hStLeftCornerIdx_ + shLocalBatch); | ||
| 462 | + uint32_t hBatchCountMask = static_cast<uint32_t>(curHBatchCount); | ||
| 459 | 463 | ||
| 460 | __VEC_SCOPE__ | 464 | __VEC_SCOPE__ |
| 461 | { | 465 | { |
| 462 | - MicroAPI::RegTensor<INDEX> idx; | 466 | + MicroAPI::RegTensor<INDEX> hIdx; |
| 463 | - MicroAPI::Arange(idx, wBaseIdx); | 467 | + MicroAPI::Arange(hIdx, hBaseIdx); |
| 464 | 468 | ||
| 465 | - MicroAPI::MaskReg allMask = MicroAPI::CreateMask<INDEX, MicroAPI::MaskPattern::ALL>(); | 469 | + MicroAPI::MaskReg hAllMask = MicroAPI::CreateMask<INDEX, MicroAPI::MaskPattern::ALL>(); |
| 466 | - MicroAPI::MaskReg batchMask = MicroAPI::UpdateMask<INDEX>(batchCountMask); | 470 | + MicroAPI::MaskReg hBatchMask = MicroAPI::UpdateMask<INDEX>(hBatchCountMask); |
| 467 | 471 | ||
| 468 | - MicroAPI::RegTensor<INDEX> regConstOutput; | 472 | + MicroAPI::RegTensor<INDEX> hRegOut; |
| 469 | - MicroAPI::Duplicate(regConstOutput, wOutput); | 473 | + MicroAPI::Duplicate(hRegOut, hOutput); |
| 470 | 474 | ||
| 471 | - MicroAPI::RegTensor<INDEX> regConstInput; | 475 | + MicroAPI::RegTensor<INDEX> hRegIn; |
| 472 | - MicroAPI::Duplicate(regConstInput, wGradInput); | 476 | + MicroAPI::Duplicate(hRegIn, hGradInput); |
| 473 | 477 | ||
| 474 | - MicroAPI::RegTensor<INDEX> stGlobal; | 478 | + MicroAPI::RegTensor<INDEX> hStGlobal; |
| 475 | - MicroAPI::Mul(stGlobal, idx, regConstOutput, allMask); | 479 | + MicroAPI::Mul(hStGlobal, hIdx, hRegOut, hAllMask); |
| 476 | - MicroAPI::Div(stGlobal, stGlobal, regConstInput, allMask); | 480 | + MicroAPI::Div(hStGlobal, hStGlobal, hRegIn, hAllMask); |
| 477 | 481 | ||
| 478 | - MicroAPI::RegTensor<INDEX> edGlobal; | 482 | + MicroAPI::RegTensor<INDEX> hEdGlobal; |
| 479 | - MicroAPI::Adds(edGlobal, idx, INDEX(1), allMask); | 483 | + MicroAPI::Adds(hEdGlobal, hIdx, INDEX(1), hAllMask); |
| 480 | - MicroAPI::Mul(edGlobal, edGlobal, regConstOutput, allMask); | 484 | + MicroAPI::Mul(hEdGlobal, hEdGlobal, hRegOut, hAllMask); |
| 481 | - MicroAPI::Add(edGlobal, edGlobal, regConstInput, allMask); | 485 | + MicroAPI::Add(hEdGlobal, hEdGlobal, hRegIn, hAllMask); |
| 482 | - MicroAPI::Adds(edGlobal, edGlobal, INDEX(-1), allMask); | 486 | + MicroAPI::Adds(hEdGlobal, hEdGlobal, INDEX(-1), hAllMask); |
| 483 | - MicroAPI::Div(edGlobal, edGlobal, regConstInput, allMask); | 487 | + MicroAPI::Div(hEdGlobal, hEdGlobal, hRegIn, hAllMask); |
| 484 | 488 | ||
| 485 | - MicroAPI::RegTensor<INDEX> cover; | 489 | + MicroAPI::RegTensor<INDEX> hCover; |
| 486 | - MicroAPI::Sub(cover, edGlobal, stGlobal, allMask); | 490 | + MicroAPI::Sub(hCover, hEdGlobal, hStGlobal, hAllMask); |
| 487 | 491 | ||
| 488 | - MicroAPI::Maxs(stGlobal, stGlobal, wTileStart, allMask); | 492 | + MicroAPI::Maxs(hStGlobal, hStGlobal, hTileStart, hAllMask); |
| 489 | - MicroAPI::Adds(stGlobal, stGlobal, INDEX(-wTileStart), allMask); | 493 | + MicroAPI::Adds(hStGlobal, hStGlobal, INDEX(-hTileStart), hAllMask); |
| 490 | 494 | ||
| 491 | - MicroAPI::Mins(edGlobal, edGlobal, wTileEnd, allMask); | 495 | + MicroAPI::Mins(hEdGlobal, hEdGlobal, hTileEnd, hAllMask); |
| 492 | - MicroAPI::Adds(edGlobal, edGlobal, INDEX(-wTileStart), allMask); | 496 | + MicroAPI::Adds(hEdGlobal, hEdGlobal, INDEX(-hTileStart), hAllMask); |
| 493 | 497 | ||
| 494 | - MicroAPI::DataCopy(stWAddr, stGlobal, batchMask); | 498 | + MicroAPI::DataCopy(stHAddr, hStGlobal, hBatchMask); |
| 495 | - MicroAPI::DataCopy(edWAddr, edGlobal, batchMask); | 499 | + MicroAPI::DataCopy(edHAddr, hEdGlobal, hBatchMask); |
| 496 | - MicroAPI::DataCopy(coverWAddr, cover, batchMask); | 500 | + MicroAPI::DataCopy(coverHAddr, hCover, hBatchMask); |
| 497 | } | 501 | } |
| 498 | 502 | ||
| 499 | - PIPE_V_S(); | 503 | + for (int64_t swLocalBatch = 0; swLocalBatch < wGradInputActual_; swLocalBatch += INDEX_VF_LEN) { |
| 504 | + int64_t curBatchCount = wGradInputActual_ - swLocalBatch; | ||
| 505 | + curBatchCount = curBatchCount > INDEX_VF_LEN ? INDEX_VF_LEN : curBatchCount; | ||
| 500 | 506 | ||
| 501 | - for (int64_t shLocal = 0; shLocal < hGradInputActual_; ++shLocal) { | 507 | + const INDEX wBaseIdx = static_cast<INDEX>(wStLeftCornerIdx_ + swLocalBatch); |
| 502 | - int64_t stH = 0; | 508 | + uint32_t batchCountMask = static_cast<uint32_t>(curBatchCount); |
| 503 | - int64_t edH = 0; | ||
| 504 | - int64_t coverH = 0; | ||
| 505 | - CalcOutputRangeFromInputIndex( | ||
| 506 | - hStLeftCornerIdx_ + shLocal, tiling_->hOutput, tiling_->hInput, hAxisIndex_, tiling_->hOutputInner, | ||
| 507 | - hOutputActual_, stH, edH, coverH); | ||
| 508 | 509 | ||
| 509 | - if (edH <= stH || coverH <= 0) { | 510 | + __VEC_SCOPE__ |
| 510 | - continue; | 511 | + { |
| 512 | + MicroAPI::RegTensor<INDEX> idx; | ||
| 513 | + MicroAPI::Arange(idx, wBaseIdx); | ||
| 514 | + | ||
| 515 | + MicroAPI::MaskReg allMask = MicroAPI::CreateMask<INDEX, MicroAPI::MaskPattern::ALL>(); | ||
| 516 | + MicroAPI::MaskReg batchMask = MicroAPI::UpdateMask<INDEX>(batchCountMask); | ||
| 517 | + | ||
| 518 | + MicroAPI::RegTensor<INDEX> regConstOutput; | ||
| 519 | + MicroAPI::Duplicate(regConstOutput, wOutput); | ||
| 520 | + | ||
| 521 | + MicroAPI::RegTensor<INDEX> regConstInput; | ||
| 522 | + MicroAPI::Duplicate(regConstInput, wGradInput); | ||
| 523 | + | ||
| 524 | + MicroAPI::RegTensor<INDEX> stGlobal; | ||
| 525 | + MicroAPI::Mul(stGlobal, idx, regConstOutput, allMask); | ||
| 526 | + MicroAPI::Div(stGlobal, stGlobal, regConstInput, allMask); | ||
| 527 | + | ||
| 528 | + MicroAPI::RegTensor<INDEX> edGlobal; | ||
| 529 | + MicroAPI::Adds(edGlobal, idx, INDEX(1), allMask); | ||
| 530 | + MicroAPI::Mul(edGlobal, edGlobal, regConstOutput, allMask); | ||
| 531 | + MicroAPI::Add(edGlobal, edGlobal, regConstInput, allMask); | ||
| 532 | + MicroAPI::Adds(edGlobal, edGlobal, INDEX(-1), allMask); | ||
| 533 | + MicroAPI::Div(edGlobal, edGlobal, regConstInput, allMask); | ||
| 534 | + | ||
| 535 | + MicroAPI::RegTensor<INDEX> cover; | ||
| 536 | + MicroAPI::Sub(cover, edGlobal, stGlobal, allMask); | ||
| 537 | + | ||
| 538 | + MicroAPI::Maxs(stGlobal, stGlobal, wTileStart, allMask); | ||
| 539 | + MicroAPI::Adds(stGlobal, stGlobal, INDEX(-wTileStart), allMask); | ||
| 540 | + | ||
| 541 | + MicroAPI::Mins(edGlobal, edGlobal, wTileEnd, allMask); | ||
| 542 | + MicroAPI::Adds(edGlobal, edGlobal, INDEX(-wTileStart), allMask); | ||
| 543 | + | ||
| 544 | + MicroAPI::DataCopy(stWAddr, stGlobal, batchMask); | ||
| 545 | + MicroAPI::DataCopy(edWAddr, edGlobal, batchMask); | ||
| 546 | + MicroAPI::DataCopy(coverWAddr, cover, batchMask); | ||
| 511 | } | 547 | } |
| 512 | 548 | ||
| 513 | - const int64_t hBase = shLocal * wGradInputAligned_; | 549 | + PIPE_V_S(); |
| 514 | 550 | ||
| 515 | - for (int64_t wInBatch = 0; wInBatch < curBatchCount; ++wInBatch) { | 551 | + for (int64_t wi = 0; wi < curBatchCount; ++wi) { |
| 516 | - const int64_t swLocal = swLocalBatch + wInBatch; | 552 | + const int64_t cw = static_cast<int64_t>(coverWAddr[wi]); |
| 517 | - const int64_t stW = static_cast<int64_t>(stWAddr[wInBatch]); | 553 | + invCoverWAddr[wi] = (cw > 0) ? static_cast<COMPUTE_TYPE>(1.0f / static_cast<float>(cw)) : |
| 518 | - const int64_t edW = static_cast<int64_t>(edWAddr[wInBatch]); | 554 | + static_cast<COMPUTE_TYPE>(0.0f); |
| 519 | - const int64_t coverW = static_cast<int64_t>(coverWAddr[wInBatch]); | 555 | + } |
| 520 | 556 | ||
| 521 | - if (edW <= stW || coverW <= 0) { | 557 | + for (int64_t shLocal = 0; shLocal < curHBatchCount; ++shLocal) { |
| 558 | + const int64_t stH = static_cast<int64_t>(stHAddr[shLocal]); | ||
| 559 | + const int64_t edH = static_cast<int64_t>(edHAddr[shLocal]); | ||
| 560 | + const int64_t coverH = static_cast<int64_t>(coverHAddr[shLocal]); | ||
| 561 | + | ||
| 562 | + if (edH <= stH || coverH <= 0) { | ||
| 522 | continue; | 563 | continue; |
| 523 | } | 564 | } |
| 524 | 565 | ||
| 525 | - const int64_t kernelSize = coverH * coverW; | 566 | + const int64_t hIdxGlobal = shLocalBatch + shLocal; |
| 526 | - if (kernelSize <= 0) { | 567 | + const int64_t hBase = hIdxGlobal * wGradInputAligned_; |
| 527 | - continue; | 568 | + const COMPUTE_TYPE invCoverH = static_cast<COMPUTE_TYPE>(1.0f / static_cast<float>(coverH)); |
| 569 | + | ||
| 570 | + for (int64_t wInBatch = 0; wInBatch < curBatchCount; ++wInBatch) { | ||
| 571 | + const int64_t swLocal = swLocalBatch + wInBatch; | ||
| 572 | + const int64_t stW = static_cast<int64_t>(stWAddr[wInBatch]); | ||
| 573 | + const int64_t edW = static_cast<int64_t>(edWAddr[wInBatch]); | ||
| 574 | + | ||
| 575 | + if (edW <= stW) { | ||
| 576 | + continue; | ||
| 577 | + } | ||
| 578 | + | ||
| 579 | + const COMPUTE_TYPE iw = invCoverWAddr[wInBatch]; | ||
| 580 | + if (iw <= static_cast<COMPUTE_TYPE>(0.0f)) { | ||
| 581 | + continue; | ||
| 582 | + } | ||
| 583 | + | ||
| 584 | + const int64_t inBase = (hBase + swLocal) * highAxisLocalStride_; | ||
| 585 | + const COMPUTE_TYPE scale = invCoverH * iw; | ||
| 586 | + AccumulateOutputRowsForInputPointRegFp32(srcLocal, dstLocal, inBase, scale, stH, edH, stW, edW); | ||
| 528 | } | 587 | } |
| 529 | - | ||
| 530 | - const int64_t inBase = (hBase + swLocal) * tiling_->highAxisInner; | ||
| 531 | - const COMPUTE_TYPE scale = static_cast<COMPUTE_TYPE>(1.0f / static_cast<float>(kernelSize)); | ||
| 532 | - | ||
| 533 | - AccumulateOutputRowsForInputPointRegFp32(srcLocal, dstLocal, inBase, scale, stH, edH, stW, edW); | ||
| 534 | } | 588 | } |
| 535 | } | 589 | } |
| 536 | } | 590 | } |
| @@ -541,8 +595,8 @@ __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::Compute() | |||
| 541 | { | 595 | { |
| 542 | LocalTensor<T> srcLocalT = transQue_.DeQue<T>(); | 596 | LocalTensor<T> srcLocalT = transQue_.DeQue<T>(); |
| 543 | 597 | ||
| 544 | - const uint32_t srcElemCount = static_cast<uint32_t>(tiling_->highAxisInner * inputColNum_); | 598 | + const uint32_t srcElemCount = static_cast<uint32_t>(highAxisLocalStride_ * inputColNum_); |
| 545 | - const uint32_t dstElemCount = static_cast<uint32_t>(outputRowNumAligned_ * tiling_->highAxisInner); | 599 | + const uint32_t dstElemCount = static_cast<uint32_t>(outputRowNumAligned_ * highAxisLocalStride_); |
| 546 | 600 | ||
| 547 | if constexpr (std::is_same_v<T, float>) { | 601 | if constexpr (std::is_same_v<T, float>) { |
| 548 | LocalTensor<T> dstLocalT = transOutQue_.AllocTensor<T>(); | 602 | LocalTensor<T> dstLocalT = transOutQue_.AllocTensor<T>(); |
| @@ -585,7 +639,7 @@ template <typename T, typename INDEX> | |||
| 585 | __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::TransOut() | 639 | __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::TransOut() |
| 586 | { | 640 | { |
| 587 | const uint32_t rowNum = static_cast<uint32_t>(outputRowNumAligned_); | 641 | const uint32_t rowNum = static_cast<uint32_t>(outputRowNumAligned_); |
| 588 | - const uint32_t colNum = static_cast<uint32_t>(tiling_->highAxisInner); | 642 | + const uint32_t colNum = static_cast<uint32_t>(highAxisLocalStride_); |
| 589 | 643 | ||
| 590 | LocalTensor<T> srcLocal = transOutQue_.DeQue<T>(); | 644 | LocalTensor<T> srcLocal = transOutQue_.DeQue<T>(); |
| 591 | LocalTensor<T> dstLocal = transQue_.AllocTensor<T>(); | 645 | LocalTensor<T> dstLocal = transQue_.AllocTensor<T>(); |
| @@ -643,7 +697,7 @@ template <typename T, typename INDEX> | |||
| 643 | __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::ProcessPerLoop() | 697 | __aicore__ inline void AdaptiveAvgPool2dGradNCHWSmallKernel<T, INDEX>::ProcessPerLoop() |
| 644 | { | 698 | { |
| 645 | CopyIn(); | 699 | CopyIn(); |
| 646 | - TransInput(static_cast<uint32_t>(tiling_->highAxisInner), static_cast<uint32_t>(inputColNum_)); | 700 | + TransInput(static_cast<uint32_t>(highAxisLocalStride_), static_cast<uint32_t>(inputColNum_)); |
| 647 | Compute(); | 701 | Compute(); |
| 648 | TransOut(); | 702 | TransOut(); |
| 649 | CopyOut(); | 703 | CopyOut(); |
| @@ -8,11 +8,12 @@ | |||
| 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. |
| 9 | */ | 9 | */ |
| 10 | 10 | ||
| 11 | -/* ! | 11 | + /*! |
| 12 | * \file adaptive_avg_pool2d_grad_simt.h | 12 | * \file adaptive_avg_pool2d_grad_simt.h |
| 13 | - * \brief adaptive_avg_pool2d_grad implied by simt | 13 | + * \brief |
| 14 | */ | 14 | */ |
| 15 | 15 | ||
| 16 | + | ||
| 16 | 17 | ||
| 17 | 18 | ||
| 18 | 19 | ||
| @@ -26,24 +27,25 @@ using namespace AscendC; | |||
| 26 | 27 | ||
| 27 | namespace AdaptiveAvgPool2dGradOp { | 28 | namespace AdaptiveAvgPool2dGradOp { |
| 28 | constexpr static uint32_t THREAD_DIM = 1024; | 29 | constexpr static uint32_t THREAD_DIM = 1024; |
| 29 | -constexpr static uint32_t TILING_DATA_NUM = 6; | 30 | +constexpr static uint32_t SIMT_PARAMS_NUM = 64; |
| 30 | -constexpr static uint32_t SIMT_PARAMS_NUM = 32; | ||
| 31 | constexpr static uint32_t MAGIC_C_IDX = 0; | 31 | constexpr static uint32_t MAGIC_C_IDX = 0; |
| 32 | constexpr static uint32_t MAGIC_IN_H_IDX = 2; | 32 | constexpr static uint32_t MAGIC_IN_H_IDX = 2; |
| 33 | constexpr static uint32_t MAGIC_IN_W_IDX = 4; | 33 | constexpr static uint32_t MAGIC_IN_W_IDX = 4; |
| 34 | constexpr static uint32_t MAGIC_OSIZE_H_IDX = 6; | 34 | constexpr static uint32_t MAGIC_OSIZE_H_IDX = 6; |
| 35 | constexpr static uint32_t MAGIC_OSIZE_W_IDX = 8; | 35 | constexpr static uint32_t MAGIC_OSIZE_W_IDX = 8; |
| 36 | +constexpr static uint32_t MAGIC_SEG_IDX = 10; | ||
| 37 | +constexpr static uint32_t SEG_INFO_IDX = 12; | ||
| 38 | +constexpr static uint32_t SEG_INFO_STRIDE = 4; | ||
| 39 | + | ||
| 40 | +template <typename OFFSET_T> | ||
| 41 | +using DivForOffset = typename std::conditional<std::is_same<OFFSET_T, int32_t>::value, uint32_t, uint64_t>::type; | ||
| 36 | 42 | ||
| 37 | template <typename VALUE_T, typename OFFSET_T, int64_t CHANNEL_LAST> | 43 | template <typename VALUE_T, typename OFFSET_T, int64_t CHANNEL_LAST> |
| 38 | class AdaptiveAvgPool2dGradSimt { | 44 | class AdaptiveAvgPool2dGradSimt { |
| 39 | public: | 45 | public: |
| 40 | - __aicore__ inline AdaptiveAvgPool2dGradSimt( | 46 | + __aicore__ inline AdaptiveAvgPool2dGradSimt(TPipe* pipe, const AdaptiveAvgPool2dGradSimtTiling* __restrict__ tilingData) : pipe_(pipe), tilingData_(tilingData) {} |
| 41 | - TPipe* pipe, const AdaptiveAvgPool2dGradSimtTiling* __restrict__ tilingData) | ||
| 42 | - : pipe_(pipe), tilingData_(tilingData) | ||
| 43 | - {} | ||
| 44 | __aicore__ inline void Init(GM_ADDR yGrad, GM_ADDR xGrad); | 47 | __aicore__ inline void Init(GM_ADDR yGrad, GM_ADDR xGrad); |
| 45 | __aicore__ inline void Process(); | 48 | __aicore__ inline void Process(); |
| 46 | - | ||
| 47 | private: | 49 | private: |
| 48 | TPipe* pipe_; | 50 | TPipe* pipe_; |
| 49 | AscendC::GlobalTensor<VALUE_T> yGrad_; | 51 | AscendC::GlobalTensor<VALUE_T> yGrad_; |
| @@ -53,97 +55,435 @@ private: | |||
| 53 | }; | 55 | }; |
| 54 | 56 | ||
| 55 | template <typename OFFSET_T, typename DIV_T> | 57 | template <typename OFFSET_T, typename DIV_T> |
| 56 | -__simt_callee__ __aicore__ inline static OFFSET_T FloorDivMul( | 58 | +__simt_callee__ __aicore__ inline static OFFSET_T FloorDivMul(OFFSET_T numerator, OFFSET_T mulFactor, DIV_T divisorMagic, DIV_T divisorShift) |
| 57 | - OFFSET_T numerator, OFFSET_T mulFactor, DIV_T divisorMagic, DIV_T divisorShift) | ||
| 58 | { | 59 | { |
| 59 | - DIV_T wideNumerator = static_cast<DIV_T>(numerator) * static_cast<DIV_T>(mulFactor); | 60 | + return static_cast<OFFSET_T>(Simt::UintDiv<DIV_T>(static_cast<DIV_T>(numerator) * static_cast<DIV_T>(mulFactor), divisorMagic, divisorShift)); |
| 60 | - DIV_T quotient = Simt::UintDiv<DIV_T>(wideNumerator, divisorMagic, divisorShift); | ||
| 61 | - return static_cast<OFFSET_T>(quotient); | ||
| 62 | } | 61 | } |
| 63 | 62 | ||
| 64 | template <typename OFFSET_T, typename DIV_T> | 63 | template <typename OFFSET_T, typename DIV_T> |
| 65 | -__simt_callee__ __aicore__ inline static OFFSET_T CeilDivMul( | 64 | +__simt_callee__ __aicore__ inline static OFFSET_T CeilDivMul(OFFSET_T numerator, OFFSET_T mulFactor, OFFSET_T ceilAddend, DIV_T divisorMagic, DIV_T divisorShift) |
| 66 | - OFFSET_T numerator, OFFSET_T mulFactor, OFFSET_T ceilAddend, DIV_T divisorMagic, DIV_T divisorShift) | ||
| 67 | { | 65 | { |
| 68 | - DIV_T wideNumerator = | 66 | + return static_cast<OFFSET_T>(Simt::UintDiv<DIV_T>(static_cast<DIV_T>(numerator) * static_cast<DIV_T>(mulFactor) + static_cast<DIV_T>(ceilAddend), divisorMagic, divisorShift)); |
| 69 | - static_cast<DIV_T>(numerator) * static_cast<DIV_T>(mulFactor) + static_cast<DIV_T>(ceilAddend); | ||
| 70 | - DIV_T quotient = Simt::UintDiv<DIV_T>(wideNumerator, divisorMagic, divisorShift); | ||
| 71 | - return static_cast<OFFSET_T>(quotient); | ||
| 72 | } | 67 | } |
| 73 | 68 | ||
| 74 | template <typename OFFSET_T, typename DIV_T> | 69 | template <typename OFFSET_T, typename DIV_T> |
| 75 | -__simt_callee__ __aicore__ inline static OFFSET_T StartIndexIn2Out( | 70 | +__simt_callee__ __aicore__ inline static OFFSET_T StartIndexIn2Out(OFFSET_T inIdx, OFFSET_T osize, DIV_T magicIsize, DIV_T shiftIsize) |
| 76 | - OFFSET_T inIdx, OFFSET_T osize, DIV_T magicIsize, DIV_T shiftIsize) | ||
| 77 | { | 71 | { |
| 78 | return FloorDivMul<OFFSET_T, DIV_T>(inIdx, osize, magicIsize, shiftIsize); | 72 | return FloorDivMul<OFFSET_T, DIV_T>(inIdx, osize, magicIsize, shiftIsize); |
| 79 | } | 73 | } |
| 80 | 74 | ||
| 81 | template <typename OFFSET_T, typename DIV_T> | 75 | template <typename OFFSET_T, typename DIV_T> |
| 82 | -__simt_callee__ __aicore__ inline static OFFSET_T EndIndexIn2Out( | 76 | +__simt_callee__ __aicore__ inline static OFFSET_T EndIndexIn2Out(OFFSET_T inIdx, OFFSET_T isize, OFFSET_T osize, DIV_T magicIsize, DIV_T shiftIsize) |
| 83 | - OFFSET_T inIdx, OFFSET_T isize, OFFSET_T osize, DIV_T magicIsize, DIV_T shiftIsize) | ||
| 84 | { | 77 | { |
| 85 | return CeilDivMul<OFFSET_T, DIV_T>(inIdx + 1, osize, isize - 1, magicIsize, shiftIsize); | 78 | return CeilDivMul<OFFSET_T, DIV_T>(inIdx + 1, osize, isize - 1, magicIsize, shiftIsize); |
| 86 | } | 79 | } |
| 87 | 80 | ||
| 88 | template <typename OFFSET_T, typename DIV_T> | 81 | template <typename OFFSET_T, typename DIV_T> |
| 89 | -__simt_callee__ __aicore__ inline static OFFSET_T StartIndexOut2In( | 82 | +__simt_callee__ __aicore__ inline static OFFSET_T StartIndexOut2In(OFFSET_T outIdx, OFFSET_T isize, DIV_T magicOsize, DIV_T shiftOsize) |
| 90 | - OFFSET_T outIdx, OFFSET_T isize, DIV_T magicOsize, DIV_T shiftOsize) | ||
| 91 | { | 83 | { |
| 92 | return FloorDivMul<OFFSET_T, DIV_T>(outIdx, isize, magicOsize, shiftOsize); | 84 | return FloorDivMul<OFFSET_T, DIV_T>(outIdx, isize, magicOsize, shiftOsize); |
| 93 | } | 85 | } |
| 94 | 86 | ||
| 95 | template <typename OFFSET_T, typename DIV_T> | 87 | template <typename OFFSET_T, typename DIV_T> |
| 96 | -__simt_callee__ __aicore__ inline static OFFSET_T EndIndexOut2In( | 88 | +__simt_callee__ __aicore__ inline static OFFSET_T EndIndexOut2In(OFFSET_T outIdx, OFFSET_T osize, OFFSET_T isize, DIV_T magicOsize, DIV_T shiftOsize) |
| 97 | - OFFSET_T outIdx, OFFSET_T osize, OFFSET_T isize, DIV_T magicOsize, DIV_T shiftOsize) | ||
| 98 | { | 89 | { |
| 99 | return CeilDivMul<OFFSET_T, DIV_T>(outIdx + 1, isize, osize - 1, magicOsize, shiftOsize); | 90 | return CeilDivMul<OFFSET_T, DIV_T>(outIdx + 1, isize, osize - 1, magicOsize, shiftOsize); |
| 100 | } | 91 | } |
| 101 | 92 | ||
| 102 | -template <typename VALUE_T, typename OFFSET_T> | 93 | +template <typename OFFSET_T, typename DIV_T> |
| 103 | -__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveAvgPool2dGradNchw( | 94 | +__simt_callee__ __aicore__ inline static void LoadSimtParams(__ubuf__ OFFSET_T* p, DIV_T& magicInH, DIV_T& shiftInH, DIV_T& magicInW, DIV_T& shiftInW, DIV_T& magicOsizeH, DIV_T& shiftOsizeH, DIV_T& magicOsizeW, DIV_T& shiftOsizeW) |
| 104 | - __ubuf__ OFFSET_T* simtParams, const __gm__ VALUE_T* gradY, const OFFSET_T nDims, const OFFSET_T cDims, | ||
| 105 | - const OFFSET_T inH, const OFFSET_T inW, const OFFSET_T outH, const OFFSET_T outW, __gm__ VALUE_T* gradX) | ||
| 106 | { | 95 | { |
| 107 | - using DIV_T = typename std::conditional<std::is_same<OFFSET_T, int32_t>::value, uint32_t, uint64_t>::type; | 96 | + magicInH = static_cast<DIV_T>(p[MAGIC_IN_H_IDX]); shiftInH = static_cast<DIV_T>(p[MAGIC_IN_H_IDX + 1]); |
| 108 | - DIV_T magicC = simtParams[MAGIC_C_IDX]; | 97 | + magicInW = static_cast<DIV_T>(p[MAGIC_IN_W_IDX]); shiftInW = static_cast<DIV_T>(p[MAGIC_IN_W_IDX + 1]); |
| 109 | - DIV_T shiftC = simtParams[MAGIC_C_IDX + 1]; | 98 | + magicOsizeH = static_cast<DIV_T>(p[MAGIC_OSIZE_H_IDX]); shiftOsizeH = static_cast<DIV_T>(p[MAGIC_OSIZE_H_IDX + 1]); |
| 110 | - DIV_T magicInH = simtParams[MAGIC_IN_H_IDX]; | 99 | + magicOsizeW = static_cast<DIV_T>(p[MAGIC_OSIZE_W_IDX]); shiftOsizeW = static_cast<DIV_T>(p[MAGIC_OSIZE_W_IDX + 1]); |
| 111 | - DIV_T shiftInH = simtParams[MAGIC_IN_H_IDX + 1]; | 100 | +} |
| 112 | - DIV_T magicInW = simtParams[MAGIC_IN_W_IDX]; | ||
| 113 | - DIV_T shiftInW = simtParams[MAGIC_IN_W_IDX + 1]; | ||
| 114 | - DIV_T magicOsizeH = simtParams[MAGIC_OSIZE_H_IDX]; | ||
| 115 | - DIV_T shiftOsizeH = simtParams[MAGIC_OSIZE_H_IDX + 1]; | ||
| 116 | - DIV_T magicOsizeW = simtParams[MAGIC_OSIZE_W_IDX]; | ||
| 117 | - DIV_T shiftOsizeW = simtParams[MAGIC_OSIZE_W_IDX + 1]; | ||
| 118 | 101 | ||
| 119 | - DIV_T count = nDims * cDims * inH * inW; | 102 | +template <typename VALUE_T, typename OFFSET_T> |
| 120 | - for (DIV_T index = blockIdx.x * blockDim.x + threadIdx.x; index < count; | 103 | +__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveAvgPool2dGradHInOne(__ubuf__ OFFSET_T* p, const __gm__ VALUE_T* gradY, OFFSET_T nDims, OFFSET_T cDims, OFFSET_T inH, OFFSET_T inW, OFFSET_T outH, OFFSET_T outW, __gm__ VALUE_T* gradX) |
| 121 | - index += gridDim.x * blockDim.x) { | 104 | +{ |
| 122 | - DIV_T temp1 = Simt::UintDiv(index, magicInW, shiftInW); | 105 | + using DIV_T = DivForOffset<OFFSET_T>; |
| 123 | - DIV_T w = index - temp1 * static_cast<DIV_T>(inW); | 106 | + DIV_T magicInH = 0, shiftInH = 0, magicInW = 0, shiftInW = 0, magicOsizeH = 0, shiftOsizeH = 0, magicOsizeW = 0, shiftOsizeW = 0; |
| 124 | - DIV_T temp2 = Simt::UintDiv(temp1, magicInH, shiftInH); | 107 | + LoadSimtParams<OFFSET_T, DIV_T>(p, magicInH, shiftInH, magicInW, shiftInW, magicOsizeH, shiftOsizeH, magicOsizeW, shiftOsizeW); |
| 125 | - DIV_T h = temp1 - temp2 * static_cast<DIV_T>(inH); | 108 | + DIV_T count = static_cast<DIV_T>(nDims) * static_cast<DIV_T>(cDims) * static_cast<DIV_T>(inW); |
| 126 | - DIV_T n = Simt::UintDiv(temp2, magicC, shiftC); | 109 | + DIV_T outHW = static_cast<DIV_T>(outH) * static_cast<DIV_T>(outW); |
| 127 | - DIV_T c = temp2 - n * static_cast<DIV_T>(cDims); | 110 | + for (DIV_T index = blockIdx.x * blockDim.x + threadIdx.x; index < count; index += gridDim.x * blockDim.x) { |
| 128 | - | 111 | + DIV_T nc = Simt::UintDiv<DIV_T>(index, magicInW, shiftInW), w = index - nc * static_cast<DIV_T>(inW), base = nc * outHW; |
| 129 | - OFFSET_T ohStarts = StartIndexIn2Out<OFFSET_T, DIV_T>(h, outH, magicInH, shiftInH); | 112 | + OFFSET_T owStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), outW, magicInW, shiftInW); |
| 130 | - OFFSET_T ohEnds = EndIndexIn2Out<OFFSET_T, DIV_T>(h, inH, outH, magicInH, shiftInH); | 113 | + OFFSET_T owEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), inW, outW, magicInW, shiftInW); |
| 131 | - OFFSET_T owStarts = StartIndexIn2Out<OFFSET_T, DIV_T>(w, outW, magicInW, shiftInW); | ||
| 132 | - OFFSET_T owEnds = EndIndexIn2Out<OFFSET_T, DIV_T>(w, inW, outW, magicInW, shiftInW); | ||
| 133 | - // 遍历所有可能覆盖这个input点的输出窗口 | ||
| 134 | float gradient = 0.0f; | 114 | float gradient = 0.0f; |
| 115 | + for (OFFSET_T ow = owStart; ow < owEnd; ++ow) { | ||
| 116 | + OFFSET_T iw0 = StartIndexOut2In<OFFSET_T, DIV_T>(ow, inW, magicOsizeW, shiftOsizeW); | ||
| 117 | + OFFSET_T iw1 = EndIndexOut2In<OFFSET_T, DIV_T>(ow, outW, inW, magicOsizeW, shiftOsizeW); | ||
| 118 | + float invKW = 1.0f / static_cast<float>(iw1 - iw0); | ||
| 119 | + for (OFFSET_T oh = 0; oh < outH; ++oh) { | ||
| 120 | + gradient += static_cast<float>(gradY[base + static_cast<DIV_T>(oh) * static_cast<DIV_T>(outW) + static_cast<DIV_T>(ow)]) * invKW; | ||
| 121 | + } | ||
| 122 | + } | ||
| 123 | + gradX[index] = static_cast<VALUE_T>(gradient); | ||
| 124 | + } | ||
| 125 | +} | ||
| 135 | 126 | ||
| 136 | - for (OFFSET_T oh = ohStarts; oh < ohEnds; ++oh) { | 127 | +template <typename VALUE_T, typename OFFSET_T, uint32_t OUT_W> |
| 128 | +__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveAvgPool2dGradSmallOutWRow(__ubuf__ OFFSET_T* p, const __gm__ VALUE_T* gradY, OFFSET_T nDims, OFFSET_T cDims, OFFSET_T inH, OFFSET_T inW, OFFSET_T outH, OFFSET_T outW, __gm__ VALUE_T* gradX) | ||
| 129 | +{ | ||
| 130 | + using DIV_T = DivForOffset<OFFSET_T>; | ||
| 131 | + DIV_T magicInH = 0, shiftInH = 0, magicInW = 0, shiftInW = 0, magicOsizeH = 0, shiftOsizeH = 0, magicOsizeW = 0, shiftOsizeW = 0; | ||
| 132 | + LoadSimtParams<OFFSET_T, DIV_T>(p, magicInH, shiftInH, magicInW, shiftInW, magicOsizeH, shiftOsizeH, magicOsizeW, shiftOsizeW); | ||
| 133 | + OFFSET_T iw0[OUT_W], iw1[OUT_W]; float invKW[OUT_W]; | ||
| 134 | + for (uint32_t ow = 0; ow < OUT_W; ++ow) { | ||
| 135 | + iw0[ow] = StartIndexOut2In<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(ow), inW, magicOsizeW, shiftOsizeW); | ||
| 136 | + iw1[ow] = EndIndexOut2In<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(ow), static_cast<OFFSET_T>(OUT_W), inW, magicOsizeW, shiftOsizeW); | ||
| 137 | + invKW[ow] = 1.0f / static_cast<float>(iw1[ow] - iw0[ow]); | ||
| 138 | + } | ||
| 139 | + DIV_T rowCount = static_cast<DIV_T>(nDims) * static_cast<DIV_T>(cDims) * static_cast<DIV_T>(inH); | ||
| 140 | + DIV_T outHW = static_cast<DIV_T>(outH) * static_cast<DIV_T>(OUT_W); | ||
| 141 | + for (DIV_T row = blockIdx.x * blockDim.x + threadIdx.x; row < rowCount; row += gridDim.x * blockDim.x) { | ||
| 142 | + DIV_T nc = Simt::UintDiv<DIV_T>(row, magicInH, shiftInH), h = row - nc * static_cast<DIV_T>(inH), base = nc * outHW; | ||
| 143 | + DIV_T xBase = (nc * static_cast<DIV_T>(inH) + h) * static_cast<DIV_T>(inW); | ||
| 144 | + OFFSET_T ohStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), outH, magicInH, shiftInH); | ||
| 145 | + OFFSET_T ohEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), inH, outH, magicInH, shiftInH); | ||
| 146 | + float owSum[OUT_W]; | ||
| 147 | + for (uint32_t ow = 0; ow < OUT_W; ++ow) { | ||
| 148 | + owSum[ow] = 0.0f; | ||
| 149 | + } | ||
| 150 | + for (OFFSET_T oh = ohStart; oh < ohEnd; ++oh) { | ||
| 137 | OFFSET_T ih0 = StartIndexOut2In<OFFSET_T, DIV_T>(oh, inH, magicOsizeH, shiftOsizeH); | 151 | OFFSET_T ih0 = StartIndexOut2In<OFFSET_T, DIV_T>(oh, inH, magicOsizeH, shiftOsizeH); |
| 138 | OFFSET_T ih1 = EndIndexOut2In<OFFSET_T, DIV_T>(oh, outH, inH, magicOsizeH, shiftOsizeH); | 152 | OFFSET_T ih1 = EndIndexOut2In<OFFSET_T, DIV_T>(oh, outH, inH, magicOsizeH, shiftOsizeH); |
| 139 | - OFFSET_T kH = ih1 - ih0; | 153 | + DIV_T yBase = base + static_cast<DIV_T>(oh) * static_cast<DIV_T>(OUT_W); |
| 140 | - for (OFFSET_T ow = owStarts; ow < owEnds; ++ow) { | 154 | + float invKH = 1.0f / static_cast<float>(ih1 - ih0); |
| 155 | + for (uint32_t ow = 0; ow < OUT_W; ++ow) { | ||
| 156 | + owSum[ow] += static_cast<float>(gradY[yBase + static_cast<DIV_T>(ow)]) * invKH * invKW[ow]; | ||
| 157 | + } | ||
| 158 | + } | ||
| 159 | + for (OFFSET_T w = 0; w < inW; ++w) { | ||
| 160 | + float gradient = 0.0f; | ||
| 161 | + for (uint32_t ow = 0; ow < OUT_W; ++ow) { | ||
| 162 | + if (w >= iw0[ow] && w < iw1[ow]) { | ||
| 163 | + gradient += owSum[ow]; | ||
| 164 | + } | ||
| 165 | + } | ||
| 166 | + gradX[xBase + static_cast<DIV_T>(w)] = static_cast<VALUE_T>(gradient); | ||
| 167 | + } | ||
| 168 | + } | ||
| 169 | +} | ||
| 170 | + | ||
| 171 | +template <typename VALUE_T, typename OFFSET_T, uint32_t OUT_W> | ||
| 172 | +__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveAvgPool2dGradSmallOutWSegFast(__ubuf__ OFFSET_T* p, const __gm__ VALUE_T* gradY, OFFSET_T nDims, OFFSET_T cDims, OFFSET_T inH, OFFSET_T inW, OFFSET_T outH, OFFSET_T outW, __gm__ VALUE_T* gradX) | ||
| 173 | +{ | ||
| 174 | + using DIV_T = DivForOffset<OFFSET_T>; | ||
| 175 | + DIV_T magicInH = 0, shiftInH = 0, magicInW = 0, shiftInW = 0, magicOsizeH = 0, shiftOsizeH = 0, magicOsizeW = 0, shiftOsizeW = 0; | ||
| 176 | + LoadSimtParams<OFFSET_T, DIV_T>(p, magicInH, shiftInH, magicInW, shiftInW, magicOsizeH, shiftOsizeH, magicOsizeW, shiftOsizeW); | ||
| 177 | + constexpr uint32_t SEG_NUM_U32 = OUT_W * 2 - 1; | ||
| 178 | + constexpr DIV_T SEG_NUM = static_cast<DIV_T>(SEG_NUM_U32); | ||
| 179 | + DIV_T magicSeg = static_cast<DIV_T>(p[MAGIC_SEG_IDX]), shiftSeg = static_cast<DIV_T>(p[MAGIC_SEG_IDX + 1]); | ||
| 180 | + DIV_T taskCount = static_cast<DIV_T>(nDims) * static_cast<DIV_T>(cDims) * static_cast<DIV_T>(inH) * SEG_NUM; | ||
| 181 | + DIV_T outHW = static_cast<DIV_T>(outH) * static_cast<DIV_T>(OUT_W); | ||
| 182 | + for (DIV_T task = blockIdx.x * blockDim.x + threadIdx.x; task < taskCount; task += gridDim.x * blockDim.x) { | ||
| 183 | + DIV_T row = Simt::UintDiv<DIV_T>(task, magicSeg, shiftSeg), seg = task - row * SEG_NUM; | ||
| 184 | + DIV_T nc = Simt::UintDiv<DIV_T>(row, magicInH, shiftInH), h = row - nc * static_cast<DIV_T>(inH); | ||
| 185 | + uint32_t meta = SEG_INFO_IDX + static_cast<uint32_t>(seg) * SEG_INFO_STRIDE; | ||
| 186 | + OFFSET_T startW = static_cast<OFFSET_T>(p[meta]), endW = static_cast<OFFSET_T>(p[meta + 1]); | ||
| 187 | + OFFSET_T kW0 = static_cast<OFFSET_T>(p[meta + 2]), kW1 = static_cast<OFFSET_T>(p[meta + 3]); | ||
| 188 | + if (startW >= endW) { | ||
| 189 | + continue; | ||
| 190 | + } | ||
| 191 | + if (kW0 <= static_cast<OFFSET_T>(0)) { | ||
| 192 | + continue; | ||
| 193 | + } | ||
| 194 | + uint32_t ow0 = static_cast<uint32_t>(seg >> 1); | ||
| 195 | + bool boundary = (seg & static_cast<DIV_T>(1)) != 0; | ||
| 196 | + if (boundary && kW1 <= static_cast<OFFSET_T>(0)) { | ||
| 197 | + continue; | ||
| 198 | + } | ||
| 199 | + OFFSET_T ohStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), outH, magicInH, shiftInH); | ||
| 200 | + OFFSET_T ohEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), inH, outH, magicInH, shiftInH); | ||
| 201 | + DIV_T hLeft = h * static_cast<DIV_T>(outH), hRight = (h + static_cast<DIV_T>(1)) * static_cast<DIV_T>(outH); | ||
| 202 | + bool leftCross = static_cast<DIV_T>(ohStart) * static_cast<DIV_T>(inH) < hLeft; | ||
| 203 | + bool rightCross = static_cast<DIV_T>(ohEnd) * static_cast<DIV_T>(inH) > hRight; | ||
| 204 | + bool singleOut = ohEnd <= ohStart + static_cast<OFFSET_T>(1); | ||
| 205 | + float firstWeight = (leftCross || (singleOut && rightCross)) ? 0.5f : 1.0f; | ||
| 206 | + float lastWeight = rightCross ? 0.5f : 1.0f; | ||
| 207 | + float invKW0 = 1.0f / static_cast<float>(kW0); | ||
| 208 | + float invKW1 = boundary ? ((kW0 == kW1) ? invKW0 : 1.0f / static_cast<float>(kW1)) : 0.0f; | ||
| 209 | + DIV_T base = nc * outHW; | ||
| 210 | + DIV_T firstBase = base + static_cast<DIV_T>(ohStart) * static_cast<DIV_T>(OUT_W); | ||
| 211 | + float gradient = static_cast<float>(gradY[firstBase + static_cast<DIV_T>(ow0)]) * firstWeight * invKW0; | ||
| 212 | + if (boundary) { | ||
| 213 | + gradient += static_cast<float>(gradY[firstBase + static_cast<DIV_T>(ow0 + 1)]) * firstWeight * invKW1; | ||
| 214 | + } | ||
| 215 | + for (OFFSET_T oh = ohStart + static_cast<OFFSET_T>(1); oh < ohEnd - static_cast<OFFSET_T>(1); ++oh) { | ||
| 216 | + DIV_T yBase = base + static_cast<DIV_T>(oh) * static_cast<DIV_T>(OUT_W); | ||
| 217 | + gradient += static_cast<float>(gradY[yBase + static_cast<DIV_T>(ow0)]) * invKW0; | ||
| 218 | + if (boundary) { | ||
| 219 | + gradient += static_cast<float>(gradY[yBase + static_cast<DIV_T>(ow0 + 1)]) * invKW1; | ||
| 220 | + } | ||
| 221 | + } | ||
| 222 | + if (ohEnd > ohStart + static_cast<OFFSET_T>(1)) { | ||
| 223 | + DIV_T lastBase = base + static_cast<DIV_T>(ohEnd - static_cast<OFFSET_T>(1)) * static_cast<DIV_T>(OUT_W); | ||
| 224 | + gradient += static_cast<float>(gradY[lastBase + static_cast<DIV_T>(ow0)]) * lastWeight * invKW0; | ||
| 225 | + if (boundary) { | ||
| 226 | + gradient += static_cast<float>(gradY[lastBase + static_cast<DIV_T>(ow0 + 1)]) * lastWeight * invKW1; | ||
| 227 | + } | ||
| 228 | + } | ||
| 229 | + VALUE_T outVal = static_cast<VALUE_T>(gradient); | ||
| 230 | + DIV_T xBase = (nc * static_cast<DIV_T>(inH) + h) * static_cast<DIV_T>(inW); | ||
| 231 | + for (OFFSET_T w = startW; w < endW; ++w) { | ||
| 232 | + gradX[xBase + static_cast<DIV_T>(w)] = outVal; | ||
| 233 | + } | ||
| 234 | + } | ||
| 235 | +} | ||
| 236 | + | ||
| 237 | +template <typename VALUE_T, typename OFFSET_T> | ||
| 238 | +__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveAvgPool2dGradOutWSmall(__ubuf__ OFFSET_T* p, const __gm__ VALUE_T* gradY, OFFSET_T nDims, OFFSET_T cDims, OFFSET_T inH, OFFSET_T inW, OFFSET_T outH, OFFSET_T outW, __gm__ VALUE_T* gradX) | ||
| 239 | +{ | ||
| 240 | + using DIV_T = DivForOffset<OFFSET_T>; | ||
| 241 | + DIV_T magicInH = 0, shiftInH = 0, magicInW = 0, shiftInW = 0, magicOsizeH = 0, shiftOsizeH = 0, magicOsizeW = 0, shiftOsizeW = 0; | ||
| 242 | + LoadSimtParams<OFFSET_T, DIV_T>(p, magicInH, shiftInH, magicInW, shiftInW, magicOsizeH, shiftOsizeH, magicOsizeW, shiftOsizeW); | ||
| 243 | + DIV_T count = static_cast<DIV_T>(nDims) * static_cast<DIV_T>(cDims) * static_cast<DIV_T>(inH) * static_cast<DIV_T>(inW); | ||
| 244 | + DIV_T outHW = static_cast<DIV_T>(outH) * static_cast<DIV_T>(outW); | ||
| 245 | + for (DIV_T index = blockIdx.x * blockDim.x + threadIdx.x; index < count; index += gridDim.x * blockDim.x) { | ||
| 246 | + DIV_T tmp = Simt::UintDiv<DIV_T>(index, magicInW, shiftInW); | ||
| 247 | + DIV_T w = index - tmp * static_cast<DIV_T>(inW), nc = Simt::UintDiv<DIV_T>(tmp, magicInH, shiftInH), h = tmp - nc * static_cast<DIV_T>(inH); | ||
| 248 | + DIV_T base = nc * outHW; | ||
| 249 | + OFFSET_T ohStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), outH, magicInH, shiftInH); | ||
| 250 | + OFFSET_T ohEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), inH, outH, magicInH, shiftInH); | ||
| 251 | + OFFSET_T owStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), outW, magicInW, shiftInW); | ||
| 252 | + OFFSET_T owEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), inW, outW, magicInW, shiftInW); | ||
| 253 | + float gradient = 0.0f; | ||
| 254 | + for (OFFSET_T ow = owStart; ow < owEnd; ++ow) { | ||
| 255 | + OFFSET_T iw0 = StartIndexOut2In<OFFSET_T, DIV_T>(ow, inW, magicOsizeW, shiftOsizeW); | ||
| 256 | + OFFSET_T iw1 = EndIndexOut2In<OFFSET_T, DIV_T>(ow, outW, inW, magicOsizeW, shiftOsizeW); | ||
| 257 | + float invKW = 1.0f / static_cast<float>(iw1 - iw0); | ||
| 258 | + for (OFFSET_T oh = ohStart; oh < ohEnd; ++oh) { | ||
| 259 | + OFFSET_T ih0 = StartIndexOut2In<OFFSET_T, DIV_T>(oh, inH, magicOsizeH, shiftOsizeH); | ||
| 260 | + OFFSET_T ih1 = EndIndexOut2In<OFFSET_T, DIV_T>(oh, outH, inH, magicOsizeH, shiftOsizeH); | ||
| 261 | + gradient += static_cast<float>(gradY[base + static_cast<DIV_T>(oh) * static_cast<DIV_T>(outW) + static_cast<DIV_T>(ow)]) * invKW / static_cast<float>(ih1 - ih0); | ||
| 262 | + } | ||
| 263 | + } | ||
| 264 | + gradX[index] = static_cast<VALUE_T>(gradient); | ||
| 265 | + } | ||
| 266 | +} | ||
| 267 | + | ||
| 268 | + | ||
| 269 | + | ||
| 270 | + | ||
| 271 | +template <typename VALUE_T, typename OFFSET_T, uint32_t H_SCALE> | ||
| 272 | +__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveAvgPool2dGradHExpandExactOutWSmallFast(__ubuf__ OFFSET_T* p, const __gm__ VALUE_T* gradY, OFFSET_T nDims, OFFSET_T cDims, OFFSET_T inH, OFFSET_T inW, OFFSET_T outH, OFFSET_T outW, __gm__ VALUE_T* gradX) | ||
| 273 | +{ | ||
| 274 | + using DIV_T = DivForOffset<OFFSET_T>; | ||
| 275 | + DIV_T magicInH = 0, shiftInH = 0, magicInW = 0, shiftInW = 0, magicOsizeH = 0, shiftOsizeH = 0, magicOsizeW = 0, shiftOsizeW = 0; | ||
| 276 | + LoadSimtParams<OFFSET_T, DIV_T>(p, magicInH, shiftInH, magicInW, shiftInW, magicOsizeH, shiftOsizeH, magicOsizeW, shiftOsizeW); | ||
| 277 | + DIV_T count = static_cast<DIV_T>(nDims) * static_cast<DIV_T>(cDims) * static_cast<DIV_T>(inH) * static_cast<DIV_T>(inW); | ||
| 278 | + DIV_T outHW = static_cast<DIV_T>(outH) * static_cast<DIV_T>(outW); | ||
| 279 | + constexpr DIV_T hScale = static_cast<DIV_T>(H_SCALE); | ||
| 280 | + for (DIV_T index = blockIdx.x * blockDim.x + threadIdx.x; index < count; index += gridDim.x * blockDim.x) { | ||
| 281 | + DIV_T tmp = Simt::UintDiv<DIV_T>(index, magicInW, shiftInW); | ||
| 282 | + DIV_T w = index - tmp * static_cast<DIV_T>(inW); | ||
| 283 | + DIV_T nc = Simt::UintDiv<DIV_T>(tmp, magicInH, shiftInH); | ||
| 284 | + DIV_T h = tmp - nc * static_cast<DIV_T>(inH); | ||
| 285 | + DIV_T base = nc * outHW + h * hScale * static_cast<DIV_T>(outW); | ||
| 286 | + OFFSET_T owStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), outW, magicInW, shiftInW); | ||
| 287 | + OFFSET_T owEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), inW, outW, magicInW, shiftInW); | ||
| 288 | + float gradient = 0.0f; | ||
| 289 | + for (OFFSET_T ow = owStart; ow < owEnd; ++ow) { | ||
| 290 | + OFFSET_T iw0 = StartIndexOut2In<OFFSET_T, DIV_T>(ow, inW, magicOsizeW, shiftOsizeW); | ||
| 291 | + OFFSET_T iw1 = EndIndexOut2In<OFFSET_T, DIV_T>(ow, outW, inW, magicOsizeW, shiftOsizeW); | ||
| 292 | + DIV_T yBase = base + static_cast<DIV_T>(ow); | ||
| 293 | + float invKW = 1.0f / static_cast<float>(iw1 - iw0); | ||
| 294 | + | ||
| 295 | + for (uint32_t i = 0; i < H_SCALE; ++i) { | ||
| 296 | + gradient += static_cast<float>(gradY[yBase + static_cast<DIV_T>(i) * static_cast<DIV_T>(outW)]) * invKW; | ||
| 297 | + } | ||
| 298 | + } | ||
| 299 | + gradX[index] = static_cast<VALUE_T>(gradient); | ||
| 300 | + } | ||
| 301 | +} | ||
| 302 | + | ||
| 303 | +template <typename VALUE_T, typename OFFSET_T> | ||
| 304 | +__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveAvgPool2dGradHReduceWExpandFast(__ubuf__ OFFSET_T* p, const __gm__ VALUE_T* gradY, OFFSET_T nDims, OFFSET_T cDims, OFFSET_T inH, OFFSET_T inW, OFFSET_T outH, OFFSET_T outW, __gm__ VALUE_T* gradX) | ||
| 305 | +{ | ||
| 306 | + using DIV_T = DivForOffset<OFFSET_T>; | ||
| 307 | + DIV_T magicInH = 0, shiftInH = 0, magicInW = 0, shiftInW = 0, magicOsizeH = 0, shiftOsizeH = 0, magicOsizeW = 0, shiftOsizeW = 0; | ||
| 308 | + LoadSimtParams<OFFSET_T, DIV_T>(p, magicInH, shiftInH, magicInW, shiftInW, magicOsizeH, shiftOsizeH, magicOsizeW, shiftOsizeW); | ||
| 309 | + DIV_T count = static_cast<DIV_T>(nDims) * static_cast<DIV_T>(cDims) * static_cast<DIV_T>(inH) * static_cast<DIV_T>(inW); | ||
| 310 | + DIV_T outHW = static_cast<DIV_T>(outH) * static_cast<DIV_T>(outW); | ||
| 311 | + for (DIV_T index = blockIdx.x * blockDim.x + threadIdx.x; index < count; index += gridDim.x * blockDim.x) { | ||
| 312 | + DIV_T tmp = Simt::UintDiv<DIV_T>(index, magicInW, shiftInW); | ||
| 313 | + DIV_T w = index - tmp * static_cast<DIV_T>(inW), nc = Simt::UintDiv<DIV_T>(tmp, magicInH, shiftInH), h = tmp - nc * static_cast<DIV_T>(inH); | ||
| 314 | + DIV_T base = nc * outHW; | ||
| 315 | + OFFSET_T ohStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), outH, magicInH, shiftInH); | ||
| 316 | + OFFSET_T ohEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), inH, outH, magicInH, shiftInH); | ||
| 317 | + OFFSET_T owStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), outW, magicInW, shiftInW); | ||
| 318 | + OFFSET_T owEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), inW, outW, magicInW, shiftInW); | ||
| 319 | + float gradient = 0.0f; | ||
| 320 | + for (OFFSET_T oh = ohStart; oh < ohEnd; ++oh) { | ||
| 321 | + OFFSET_T ih0 = StartIndexOut2In<OFFSET_T, DIV_T>(oh, inH, magicOsizeH, shiftOsizeH); | ||
| 322 | + OFFSET_T ih1 = EndIndexOut2In<OFFSET_T, DIV_T>(oh, outH, inH, magicOsizeH, shiftOsizeH); | ||
| 323 | + float invKH = 1.0f / static_cast<float>(ih1 - ih0); | ||
| 324 | + DIV_T yBase = base + static_cast<DIV_T>(oh) * static_cast<DIV_T>(outW); | ||
| 325 | + if (owEnd <= owStart) { | ||
| 326 | + continue; | ||
| 327 | + } | ||
| 328 | + if (owEnd == owStart + static_cast<OFFSET_T>(1)) { | ||
| 329 | + OFFSET_T iw0 = StartIndexOut2In<OFFSET_T, DIV_T>(owStart, inW, magicOsizeW, shiftOsizeW); | ||
| 330 | + OFFSET_T iw1 = EndIndexOut2In<OFFSET_T, DIV_T>(owStart, outW, inW, magicOsizeW, shiftOsizeW); | ||
| 331 | + gradient += static_cast<float>(gradY[yBase + static_cast<DIV_T>(owStart)]) * invKH / static_cast<float>(iw1 - iw0); | ||
| 332 | + continue; | ||
| 333 | + } | ||
| 334 | + OFFSET_T firstIw0 = StartIndexOut2In<OFFSET_T, DIV_T>(owStart, inW, magicOsizeW, shiftOsizeW); | ||
| 335 | + OFFSET_T firstIw1 = EndIndexOut2In<OFFSET_T, DIV_T>(owStart, outW, inW, magicOsizeW, shiftOsizeW); | ||
| 336 | + gradient += static_cast<float>(gradY[yBase + static_cast<DIV_T>(owStart)]) * invKH / static_cast<float>(firstIw1 - firstIw0); | ||
| 337 | + OFFSET_T middleEnd = owEnd - static_cast<OFFSET_T>(1); | ||
| 338 | + for (OFFSET_T ow = owStart + static_cast<OFFSET_T>(1); ow < middleEnd; ++ow) { | ||
| 339 | + gradient += static_cast<float>(gradY[yBase + static_cast<DIV_T>(ow)]) * invKH; | ||
| 340 | + } | ||
| 341 | + OFFSET_T lastOw = owEnd - static_cast<OFFSET_T>(1); | ||
| 342 | + OFFSET_T lastIw0 = StartIndexOut2In<OFFSET_T, DIV_T>(lastOw, inW, magicOsizeW, shiftOsizeW); | ||
| 343 | + OFFSET_T lastIw1 = EndIndexOut2In<OFFSET_T, DIV_T>(lastOw, outW, inW, magicOsizeW, shiftOsizeW); | ||
| 344 | + gradient += static_cast<float>(gradY[yBase + static_cast<DIV_T>(lastOw)]) * invKH / static_cast<float>(lastIw1 - lastIw0); | ||
| 345 | + } | ||
| 346 | + gradX[index] = static_cast<VALUE_T>(gradient); | ||
| 347 | + } | ||
| 348 | +} | ||
| 349 | + | ||
| 350 | + | ||
| 351 | + | ||
| 352 | +template <typename VALUE_T, typename OFFSET_T> | ||
| 353 | +__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveAvgPool2dGradHExpandW2SmallFast(__ubuf__ OFFSET_T* p, const __gm__ VALUE_T* gradY, OFFSET_T nDims, OFFSET_T cDims, OFFSET_T inH, OFFSET_T inW, OFFSET_T outH, OFFSET_T outW, __gm__ VALUE_T* gradX) | ||
| 354 | +{ | ||
| 355 | + using DIV_T = DivForOffset<OFFSET_T>; | ||
| 356 | + DIV_T magicInH = 0, shiftInH = 0, magicInW = 0, shiftInW = 0, magicOsizeH = 0, shiftOsizeH = 0, magicOsizeW = 0, shiftOsizeW = 0; | ||
| 357 | + LoadSimtParams<OFFSET_T, DIV_T>(p, magicInH, shiftInH, magicInW, shiftInW, magicOsizeH, shiftOsizeH, magicOsizeW, shiftOsizeW); | ||
| 358 | + static_cast<void>(outW); | ||
| 359 | + | ||
| 360 | + constexpr DIV_T OUT_W = static_cast<DIV_T>(2); | ||
| 361 | + DIV_T count = static_cast<DIV_T>(nDims) * static_cast<DIV_T>(cDims) * static_cast<DIV_T>(inH) * static_cast<DIV_T>(inW); | ||
| 362 | + DIV_T outHW = static_cast<DIV_T>(outH) * OUT_W; | ||
| 363 | + DIV_T kW = (static_cast<DIV_T>(inW) + static_cast<DIV_T>(1)) >> static_cast<DIV_T>(1); | ||
| 364 | + DIV_T rightStart = static_cast<DIV_T>(inW) >> static_cast<DIV_T>(1); | ||
| 365 | + float invKW = 1.0f / static_cast<float>(kW); | ||
| 366 | + | ||
| 367 | + for (DIV_T index = blockIdx.x * blockDim.x + threadIdx.x; index < count; index += gridDim.x * blockDim.x) { | ||
| 368 | + DIV_T tmp = Simt::UintDiv<DIV_T>(index, magicInW, shiftInW); | ||
| 369 | + DIV_T w = index - tmp * static_cast<DIV_T>(inW); | ||
| 370 | + DIV_T nc = Simt::UintDiv<DIV_T>(tmp, magicInH, shiftInH); | ||
| 371 | + DIV_T h = tmp - nc * static_cast<DIV_T>(inH); | ||
| 372 | + | ||
| 373 | + OFFSET_T ohStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), outH, magicInH, shiftInH); | ||
| 374 | + OFFSET_T ohEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), inH, outH, magicInH, shiftInH); | ||
| 375 | + DIV_T hLeft = h * static_cast<DIV_T>(outH); | ||
| 376 | + DIV_T hRight = (h + static_cast<DIV_T>(1)) * static_cast<DIV_T>(outH); | ||
| 377 | + bool leftCross = (h != static_cast<DIV_T>(0)) && (static_cast<DIV_T>(ohStart) * static_cast<DIV_T>(inH) < hLeft); | ||
| 378 | + bool rightCross = (h + static_cast<DIV_T>(1) < static_cast<DIV_T>(inH)) && (static_cast<DIV_T>(ohEnd) * static_cast<DIV_T>(inH) > hRight); | ||
| 379 | + OFFSET_T midBegin = leftCross ? (ohStart + static_cast<OFFSET_T>(1)) : ohStart; | ||
| 380 | + OFFSET_T midEnd = rightCross ? (ohEnd - static_cast<OFFSET_T>(1)) : ohEnd; | ||
| 381 | + | ||
| 382 | + DIV_T base = nc * outHW; | ||
| 383 | + bool useOw0 = w < kW; | ||
| 384 | + bool useOw1 = w >= rightStart; | ||
| 385 | + | ||
| 386 | + if (useOw0 && !useOw1) { | ||
| 387 | + float sum0 = 0.0f; | ||
| 388 | + if (leftCross) { | ||
| 389 | + sum0 += static_cast<float>(gradY[base + static_cast<DIV_T>(ohStart) * OUT_W]) * 0.5f * invKW; | ||
| 390 | + } | ||
| 391 | + OFFSET_T oh = midBegin; | ||
| 392 | + for (; oh + static_cast<OFFSET_T>(3) < midEnd; oh += static_cast<OFFSET_T>(4)) { | ||
| 393 | + DIV_T yBase = base + static_cast<DIV_T>(oh) * OUT_W; | ||
| 394 | + sum0 += static_cast<float>(gradY[yBase]) * invKW + | ||
| 395 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(2)]) * invKW + | ||
| 396 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(4)]) * invKW + | ||
| 397 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(6)]) * invKW; | ||
| 398 | + } | ||
| 399 | + for (; oh < midEnd; ++oh) { | ||
| 400 | + sum0 += static_cast<float>(gradY[base + static_cast<DIV_T>(oh) * OUT_W]) * invKW; | ||
| 401 | + } | ||
| 402 | + if (rightCross) { | ||
| 403 | + sum0 += static_cast<float>(gradY[base + static_cast<DIV_T>(ohEnd - static_cast<OFFSET_T>(1)) * OUT_W]) * 0.5f * invKW; | ||
| 404 | + } | ||
| 405 | + gradX[index] = static_cast<VALUE_T>(sum0); | ||
| 406 | + } else if (!useOw0 && useOw1) { | ||
| 407 | + float sum1 = 0.0f; | ||
| 408 | + if (leftCross) { | ||
| 409 | + sum1 += static_cast<float>(gradY[base + static_cast<DIV_T>(ohStart) * OUT_W + static_cast<DIV_T>(1)]) * 0.5f * invKW; | ||
| 410 | + } | ||
| 411 | + OFFSET_T oh = midBegin; | ||
| 412 | + for (; oh + static_cast<OFFSET_T>(3) < midEnd; oh += static_cast<OFFSET_T>(4)) { | ||
| 413 | + DIV_T yBase = base + static_cast<DIV_T>(oh) * OUT_W + static_cast<DIV_T>(1); | ||
| 414 | + sum1 += static_cast<float>(gradY[yBase]) * invKW + | ||
| 415 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(2)]) * invKW + | ||
| 416 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(4)]) * invKW + | ||
| 417 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(6)]) * invKW; | ||
| 418 | + } | ||
| 419 | + for (; oh < midEnd; ++oh) { | ||
| 420 | + sum1 += static_cast<float>(gradY[base + static_cast<DIV_T>(oh) * OUT_W + static_cast<DIV_T>(1)]) * invKW; | ||
| 421 | + } | ||
| 422 | + if (rightCross) { | ||
| 423 | + sum1 += static_cast<float>(gradY[base + static_cast<DIV_T>(ohEnd - static_cast<OFFSET_T>(1)) * OUT_W + static_cast<DIV_T>(1)]) * 0.5f * invKW; | ||
| 424 | + } | ||
| 425 | + gradX[index] = static_cast<VALUE_T>(sum1); | ||
| 426 | + } else { | ||
| 427 | + float sum0 = 0.0f; | ||
| 428 | + float sum1 = 0.0f; | ||
| 429 | + if (leftCross) { | ||
| 430 | + DIV_T yBase = base + static_cast<DIV_T>(ohStart) * OUT_W; | ||
| 431 | + sum0 += static_cast<float>(gradY[yBase]) * 0.5f * invKW; | ||
| 432 | + sum1 += static_cast<float>(gradY[yBase + static_cast<DIV_T>(1)]) * 0.5f * invKW; | ||
| 433 | + } | ||
| 434 | + OFFSET_T oh = midBegin; | ||
| 435 | + for (; oh + static_cast<OFFSET_T>(3) < midEnd; oh += static_cast<OFFSET_T>(4)) { | ||
| 436 | + DIV_T yBase = base + static_cast<DIV_T>(oh) * OUT_W; | ||
| 437 | + sum0 += static_cast<float>(gradY[yBase]) * invKW + | ||
| 438 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(2)]) * invKW + | ||
| 439 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(4)]) * invKW + | ||
| 440 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(6)]) * invKW; | ||
| 441 | + sum1 += static_cast<float>(gradY[yBase + static_cast<DIV_T>(1)]) * invKW + | ||
| 442 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(3)]) * invKW + | ||
| 443 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(5)]) * invKW + | ||
| 444 | + static_cast<float>(gradY[yBase + static_cast<DIV_T>(7)]) * invKW; | ||
| 445 | + } | ||
| 446 | + for (; oh < midEnd; ++oh) { | ||
| 447 | + DIV_T yBase = base + static_cast<DIV_T>(oh) * OUT_W; | ||
| 448 | + sum0 += static_cast<float>(gradY[yBase]) * invKW; | ||
| 449 | + sum1 += static_cast<float>(gradY[yBase + static_cast<DIV_T>(1)]) * invKW; | ||
| 450 | + } | ||
| 451 | + if (rightCross) { | ||
| 452 | + DIV_T yBase = base + static_cast<DIV_T>(ohEnd - static_cast<OFFSET_T>(1)) * OUT_W; | ||
| 453 | + sum0 += static_cast<float>(gradY[yBase]) * 0.5f * invKW; | ||
| 454 | + sum1 += static_cast<float>(gradY[yBase + static_cast<DIV_T>(1)]) * 0.5f * invKW; | ||
| 455 | + } | ||
| 456 | + gradX[index] = static_cast<VALUE_T>(sum0 + sum1); | ||
| 457 | + } | ||
| 458 | + } | ||
| 459 | +} | ||
| 460 | + | ||
| 461 | + | ||
| 462 | +template <typename VALUE_T, typename OFFSET_T> | ||
| 463 | +__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveAvgPool2dGradNchw(__ubuf__ OFFSET_T* p, const __gm__ VALUE_T* gradY, OFFSET_T nDims, OFFSET_T cDims, OFFSET_T inH, OFFSET_T inW, OFFSET_T outH, OFFSET_T outW, __gm__ VALUE_T* gradX) | ||
| 464 | +{ | ||
| 465 | + using DIV_T = DivForOffset<OFFSET_T>; | ||
| 466 | + DIV_T magicInH = 0, shiftInH = 0, magicInW = 0, shiftInW = 0, magicOsizeH = 0, shiftOsizeH = 0, magicOsizeW = 0, shiftOsizeW = 0; | ||
| 467 | + LoadSimtParams<OFFSET_T, DIV_T>(p, magicInH, shiftInH, magicInW, shiftInW, magicOsizeH, shiftOsizeH, magicOsizeW, shiftOsizeW); | ||
| 468 | + DIV_T count = static_cast<DIV_T>(nDims) * static_cast<DIV_T>(cDims) * static_cast<DIV_T>(inH) * static_cast<DIV_T>(inW); | ||
| 469 | + DIV_T outHW = static_cast<DIV_T>(outH) * static_cast<DIV_T>(outW); | ||
| 470 | + for (DIV_T index = blockIdx.x * blockDim.x + threadIdx.x; index < count; index += gridDim.x * blockDim.x) { | ||
| 471 | + DIV_T tmp = Simt::UintDiv<DIV_T>(index, magicInW, shiftInW); | ||
| 472 | + DIV_T w = index - tmp * static_cast<DIV_T>(inW), nc = Simt::UintDiv<DIV_T>(tmp, magicInH, shiftInH), h = tmp - nc * static_cast<DIV_T>(inH); | ||
| 473 | + DIV_T base = nc * outHW; | ||
| 474 | + OFFSET_T ohStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), outH, magicInH, shiftInH); | ||
| 475 | + OFFSET_T ohEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(h), inH, outH, magicInH, shiftInH); | ||
| 476 | + OFFSET_T owStart = StartIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), outW, magicInW, shiftInW); | ||
| 477 | + OFFSET_T owEnd = EndIndexIn2Out<OFFSET_T, DIV_T>(static_cast<OFFSET_T>(w), inW, outW, magicInW, shiftInW); | ||
| 478 | + float gradient = 0.0f; | ||
| 479 | + for (OFFSET_T oh = ohStart; oh < ohEnd; ++oh) { | ||
| 480 | + OFFSET_T ih0 = StartIndexOut2In<OFFSET_T, DIV_T>(oh, inH, magicOsizeH, shiftOsizeH); | ||
| 481 | + OFFSET_T ih1 = EndIndexOut2In<OFFSET_T, DIV_T>(oh, outH, inH, magicOsizeH, shiftOsizeH); | ||
| 482 | + float invKH = 1.0f / static_cast<float>(ih1 - ih0); | ||
| 483 | + for (OFFSET_T ow = owStart; ow < owEnd; ++ow) { | ||
| 141 | OFFSET_T iw0 = StartIndexOut2In<OFFSET_T, DIV_T>(ow, inW, magicOsizeW, shiftOsizeW); | 484 | OFFSET_T iw0 = StartIndexOut2In<OFFSET_T, DIV_T>(ow, inW, magicOsizeW, shiftOsizeW); |
| 142 | OFFSET_T iw1 = EndIndexOut2In<OFFSET_T, DIV_T>(ow, outW, inW, magicOsizeW, shiftOsizeW); | 485 | OFFSET_T iw1 = EndIndexOut2In<OFFSET_T, DIV_T>(ow, outW, inW, magicOsizeW, shiftOsizeW); |
| 143 | - OFFSET_T kW = iw1 - iw0; | 486 | + gradient += static_cast<float>(gradY[base + static_cast<DIV_T>(oh) * static_cast<DIV_T>(outW) + static_cast<DIV_T>(ow)]) * invKH / static_cast<float>(iw1 - iw0); |
| 144 | - OFFSET_T div = kH * kW; | ||
| 145 | - DIV_T outputIdx = n * cDims * outH * outW + c * outH * outW + oh * outW + ow; | ||
| 146 | - gradient += static_cast<float>(gradY[outputIdx]) / static_cast<float>(div); | ||
| 147 | } | 487 | } |
| 148 | } | 488 | } |
| 149 | gradX[index] = static_cast<VALUE_T>(gradient); | 489 | gradX[index] = static_cast<VALUE_T>(gradient); |
| @@ -161,48 +501,74 @@ __aicore__ inline void AdaptiveAvgPool2dGradSimt<VALUE_T, OFFSET_T, CHANNEL_LAST | |||
| 161 | template <typename VALUE_T, typename OFFSET_T, int64_t CHANNEL_LAST> | 501 | template <typename VALUE_T, typename OFFSET_T, int64_t CHANNEL_LAST> |
| 162 | __aicore__ inline void AdaptiveAvgPool2dGradSimt<VALUE_T, OFFSET_T, CHANNEL_LAST>::Process() | 502 | __aicore__ inline void AdaptiveAvgPool2dGradSimt<VALUE_T, OFFSET_T, CHANNEL_LAST>::Process() |
| 163 | { | 503 | { |
| 164 | - using DIV_T = typename std::conditional<std::is_same<OFFSET_T, int32_t>::value, uint32_t, uint64_t>::type; | 504 | + using DIV_T = DivForOffset<OFFSET_T>; |
| 505 | + OFFSET_T n = static_cast<OFFSET_T>(tilingData_->nDim), c = static_cast<OFFSET_T>(tilingData_->cDim), hIn = static_cast<OFFSET_T>(tilingData_->hInDim); | ||
| 506 | + OFFSET_T wIn = static_cast<OFFSET_T>(tilingData_->wInDim), hOut = static_cast<OFFSET_T>(tilingData_->hOutDim), wOut = static_cast<OFFSET_T>(tilingData_->wOutDim); | ||
| 165 | LocalTensor<DIV_T> simtParam = paramBuf_.Get<DIV_T>(); | 507 | LocalTensor<DIV_T> simtParam = paramBuf_.Get<DIV_T>(); |
| 166 | - DIV_T magicC = 0; | 508 | + DIV_T magicC = 0, shiftC = 0, magicInH = 0, shiftInH = 0, magicInW = 0, shiftInW = 0, magicOsizeH = 0, shiftOsizeH = 0, magicOsizeW = 0, shiftOsizeW = 0, magicSeg = 0, shiftSeg = 0; |
| 167 | - DIV_T shiftC = 0; | 509 | + DIV_T segNum = (wOut >= static_cast<OFFSET_T>(2) && wOut <= static_cast<OFFSET_T>(4)) ? static_cast<DIV_T>(wOut) * static_cast<DIV_T>(2) - static_cast<DIV_T>(1) : static_cast<DIV_T>(7); |
| 168 | - DIV_T magicInH = 0; | 510 | + GetUintDivMagicAndShift<DIV_T>(magicC, shiftC, static_cast<DIV_T>(c)); |
| 169 | - DIV_T shiftInH = 0; | 511 | + GetUintDivMagicAndShift<DIV_T>(magicInH, shiftInH, static_cast<DIV_T>(hIn)); |
| 170 | - DIV_T magicInW = 0; | 512 | + GetUintDivMagicAndShift<DIV_T>(magicInW, shiftInW, static_cast<DIV_T>(wIn)); |
| 171 | - DIV_T shiftInW = 0; | 513 | + GetUintDivMagicAndShift<DIV_T>(magicOsizeH, shiftOsizeH, static_cast<DIV_T>(hOut)); |
| 172 | - DIV_T magicOsizeH = 0; | 514 | + GetUintDivMagicAndShift<DIV_T>(magicOsizeW, shiftOsizeW, static_cast<DIV_T>(wOut)); |
| 173 | - DIV_T shiftOsizeH = 0; | 515 | + GetUintDivMagicAndShift<DIV_T>(magicSeg, shiftSeg, segNum); |
| 174 | - DIV_T magicOsizeW = 0; | 516 | + simtParam.SetValue(MAGIC_C_IDX, magicC); simtParam.SetValue(MAGIC_C_IDX + 1, shiftC); |
| 175 | - DIV_T shiftOsizeW = 0; | 517 | + simtParam.SetValue(MAGIC_IN_H_IDX, magicInH); simtParam.SetValue(MAGIC_IN_H_IDX + 1, shiftInH); |
| 176 | - | 518 | + simtParam.SetValue(MAGIC_IN_W_IDX, magicInW); simtParam.SetValue(MAGIC_IN_W_IDX + 1, shiftInW); |
| 177 | - GetUintDivMagicAndShift<DIV_T>(magicC, shiftC, static_cast<DIV_T>(tilingData_->cDim)); | 519 | + simtParam.SetValue(MAGIC_OSIZE_H_IDX, magicOsizeH); simtParam.SetValue(MAGIC_OSIZE_H_IDX + 1, shiftOsizeH); |
| 178 | - GetUintDivMagicAndShift<DIV_T>(magicInH, shiftInH, static_cast<DIV_T>(tilingData_->hInDim)); | 520 | + simtParam.SetValue(MAGIC_OSIZE_W_IDX, magicOsizeW); simtParam.SetValue(MAGIC_OSIZE_W_IDX + 1, shiftOsizeW); |
| 179 | - GetUintDivMagicAndShift<DIV_T>(magicInW, shiftInW, static_cast<DIV_T>(tilingData_->wInDim)); | 521 | + simtParam.SetValue(MAGIC_SEG_IDX, magicSeg); simtParam.SetValue(MAGIC_SEG_IDX + 1, shiftSeg); |
| 180 | - GetUintDivMagicAndShift<DIV_T>(magicOsizeH, shiftOsizeH, static_cast<DIV_T>(tilingData_->hOutDim)); | 522 | + if (wOut >= static_cast<OFFSET_T>(2) && wOut <= static_cast<OFFSET_T>(4)) { |
| 181 | - GetUintDivMagicAndShift<DIV_T>(magicOsizeW, shiftOsizeW, static_cast<DIV_T>(tilingData_->wOutDim)); | 523 | + for (uint32_t seg = 0; seg < 7; ++seg) { |
| 182 | - | 524 | + DIV_T startW = 0, endW = 0, kW0 = 1, kW1 = 1; |
| 183 | - simtParam.SetValue(MAGIC_C_IDX, magicC); | 525 | + if (static_cast<DIV_T>(seg) < segNum) { |
| 184 | - simtParam.SetValue(MAGIC_C_IDX + 1, shiftC); | 526 | + DIV_T ow = static_cast<DIV_T>(seg >> 1), wInDiv = static_cast<DIV_T>(wIn), wOutDiv = static_cast<DIV_T>(wOut); |
| 185 | - simtParam.SetValue(MAGIC_IN_H_IDX, magicInH); | 527 | + DIV_T left0 = ow * wInDiv, right0 = (ow + static_cast<DIV_T>(1)) * wInDiv; |
| 186 | - simtParam.SetValue(MAGIC_IN_H_IDX + 1, shiftInH); | 528 | + DIV_T iw0 = left0 / wOutDiv, iw1 = (right0 + wOutDiv - static_cast<DIV_T>(1)) / wOutDiv; |
| 187 | - simtParam.SetValue(MAGIC_IN_W_IDX, magicInW); | 529 | + kW0 = iw1 - iw0; |
| 188 | - simtParam.SetValue(MAGIC_IN_W_IDX + 1, shiftInW); | 530 | + if ((seg & 1) != 0) { |
| 189 | - simtParam.SetValue(MAGIC_OSIZE_H_IDX, magicOsizeH); | 531 | + DIV_T boundaryPos = (ow + static_cast<DIV_T>(1)) * wInDiv; |
| 190 | - simtParam.SetValue(MAGIC_OSIZE_H_IDX + 1, shiftOsizeH); | 532 | + startW = boundaryPos / wOutDiv; endW = (boundaryPos + wOutDiv - static_cast<DIV_T>(1)) / wOutDiv; |
| 191 | - simtParam.SetValue(MAGIC_OSIZE_W_IDX, magicOsizeW); | 533 | + DIV_T ow1 = ow + static_cast<DIV_T>(1), left1 = ow1 * wInDiv, right1 = (ow1 + static_cast<DIV_T>(1)) * wInDiv; |
| 192 | - simtParam.SetValue(MAGIC_OSIZE_W_IDX + 1, shiftOsizeW); | 534 | + kW1 = (right1 + wOutDiv - static_cast<DIV_T>(1)) / wOutDiv - left1 / wOutDiv; |
| 193 | - | 535 | + } else { |
| 536 | + startW = (ow == static_cast<DIV_T>(0)) ? static_cast<DIV_T>(0) : (left0 + wOutDiv - static_cast<DIV_T>(1)) / wOutDiv; | ||
| 537 | + endW = (ow == wOutDiv - static_cast<DIV_T>(1)) ? wInDiv : right0 / wOutDiv; | ||
| 538 | + } | ||
| 539 | + } | ||
| 540 | + uint32_t meta = SEG_INFO_IDX + seg * SEG_INFO_STRIDE; | ||
| 541 | + simtParam.SetValue(meta, startW); simtParam.SetValue(meta + 1, endW); simtParam.SetValue(meta + 2, kW0); simtParam.SetValue(meta + 3, kW1); | ||
| 542 | + } | ||
| 543 | + } | ||
| 194 | DataSyncBarrier<MemDsbT::UB>(); | 544 | DataSyncBarrier<MemDsbT::UB>(); |
| 195 | - | ||
| 196 | auto gradData = (__gm__ VALUE_T*)yGrad_.GetPhyAddr(); | 545 | auto gradData = (__gm__ VALUE_T*)yGrad_.GetPhyAddr(); |
| 197 | auto outputData = (__gm__ VALUE_T*)xGrad_.GetPhyAddr(); | 546 | auto outputData = (__gm__ VALUE_T*)xGrad_.GetPhyAddr(); |
| 198 | - | 547 | + auto params = (__ubuf__ OFFSET_T*)simtParam.GetPhyAddr(); |
| 199 | - asc_vf_call<AdaptiveAvgPool2dGradNchw<VALUE_T, OFFSET_T>>( | 548 | + if (hIn == static_cast<OFFSET_T>(1)) { |
| 200 | - dim3(THREAD_DIM), (__ubuf__ OFFSET_T*)simtParam.GetPhyAddr(), gradData, | 549 | + asc_vf_call<AdaptiveAvgPool2dGradHInOne<VALUE_T, OFFSET_T>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); |
| 201 | - static_cast<OFFSET_T>(tilingData_->nDim), static_cast<OFFSET_T>(tilingData_->cDim), | 550 | + } else if (hOut > hIn && wOut == static_cast<OFFSET_T>(4) && wOut <= wIn && wIn >= static_cast<OFFSET_T>(32)) { |
| 202 | - static_cast<OFFSET_T>(tilingData_->hInDim), static_cast<OFFSET_T>(tilingData_->wInDim), | 551 | + asc_vf_call<AdaptiveAvgPool2dGradSmallOutWSegFast<VALUE_T, OFFSET_T, 4>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); |
| 203 | - static_cast<OFFSET_T>(tilingData_->hOutDim), static_cast<OFFSET_T>(tilingData_->wOutDim), outputData); | 552 | + } else if (hOut > hIn && wOut == static_cast<OFFSET_T>(3) && wOut <= wIn && wIn >= static_cast<OFFSET_T>(24)) { |
| 553 | + asc_vf_call<AdaptiveAvgPool2dGradSmallOutWSegFast<VALUE_T, OFFSET_T, 3>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); | ||
| 554 | + } else if (hOut > hIn && wOut == static_cast<OFFSET_T>(2) && wOut <= wIn && wIn >= static_cast<OFFSET_T>(32)) { | ||
| 555 | + asc_vf_call<AdaptiveAvgPool2dGradSmallOutWSegFast<VALUE_T, OFFSET_T, 2>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); | ||
| 556 | + } else if (hOut > hIn && wOut == static_cast<OFFSET_T>(2) && wOut <= wIn && wIn <= static_cast<OFFSET_T>(16)) { | ||
| 557 | + if (hOut >= hIn * static_cast<OFFSET_T>(1024)) { | ||
| 558 | + asc_vf_call<AdaptiveAvgPool2dGradHExpandW2SmallFast<VALUE_T, OFFSET_T>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); | ||
| 559 | + } else { | ||
| 560 | + asc_vf_call<AdaptiveAvgPool2dGradSmallOutWRow<VALUE_T, OFFSET_T, 2>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); | ||
| 561 | + } | ||
| 562 | + } else if (hOut > hIn && wOut <= wIn && hIn == static_cast<OFFSET_T>(2) && hOut == static_cast<OFFSET_T>(24)) { | ||
| 563 | + asc_vf_call<AdaptiveAvgPool2dGradHExpandExactOutWSmallFast<VALUE_T, OFFSET_T, 12>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); | ||
| 564 | + } else if (hOut > hIn && wOut <= wIn) { | ||
| 565 | + asc_vf_call<AdaptiveAvgPool2dGradOutWSmall<VALUE_T, OFFSET_T>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); | ||
| 566 | + } else if (hIn > hOut && wOut >= wIn * static_cast<OFFSET_T>(4)) { | ||
| 567 | + asc_vf_call<AdaptiveAvgPool2dGradHReduceWExpandFast<VALUE_T, OFFSET_T>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); | ||
| 568 | + } else { | ||
| 569 | + asc_vf_call<AdaptiveAvgPool2dGradNchw<VALUE_T, OFFSET_T>>(dim3(THREAD_DIM), params, gradData, n, c, hIn, wIn, hOut, wOut, outputData); | ||
| 570 | + } | ||
| 204 | } | 571 | } |
| 205 | 572 | ||
| 206 | } // namespace AdaptiveAvgPool2dGradOp | 573 | } // namespace AdaptiveAvgPool2dGradOp |
| 207 | - | ||
| 208 | 574 | ||
Mpooling/adaptive_avg_pool2d_grad/tests/ut/op_host/arch35/test_adaptive_avg_pool2d_grad_tiling.cpp+4-3
| @@ -296,6 +296,7 @@ TEST_F(AdaptiveAvgPool2dGradTilingTest, small_kernel_positive_branch_matrix) | |||
| 296 | { | 296 | { |
| 297 | const uint64_t smallInt32Key = GET_TPL_TILING_KEY(TPL_SMALL_KERNEL, TPL_INT32, 0); | 297 | const uint64_t smallInt32Key = GET_TPL_TILING_KEY(TPL_SMALL_KERNEL, TPL_INT32, 0); |
| 298 | const uint64_t smallInt64Key = GET_TPL_TILING_KEY(TPL_SMALL_KERNEL, TPL_INT64, 0); | 298 | const uint64_t smallInt64Key = GET_TPL_TILING_KEY(TPL_SMALL_KERNEL, TPL_INT64, 0); |
| 299 | + const uint64_t simtInt32Key = GET_TPL_TILING_KEY(TPL_SIMT_KERNEL, TPL_INT32, 0); | ||
| 299 | 300 | ||
| 300 | std::vector<SmallKernelPositiveCase> cases = { | 301 | std::vector<SmallKernelPositiveCase> cases = { |
| 301 | // TrySplitNC success, fp32 branch, NC exactly aligned by highAxisInner. | 302 | // TrySplitNC success, fp32 branch, NC exactly aligned by highAxisInner. |
| @@ -344,7 +345,7 @@ TEST_F(AdaptiveAvgPool2dGradTilingTest, small_kernel_positive_branch_matrix) | |||
| 344 | {1, 128, 11, 29}, | 345 | {1, 128, 11, 29}, |
| 345 | ge::DT_FLOAT, | 346 | ge::DT_FLOAT, |
| 346 | ge::FORMAT_NCHW, | 347 | ge::FORMAT_NCHW, |
| 347 | - smallInt32Key}, | 348 | + simtInt32Key}, |
| 348 | 349 | ||
| 349 | // W dynamic adjustment, but final wOutputTail aligned. | 350 | // W dynamic adjustment, but final wOutputTail aligned. |
| 350 | {"fp32_split_h_then_split_w_aligned_tail", | 351 | {"fp32_split_h_then_split_w_aligned_tail", |
| @@ -352,7 +353,7 @@ TEST_F(AdaptiveAvgPool2dGradTilingTest, small_kernel_positive_branch_matrix) | |||
| 352 | {1, 128, 11, 30}, | 353 | {1, 128, 11, 30}, |
| 353 | ge::DT_FLOAT, | 354 | ge::DT_FLOAT, |
| 354 | ge::FORMAT_NCHW, | 355 | ge::FORMAT_NCHW, |
| 355 | - smallInt32Key}, | 356 | + simtInt32Key}, |
| 356 | 357 | ||
| 357 | // Split W plus highAxisTail unaligned. | 358 | // Split W plus highAxisTail unaligned. |
| 358 | {"fp32_split_w_high_axis_tail_unaligned", | 359 | {"fp32_split_w_high_axis_tail_unaligned", |
| @@ -363,7 +364,7 @@ TEST_F(AdaptiveAvgPool2dGradTilingTest, small_kernel_positive_branch_matrix) | |||
| 363 | smallInt32Key}, | 364 | smallInt32Key}, |
| 364 | 365 | ||
| 365 | // Non-fp32 branch, SplitUnalignHW, W branch. | 366 | // Non-fp32 branch, SplitUnalignHW, W branch. |
| 366 | - {"fp16_split_w_branch", {1, 257, 5, 16}, {1, 257, 11, 31}, ge::DT_FLOAT16, ge::FORMAT_NCHW, smallInt32Key}, | 367 | + {"fp16_split_w_branch", {1, 257, 5, 16}, {1, 257, 11, 31}, ge::DT_FLOAT16, ge::FORMAT_NCHW, simtInt32Key}, |
| 367 | 368 | ||
| 368 | // CHW / FORMAT_NCL small kernel path, fp32 branch. | 369 | // CHW / FORMAT_NCL small kernel path, fp32 branch. |
| 369 | {"chw_fp32_small_kernel_split_hw", {257, 7, 19}, {257, 13, 29}, ge::DT_FLOAT, ge::FORMAT_NCL, smallInt32Key}, | 370 | {"chw_fp32_small_kernel_split_hw", {257, 7, 19}, {257, 13, 29}, ge::DT_FLOAT, ge::FORMAT_NCL, smallInt32Key}, |
🟡 Medium Priority
变更行 611–622：
PrintSplitData()原先通过OP_LOGI输出 splitData 中各字段的详细调试信息（包括 highAxisInner、hOutputInner、wOutputInner、totalBufferSize 等十余项关键数据）。新代码只计算了highAxisPadding、highAxisValidRate、ubUseRate、coreUseRate四个派生指标，但末尾的OP_LOGI/OP_LOGD调用被完全删除，这些计算结果未输出到任何地方。影响：该函数在
DoOpTiling()(第 629 行) 被调用，是算子运行时唯一的调试/可观测性输出路径。移除后，当出现 UB 超限、core 使用率异常等问题时，开发者无法从日志中获取 split 数据的运行时快照，调试能力严重退化。建议：恢复对 splitData 各字段的日志输出，并将新计算的 highAxisPadding / highAxisValidRate / ubUseRate / coreUseRate 也一并输出。例如在函数末尾添加：
OP_LOGI("AdaptiveAvgPool2dGradNCHW", "highAxisPadding=%ld validRate=%.2f ubUseRate=%.2f coreUseRate=%.2f", ...);
同时至少保留旧代码中总 bufferSize、usedCoreNum 等核心字段的输出。