New high level interface for cuDNN #523

denizyuret · 2020-11-06T11:37:02Z

Please see README.md for some design decisions, comments welcome. I tried not to break any existing code by defining deprecated methods. Here is a todo list.

lib/cudnn/util.jl

denizyuret · 2020-11-07T16:18:53Z

@maleadt I need some advice on https://github.com/denizyuret/Knet.jl/blob/2dec7d58a9d40899c5c38d6f4f17af463163409f/src/cudnn/dropout.jl#L66-L88

As far as I can understand from cudnn docs, cudnnDropoutState should be similar to Random.GLOBAL_RNG. There is this warning about concurrently running dropouts using the same state. I don't know if this concerns threads on the CPU or streams on the GPU or both. How should we handle this? Did you have a similar issue with CURAND?

maleadt · 2020-11-10T15:42:39Z

If you rebase this on master, CI will work.

denizyuret · 2020-11-10T19:53:27Z

If you rebase this on master, CI will work.

I occasionally merge master. I don't know how/why to rebase.

maleadt · 2020-11-10T20:02:10Z

I occasionally merge master. I don't know how/why to rebase.

git rebase master (after pulling it, else git fetch && git rebase origin/master). Better do so, to avoid ugly merge commits.

denizyuret · 2020-11-14T08:35:04Z

I occasionally merge master. I don't know how/why to rebase.

git rebase master (after pulling it, else git fetch && git rebase origin/master). Better do so, to avoid ugly merge commits.

I was able to do a successful git rebase master once but now I get scary conflict messages when I try. We'll clean up the history when it is time to review.

simeonschaub · 2020-11-18T16:35:03Z

lib/cudnn/nnlib.jl

-                      algo=CUDNN_SOFTMAX_FAST, mode=cudnnSoftmaxMode_t(dims-1))
-  return y
+    mode, xsize = softmaxhelper(size(x), dims)
+    cudnnSoftmaxForward!(reshape(y,xsize), reshape(x,xsize); algo=CUDNN_SOFTMAX_FAST, mode)


Does this always use fast mode? Would it be possible to respect the default math mode here?

It does right now, but I agree it should be consistent. I am not sure what the "default math mode" should be here: there is a CUDNN.math_mode, but that's about tensor optimizations.

CUDA.math_mode is a more generic property, which CUDNN.math_mode copies for handle creation. But the former should be used for other purposes:

CUDA.jl/lib/cublas/wrappers.jl

Lines 757 to 763 in d16a4b3

math_mode = CUDA.math_mode()

reduced_precision = CUDA.math_precision()

if sig === (Float16, Float16)

# NOTE: Float16=Float16*Float16 can also happen in 32-bit compute

return math_mode==CUDA.PEDANTIC_MATH ? CUBLAS_COMPUTE_16F_PEDANTIC : CUBLAS_COMPUTE_16F

end

I think we should have the current math mode as the default and have libraries request the fast paths where possible.

codecov · 2020-11-20T11:46:16Z

Codecov Report

Merging #523 (3d016ba) into master (df9b4df) will increase coverage by 1.88%.
The diff coverage is 73.72%.

@@            Coverage Diff             @@
##           master     #523      +/-   ##
==========================================
+ Coverage   77.91%   79.79%   +1.88%     
==========================================
  Files         118      122       +4     
  Lines        7117     7329     +212     
==========================================
+ Hits         5545     5848     +303     
+ Misses       1572     1481      -91

Impacted Files	Coverage Δ
lib/cudnn/CUDNN.jl	`66.03% <ø> (ø)`
lib/cudnn/tensor.jl	`34.48% <34.48%> (-22.67%)`	⬇️
lib/cudnn/batchnorm.jl	`35.29% <42.85%> (ø)`
lib/cudnn/util.jl	`44.44% <50.00%> (-8.50%)`	⬇️
lib/cudnn/reduce.jl	`59.25% <59.25%> (ø)`
lib/cudnn/pooling.jl	`63.33% <62.06%> (-32.32%)`	⬇️
lib/cudnn/softmax.jl	`69.23% <66.66%> (-30.77%)`	⬇️
lib/cudnn/activation.jl	`71.42% <71.42%> (-22.33%)`	⬇️
lib/cudnn/rnn.jl	`71.92% <71.92%> (+71.92%)`	⬆️
lib/cudnn/nnlib.jl	`72.94% <72.22%> (+3.51%)`	⬆️
... and 19 more

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update df9b4df...0b57f52. Read the comment docs.

maleadt · 2020-12-09T08:21:44Z

CI failure is likely a regression in Julia fixed with LLVM#master.

maleadt · 2020-12-09T16:57:32Z

Rebasing on master should fix that.

denizyuret · 2020-12-16T13:41:28Z

