Add vectorized cub::DeviceTransform algorithm - #4815
Conversation
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
cub::DeviceTransform algorithm
3ca5ba0 to
75524e0
Compare
75524e0 to
2b05de2
Compare
0f0cb45 to
a6d1fdd
Compare
🟨 CI finished in 3h 03m: Pass: 84%/142 | Total: 3d 06h | Avg: 33m 12s | Max: 3h 03m | Hits: 78%/130051
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 142)
| # | Runner |
|---|---|
| 95 | linux-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 4 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
a6d1fdd to
ee55dbc
Compare
🟨 CI finished in 1h 57m: Pass: 84%/142 | Total: 1d 02h | Avg: 11m 04s | Max: 43m 20s | Hits: 99%/130051
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| stdpar | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 142)
| # | Runner |
|---|---|
| 95 | linux-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 4 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
ee55dbc to
9836026
Compare
🟨 CI finished in 2h 52m: Pass: 84%/142 | Total: 1d 09h | Avg: 14m 00s | Max: 2h 24m | Hits: 98%/130315
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 142)
| # | Runner |
|---|---|
| 95 | linux-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 4 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
9836026 to
1b4d642
Compare
🟨 CI finished in 2h 51m: Pass: 61%/142 | Total: 3d 00h | Avg: 30m 39s | Max: 2h 50m | Hits: 83%/92005
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 142)
| # | Runner |
|---|---|
| 95 | linux-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 4 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
5556f5c to
7e36771
Compare
🟨 CI finished in 5h 55m: Pass: 97%/138 | Total: 3d 05h | Avg: 33m 37s | Max: 1h 27m | Hits: 80%/158790
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| python | |
| +/- | CCCL C Parallel Library |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 138)
| # | Runner |
|---|---|
| 91 | linux-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 4 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
e18959b to
f81d7e6
Compare
🟨 CI finished in 3h 52m: Pass: 97%/138 | Total: 2d 11h | Avg: 26m 02s | Max: 2h 35m | Hits: 88%/158861
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| python | |
| +/- | CCCL C Parallel Library |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 138)
| # | Runner |
|---|---|
| 91 | linux-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 4 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
🟨 CI finished in 1h 08m: Pass: 97%/138 | Total: 1d 01h | Avg: 11m 06s | Max: 59m 28s | Hits: 99%/158861
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| python | |
| +/- | CCCL C Parallel Library |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 138)
| # | Runner |
|---|---|
| 91 | linux-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 4 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
| }; | ||
|
|
||
| template <int Bytes> | ||
| _CCCL_HOST_DEVICE _CCCL_CONSTEVAL auto load_store_type() |
There was a problem hiding this comment.
this is related to #4003.
In this PR, I used a template type with a given alignment
There was a problem hiding this comment.
You mean AlignedData? I tried that and the codegen was worse.
| @@ -274,11 +430,11 @@ _CCCL_DEVICE void transform_kernel_ublkcp( | |||
| _CCCL_ASSERT(reinterpret_cast<uintptr_t>(src) % alignof(T) == 0, ""); | |||
There was a problem hiding this comment.
not related to this PR, but I'm realizing that having a function to check the alignment would be more readable
There was a problem hiding this comment.
We actually do have thrust::detail::util::is_aligned, but it should be exposed more centrally.
b7051b4 to
e1816f5
Compare
🟩 CI finished in 3h 13m: Pass: 100%/145 | Total: 3d 06h | Avg: 32m 27s | Max: 1h 19m | Hits: 79%/163306
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| python | |
| +/- | CCCL C Parallel Library |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 145)
| # | Runner |
|---|---|
| 91 | linux-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | linux-amd64-gpu-h100-latest-1 |
| 11 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
| else | ||
| { | ||
| return true; // fancy iterators are aligned, since the vectorized kernel chooses a different code path | ||
| } |
There was a problem hiding this comment.
question: is this code unreacheable? We only ever select vectorized code path when all inputs are contiguous. Also, the vectorized kernel doesn't seem to have a different code path for fancy inputs. Is this comment a leftover?
There was a problem hiding this comment.
Yes. I also call this function on the output iterator below, which can be a fancy iterator.
e1816f5 to
4dab503
Compare
🟩 CI finished in 6h 56m: Pass: 100%/153 | Total: 3d 15h | Avg: 34m 17s | Max: 2h 52m | Hits: 77%/173216
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| stdpar | |
| python | |
| +/- | CCCL C Parallel Library |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 153)
| # | Runner |
|---|---|
| 93 | linux-amd64-cpu16 |
| 17 | windows-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | linux-amd64-gpu-h100-latest-1 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
7e0d852 to
7e85b2d
Compare
🟩 CI finished in 2h 19m: Pass: 100%/153 | Total: 3d 18h | Avg: 35m 32s | Max: 1h 46m | Hits: 64%/173216
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| stdpar | |
| python | |
| +/- | CCCL C Parallel Library |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 153)
| # | Runner |
|---|---|
| 93 | linux-amd64-cpu16 |
| 17 | windows-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | linux-amd64-gpu-h100-latest-1 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
|
I would like to squash the commits on this branch to make the rebase easier. @gevtushenko please flush any comments you still have. Thx! |
7e85b2d to
1518653
Compare
|
After the rebase, the vectorized kernel will be preferred over the prefetch kernel before Ampere (where the LDGSTS kernel will be preferred). I will file a separate PR with benchmark to decide between vectorized and LDGSTS kernel on Ampere. |
1518653 to
f469ae2
Compare
🟨 CI finished in 2h 04m: Pass: 99%/153 | Total: 3d 13h | Avg: 33m 36s | Max: 1h 24m | Hits: 77%/171303
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| stdpar | |
| python | |
| +/- | CCCL C Parallel Library |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 153)
| # | Runner |
|---|---|
| 93 | linux-amd64-cpu16 |
| 17 | windows-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | linux-amd64-gpu-h100-latest-1 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
🟩 CI finished in 2h 50m: Pass: 100%/153 | Total: 3d 13h | Avg: 33m 27s | Max: 1h 24m | Hits: 78%/173216
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| CCCL Packaging | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| stdpar | |
| python | |
| +/- | CCCL C Parallel Library |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | CCCL Packaging |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | stdpar |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 153)
| # | Runner |
|---|---|
| 93 | linux-amd64-cpu16 |
| 17 | windows-amd64-cpu16 |
| 12 | linux-amd64-gpu-rtxa6000-latest-1 |
| 11 | linux-amd64-gpu-h100-latest-1 |
| 10 | linux-arm64-cpu16 |
| 7 | linux-amd64-gpu-rtx2080-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
Fixes: #4837
Review and merge the test reorganizatin first:
prefetch (base) vs. vectorized (new)