Speedup bernoulli_scalar_cuda_kernel with grid-stride loop#20626
Closed
syed-ahmed wants to merge 20 commits intogh/syed-ahmed/7/basefrom
Closed
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop#20626syed-ahmed wants to merge 20 commits intogh/syed-ahmed/7/basefrom
syed-ahmed wants to merge 20 commits intogh/syed-ahmed/7/basefrom
Conversation
This was referenced May 17, 2019
Member
|
After these PRs, will the CUDA RNGs be thread and stream-safe in PyTorch? Does the |
ssnl
reviewed
May 17, 2019
| if (std::is_same<scalar_t, double>::value) { | ||
| distribution_nullary_kernel<scalar_t, accscalar_t, curand4_engine_calls/2>(iter, | ||
| gen, | ||
| [] __device__ (curandStatePhilox4_32_10_t* state) { return curand_uniform2_double(state); }, |
Collaborator
There was a problem hiding this comment.
so @ngimel told me that curand's uniform with double is actually a lie because there is not enough precision. I don't know if things have changed or not.
Member
There was a problem hiding this comment.
How many bits of randomness does it provide? 53 bits is standard so that it generates all the rationals x / 2^53 for x in [0,1, ..., 2^53-1].
That doesn't generate all valid doubles in [0, 1), but basically nobody does that outside of toy programs.
Collaborator
There was a problem hiding this comment.
curand_uniform_double is a lie (uses 32 bits), curand_uniform2_double is not (uses 53 bits per value).
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Collaborator
Author
That is the intention. These PRs achieve thread and stream-safety by replacing curandStateMTGP with curandStatePhilox: #19508 (comment)
I checked locally and got the following ouput: So may be it fixed the issue? I didn't triage and find the source of that bug. May be the tensor iterator magic in the normal PR has fixed this. |
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
ezyang
approved these changes
May 29, 2019
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
Speedup bernoulli_scalar_cuda_kernel with grid-stride loop gh-metadata: pytorch pytorch 20626 gh/syed-ahmed/7/head
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
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.
Stack from ghstack:
Differential Revision: D15454046
Effective Bandwidth Benchmark
Float Type
Before:
After:
Double Type
Before:
After: