Fixes deform_conv issue with large input/output #4351
Merged
Conversation
This file contains 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
vfdev-5
force-pushed
the
fix-4269-deform-conv-index-overflow
branch
from
30a2324 to
3498de6
Compare
vfdev-5
commented
Sep 2, 2021
Comment on lines
+246
to
+299
| if (use_64bits_indexing) { | ||
| AT_DISPATCH_FLOATING_TYPES_AND_HALF( | ||
| input.scalar_type(), "deformable_im2col", ([&] { | ||
| deformable_im2col_kernel<scalar_t, int64_t><<<blocks, threads>>>( | ||
| num_kernels, | ||
| input.data_ptr<scalar_t>(), | ||
| data_offset.data_ptr<scalar_t>(), | ||
| data_mask.data_ptr<scalar_t>(), | ||
| height, | ||
| width, | ||
| weight_h, | ||
| weight_w, | ||
| pad_h, | ||
| pad_w, | ||
| stride_h, | ||
| stride_w, | ||
| dilation_h, | ||
| dilation_w, | ||
| parallel_imgs, | ||
| n_in_channels, | ||
| deformable_group, | ||
| out_h, | ||
| out_w, | ||
| use_mask, | ||
| data_col.data_ptr<scalar_t>()); | ||
| })); | ||
|
|
||
| } else { | ||
| AT_DISPATCH_FLOATING_TYPES_AND_HALF( | ||
| input.scalar_type(), "deformable_im2col", ([&] { | ||
| deformable_im2col_kernel<scalar_t, int><<<blocks, threads>>>( | ||
| num_kernels, | ||
| input.data_ptr<scalar_t>(), | ||
| data_offset.data_ptr<scalar_t>(), | ||
| data_mask.data_ptr<scalar_t>(), | ||
| height, | ||
| width, | ||
| weight_h, | ||
| weight_w, | ||
| pad_h, | ||
| pad_w, | ||
| stride_h, | ||
| stride_w, | ||
| dilation_h, | ||
| dilation_w, | ||
| parallel_imgs, | ||
| n_in_channels, | ||
| deformable_group, | ||
| out_h, | ||
| out_w, | ||
| use_mask, | ||
| data_col.data_ptr<scalar_t>()); | ||
| })); | ||
| } |
There was a problem hiding this comment.
@fmassa any clever ideas on how to reduce code duplication ?
There was a problem hiding this comment.
I think this is fine. You could write a macro here to call this only once, like what PyTorch does here for example, and then is dispatched here.
If you do this, don't forget to #undef the macros after they are used
fmassa
approved these changes
Sep 6, 2021
Comment on lines
+246
to
+299
| if (use_64bits_indexing) { | ||
| AT_DISPATCH_FLOATING_TYPES_AND_HALF( | ||
| input.scalar_type(), "deformable_im2col", ([&] { | ||
| deformable_im2col_kernel<scalar_t, int64_t><<<blocks, threads>>>( | ||
| num_kernels, | ||
| input.data_ptr<scalar_t>(), | ||
| data_offset.data_ptr<scalar_t>(), | ||
| data_mask.data_ptr<scalar_t>(), | ||
| height, | ||
| width, | ||
| weight_h, | ||
| weight_w, | ||
| pad_h, | ||
| pad_w, | ||
| stride_h, | ||
| stride_w, | ||
| dilation_h, | ||
| dilation_w, | ||
| parallel_imgs, | ||
| n_in_channels, | ||
| deformable_group, | ||
| out_h, | ||
| out_w, | ||
| use_mask, | ||
| data_col.data_ptr<scalar_t>()); | ||
| })); | ||
|
|
||
| } else { | ||
| AT_DISPATCH_FLOATING_TYPES_AND_HALF( | ||
| input.scalar_type(), "deformable_im2col", ([&] { | ||
| deformable_im2col_kernel<scalar_t, int><<<blocks, threads>>>( | ||
| num_kernels, | ||
| input.data_ptr<scalar_t>(), | ||
| data_offset.data_ptr<scalar_t>(), | ||
| data_mask.data_ptr<scalar_t>(), | ||
| height, | ||
| width, | ||
| weight_h, | ||
| weight_w, | ||
| pad_h, | ||
| pad_w, | ||
| stride_h, | ||
| stride_w, | ||
| dilation_h, | ||
| dilation_w, | ||
| parallel_imgs, | ||
| n_in_channels, | ||
| deformable_group, | ||
| out_h, | ||
| out_w, | ||
| use_mask, | ||
| data_col.data_ptr<scalar_t>()); | ||
| })); | ||
| } |
There was a problem hiding this comment.
I think this is fine. You could write a macro here to call this only once, like what PyTorch does here for example, and then is dispatched here.
If you do this, don't forget to #undef the macros after they are used
facebook-github-bot
pushed a commit
that referenced
this pull request
Sep 9, 2021
Summary: * WIP on fixing index overflow issue * Fixed backward pass for large num_kernels * Fixed clang formatting * Fixed GET_BLOCKS int/int64_t types issue Reviewed By: fmassa Differential Revision: D30793320 fbshipit-source-id: ce99a6c2c0f859b32d2c565da451640331f935f8 Co-authored-by: vfdev-5 <vfdev-5@gmail.com> Co-authored-by: Francisco Massa <fvsmassa@gmail.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
Fixes #4269
Description:
!!! Currently, no tests provided as reproducible example allocates large tensors - any suggestions on that ?