Optimize the DCT GPU kernel. #2471
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
Signed-off-by: Rafal <Banas.Rafal97@gmail.com>
|
!build |
|
CI MESSAGE: [1805832]: BUILD STARTED |
|
What is the speed now, and what used to be before that optimization? |
|
CI MESSAGE: [1805832]: BUILD FAILED |
mzient
reviewed
Nov 17, 2020
dali/kernels/signal/dct/dct_gpu.cu
Outdated
| __global__ void ApplyDctInner(const typename Dct1DGpu<OutputType, InputType>::SampleDesc *samples, | ||
| const BlockSetupInner::BlockDesc *blocks, | ||
| const float *lifter_coeffs) { | ||
| extern __shared__ char *shm[]; |
There was a problem hiding this comment.
Suggested change
| extern __shared__ char *shm[]; | |
| extern __shared__ char shm[]; |
Signed-off-by: Rafal <Banas.Rafal97@gmail.com>
mzient
reviewed
Nov 17, 2020
dali/kernels/signal/dct/dct_gpu.h
Outdated
| struct BlockDesc { | ||
| int64_t sample_idx; | ||
| int64_t frame_start; | ||
| int64_t frames_num; |
There was a problem hiding this comment.
Suggested change
| int64_t frames_num; | |
| int64_t num_frames; |
or
Suggested change
| int64_t frames_num; | |
| int64_t frame_count; |
banasraf
force-pushed
the
dct-kernel-optimization
branch
from
be2c955
to
1e59882
Compare
Signed-off-by: Rafal <Banas.Rafal97@gmail.com>
banasraf
force-pushed
the
dct-kernel-optimization
branch
from
1e59882
to
f2ac7a0
Compare
|
!build |
For the planar layout it's ~550 GFLOPS -> ~630 GFLOPS and for interleaved it's ~30 GFLOPS -> ~315 GFLOPS |
|
CI MESSAGE: [1806677]: BUILD STARTED |
JanuszL
approved these changes
Nov 17, 2020
Signed-off-by: Rafal <Banas.Rafal97@gmail.com>
JanuszL
reviewed
Nov 17, 2020
| @@ -40,7 +40,7 @@ def define_graph(self): | |||
| test_data_root = get_dali_extra_path() | |||
| good_path = 'db/single' | |||
| missnamed_path = 'db/single/missnamed' | |||
| test_good_path = {'jpeg', 'mixed', 'png', 'tiff', 'pnm', 'bmp', 'jpeg2k'} | |||
| test_good_path = {'jpeg2k'} | |||
There was a problem hiding this comment.
eeeh, good catch. I've committed too much
|
CI MESSAGE: [1806677]: BUILD FAILED |
Signed-off-by: Rafal <Banas.Rafal97@gmail.com>
JanuszL
approved these changes
Nov 18, 2020
Signed-off-by: Rafal <Banas.Rafal97@gmail.com>
|
!build |
|
CI MESSAGE: [1809572]: BUILD STARTED |
|
CI MESSAGE: [1809572]: BUILD PASSED |
mzient
approved these changes
Nov 20, 2020
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
Why we need this PR?
What happened in this PR?
The case with the transform done over the inner axis is handled with a separate CUDA kernel. The existing kernel was optimized by employing shared memory.
GPU DCT kernel.
A new CUDA kernel and the changes in the old one.
Existing tests still apply. I added a performance test.
N/A
JIRA TASK: DALI-1690