Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Fix bug when dtype=fp16 in deformable conv #46975

Closed
wants to merge 8 commits into from
Closed

Fix bug when dtype=fp16 in deformable conv #46975

wants to merge 8 commits into from

Conversation

Rayman96
Copy link
Contributor

PR types

New features

PR changes

OPs

Describe

修复FP16中dx精度bug，基于#46111

@paddle-bot
Copy link

paddle-bot bot commented Oct 13, 2022

你的PR提交成功，感谢你对开源项目的贡献!
请关注后续CI自动化测试结果，详情请参考Paddle-CI手册。
Your PR has been submitted. Thanks for your contribution!
Please wait for the result of CI firstly. See Paddle CI Manual for details.

@paddle-bot paddle-bot bot added contributor External developers status: proposed labels Oct 13, 2022
@Rayman96 Rayman96 mentioned this pull request Oct 13, 2022
Merged
@Rayman96 Rayman96 marked this pull request as ready for review October 13, 2022 06:09
Rayman96
Rayman96 commented Oct 13, 2022
View reviewed changes
paddle/phi/kernels/impl/deformable_conv_grad_kernel_impl.h Outdated
Comment on lines 305 to 308
DenseTensor mt_dx = phi::EmptyLike<MT, Context>(dev_ctx, *dx);
MT* mt_dx_ptr = (x.dtype() == DataType::FLOAT16)
? (dev_ctx.template Alloc<MT>(&mt_dx))
: (dev_ctx.template Alloc<MT>(dx));
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@zhangting2020 这里已修改 从L305 ~ L324，对于fp16的情况做了结果转换。本地测试结果符合预期
image

image

@Rayman96
Copy link
Contributor Author

由于前向加速的先合入了，导致这个代码有一行conflict。已修复重新跑CI流程

@Ligoml Ligoml mentioned this pull request Oct 14, 2022
Closed
@zhangting2020
Copy link
Contributor

目前这个实现方式你创建了一个新的tensor，会增加内存，同时又插入了一个cast op，也会引入额外的代价。这个方案需要再斟酌，你可以借鉴下lookup_table_v2算子的梯度实现，其中也使用了原子操作，但我们过去也为它做了fp16的优化。

@Rayman96
Copy link
Contributor Author

目前这个实现方式你创建了一个新的tensor，会增加内存，同时又插入了一个cast op，也会引入额外的代价。这个方案需要再斟酌，你可以借鉴下lookup_table_v2算子的梯度实现，其中也使用了原子操作，但我们过去也为它做了fp16的优化。

好的 我参考下实现，试一下速度是否可以满足要求

@Rayman96
Copy link
Contributor Author

@zhangting2020 paddle/fluid/operators/lookup_table_v2_op.cu
image
是这个paddle::platform::VectorizedAtomicAddPerBlock为fp16做过优化吗？
其他类似的操作找到比如paddle::platform::fastAtomicAdd我尝试过速度会很慢

@Rayman96
Copy link
Contributor Author

@zhangting2020 paddle/fluid/operators/lookup_table_v2_op.cu image 是这个paddle::platform::VectorizedAtomicAddPerBlock为fp16做过优化吗？ 其他类似的操作找到比如paddle::platform::fastAtomicAdd我尝试过速度会很慢

VectorizedAtomicAddPerBlock看了下实现也是调用了CudaAtomicAdd，在fp16时速度和fp32差距还是会比较大。

@Rayman96
Copy link
Contributor Author

Rayman96 commented Oct 14, 2022 •
edited
Loading

@zhangting2020 现在的代码会在fp32和fp64的时候也多增加一个新的tensor内存，这部分确实是浪费的。如果只在fp16是增加新的tensor和Cast_op，这样用内存换时间是可以接受的吗？
目前已这样修改

@Rayman96
Copy link
Contributor Author

Rayman96 commented Oct 15, 2022 •
edited
Loading

