已合并
fix(softsign)：修复精度问题，修复空tensor问题 #8479
yulianjie创建于 26 天前
fix(softsign)：修复精度问题，修复空tensor问题 #8479
已合并
共 3 个文件变更+44-8
| @@ -46,11 +46,11 @@ static ge::graphStatus RunEleTiling(gert::TilingContext* context, SoftsignTiling | |||
| 46 | static ge::graphStatus DoTilingByDtype(gert::TilingContext* context, ge::DataType dtype, SoftsignTilingData* tilingData) | 46 | static ge::graphStatus DoTilingByDtype(gert::TilingContext* context, ge::DataType dtype, SoftsignTilingData* tilingData) |
| 47 | { | 47 | { |
| 48 | if (dtype == ge::DT_FLOAT) { | 48 | if (dtype == ge::DT_FLOAT) { |
| 49 | - return RunEleTiling<SoftsignOp::GraphSoftsign<float, float>::OpDag>(context, tilingData); | 49 | + return RunEleTiling<SoftsignOp::GraphSoftsignByDtype<float>::OpDag>(context, tilingData); |
| 50 | } else if (dtype == ge::DT_FLOAT16) { | 50 | } else if (dtype == ge::DT_FLOAT16) { |
| 51 | - return RunEleTiling<SoftsignOp::GraphSoftsign<half, float>::OpDag>(context, tilingData); | 51 | + return RunEleTiling<SoftsignOp::GraphSoftsignByDtype<half>::OpDag>(context, tilingData); |
| 52 | } else if (dtype == ge::DT_BF16) { | 52 | } else if (dtype == ge::DT_BF16) { |
| 53 | - return RunEleTiling<SoftsignOp::GraphSoftsign<bfloat16_t, float>::OpDag>(context, tilingData); | 53 | + return RunEleTiling<SoftsignOp::GraphSoftsignByDtype<bfloat16_t>::OpDag>(context, tilingData); |
| 54 | } | 54 | } |
| 55 | OP_LOGE(context, "Softsign: unsupported dtype=%d", static_cast<int>(dtype)); | 55 | OP_LOGE(context, "Softsign: unsupported dtype=%d", static_cast<int>(dtype)); |
| 56 | return ge::GRAPH_FAILED; | 56 | return ge::GRAPH_FAILED; |
| @@ -21,10 +21,14 @@ | |||
| 21 | * | -> Div -> CopyOut -> y (GM) | 21 | * | -> Div -> CopyOut -> y (GM) |
| 22 | * -> Abs -> Adds(+1) -----/ | 22 | * -> Abs -> Adds(+1) -----/ |
| 23 | * | 23 | * |
| 24 | - * 数据流 (FP16/BF16): | 24 | + * 数据流 (FP16/BF16, 与 TensorFlow 的原 dtype 计算语义对齐): |
| 25 | - * x (GM) -> CopyIn -> CastIn(→FP32) -> SaveX(Vec::Copy) -----\ | 25 | + * x (GM) -> CopyIn -> CastIn(→FP32) -> SaveX(Vec::Copy) ----------------------------\ |
| 26 | - * | -> Div -> CastOut(→U) -> CopyOut -> y (GM) | 26 | + * | -> Div |
| 27 | - * -> Abs -> Adds(+1) ---------/ | 27 | + * -> Abs -> Adds(+1) -> CastDenomOut(→U) -> CastDenomIn(→FP32) -/ |
| 28 | + * Div -> CastOut(→U) -> CopyOut -> y (GM) | ||
| 29 | + * | ||
| 30 | + * TensorFlow 在 FP16/BF16 下会先将 (abs(x) + 1) 舍入到原 dtype。因此分母需要显式窄化后再拓宽， | ||
| 31 | + * 不能在 FP32 中一次性完成所有中间计算。 | ||
| 28 | */ | 32 | */ |
| 29 | 33 | ||
| 30 | 34 | ||
| @@ -62,6 +66,35 @@ struct GraphSoftsign { | |||
| 62 | using OpDag = DAGSch<Outputs>; | 66 | using OpDag = DAGSch<Outputs>; |
| 63 | }; | 67 | }; |
| 64 | 68 | ||
| 69 | +template <typename U, typename T = float> | ||
| 70 | +struct GraphSoftsignLowPrecision { | ||
| 71 | + using ConstOne = MAKE_CONST(float, 1); | ||
| 72 | + | ||
| 73 | + using OpCopyIn = Bind<Vec::CopyIn<U>, Placeholder::In0<U>>; | ||
| 74 | + using OpCastIn = Bind<Vec::Cast<T, U, 0>, OpCopyIn>; | ||
| 75 | + using OpSaveX = Bind<Vec::Copy<T>, OpCastIn>; | ||
| 76 | + using OpAbs = Bind<Vec::Abs<T>, OpCastIn>; | ||
| 77 | + using OpAdds = Bind<Vec::Adds<T>, OpAbs, ConstOne>; | ||
| 78 | + using OpCastDenomOut = Bind<Vec::Cast<U, T, 1>, OpAdds>; | ||
| 79 | + using OpCastDenomIn = Bind<Vec::Cast<T, U, 0>, OpCastDenomOut>; | ||
| 80 | + using OpDiv = Bind<Vec::Div<T>, OpSaveX, OpCastDenomIn>; | ||
| 81 | + using OpCastOut = Bind<Vec::Cast<U, T, 1>, OpDiv>; | ||
| 82 | + using OpCopyOut = Bind<Vec::CopyOut<U>, Placeholder::Out0<U>, OpCastOut>; | ||
| 83 | + | ||
| 84 | + using Outputs = Elems<OpCopyOut>; | ||
| 85 | + using OpDag = DAGSch<Outputs>; | ||
| 86 | +}; | ||
| 87 | + | ||
| 88 | +template <typename U> | ||
| 89 | +struct GraphSoftsignByDtype { | ||
| 90 | + using OpDag = typename GraphSoftsignLowPrecision<U, float>::OpDag; | ||
| 91 | +}; | ||
| 92 | + | ||
| 93 | +template <> | ||
| 94 | +struct GraphSoftsignByDtype<float> { | ||
| 95 | + using OpDag = typename GraphSoftsign<float, float>::OpDag; | ||
| 96 | +}; | ||
| 97 | + | ||
| 65 | } // namespace SoftsignOp | 98 | } // namespace SoftsignOp |
| 66 | 99 | ||
| 67 | 100 | ||
| @@ -33,9 +33,12 @@ __global__ __aicore__ void softsign(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ | |||
| 33 | KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); | 33 | KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); |
| 34 | REGISTER_TILING_DEFAULT(SoftsignTilingData); | 34 | REGISTER_TILING_DEFAULT(SoftsignTilingData); |
| 35 | GET_TILING_DATA_WITH_STRUCT(SoftsignTilingData, tilingData, tiling); | 35 | GET_TILING_DATA_WITH_STRUCT(SoftsignTilingData, tilingData, tiling); |
| 36 | + if (tilingData.baseTiling.blockNum == 0) { | ||
| 37 | + return; | ||
| 38 | + } | ||
| 36 | TPipe pipe; | 39 | TPipe pipe; |
| 37 | using InputT = DTYPE_X; | 40 | using InputT = DTYPE_X; |
| 38 | - using OpDag = GraphSoftsign<InputT, float>::OpDag; | 41 | + using OpDag = GraphSoftsignByDtype<InputT>::OpDag; |
| 39 | ElementwiseSch<0UL, OpDag> sch(&(tilingData.baseTiling), &pipe); | 42 | ElementwiseSch<0UL, OpDag> sch(&(tilingData.baseTiling), &pipe); |
| 40 | sch.Init(x, y); | 43 | sch.Init(x, y); |
| 41 | sch.Process(); | 44 | sch.Process(); |