Coordinate Flip GPU operator #1895
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
JanuszL
approved these changes
Apr 23, 2020
klecki
reviewed
Apr 27, 2020
dali/operators/coord/coord_flip.cu
Outdated
|
|
||
| void CoordFlipGPU::RunImpl(workspace_t<GPUBackend> &ws) { | ||
| const auto &input = ws.InputRef<GPUBackend>(0); | ||
| DALI_ENFORCE(input.type().id() == DALI_FLOAT, "Input is expected to be float"); |
There was a problem hiding this comment.
Please move to SetupImpl (same as CPU).
klecki
reviewed
Apr 27, 2020
| cudaMemcpyAsync(sample_descs_gpu_, sample_descs_.data(), sz, cudaMemcpyHostToDevice, stream)); | ||
|
|
||
| dim3 block(32, 32); | ||
| auto blocks_per_sample = std::max(32, 1024 / batch_size_); |
There was a problem hiding this comment.
What's the reason for this gridDim.x?
There was a problem hiding this comment.
it's just a good enough number. The second term is meant to reduce the number of total blocks in case of a batch with many samples
klecki
reviewed
Apr 27, 2020
dali/operators/coord/coord_flip.cu
Outdated
| CUDA_CALL( | ||
| cudaMemcpyAsync(sample_descs_gpu_, sample_descs_.data(), sz, cudaMemcpyHostToDevice, stream)); | ||
|
|
||
| dim3 block(32, 32); |
There was a problem hiding this comment.
Is there any benefit in using 2D blocks instead of 1D of the same volume?
That would probably simplify the addressing a bit.
jantonguirao
force-pushed
the
coord_flip_gpu
branch
4 times, most recently
from
819bf25
to
57cee51
Compare
mzient
reviewed
Apr 28, 2020
dali/operators/coord/coord_flip.cu
Outdated
| int64_t tid = threadIdx.y * blockDim.x + threadIdx.x; | ||
| for (int64_t idx = offset + tid; idx < sample.size; idx += grid_size) { | ||
| int d = idx % ndim; | ||
| bool flip = static_cast<bool>(sample.flip_dim_mask & (1 << d)); |
There was a problem hiding this comment.
I think implicit conversion would do just fine, like it would in if statement.
mzient
reviewed
Apr 28, 2020
dali/operators/coord/coord_flip.cu
Outdated
| for (int64_t idx = offset + tid; idx < sample.size; idx += grid_size) { | ||
| int d = idx % ndim; | ||
| bool flip = static_cast<bool>(sample.flip_dim_mask & (1 << d)); | ||
| sample.out[idx] = flip ? T(1) - sample.in[idx] : sample.in[idx]; |
jantonguirao
force-pushed
the
coord_flip_gpu
branch
from
57cee51
to
41d28d1
Compare
Signed-off-by: Joaquin Anton <janton@nvidia.com>
Signed-off-by: Joaquin Anton <janton@nvidia.com>
Signed-off-by: Joaquin Anton <janton@nvidia.com>
jantonguirao
force-pushed
the
coord_flip_gpu
branch
2 times, most recently
from
09c81ee
to
6c6834a
Compare
jantonguirao
force-pushed
the
coord_flip_gpu
branch
from
6c6834a
to
561c55c
Compare
Signed-off-by: Joaquin Anton <janton@nvidia.com>
mzient
approved these changes
Apr 29, 2020
|
!build |
|
CI MESSAGE: [1291137]: BUILD STARTED |
|
CI MESSAGE: [1291137]: BUILD PASSED |
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.
Signed-off-by: Joaquin Anton janton@nvidia.com
Why we need this PR?
What happened in this PR?
Fill relevant points, put NA otherwise. Replace anything inside []
Added Coordinate Flip GPU operator
New operator
The operator implementation
Python operator tests added
NA
JIRA TASK: [DALI-1392]