This repository has been archived by the owner on Nov 17, 2023. It is now read-only.
Performance improvement in Normalize GPU Kernel #14139
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
sandeep-krishnamurthy
added
Operator
pr-awaiting-review
zhreshold
reviewed
Feb 13, 2019
zhreshold
reviewed
Feb 13, 2019
zhreshold
suggested changes
Feb 13, 2019
zhreshold
suggested changes
Feb 13, 2019
src/operator/image/image_random.cu
Outdated
| ToTensorCudaKernel<gpu, DType> | ||
| <<<blocks, dim3(32, 32), 0, stream>>>(input, output, | ||
| <<<blocks, dim3(H, cuda::kMaxThreadsPerBlock / H), 0, stream>>>(input, output, |
There was a problem hiding this comment.
please fix ToTensor similarly in a separate PR.
There was a problem hiding this comment.
Sure already work in progress. Should this PR wait till then or undo dim3(H, cuda::kMaxThreadsPerBlock / H) back to dim3(32, 32). Let me know next steps for this PR. Thanks!
There was a problem hiding this comment.
back to (32, 32), we can address it later.
|
Basically lgtm, please make minor revision and once CI passes ,we can merge |
sandeep-krishnamurthy
force-pushed
the
normalize_gpu_perf
branch
from
9ca1aec
to
26a0532
Compare
Done. Will create ToTensor refactoring PR in a day or two. Thanks again for your time and fast turn around time for all PR reviews. |
sandeep-krishnamurthy
force-pushed
the
normalize_gpu_perf
branch
from
cf05dd7
to
f595dff
Compare
zhreshold
approved these changes
Feb 14, 2019
stephenrawls
pushed a commit
to stephenrawls/incubator-mxnet
that referenced
this pull request
Feb 16, 2019
* New CPU kernel for normalize * New GPU kernel for Normalize * Add launch bounds and increase threads to 32*32 * do not hardcode number of threads * Try fix windows build failure * make channels as int to fix windows build issues with omp * Simplify cuda kernels with 1 D thread block * Minor refactoring * Revert thread dim for ToTensor operator
jessr92
pushed a commit
to jessr92/incubator-mxnet
that referenced
this pull request
Feb 19, 2019
* New CPU kernel for normalize * New GPU kernel for Normalize * Add launch bounds and increase threads to 32*32 * do not hardcode number of threads * Try fix windows build failure * make channels as int to fix windows build issues with omp * Simplify cuda kernels with 1 D thread block * Minor refactoring * Revert thread dim for ToTensor operator
drivanov
pushed a commit
to drivanov/incubator-mxnet
that referenced
this pull request
Mar 4, 2019
* New CPU kernel for normalize * New GPU kernel for Normalize * Add launch bounds and increase threads to 32*32 * do not hardcode number of threads * Try fix windows build failure * make channels as int to fix windows build issues with omp * Simplify cuda kernels with 1 D thread block * Minor refactoring * Revert thread dim for ToTensor operator
vdantu
pushed a commit
to vdantu/incubator-mxnet
that referenced
this pull request
Mar 31, 2019
* New CPU kernel for normalize * New GPU kernel for Normalize * Add launch bounds and increase threads to 32*32 * do not hardcode number of threads * Try fix windows build failure * make channels as int to fix windows build issues with omp * Simplify cuda kernels with 1 D thread block * Minor refactoring * Revert thread dim for ToTensor operator
haohuanw
pushed a commit
to haohuanw/incubator-mxnet
that referenced
this pull request
Jun 23, 2019
* New CPU kernel for normalize * New GPU kernel for Normalize * Add launch bounds and increase threads to 32*32 * do not hardcode number of threads * Try fix windows build failure * make channels as int to fix windows build issues with omp * Simplify cuda kernels with 1 D thread block * Minor refactoring * Revert thread dim for ToTensor operator
Sign up for free
to subscribe to this conversation on GitHub.
Already have an account?
Sign in.
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.
Description
Similar to perf improvements of ToTensor GPU kernel in PR - #14099
In this PR, I wrote a separate CUDA kernel for GPU and moved out of Kernel launch/map.
Benchmarks below.
Benchmarks
Ran 500 Normalize operation on (3, 512, 512) sample input.
GPU
Before: ('Average time per Normalize 3,512,512 - ', 38.19581985473633)
After: ('Average time per Normalize 3,512,512 - ', 0.5398507118225098)
CPU
Before: ('Average time per Normalize 3,512,512 - ', 1.8209707736968994)
After: ('Average time per Normalize 3,512,512 - ', 1.2644755840301514)
('Total time for CPU ToTensor - ', 1264.4755840301514)
('Average time per Normalize 3,512,512 - ', 1.2644755840301514)
Checklist
Essentials
Please feel free to remove inapplicable items for your PR.
Changes
@stu1130 @zhreshold