@maleadt I think this PR is ready to merge. We can meet up to take care of squash/rebase etc. if you'd like.

lib/cudnn/README.md

DhairyaLGandhi · 2020-12-23T06:28:25Z

I've been playing with the PR for a couple days now, and I was noticing some performance regressions training with several models, but resnet shows it fairly simply

Current:

BenchmarkTools.Trial:
  memory estimate:  676.37 KiB
  allocs estimate:  22051
  --------------
  minimum time:     413.535 ms (0.99% GC)
  median time:      496.596 ms (0.82% GC)
  mean time:        571.457 ms (0.78% GC)
  maximum time:     952.702 ms (0.63% GC)
  --------------
  samples:          5
  evals/sample:     1

This:

BenchmarkTools.Trial:
  memory estimate:  348.56 KiB
  allocs estimate:  9584
  --------------
  minimum time:     619.104 ms (0.00% GC)
  median time:      685.281 ms (0.00% GC)
  mean time:        868.306 ms (0.00% GC)
  maximum time:     1.484 s (0.00% GC)
  --------------
  samples:          4
  evals/sample:     1

This is with an input of (224,224,3,128) using Resnet from Metalhead.jl. The problem seems to rise with batch sizes, although in some cudnn bound cases, I see smaller batch outperform the current 🎉

DhairyaLGandhi · 2021-01-16T13:54:04Z

We'll probably want to sync releases and compat accordingly.

DhairyaLGandhi · 2021-01-20T07:30:07Z

@denizyuret what are your thoughts on using the RNN primitives from CUDNN vs naive Julia ones. Some frameworks opt to not use the ones in CUDNN because of poorer performance

denizyuret · 2021-01-20T11:57:34Z

@denizyuret what are your thoughts on using the RNN primitives from CUDNN vs naive Julia ones. Some frameworks opt to not use the ones in CUDNN because of poorer performance

