emulate hip/cuda-Memcpy3D with a kernel #1014
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
psychocoderHPC
force-pushed
the
topic-emulate3DMemcopy
branch
from
5e529c7
to
08ba662
Compare
psychocoderHPC
commented
May 29, 2020
| /// It is required to start `height * depth` HIP/CUDA blocks. | ||
| /// The kernel loops over the memory rows. | ||
| template<typename T> | ||
| __global__ void hipMemcpy3DEmulatedKernelD2D( |
There was a problem hiding this comment.
note: a native kernel is used to avoid cyclic dependencies within alpaka.
BenjaminW3
reviewed
May 29, 2020
cmake/alpakaCommon.cmake
Outdated
| @@ -18,7 +19,10 @@ set(ALPAKA_ACC_CPU_B_SEQ_T_FIBERS_ENABLE_DEFAULT ON) | |||
| set(ALPAKA_ACC_CPU_B_TBB_T_SEQ_ENABLE_DEFAULT ON) | |||
| set(ALPAKA_ACC_CPU_B_OMP2_T_SEQ_ENABLE_DEFAULT ON) | |||
| set(ALPAKA_ACC_CPU_B_SEQ_T_OMP2_ENABLE_DEFAULT ON) | |||
| set(ALPAKA_ACC_CPU_BT_OMP4_ENABLE_DEFAULT ON) | |||
| set(ALPAKA_ACC_ANY_BT_OMP5_ENABLE_DEFAULT ON) | |||
| set(ALPAKA_ACC_ANY_BT_OACC_ENABLE_DEFAULT OFF) | |||
There was a problem hiding this comment.
I think this should not be part of this PR.
psychocoderHPC
force-pushed
the
topic-emulate3DMemcopy
branch
2 times, most recently
from
e5889d1
to
d7a4c63
Compare
- add kernel to emulate hip/cuda-Memcpy3D - add CMake option to enable/disable emulated memory copy (by default only for HIP enabled)
psychocoderHPC
force-pushed
the
topic-emulate3DMemcopy
branch
from
d7a4c63
to
e503b7f
Compare
sbastrakov
approved these changes
Jun 2, 2020
|
@BenjaminW3 Do you know what we can do if a CI test fails. Actions do not allow to restart single tests. |
|
Github is working on it. I do not know a solution despite rebuilding all or merging it nevertheless. |
|
Feel free to merge it even if the CI is not passing. The CI passed before but I fixed a indention issue in CMake and today not all tests passing. |
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.
only for HIP enabled)
ALPAKA_EMU_MEMCPY3Dfor one HIP and one CUDA CI testThis optimization based on my HIP issue and will increase the memory copy performance for device to device copies on the same device.
I enabled the emulated copy only for HIP, for CUDA it can be optional enabled but is not showing any improvement. I assume the CUDA driver is already using a kernel instead of looping over the rows and call 1D mem-copies.