[libcu++] Implement internal resizable buffer - #10221
Conversation
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughWalkthroughChangesAdds internal Resizable buffer
Assessment against linked issues
Suggested reviewers: Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
🧹 Nitpick comments (1)
libcudacxx/include/cuda/__container/buffer.h (1)
829-846: 🎯 Functional Correctness | 🔵 Trivial | 💤 Low valuesuggestion: Add
_CCCL_HOST_APIto both__copy_cross_buffersoverloads; these helpers are the only functions in this section missing the visibility macro.Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: b3e156b4-7280-4f44-b335-ff2b626a3807
📒 Files selected for processing (4)
libcudacxx/include/cuda/__container/buffer.hlibcudacxx/include/cuda/__container/resizable_buffer.hlibcudacxx/include/cuda/__container/uninitialized_async_buffer.hlibcudacxx/test/libcudacxx/cuda/containers/buffer/resizable_buffer.cu
This comment has been minimized.
This comment has been minimized.
| _CCCL_HOST_API void resize(size_type __new_size, ::cuda::no_init_t) | ||
| { | ||
| if (__new_size > __capacity_) | ||
| { | ||
| _CCCL_THROW(::std::invalid_argument, | ||
| "cuda::__resizable_buffer::resize requires an explicit stream to grow beyond capacity"); | ||
| } | ||
| __base_t::__set_size_unsynchronized(__new_size); | ||
| } |
There was a problem hiding this comment.
We can delete this overload and just require the stream always.
There was a problem hiding this comment.
I would like to keep at least one of the non-stream overloads, I think as long as reallocation is not triggered and the user knows the reallocation won't happen, its good to have an option to do it in non-stream order
There was a problem hiding this comment.
The is should be called resize_unsynchronized() or something like that. I also don't really see the utility of it over resize_discard(). Both throw the old end elements into the oubliette, but this one will throw if it tries to grow. I guess it is useful if you want to ensure you aren't reallocating? But then it should have a different name.
There was a problem hiding this comment.
I think the main difference is semantics. The one that takes a stream is stream ordered and if there is any work in flight that uses the buffer you have to use a stream that captures that work in the resize call, otherwise you risk reallocation racing with the work. I think in most cases you should be able to provide a stream like that, but I thought we could have a variant that is non-stream ordered and only updates the CPU view of the memory. I don't know how useful it will be but for internal version I wanted to try it
There was a problem hiding this comment.
🧹 Nitpick comments (1)
libcudacxx/test/libcudacxx/cuda/containers/buffer/resizable_buffer.cu (1)
28-49: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winsuggestion:
check_prefix's second parameter (the expected-values array) is never used —check_prefix_kernelonly compares against the globaldevice_data, not the array passed in. All current call sites happen to pass a prefix ofdevice_data, so this coincidentally works, but the function doesn't actually verify the values callers pass; a future call with different expected data would silently pass incorrectly.Proposed fix: actually compare against the passed values
-template <typename _Iter> -__global__ void check_prefix_kernel(_Iter ptr, cuda::std::size_t size) -{ - for (cuda::std::size_t i = 0; i != size; ++i) - { - if (ptr[i] != device_data[i]) - { - __trap(); - } - } -} - -template <cuda::std::size_t _Size> -void check_prefix(const buffer_t& buf, const cuda::std::array<int, _Size>&) +template <typename _Iter, typename _Expected> +__global__ void check_prefix_kernel(_Iter ptr, _Expected expected, cuda::std::size_t size) +{ + for (cuda::std::size_t i = 0; i != size; ++i) + { + if (ptr[i] != expected[i]) + { + __trap(); + } + } +} + +template <cuda::std::size_t _Size> +void check_prefix(const buffer_t& buf, const cuda::std::array<int, _Size>& expected) { REQUIRE(buf.size() >= _Size); cuda::__ensure_current_context guard{buf.stream()}; - check_prefix_kernel<<<1, 1, 0, buf.stream().get()>>>(buf.begin(), _Size); + check_prefix_kernel<<<1, 1, 0, buf.stream().get()>>>(buf.begin(), expected, _Size); REQUIRE(::cudaGetLastError() == ::cudaSuccess); buf.stream().sync(); }
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: bc5d2c26-299c-4a09-aef1-30fc39b5f427
📒 Files selected for processing (4)
libcudacxx/include/cuda/__container/buffer.hlibcudacxx/include/cuda/__container/resizable_buffer.hlibcudacxx/include/cuda/__container/uninitialized_async_buffer.hlibcudacxx/test/libcudacxx/cuda/containers/buffer/resizable_buffer.cu
🚧 Files skipped from review as they are similar to previous changes (1)
- libcudacxx/include/cuda/__container/uninitialized_async_buffer.h
| //! @brief Helper to check container is compatible with this buffer | ||
| template <class _Range> | ||
| static constexpr bool __compatible_range = (::cuda::std::ranges::__container_compatible_range<_Range, _Tp>); | ||
| [[nodiscard]] static constexpr bool __compatible_range() noexcept |
There was a problem hiding this comment.
Why make these functions?
| const auto __old_capacity = __capacity_; | ||
|
|
||
| __capacity_ = 0; |
There was a problem hiding this comment.
| const auto __old_capacity = __capacity_; | |
| __capacity_ = 0; | |
| const auto __old_capacity = ::cuda::std::exchange(__capacity_, 0); |
| { | ||
| const auto __old_size = this->size(); | ||
| const auto __old_capacity = __capacity_; | ||
| _CCCL_ASSERT(__old_size <= __new_capacity, "cuda::__resizable_buffer cannot reallocate below size"); |
There was a problem hiding this comment.
Nit: the error message is a little confusing here
| _CCCL_HOST_API void resize(size_type __new_size, ::cuda::no_init_t) | ||
| { | ||
| if (__new_size > __capacity_) | ||
| { | ||
| _CCCL_THROW(::std::invalid_argument, | ||
| "cuda::__resizable_buffer::resize requires an explicit stream to grow beyond capacity"); | ||
| } | ||
| __base_t::__set_size_unsynchronized(__new_size); | ||
| } |
There was a problem hiding this comment.
The is should be called resize_unsynchronized() or something like that. I also don't really see the utility of it over resize_discard(). Both throw the old end elements into the oubliette, but this one will throw if it tries to grow. I guess it is useful if you want to ensure you aren't reallocating? But then it should have a different name.
This comment has been minimized.
This comment has been minimized.
😬 CI Workflow Results🟥 Finished in 1h 46m: Pass: 95%/120 | Total: 1d 11h | Max: 1h 31m | Hits: 99%/336859See results here. |
be592cf to
7940e31
Compare
| const auto __old_stream = __stream_; | ||
| __uninitialized_async_buffer __ret{ | ||
| __fake_resource_ref{::cuda::std::addressof(__mr_)}, __stream_, __count, __alignment_}; | ||
| __fake_resource_ref{::cuda::std::addressof(__mr_)}, __stream, __count, __alignment_}; | ||
| ::cuda::std::swap(__count_, __ret.__count_); | ||
| ::cuda::std::swap(__alignment_, __ret.__alignment_); | ||
| ::cuda::std::swap(__buf_, __ret.__buf_); | ||
| __stream_ = __stream; | ||
| __ret.__stream_ = __old_stream; |
There was a problem hiding this comment.
Couldnt we just use cuda::std::exchange
There was a problem hiding this comment.
Also why is this creating the new allocation on __stream and then destroying it on __old_stream afterwards
| template <class _Range> | ||
| static constexpr bool __compatible_range = (::cuda::std::ranges::__container_compatible_range<_Range, _Tp>); | ||
| [[nodiscard]] static constexpr bool __compatible_range() noexcept | ||
| { | ||
| return ::cuda::std::ranges::__container_compatible_range<_Range, _Tp>; | ||
| } | ||
|
|
||
| //! @brief Helper to check whether a different buffer still satisfies all | ||
| //! properties of this one | ||
| template <class... _OtherProperties> | ||
| static constexpr bool __properties_match = | ||
| ::cuda::std::__type_set_contains_v<::cuda::std::__make_type_set<_OtherProperties...>, _Properties...>; | ||
| [[nodiscard]] static constexpr bool __properties_match() noexcept | ||
| { | ||
| return ::cuda::std::__type_set_contains_v<::cuda::std::__make_type_set<_OtherProperties...>, _Properties...>; | ||
| } |
There was a problem hiding this comment.
I am not sure about those, they should stay static constexpr bool variables afaik
3 participants
closes #10213
Adds this functionality as internal separate class because of ABI constraints. Instead of adding a bunch of
make_*functions it exposes rvalue constructor from normal buffer, but exposes the constructors too since that part was easy.We can see how internal usage of it goes and consider folding it into the buffer in the next major version when we can break ABI.
The new buffer class has a capacity separate from size and size has a few variants, the less arguments provided the less cases are supported by it.
We could consider adding
appendthat actually initializes growing the buffer, but I think its not necessary right now.We could also think about exposing this type in
cuda::experimental, but I want to start as internal only