@DhairyaLGandhi, do you have any examples where CUDNN has poorer performance than an alternative implementation of RNN? In my experience a manually programmed RNN (e.g. https://github.com/denizyuret/Knet.jl/blob/master/tutorial/90.s2s.ipynb) is significantly slower than a CUDNN call (e.g. https://github.com/denizyuret/Knet.jl/blob/master/tutorial/60.rnn.ipynb).

This is expected because CUDNN parallelizes certain operations (e.g. matmul of the whole input sequence rather than individual time steps) that cannot be done with a single time-step implementation.

lib/cudnn/descriptors.jl

DhairyaLGandhi · 2021-01-21T07:52:08Z

lib/cudnn/convolution.jl

+# Allocate the maximum reasonable amount of memory for algorithm discovery
+function cudnnFindConvolutionAlgorithmWorkspaceSize(x)
+    gpufree = Mem.info()[1] + CUDA.cached_memory()
+    min(gpufree ÷ 10, sizeof(x) * 100)


100 seems a bit aggressive, it might mean we start going to slower algos as we get more memory constrained, when tensors should have been freed. Maybe @maleadt can clarify if my thinking is incorrect here

Mem.available_memory() instead of Mem.info()[1].

100x seems aggressive indeed; is there any guide in the documentation, or how did you come up with that?

DhairyaLGandhi · 2021-01-21T10:54:47Z

Still noticing some slowness;

# Release
julia> benchmark_bw_cu(64)
BenchmarkTools.Trial:
  memory estimate:  2.23 MiB
  allocs estimate:  48336
  --------------
  minimum time:     552.016 ms (1.72% GC)
  median time:      601.561 ms (1.44% GC)
  mean time:        598.092 ms (1.45% GC)
  maximum time:     637.229 ms (1.47% GC)
  --------------
  samples:          4
  evals/sample:     1

# This PR
julia> benchmark_bw_cu(64)
BenchmarkTools.Trial:
  memory estimate:  3.58 MiB
  allocs estimate:  57069
  --------------
  minimum time:     730.687 ms (2.78% GC)
  median time:      753.419 ms (5.19% GC)
  mean time:        809.694 ms (10.59% GC)
  maximum time:     944.977 ms (20.92% GC)
  --------------
  samples:          3
  evals/sample:     1

And I am OOMing at higher batch sizes, but I haven't tried TF to see how they do with a 32 GB V100

maleadt · 2021-01-27T07:15:34Z

There's some conflicts because of #672.

denizyuret · 2021-01-27T11:21:15Z

I will take a look. @DhairyaLGandhi when can we merge this to avoid future conflicts?

maleadt · 2021-01-28T09:41:35Z

Let's go ahead and merge this once the conflicts are resolved.

DhairyaLGandhi · 2021-01-28T10:42:32Z

Yes, we were discussing it for a couple days. We will need to change some nnlib APIs and get 3D recurrent layers with cudnn, and look into some regressions, but this is the base to grow from.

denizyuret · 2021-01-30T13:21:54Z

@DhairyaLGandhi @maleadt, I checked the conflict with the #672 merge. This is practically impossible to merge with #523 -- it would be easier for me to rewrite whatever functionality was intended. I spent a month on #523 redesigning all of lib/cudnn and test/cudnn. #672 used the master versions, e.g. the test file it is using has been moved, renamed and redesigned, nnlib.jl has been redesigned. If all that is needed is adding a beta keyword to the nnlib conv functions why don't I just add this to #523 instead of trying to merge?

DhairyaLGandhi · 2021-01-30T13:28:31Z

That should be fine, we could also revert #672, and rebase the functionality over on top of this. Whatever would be easier I think.

denizyuret · 2021-01-30T15:02:36Z

I can add alpha/beta keyword args in CUDA/lib/cudnn/nnlib.jl in this PR. I assume these have the same semantics as cuDNN and the default is alpha=1, beta=0. @DhairyaLGandhi you mentioned that these were supported in NNlib.jl but I did not see them in the master, am I looking at the wrong branch or did I misunderstand what #672 is trying to do?

DhairyaLGandhi · 2021-01-30T15:06:23Z

Yes, these are implemented here https://github.com/FluxML/NNlib.jl/blob/d9aaaf7ac4df683f1d3361efb3cbaef9aa8112a1/src/impl/conv_im2col.jl#L25

And the defaults are consistent with CuDNN

denizyuret · 2021-01-31T18:11:31Z

@DhairyaLGandhi I added explicit alpha/beta kwargs to all NNlib.conv related functions in this branch (lib/cudnn/nnlib.jl) along with associated tests in (test/cudnn/nnlib.jl). All tests pass except for NNlib.∇conv_data! with nonzero beta. The cpu version does nothing, the gpu version adds the output.

At this point we need to decide the semantics of alpha/beta in backward functions NNlib.∇conv*: there are two possibilities:
(1) They simply scale / add to the output as in the forward call and are independent of whether the forward call had alpha/beta.
(2) They are assumed to be the alpha/beta from the forward call and do the right thing for the gradient.

If you choose (2) and I remember correctly:
(a) both backward-filter and backward-data should use beta=0, alpha=forward-alpha.
(b) backward-bias should use alpha=1, beta=0 regardless of the forward params.
(c) Finally the gradient for the residual (the cudnn z argument, or the original y argument that got added to) should multiply dy by beta.

Right now NNlib seems inconsistent in this regard: backward-filter adds when beta is non-zero, backward-data doesn't.

I also noticed that conv_bias_act! does not use alpha/beta consistently in nnlib.

In any case, these are NNlib issues that can be addressed later. Can we revert #672 and merge this PR?

maleadt · 2021-02-01T08:45:21Z

I'll revert/rebase/push, we can merge this once it passes a final round of CI.

This reverts commit f7ea311.

DhairyaLGandhi · 2021-02-01T13:01:13Z

Yeah, I'm thinking that conv_bias_ act needs to be compatible with conv.

I think alpha and beta should be consistent with the cpu case. Do you think using the cudnn default as in 2 should be triggered only when these conditions are met?

maleadt · 2021-02-02T07:27:56Z

@denizyuret Why did you push a merge commit? This was ready to go.

denizyuret · 2021-02-02T08:19:48Z

@maleadt sorry about that, it was automatic reflex (open laptop, git pull, merge if new things). How can I fix it?

maleadt · 2021-02-02T08:20:47Z

I'll drop the merge commit.

New high level interface for cuDNN

denizyuret mentioned this pull request Nov 6, 2020

Incomplete CUDNN wrappers #343

Closed

maleadt marked this pull request as draft November 6, 2020 13:09

maleadt added cuda libraries Stuff about CUDA library wrappers. enhancement New feature or request labels Nov 6, 2020

maleadt reviewed Nov 6, 2020

View reviewed changes

lib/cudnn/util.jl Outdated Show resolved Hide resolved

denizyuret self-assigned this Nov 7, 2020

simeonschaub reviewed Nov 18, 2020

View reviewed changes

denizyuret mentioned this pull request Dec 7, 2020

softmax(x) and logsoftmax(x) update their arguments #592

Closed

This was linked to issues Dec 11, 2020

Incomplete CUDNN wrappers #343

Closed

CuPtr should be Ptr in cudnnGetDropoutDescriptor #397

Closed

softmax(x) and logsoftmax(x) update their arguments #592

Closed

denizyuret removed a link to an issue Dec 11, 2020

CuPtr should be Ptr in cudnnGetDropoutDescriptor #397

Closed

This was linked to issues Dec 11, 2020

dims support for softmax #226

Closed

softmax has problem with dim parameter #599

Closed

This was referenced Dec 14, 2020

softmax has problem with dim parameter #599

Closed

added a softmax helper. #600

Closed

maleadt force-pushed the dy/cudnn branch from ff71ad2 to 9234ffc Compare December 17, 2020 14:46

maleadt marked this pull request as ready for review December 17, 2020 14:51

maleadt changed the title ~~dy/cudnn: New high level interface for cuDNN [WIP]~~ New high level interface for cuDNN Dec 17, 2020

CarloLucibello reviewed Dec 21, 2020

View reviewed changes

lib/cudnn/README.md Outdated Show resolved Hide resolved

maleadt force-pushed the dy/cudnn branch from 95a7af8 to a3bf7ee Compare January 20, 2021 07:28

DhairyaLGandhi reviewed Jan 21, 2021

View reviewed changes

maleadt and others added 3 commits February 1, 2021 10:09

Revert "Add beta keyword to conv (JuliaGPU#672)"

962dc02

This reverts commit f7ea311.

New high level interface for cuDNN

588fdb8

dy/cudnn beta kwarg for conv

0b57f52

maleadt force-pushed the dy/cudnn branch from 15565f0 to 0b57f52 Compare February 1, 2021 09:09

maleadt force-pushed the dy/cudnn branch 2 times, most recently from 15565f0 to 0b57f52 Compare February 2, 2021 08:22

maleadt merged commit f1caf8d into JuliaGPU:master Feb 2, 2021

maleadt added a commit that referenced this pull request Feb 10, 2021

Merge pull request #523 from denizyuret/dy/cudnn

0facc55

New high level interface for cuDNN

DhairyaLGandhi mentioned this pull request Feb 22, 2021

PyTorch feature parity FluxML/Flux.jl#1431

Open

92 tasks

maleadt added a commit that referenced this pull request Mar 16, 2021

Merge pull request #523 from denizyuret/dy/cudnn

b86bc4c

New high level interface for cuDNN

	math_mode = CUDA.math_mode()
	reduced_precision = CUDA.math_precision()

	if sig === (Float16, Float16)
	# NOTE: Float16=Float16*Float16 can also happen in 32-bit compute
	return math_mode==CUDA.PEDANTIC_MATH ? CUBLAS_COMPUTE_16F_PEDANTIC : CUBLAS_COMPUTE_16F
	end

New high level interface for cuDNN #523

New high level interface for cuDNN #523

Uh oh!

Conversation

denizyuret commented Nov 6, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

denizyuret commented Nov 7, 2020

Uh oh!

maleadt commented Nov 10, 2020

Uh oh!

denizyuret commented Nov 10, 2020

Uh oh!

maleadt commented Nov 10, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

denizyuret commented Nov 14, 2020

Uh oh!

simeonschaub Nov 18, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

denizyuret Nov 18, 2020

Choose a reason for hiding this comment

Uh oh!

maleadt Nov 18, 2020

Choose a reason for hiding this comment

Uh oh!

DhairyaLGandhi Dec 21, 2020

Choose a reason for hiding this comment

Uh oh!

codecov bot commented Nov 20, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

maleadt commented Dec 9, 2020

Uh oh!

maleadt commented Dec 9, 2020

Uh oh!

denizyuret commented Dec 16, 2020

Uh oh!

Uh oh!

DhairyaLGandhi commented Dec 23, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

DhairyaLGandhi commented Jan 16, 2021

Uh oh!

DhairyaLGandhi commented Jan 20, 2021

Uh oh!

denizyuret commented Jan 20, 2021

Uh oh!

Uh oh!

DhairyaLGandhi Jan 21, 2021

Choose a reason for hiding this comment

Uh oh!

maleadt Jan 21, 2021

Choose a reason for hiding this comment

Uh oh!

DhairyaLGandhi commented Jan 21, 2021

Uh oh!

maleadt commented Jan 27, 2021

Uh oh!

denizyuret commented Jan 27, 2021

Uh oh!

maleadt commented Jan 28, 2021

Uh oh!

DhairyaLGandhi commented Jan 28, 2021

Uh oh!

denizyuret commented Jan 30, 2021

Uh oh!

DhairyaLGandhi commented Jan 30, 2021

Uh oh!

denizyuret commented Jan 30, 2021

Uh oh!

DhairyaLGandhi commented Jan 30, 2021

Uh oh!

denizyuret commented Jan 31, 2021

Uh oh!

maleadt commented Feb 1, 2021

denizyuret commented Nov 6, 2020 •

edited

Loading

maleadt commented Nov 10, 2020 •

edited

Loading

simeonschaub Nov 18, 2020 •

edited

Loading

codecov bot commented Nov 20, 2020 •

edited

Loading

DhairyaLGandhi commented Dec 23, 2020 •

edited

Loading