Fix FP16 fastAtomicAdd for one case where tensor start address is not 32 bit aligned #44642
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
added 3 commits
September 14, 2020 10:20
mcarilli
changed the title
ngimel
approved these changes
Sep 14, 2020
Codecov Report
@@ Coverage Diff @@
## master #44642 +/- ##
==========================================
- Coverage 67.98% 67.98% -0.01%
==========================================
Files 384 384
Lines 49567 49567
==========================================
- Hits 33697 33696 -1
- Misses 15870 15871 +1
Continue to review full report at Codecov.
|
Contributor
facebook-github-bot
left a comment
facebook-github-bot
left a comment
There was a problem hiding this comment.
@ngimel has imported this pull request. If you are a Facebook employee, you can view this diff on Phabricator.
Contributor
xuzhao9
pushed a commit
that referenced
this pull request
Sep 18, 2020
… 32 bit aligned (#44642) Summary: For #44206 and #42218, I'd like to update trilinear interpolate backward and grid_sample backward to use `fastAtomicAdd`. As a prelude, I spotted a UB risk in `fastAtomicAdd`. I think existing code incurs a misaligned `__half2` atomicAdd when `index` is odd and `tensor` is not 32-bit aligned (`index % 2 == 1` and `(reinterpret_cast<std::uintptr_t>(tensor) % sizeof(__half2) == 1`). In this case we think we're `!low_bit` and go down the `!low_bit` code path, but in fact we are `low_bit`. It appears the original [fastAtomicAdd PR](#21879 (comment) discussion did not consider that case explicitly. I wanted to push my tentative fix for discussion ASAP. jjsjann123 and mkolod as original authors of `fastAtomicAdd`. (I'm also curious why we need to `reinterpret_cast<std::uintptr_t>(tensor...` for the address modding, but that's minor.) Pull Request resolved: #44642 Reviewed By: mruberry Differential Revision: D23699820 Pulled By: ngimel fbshipit-source-id: 0db57150715ebb45e6a1fb36897e46f00d61defd
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
6 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
For #44206 and #42218, I'd like to update trilinear interpolate backward and grid_sample backward to use
fastAtomicAdd.As a prelude, I spotted a UB risk in
fastAtomicAdd. I think existing code incurs a misaligned__half2atomicAdd whenindexis odd andtensoris not 32-bit aligned (index % 2 == 1and(reinterpret_cast<std::uintptr_t>(tensor) % sizeof(__half2) == 1). In this case we think we're!low_bitand go down the!low_bitcode path, but in fact we arelow_bit. It appears the original fastAtomicAdd PR's discussion did not consider that case explicitly.I wanted to push my tentative fix for discussion ASAP. @jjsjann123 and @mkolod as original authors of
fastAtomicAdd. (I'm also curious why we need toreinterpret_cast<std::uintptr_t>(tensor...for the address modding, but that's minor.)