已合并
[fix] resolve bool type 'repeatTimes' overflow issue #3564
nextyale创建于 6月26日
[fix] resolve bool type 'repeatTimes' overflow issue #3564
已合并
共 4 个文件变更+54-5
| @@ -59,6 +59,35 @@ | |||
| 59 | } | 59 | } |
| 60 | ] | 60 | ] |
| 61 | }, | 61 | }, |
| 62 | + { | ||
| 63 | + "bin_filename": "AngleV2_bf16", | ||
| 64 | + "inputs": [ | ||
| 65 | + { | ||
| 66 | + "name": "x", | ||
| 67 | + "index": 0, | ||
| 68 | + "dtype": "bfloat16", | ||
| 69 | + "format": "ND", | ||
| 70 | + "paramType": "required", | ||
| 71 | + "shape": [ | ||
| 72 | + -2 | ||
| 73 | + ], | ||
| 74 | + "format_match_mode": "FormatAgnostic" | ||
| 75 | + } | ||
| 76 | + ], | ||
| 77 | + "outputs": [ | ||
| 78 | + { | ||
| 79 | + "name": "y", | ||
| 80 | + "index": 0, | ||
| 81 | + "dtype": "float32", | ||
| 82 | + "format": "ND", | ||
| 83 | + "paramType": "required", | ||
| 84 | + "shape": [ | ||
| 85 | + -2 | ||
| 86 | + ], | ||
| 87 | + "format_match_mode": "FormatAgnostic" | ||
| 88 | + } | ||
| 89 | + ] | ||
| 90 | + }, | ||
| 62 | { | 91 | { |
| 63 | "bin_filename": "AngleV2_a8c107f79de7b47ce419f0ec015539f7", | 92 | "bin_filename": "AngleV2_a8c107f79de7b47ce419f0ec015539f7", |
| 64 | "inputs": [ | 93 | "inputs": [ |
| @@ -70,7 +70,7 @@ extern "C" __global__ __aicore__ void angle_v2(GM_ADDR x, GM_ADDR y, GM_ADDR wor | |||
| 70 | op.Init(x, y, &tilingData, &pipe); | 70 | op.Init(x, y, &tilingData, &pipe); |
| 71 | op.Process(); | 71 | op.Process(); |
| 72 | } | 72 | } |
| 73 | -#if (__CCE_AICORE__ >= 320) // bf16 is only supported by 950 | 73 | +#if (__CCE_AICORE__ >= 310) // bf16 is only supported by 950 |
| 74 | else if (TILING_KEY_IS(KEY_DTYPE_BF16)) { | 74 | else if (TILING_KEY_IS(KEY_DTYPE_BF16)) { |
| 75 | AngleV2N::AngleV2<bfloat16_t> op; | 75 | AngleV2N::AngleV2<bfloat16_t> op; |
| 76 | op.Init(x, y, &tilingData, &pipe); | 76 | op.Init(x, y, &tilingData, &pipe); |
| @@ -23,6 +23,7 @@ using namespace AscendC; | |||
| 23 | 23 | ||
| 24 | constexpr int32_t BUFFER_NUM = 2; | 24 | constexpr int32_t BUFFER_NUM = 2; |
| 25 | constexpr int32_t COEFFICENT = 2; | 25 | constexpr int32_t COEFFICENT = 2; |
| 26 | +constexpr uint8_t UINT8_MAX_VALUE = 255; | ||
| 26 | const double PI = 3.14159265358979323846; | 27 | const double PI = 3.14159265358979323846; |
| 27 | 28 | ||
| 28 | 29 | ||
| @@ -30,7 +30,15 @@ public: | |||
| 30 | { | 30 | { |
| 31 | pipe = inputPipe; | 31 | pipe = inputPipe; |
| 32 | this->BaseMemberDataInit(tilingData); | 32 | this->BaseMemberDataInit(tilingData); |
| 33 | - repeatTimes = (this->tileLength + this->mask - 1) / this->mask; | 33 | + uint32_t actualRepeatTimes = (this->tileLength + this->mask - 1) / this->mask; |
| 34 | + if (actualRepeatTimes > UINT8_MAX_VALUE) { | ||
| 35 | + repeatTimes = UINT8_MAX_VALUE; | ||
| 36 | + loopCnt = actualRepeatTimes / repeatTimes; | ||
| 37 | + repeatLeft = actualRepeatTimes % repeatTimes; | ||
| 38 | + } else { | ||
| 39 | + repeatTimes = static_cast<uint8_t>(actualRepeatTimes); | ||
| 40 | + repeatLeft = 0; | ||
| 41 | + } | ||
 | |||
| 34 | blockLen = this->tileLength / dataPerBlock; | 42 | blockLen = this->tileLength / dataPerBlock; |
| 35 | 43 | ||
| 36 | yGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(y) + this->offset, this->blockLength); | 44 | yGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(y) + this->offset, this->blockLength); |
| @@ -42,10 +50,19 @@ public: | |||
| 42 | __aicore__ inline void Process() | 50 | __aicore__ inline void Process() |
| 43 | { | 51 | { |
| 44 | LocalTensor<yType> zeroTensor = outQueue.AllocTensor<yType>(); | 52 | LocalTensor<yType> zeroTensor = outQueue.AllocTensor<yType>(); |
| 45 | - Duplicate( | ||
| 46 | - zeroTensor, static_cast<yType>(0.0), this->mask, repeatTimes, this->dupDstBlockStride, | ||
| 47 | - this->dupDstRepeatStride); | ||
| 48 | 53 | ||
| 54 | + for (uint32_t i = 0; i < loopCnt; i++) { | ||
| 55 | + Duplicate( | ||
| 56 | + zeroTensor, static_cast<yType>(0.0), this->mask, repeatTimes, this->dupDstBlockStride, | ||
| 57 | + this->dupDstRepeatStride); | ||
| 58 | + } | ||
| 59 | + if (repeatLeft >= 1) { | ||
| 60 | + Duplicate( | ||
| 61 | + zeroTensor[repeatTimes * loopCnt * this->mask], static_cast<yType>(0.0), this->mask, repeatLeft, this->dupDstBlockStride, | ||
| 62 | + this->dupDstRepeatStride); | ||
| 63 | + } | ||
| 64 | + | ||
| 65 | + PipeBarrier<PIPE_ALL>(); | ||
| 49 | // loop count need to be doubled, due to double buffer | 66 | // loop count need to be doubled, due to double buffer |
| 50 | for (int64_t i = 0; i < this->tileNum; i++) { | 67 | for (int64_t i = 0; i < this->tileNum; i++) { |
| 51 | int64_t coreOffset = i * this->tileLength; | 68 | int64_t coreOffset = i * this->tileLength; |
| @@ -66,6 +83,8 @@ private: | |||
| 66 | TQue<QuePosition::VECOUT, BUFFER_NUM> outQueue; | 83 | TQue<QuePosition::VECOUT, BUFFER_NUM> outQueue; |
| 67 | uint8_t repeatTimes; | 84 | uint8_t repeatTimes; |
| 68 | int32_t dataPerBlock = 32 / sizeof(yType); | 85 | int32_t dataPerBlock = 32 / sizeof(yType); |
| 86 | + uint32_t loopCnt = 1; | ||
| 87 | + uint8_t repeatLeft = 0; | ||
| 69 | uint16_t blockLen = 1; | 88 | uint16_t blockLen = 1; |
| 70 | }; | 89 | }; |
| 71 | } // namespace AngleV2N | 90 | } // namespace AngleV2N |
🟠 High Priority
在
Init()的else分支（即actualRepeatTimes <= 255的非溢出路径）中，仅设置了repeatTimes = static_cast<uint8_t>(actualRepeatTimes)，但成员变量repeatLeft保留了类内默认值1，loopCnt保留了默认值1。进入
Process()后：总重复次数 =
repeatTimes + 1，而非预期的actualRepeatTimes。例如actualRepeatTimes = 5时，实际写出 6 个 repeat；actualRepeatTimes = 0（极端边界情况）时，实际写出 1 个 repeat。额外写入的
mask个零值元素可能超出zeroTensor的分配边界（tileLength个元素），存在越界写入风险。这是本次 diff 引入的逻辑回归——原代码只有一次Duplicate，正确做了repeatTimes次 repeat。建议：在 else 分支中将
repeatLeft设为 0（或同时将loopCnt保持为 1），确保Process()中repeatLeft >= 1的分支不被触发，总重复次数恰好等于repeatTimes。