UpSample GPU Porting

[xmnlab]

resolves #16158

[rgommers]

Fewer errors than before, 4 instead of 8 with the same message:

RuntimeError: CUDA error: too many resources requested for launch

This seems to be an existing problem: gh-8103.

@skrah could you have a look at this and tell us what you think?

[skrah]

Looking at it very briefly, the Jetson TX referenced in the issue only has 256 cores.

So while the issue looks the same, I'd probably not expect it on the CI platforms. FWIW, some recent CI tests in other issues are green.

[skrah]

@pytorchbot retest this please.

[rgommers]

same failures.

[skrah]

Is this rebased on the latest master (you can also ask the bot to rebase)?

[xmnlab]

I rebased that yesterday ... also needed to resolve conflicts because 7 days ago upsample files were changed.

[xmnlab]

@skrah do you have any idea about what could be this problem? or a way to debug that?
also it seems some jobs is taking a lot of time to build, for example, for some job the estimate time is 19h .. not sure if it the regular estimate ..

[skrah]

@xmnlab If you can't reproduce it at home it seems hard to debug other than reading the diffs again. I can take a look on Monday, it's getting a bit late here.

[skrah]

Also, has anyone found a way to show the actual hardware used on the CI in detail?

[xmnlab]

@skrah I am working in parallel on a paperspace environment .. it tooks a lot of time it is running with 8 cpu cores.

[skrah]

On Fri, May 03, 2019 at 07:12:35PM +0000, Ivan Ogasawara wrote: @skrah I am working in parallel on a paperspace environment .. it tooks a lot of time it is running with 8 cpu cores.

But have you ever been able to reproduce this issue on paperspace?

[rgommers]

But have you ever been able to reproduce this issue on paperspace?

I can reproduce it locally: Arch Linux, CUDA 10.0, RTX2070 GPU. Given the CI failures, I think it should be reproducible for multiple CUDA and GPU versions. I just haven't worked on any CUDA code before, and am short on time, so I'd rather not dig too deep.

[xmnlab]

not yet .. my last building was with a my previous commit ... I will work again in this task in some minutes :)

[rgommers]

@xmnlab your previous commit had the same failure though (except for the less clear exception message), and it seems quite reproducible. So I think you'll see it now.

[skrah]

On Fri, May 03, 2019 at 12:19:42PM -0700, Ralf Gommers wrote: > But have you ever been able to reproduce this issue on paperspace? I can reproduce it locally: Arch Linux, CUDA 10.0, RTX2070 GPU. Given the CI failures, I think it should be reproducible for multiple CUDA and GPU versions. I just haven't worked on any CUDA code before, and am short on time, so I'd rather not dig too deep.

Ah, thanks. @xmnlab, then I guess just compiling with "DEBUG=1" and stepping through the offending code with gdb may be the fastest way.

[xmnlab]

@skrah thanks! I will try that!

[xmnlab]

it seems just UpSampleBicubic2d is using upsample_get_value_bounded (https://github.com/pytorch/pytorch/pull/19630/files#diff-5092da792c30694ee4adf0d0ae2a37c6R171) and upsample_increment_value_bounded (https://github.com/pytorch/pytorch/pull/19630/files#diff-5092da792c30694ee4adf0d0ae2a37c6R191)

maybe the problem could be inside one of these functions ... maybe related to the order of indexes (x, y) ...

but the problem seems to be related to cuda block/threads ... so not sure if it is really related to these functions.

[skrah]

The new code uses far more registers than the existing one. I verified that the both versions actually call the offending test case with 1024 in blockDim. So it's very likely a register issue.

Existing uses 64 registers, which seems to be optimal for my card:

ptxas info    : 77696 bytes gmem, 72 bytes cmem[3]
ptxas info    : Compiling entry function '_Z23bicubic_interp2d_kerneliddb15THCDeviceTensorIdLi4Ei16DefaultPtrTraitsES1_' for 'sm_61'
ptxas info    : Function properties for _Z23bicubic_interp2d_kerneliddb15THCDeviceTensorIdLi4Ei16DefaultPtrTraitsES1_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 64 registers, 432 bytes cmem[0], 24 bytes cmem[2]

New uses 124 registers, which is too much for my card:

ptxas info    : 77696 bytes gmem, 72 bytes cmem[3]
ptxas info    : Compiling entry function '_ZN2at6native76_GLOBAL__N__52_tmpxft_0000626d_00000000_6_UpSampleBicubic2d_cpp1_ii_b4c1e1f328upsample_bicubic2d_out_frameElddbNS_20PackedTensorAccessorIdLm4ENS_16DefaultPtrTraitsElEES4_' for 'sm_61'
ptxas info    : Function properties for _ZN2at6native76_GLOBAL__N__52_tmpxft_0000626d_00000000_6_UpSampleBicubic2d_cpp1_ii_b4c1e1f328upsample_bicubic2d_out_frameElddbNS_20PackedTensorAccessorIdLm4ENS_16DefaultPtrTraitsElEES4_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 124 registers, 496 bytes cmem[0], 24 bytes cmem[2]

With __launch_bounds__(1024), the code again uses 64 registers.

If you use C10_LAUNCH_BOUNDS_1(1024) for both kernels, the tests pass here.

Now why is the regcount higher in the new code? It could be PackedTensorAccessor, it could be the fact that many instances of int have been changed to int64_t. :)

You could experiment or just use the launch bounds. Other code in native/cuda seems to use lower bounds for blockDim, too. 1024 seems to be an outlier.

[xmnlab]

@skrah

it seems it worked locally! thank you so much! I really appreciate that!

[xmnlab]

thanks @rgommers and @skrah for all the help!

@ezyang it is done for review!

[xmnlab]

thanks so much @ezyang!
I will let you know when it is ready again for a new review! thanks!

[xmnlab]

@skrah @ezyang
not sure but the errors on CI seems to be related to jenkins ... it seems all these jobs that failed (for building) ran by 01h 01min ... not sure if it is a coincidence ...

[rgommers]

@xmnlab indeed it looks like all jobs were aborted at the same time. no obvious issues in the build log related to your code. Comparing e.g. the cuda9-cudnn7 build with a successful one from another PR, it takes 1hr 14min there and your build is about at the place where the other one is after an hour.

I suggest to just push a new commit to rebuild. Probably a temporary CI hiccup.

[ezyang]

I accidentally rebooted Jenkins yesterday which is the likely cause, my apologies.

@pytorchbot retest this please

[xmnlab]

@ezyang @skrah @rgommers

all tests passed except pr/caffe2-py2-cuda9.0-cudnn7-windows-build:

14:43:49 Build timed out (after 180 minutes). Marking the build as failed.
14:43:49 Build was aborted
14:43:49 [BFA] Scanning build for known causes...
14:43:49 [BFA] No failure causes found
14:43:49 [BFA] Done. 0s
14:43:49 Finished: FAILURE

not sure if this timeout means that the code now is slower.

what do you think? do you have any suggestion?

[ezyang]

Sometimes the Windows build flakes out like that. It didn't timeout while running a relevant test, so I judge it to be not your problem.

[ezyang]

The launch bounds logic is wrong, but I acknowledge that this is a big patch already; just fix it in a follow up. I am going to go ahead and land this.

[facebook-github-bot]

@ezyang is landing this pull request. If you are a Facebook employee, you can view this diff on Phabricator.

[xmnlab]

sounds good @ezyang thanks!

[facebook-github-bot]

@ezyang merged this pull request in 3479777.