Fix launch bounds in spatial transformer

[ptrendx]

Description

Without launch_bounds compiler is not required to use small enough number of registers to fit 1024 threads per block. Our internal CI with CUDA 10 build was failing on V100 because of this.

Checklist

Essentials

Please feel free to remove inapplicable items for your PR:

• To the my best knowledge, examples are either not affected by this change, or have been fixed to be compatible with this change

Changes

• Added launch_bounds guards around BilinearSampling[Forward,Backward]Kernel to ensure that the compiled operator works on each supported GPU.

Comments

[harshp8l]

@mxnet-label-bot add [pr-awaiting-review]

[stu1130]

@samskalicky @access2rohit @anirudh2290 could you please review it thanks!

[vandanavk]

@mxnet-label-bot add [Operator]

@apeforest for review

[Roshrini]

@eric-haibin-lin @anirudh2290 @samskalicky can you please review this PR?

[samskalicky]

@ptrendx was there a corresponding issue filed for the V100 CI failure?

[ptrendx]

It was our internal (NVIDIA's) CI failure, not external MXNet's CI, so no, there was not an issue for it.

[samskalicky]

@ptrendx Changes look straightfoward, no issues but I dont quite understand the fix.

Is 1 thread the correct value? Why choose that value (ie. how about 2)?

Please add some comment regarding the reasoning for the 1 thread value.

Are there unit tests to check this change, how can we validate the correctness?

[ptrendx]

It is 1 block, not 1 thread. Basically what this change does is tell the compiler: I'm going to run this kernel with cuda::maxThreadsPerBlock threads and so make sure that at least 1 block fits on SM. It does not limit the kernel to run only 1 block, this is minimum value.
Basically without this change compiler does not know how many threads will be used to run the kernel and so is free to generate code that uses as many registers as it wants. However, the size of register file is not infinite and so the more registers are used by the kernel, the less threads can run at the same time and trying to run the kernel with more threads results in a failure to launch (with too many resources requested for launch error). Since the kernel is launched with cuda::maxThreadsPerBlock threads, we need to make sure that it can be launched on every architecture.

Actually, I just found an issue about spatial transformer giving exactly this error (that was "fixed" by changing the kernel and so luckily being below the threshold again): #11568

[anirudh2290]

@ptrendx makes sense, #11568 had two issues of which assert issue was fixed but the too many resources requested for launch was hidden.

[Roshrini]

@samskalicky Can you check if your comment is addressed? Thanks

[samskalicky]

@ptrendx Can you add a comment in the code to make it clear that these launch bounds are required to be set this way to avoid the "too many resources requested for launch" error?

Maybe something like:

/*
 * __launch_bounds__(cuda::kMaxThreadsPerBlock, 1)
 * This sets the number of threads per block to 1 which
 * reduces the number of register resources needed.
 * Running the kernel with more threads results in a 
 * failure to launch (too many resources requested error).
 */ 

[ptrendx]

I added a comment with explanation.

[samskalicky]

Thanks @ptrendx for your help with this issue!

[ptrendx]

Is there anything else needed for this PR?

[Roshrini]

@anirudh2290 Can you merge this PR if it looks good?
@mxnet-label-bot Update [Operator, pr-awaiting-merge]

[KellenSunderland]

Looks similar to a few fixes we've provided in the past when we had too many registers for a few kernels to run on a TX1. LGTM.