forked from cms-sw/cmssw
-
Notifications
You must be signed in to change notification settings - Fork 5
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Replace CUDA API wrapper memory operations with native CUDA calls #395
Merged
fwyzard
merged 15 commits into
cms-patatrack:CMSSW_11_0_X_Patatrack
from
waredjeb:replace_cuda_memory
Oct 29, 2019
Merged
Changes from 14 commits
Commits
Show all changes
15 commits
Select commit
Hold shift + click to select a range
2ee863c
Solve conflicts with #389
waredjeb ede0cdd
Solve conflicts with #389
waredjeb a731751
Delete spurious file
fwyzard c00e7bd
Delete spurious file
fwyzard 3b7e845
Delete spurious file
fwyzard 61d12f5
Delete spurious file
fwyzard b17e7f9
Whitespaces
fwyzard 6e103c8
Merge branch 'CMSSW_11_0_X_Patatrack' of https://github.com/cms-patat…
waredjeb 51d5cc3
Merge branch 'replace_cuda_memory' of https://github.com/waredjeb/cms…
waredjeb 01bb995
Wrap cudaMem calls in call to cudaCheck
waredjeb c7b7f03
Fix errors, missing include of launch.h
waredjeb e007a5e
Apply code-format
waredjeb 9a1ca24
Reorders memory copy operations
waredjeb 463b495
Reoders memory copy in Device_to_Host section
waredjeb 49d83d6
Fix direction of the copies from device to host
fwyzard File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
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
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -1,8 +1,9 @@ | ||
#include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" | ||
|
||
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" | ||
#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" | ||
|
||
BeamSpotCUDA::BeamSpotCUDA(Data const* data_h, cudaStream_t stream) { | ||
data_d_ = cudautils::make_device_unique<Data>(stream); | ||
cuda::memory::async::copy(data_d_.get(), data_h, sizeof(Data), stream); | ||
cudaCheck(cudaMemcpyAsync(data_d_.get(), data_h, sizeof(Data), cudaMemcpyHostToDevice, stream)); | ||
} |
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
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
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
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
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
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
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
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
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
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
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
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
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
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
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -1,6 +1,7 @@ | ||
#ifndef HeterogeneousCore_CUDAUtilities_copyAsync_h | ||
#define HeterogeneousCore_CUDAUtilities_copyAsync_h | ||
|
||
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" | ||
#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" | ||
#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" | ||
|
||
|
@@ -17,15 +18,15 @@ namespace cudautils { | |
// Shouldn't compile for array types because of sizeof(T), but | ||
// let's add an assert with a more helpful message | ||
static_assert(std::is_array<T>::value == false, "For array types, use the other overload with the size parameter"); | ||
cuda::memory::async::copy(dst.get(), src.get(), sizeof(T), stream); | ||
cudaCheck(cudaMemcpyAsync(dst.get(), src.get(), sizeof(T), cudaMemcpyHostToDevice, stream)); | ||
} | ||
|
||
template <typename T> | ||
inline void copyAsync(cudautils::host::unique_ptr<T>& dst, | ||
const cudautils::device::unique_ptr<T>& src, | ||
cudaStream_t stream) { | ||
static_assert(std::is_array<T>::value == false, "For array types, use the other overload with the size parameter"); | ||
cuda::memory::async::copy(dst.get(), src.get(), sizeof(T), stream); | ||
cudaCheck(cudaMemcpyAsync(dst.get(), src.get(), sizeof(T), cudaMemcpyHostToDevice, stream)); | ||
} | ||
|
||
// Multiple elements | ||
|
@@ -34,15 +35,15 @@ namespace cudautils { | |
const cudautils::host::unique_ptr<T[]>& src, | ||
size_t nelements, | ||
cudaStream_t stream) { | ||
cuda::memory::async::copy(dst.get(), src.get(), nelements * sizeof(T), stream); | ||
cudaCheck(cudaMemcpyAsync(dst.get(), src.get(), nelements * sizeof(T), cudaMemcpyHostToDevice, stream)); | ||
} | ||
|
||
template <typename T> | ||
inline void copyAsync(cudautils::host::unique_ptr<T[]>& dst, | ||
const cudautils::device::unique_ptr<T[]>& src, | ||
size_t nelements, | ||
cudaStream_t stream) { | ||
cuda::memory::async::copy(dst.get(), src.get(), nelements * sizeof(T), stream); | ||
cudaCheck(cudaMemcpyAsync(dst.get(), src.get(), nelements * sizeof(T), cudaMemcpyHostToDevice, stream)); | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. ditto |
||
} | ||
} // namespace cudautils | ||
|
||
|
Oops, something went wrong.
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.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
this is device2host
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
And "Calling
cudaMemcpyAsync()
withdst
andsrc
pointers that do not match the direction of the copy results in an undefined behavior." (*), so specifying the direction explicitly is actually harmful?(*) https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1g85073372f776b4c4d5f89f7124b7bf79
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Indeed. I think we agreed to remove all explicit directions.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I thought it was supposed to crash...