lookup_table_v2的优化方式我理解是在将AtomAdd的过程单独作为一个kernel，且总数据量是偶数时，将两个数据构造为__half2，利用一次计算实现。
这个思路对于deform的dx计算不太适用，dx_grad的每个位置目前都是增加了一个在kernel中计算的weight结果，在每个thread中只需要有一个 sizeof(T) 的大小占用。而如果使用“将两个数据构造为__half2”的方式，需要单独再开辟一块儿内存将全部weight的值存下来（无法消除新的内存使用），并且将weight的计算和Atomadd的过程拆离成两个kernel。
另外由于理论上使用__half2会降低至多一半的时间，但目前测试使用fp16的atomAdd的速度会比fp32的二倍加cast的速度慢。（无法提高速度）
所以觉得使用当前实现是目前较优解。

@zhangting2020
Copy link
Contributor

lookup_table_v2的优化方式我理解是在将AtomAdd的过程单独作为一个kernel，且总数据量是偶数时，将两个数据构造为__half2，利用一次计算实现。 这个思路对于deform的dx计算不太适用，dx_grad的每个位置目前都是增加了一个在kernel中计算的weight结果，在每个thread中只需要有一个 sizeof(T) 的大小占用。而如果使用“将两个数据构造为__half2”的方式，需要单独再开辟一块儿内存将全部weight的值存下来（无法消除新的内存使用），并且将weight的计算和Atomadd的过程拆离成两个kernel。 另外由于理论上使用__half2会降低至多一半的时间，但目前测试使用fp16的atomAdd的速度会比fp32的二倍加cast的速度慢。（无法提高速度） 所以觉得使用当前实现是目前较优解。

在Paddle框架中混合精度训练机制会在算子计算前将输入cast到fp32，计算用fp32，计算结果cast为fp16，但这样会引入额外的cast算子的开销，并且fp16相比fp32没有显著加速。正是为了提升混合精度训练的性能，才会设计这项任务。目前这版的方案恐怕无法满足预期，因为这个任务的重点之一就是对fp16性能完成优化

@Rayman96
Copy link
Contributor Author

lookup_table_v2的优化方式我理解是在将AtomAdd的过程单独作为一个kernel，且总数据量是偶数时，将两个数据构造为__half2，利用一次计算实现。 这个思路对于deform的dx计算不太适用，dx_grad的每个位置目前都是增加了一个在kernel中计算的weight结果，在每个thread中只需要有一个 sizeof(T) 的大小占用。而如果使用“将两个数据构造为__half2”的方式，需要单独再开辟一块儿内存将全部weight的值存下来（无法消除新的内存使用），并且将weight的计算和Atomadd的过程拆离成两个kernel。 另外由于理论上使用__half2会降低至多一半的时间，但目前测试使用fp16的atomAdd的速度会比fp32的二倍加cast的速度慢。（无法提高速度） 所以觉得使用当前实现是目前较优解。

在Paddle框架中混合精度训练机制会在算子计算前将输入cast到fp32，计算用fp32，计算结果cast为fp16，但这样会引入额外的cast算子的开销，并且fp16相比fp32没有显著加速。正是为了提升混合精度训练的性能，才会设计这项任务。目前这版的方案恐怕无法满足预期，因为这个任务的重点之一就是对fp16性能完成优化

这项任务总共设计四部分计算，前向output使用fp16速度有提升，后向grad_offset, grad_filter的速度和fp32持平，只有grad_input由于计算的特殊引入cast增加了大概1.6%的计算时间。四部分整体上是和fp32持平的，测试中各有胜负，当时题目中要求也是性能不差于fp32。

@Rayman96
Copy link
Contributor Author

benchmark中例子速度基本一致，但我换用更大的测试用例后cast过程确实很耗时无法忽略

@Rayman96
Copy link
Contributor Author

@zhangting2020 之前只依赖benchmark中唯一的例子我以为cast的开销是很小的，换更大用例后发现还是挺高的。dx这里的计算我后边再继续想想办法，参考其他算子只使用更快的atomAdd也无法满足优于fp32的需要，可能需要改下dx的计算过程，今天比赛截止前应该无法满足更优实现。
目前这个题应该只有我提交，除dx外其余部分应该没什么问题，如果这道题可以算点分的话就更好了😂

@luotao1 luotao1 closed this Dec 6, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers
No reviews
Assignees

@luotao1 luotao1

@zhangting2020 zhangting2020

Labels
contributor External developers
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
4 participants
@Rayman96 @zhangting2020 @luotao1 @Ligoml