-
Notifications
You must be signed in to change notification settings - Fork 205
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
Implement P0843 inplace_vector
#1936
Conversation
907c81c
to
a5e64cd
Compare
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.
Started reviewing test, here are a few first comments
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assign.pass.cpp
Outdated
Show resolved
Hide resolved
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assign.pass.cpp
Outdated
Show resolved
Hide resolved
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assign.pass.cpp
Outdated
Show resolved
Hide resolved
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assign.pass.cpp
Outdated
Show resolved
Hide resolved
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assign.pass.cpp
Show resolved
Hide resolved
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 reviewed until the beginning of class inplace_vector
. Here is what I have so far:
//[container.intro.reqmts] | ||
# if _CCCL_STD_VER >= 2020 | ||
template <class _Range, class _Tp> | ||
concept __container_compatible_range = input_range<_Range> && convertible_to<range_reference_t<_Range>, _Tp>; | ||
# else // ^^^ C++20 ^^^ / vvv C++17 vvv | ||
template <class _Range, class _Tp> | ||
_LIBCUDACXX_CONCEPT_FRAGMENT( | ||
__container_compatible_range_, | ||
requires()(requires(input_range<_Range>), requires(convertible_to<range_reference_t<_Range>, _Tp>))); | ||
|
||
template <class _Range, class _Tp> | ||
_LIBCUDACXX_CONCEPT __container_compatible_range = _LIBCUDACXX_FRAGMENT(__container_compatible_range_, _Range, _Tp); | ||
# endif // _CCCL_STD_VER <= 2017 | ||
|
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: That does not seem to relate to inplace_vector
? Furthermore, the comment ([container.intro.reqmts]
) suggests this comes from a specification, but then all concepts are internal (starting with __
). I see the term used in P0843, but I guess this was standardized elsewhere? Where does this come from? Would it make sense to split it from this PR?
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 the exposition only concept container-compatible-range
that was introduced with ranges::to
and is used within inplace_vector
As a convention we translate all exposition only entities some-standard-name
to __some_standard_name
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assign.pass.cpp
Outdated
Show resolved
Hide resolved
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 finished a first pass over inplace_vector
. Here is what I have so far:
🟨 CI finished in 4h 57m: Pass: 98%/416 | Total: 2d 17h | Avg: 9m 27s | Max: 42m 58s | Hits: 85%/505892
|
Project | |
---|---|
CCCL Infrastructure | |
+/- | libcu++ |
CUB | |
Thrust | |
CUDA Experimental |
Modifications in project or dependencies?
Project | |
---|---|
CCCL Infrastructure | |
+/- | libcu++ |
+/- | CUB |
+/- | Thrust |
+/- | CUDA Experimental |
🏃 Runner counts (total jobs: 416)
# | Runner |
---|---|
305 | linux-amd64-cpu16 |
60 | linux-amd64-gpu-v100-latest-1 |
28 | linux-arm64-cpu16 |
23 | windows-amd64-cpu16 |
🟨 CI finished in 6h 27m: Pass: 98%/417 | Total: 2d 11h | Avg: 8m 37s | Max: 45m 38s | Hits: 83%/507125
|
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 |
🟨 CI finished in 17h 17m: Pass: 99%/417 | Total: 2d 12h | Avg: 8m 37s | Max: 44m 31s | Hits: 82%/518911
|
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 |
🟨 CI finished in 7h 11m: Pass: 96%/417 | Total: 2d 08h | Avg: 8m 05s | Max: 2h 12m | Hits: 90%/485967
|
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 |
🟨 CI finished in 1h 28m: Pass: 97%/417 | Total: 2d 04h | Avg: 7m 32s | Max: 44m 39s | Hits: 91%/516344
|
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 |
🟨 CI finished in 16h 51m: Pass: 96%/421 | Total: 7d 09h | Avg: 25m 20s | Max: 1h 14m | Hits: 14%/34188
|
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: 421)
# | Runner |
---|---|
305 | linux-amd64-cpu16 |
65 | linux-amd64-gpu-v100-latest-1 |
28 | linux-arm64-cpu16 |
23 | windows-amd64-cpu16 |
🟨 CI finished in 2d 22h: Pass: 99%/421 | Total: 7d 10h | Avg: 25m 27s | Max: 1h 14m | Hits: 14%/34188
|
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: 421)
# | Runner |
---|---|
305 | linux-amd64-cpu16 |
65 | linux-amd64-gpu-v100-latest-1 |
28 | linux-arm64-cpu16 |
23 | windows-amd64-cpu16 |
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.
Thank you Michael, this LGTM!
🟨 CI finished in 3d 13h: Pass: 99%/421 | Total: 2d 13h | Avg: 8m 43s | Max: 1h 33m | Hits: 71%/34188
|
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: 421)
# | Runner |
---|---|
305 | linux-amd64-cpu16 |
65 | linux-amd64-gpu-v100-latest-1 |
28 | linux-arm64-cpu16 |
23 | windows-amd64-cpu16 |
🟨 CI finished in 4d 03h: Pass: 99%/421 | Total: 7d 11h | Avg: 25m 32s | Max: 1h 14m | Hits: 14%/34188
|
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: 421)
# | Runner |
---|---|
305 | linux-amd64-cpu16 |
65 | linux-amd64-gpu-v100-latest-1 |
28 | linux-arm64-cpu16 |
23 | windows-amd64-cpu16 |
🟨 CI finished in 14h 24m: Pass: 99%/417 | Total: 8d 00h | Avg: 27m 41s | Max: 1h 13m | Hits: 14%/34324
|
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 |
🟩 CI finished in 19h 26m: Pass: 100%/417 | Total: 8d 00h | Avg: 27m 41s | Max: 1h 13m | Hits: 14%/34324
|
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 |
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 reviewed parts of the tests and have a few more comments. Otherwise, I think this is good to go! Well done!
{ // inplace_vector<T, N>::assign(iter, iter), with forward iterators empty range | ||
const cuda::std::array<T, 0> expected = {}; | ||
inplace_vector vec{}; | ||
vec.assign(expected.begin(), expected.end()); |
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.
The comment mentions forward iterators
but the code uses array::iterator
which is contiguous. Did you miss wrapping the array's iterators like you did above for the input iterators?
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.
That test requires "at least" forward iterator, because that is the only relevant information we need here, because we can just call distance
and that will make the right thing
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assign.pass.cpp
Outdated
Show resolved
Hide resolved
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assignment.pass.cpp
Show resolved
Hide resolved
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assignment.pass.cpp
Show resolved
Hide resolved
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/assignment.pass.cpp
Show resolved
Hide resolved
libcudacxx/test/libcudacxx/std/containers/sequences/inplace_vector/comparison.pass.cpp
Outdated
Show resolved
Hide resolved
template <class T> | ||
__host__ __device__ constexpr void test_relation() | ||
{ |
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.
Important: the specification contains operator<=>
, so we should also add a test for it.
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.
We do not have proper support for <=> on device, so I cannot add that as easily
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.
Then we should at least leave a TODO comment that we should add a test for <=> once device compiler support is major enough.
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.
Oh believe me there is going to be all kinds of fun when this ever happens
This was added after the final paper was merged
* Move enum `endian` to its own file * Move `std::rotl` and `std::rotr` to their own file * Move `std::has_single_bit` to its own file * Move `countr_{one, zero}` to its ownn file * Move `countl_{one, zero}` to their own file * Move `bit_ceil`, `bit_floor` and `bit_width` to their own file * Cleanup the `<bit>` header
🟥 CI finished in 2h 50m: Pass: 0%/417 | Total: 2d 21h | Avg: 10m 03s | Max: 1h 22m
|
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 |
🟩 CI finished in 2h 50m: Pass: 100%/417 | Total: 2d 21h | Avg: 10m 03s | Max: 1h 22m | Hits: 71%/34420
|
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 |
This implements
inplace_vector
a resizable container with a fixed capacity that stores its elements in a local array.Due to the fact that everything is local we are more or less save regarding host device issues, as long as users do not pass around references of it.
The exception guarantees are not 100% clear yet, so I implemented them on a best effort basis. We might need to revisit what we guarantee in case of an exception.