-
Notifications
You must be signed in to change notification settings - Fork 188
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
[CUDAX] Add experimental owning abstraction for cudaStream_t #2093
Conversation
🟨 CI finished in 1h 44m: Pass: 99%/417 | Total: 2d 02h | Avg: 7m 16s | Max: 1h 19m | Hits: 97%/524552
|
Project | |
---|---|
CCCL Infrastructure | |
+/- | libcu++ |
CUB | |
Thrust | |
+/- | CUDA Experimental |
pycuda |
Modifications in project or dependencies?
Project | |
---|---|
CCCL Infrastructure | |
+/- | libcu++ |
+/- | CUB |
+/- | Thrust |
+/- | CUDA Experimental |
+/- | pycuda |
🏃 Runner counts (total jobs: 417)
# | Runner |
---|---|
305 | linux-amd64-cpu16 |
61 | linux-amd64-gpu-v100-latest-1 |
28 | linux-arm64-cpu16 |
23 | windows-amd64-cpu16 |
//! @brief Constructs a stream on a specified device and with specified priority | ||
//! | ||
//! @throws cuda_error if stream creation fails | ||
explicit stream(device __dev, int __priority) |
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 would love if in new APIs we could try to be const correct:
explicit stream(device __dev, int __priority) | |
explicit stream(const device __dev, const int __priority) |
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.
If these are passed by value, is there any value in having them const?
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.
From the perspective of the user these are identical. From our perspective it would guard against accidentally changing them. I'm generally not a fan of top level qualifiers on function arguments, and they are very much an implementation detail, so we can add or remove them in any API at any time.
//! @brief Construct a new `stream` object into the moved-from state. | ||
//! | ||
//! @post `stream()` returns an invalid stream handle | ||
// Can't be constexpr because invalid_stream isn't | ||
explicit stream(uninit_t) noexcept | ||
: stream_ref(detail::invalid_stream) | ||
{} |
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.
Should this be a public constructor?
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 aligns with the event constructor from uninit and works as an opt-in to create a stream that will be assigned into later
_CCCL_NODISCARD static stream from_native_handle(::cudaStream_t __handle) | ||
{ | ||
return stream(__handle); | ||
} | ||
|
||
// Disallow construction from an `int`, e.g., `0`. | ||
static stream from_native_handle(int) = delete; | ||
|
||
// Disallow construction from `nullptr`. | ||
static stream from_native_handle(_CUDA_VSTD::nullptr_t) = delete; |
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.
Question: why shouldn't those be ctors? What problem are factory functions solving here that ctors cannot?
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 one aligns with event factory function, we can discuss it as a broader design question for cudax. These are taking the ownership of the stream, so I like the explicitness of the function.
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.
@ericniebler points out that these functions take ownership over the passed in stream, so he wants them to stand out in the code
* | ||
* \return value representing the priority of the wrapped stream. | ||
*/ | ||
_CCCL_NODISCARD int priority() const |
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.
@ericniebler Do we want to return a strong type?
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.
Question: I've always thought an owning stream abstraction should just be a unique_ptr
with a custom deleter that overrides the pointer type like:
struct stream_deleter{
using pointer = cudaStream_t;
void operator()(cudaStream_t s){ cudaStreamDestroy(s); }
};
struct owning_stream{
// other stuff
private:
std::unique_ptr<cudaStream_t, stream_deleter> __s;
};
Does that not work?
{ | ||
// TODO consider an optimization to not create an event every time and instead have one persistent event or one per | ||
// stream | ||
assert(__stream.get() != nullptr); |
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.
isn't __stream
a cudaStream_t
here?
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.
Should be check for invalid_stream, corrected
🟩 CI finished in 8h 04m: Pass: 100%/417 | Total: 1d 20h | Avg: 6m 24s | Max: 37m 49s | Hits: 98%/525419
|
Project | |
---|---|
CCCL Infrastructure | |
+/- | libcu++ |
CUB | |
Thrust | |
+/- | CUDA Experimental |
pycuda |
Modifications in project or dependencies?
Project | |
---|---|
CCCL Infrastructure | |
+/- | libcu++ |
+/- | CUB |
+/- | Thrust |
+/- | CUDA Experimental |
+/- | pycuda |
🏃 Runner counts (total jobs: 417)
# | Runner |
---|---|
305 | linux-amd64-cpu16 |
61 | linux-amd64-gpu-v100-latest-1 |
28 | linux-arm64-cpu16 |
23 | windows-amd64-cpu16 |
…2093) * construct with a stream_ref and record the event on construction --------- Co-authored-by: Eric Niebler <[email protected]>
…2093) * construct with a stream_ref and record the event on construction --------- Co-authored-by: Eric Niebler <[email protected]>
This pull request adds an owning type
cudax::stream
for cudaStream_t.Some functions in
cudax::stream
should go tocuda::stream_ref
, likerecord
andwait
, but libcu++ can't depend on cudax. We could consider having two versions ofstream_ref
, once incuda::
and second one incudax::