Add inclusive_scan with initial value support (warp/block) #1749
Conversation
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
🟩 CI Results [ Failed: 0 | Passed: 198 | Total: 198 ]
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
gonidelis
force-pushed
the
inclusive_scan_init
branch
from
da78a11
to
0ca0b3c
Compare
🟩 CI Results [ Failed: 0 | Passed: 198 | Total: 198 ]
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
bernhardmgruber
approved these changes
May 20, 2024
gonidelis
force-pushed
the
inclusive_scan_init
branch
from
0ca0b3c
to
f3142fa
Compare
gonidelis
changed the title
🟨 CI Results [ Failed: 3 | Passed: 195 | Total: 198 ]
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
🟩 CI Results [ Failed: 0 | Passed: 198 | Total: 198 ]
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
bernhardmgruber
approved these changes
May 21, 2024
gevtushenko
reviewed
May 21, 2024
gevtushenko
reviewed
May 21, 2024
gevtushenko
requested changes
May 21, 2024
🟩 CI Results [ Failed: 0 | Passed: 198 | Total: 198 ]
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
1 similar comment
🟩 CI Results [ Failed: 0 | Passed: 198 | Total: 198 ]
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
gevtushenko
requested changes
May 23, 2024
🟩 CI Results [ Failed: 0 | Passed: 198 | Total: 198 ]
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
1 similar comment
🟩 CI Results [ Failed: 0 | Passed: 198 | Total: 198 ]
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
gonidelis
force-pushed
the
inclusive_scan_init
branch
2 times, most recently
from
a932030
to
726e988
Compare
🟨 CI Results: Pass: 97%/249 | Total Time: 4d 16h | Avg Time: 27m 13s | Hits: 60%/244265
|
| 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: 249)
| # | Runner |
|---|---|
| 178 | linux-amd64-cpu16 |
| 40 | linux-amd64-gpu-v100-latest-1 |
| 16 | linux-arm64-cpu16 |
| 15 | windows-amd64-cpu16 |
gonidelis
force-pushed
the
inclusive_scan_init
branch
3 times, most recently
from
f38095f
to
5bf71e0
Compare
🟩 CI Results: Pass: 100%/249 | Total Time: 1d 05h | Avg Time: 7m 08s | Hits: 99%/248441
|
| 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: 249)
| # | Runner |
|---|---|
| 178 | linux-amd64-cpu16 |
| 40 | linux-amd64-gpu-v100-latest-1 |
| 16 | linux-arm64-cpu16 |
| 15 | windows-amd64-cpu16 |
gonidelis
force-pushed
the
inclusive_scan_init
branch
from
c5fd4e9
to
a2c3772
Compare
gonidelis
force-pushed
the
inclusive_scan_init
branch
from
a2c3772
to
e41e85a
Compare
gevtushenko
approved these changes
Jun 10, 2024
cub/cub/block/block_scan.cuh
Outdated
| //! Snippet | ||
| //! +++++++ | ||
| //! | ||
| //! The code snippet below illustrates an inclusive prefix max scan of 512 integer items that |
There was a problem hiding this comment.
Suggested change
| //! The code snippet below illustrates an inclusive prefix max scan of 512 integer items that | |
| //! The code snippet below illustrates an inclusive prefix max scan of 128 integer items that |
cub/cub/block/block_scan.cuh
Outdated
| //! +++++++ | ||
| //! | ||
| //! The code snippet below illustrates an inclusive prefix max scan of 512 integer items that | ||
| //! are partitioned in a :ref:`blocked arrangement <flexible-data-arrangement>` across 128 threads |
There was a problem hiding this comment.
Suggested change
| //! are partitioned in a :ref:`blocked arrangement <flexible-data-arrangement>` across 128 threads | |
| //! are partitioned in a :ref:`blocked arrangement <flexible-data-arrangement>` across 64 threads |
cub/cub/block/block_scan.cuh
Outdated
| //! | ||
| //! The code snippet below illustrates an inclusive prefix max scan of 512 integer items that | ||
| //! are partitioned in a :ref:`blocked arrangement <flexible-data-arrangement>` across 128 threads | ||
| //! where each thread owns 4 consecutive items. |
There was a problem hiding this comment.
Suggested change
| //! where each thread owns 4 consecutive items. | |
| //! where each thread owns 2 consecutive items. |
cub/cub/block/block_scan.cuh
Outdated
Comment on lines
2268
to
2271
| //! Suppose the set of input ``thread_data`` across the block of threads is | ||
| //! ``{[0, -1], [2, -3],[4, -5], ... [126, -127]}``. | ||
| //! The corresponding output ``thread_data`` in those threads will be | ||
| //! ``{[1, 1], [2, 2],[3, 3], ... [126, 126]}``. |
There was a problem hiding this comment.
suggestion: I think we cover this in the kernel itself. I think seeing input and expected output in the docs might be confusing, but feel free to disagree:
Suggested change
| //! Suppose the set of input ``thread_data`` across the block of threads is | |
| //! ``{[0, -1], [2, -3],[4, -5], ... [126, -127]}``. | |
| //! The corresponding output ``thread_data`` in those threads will be | |
| //! ``{[1, 1], [2, 2],[3, 3], ... [126, 126]}``. |
cub/cub/block/block_scan.cuh
Outdated
| //! | ||
| //! The code snippet below illustrates an inclusive prefix max scan of 128 integer items that | ||
| //! are partitioned in a :ref:`blocked arrangement <flexible-data-arrangement>` across 64 threads | ||
| //! where each thread owns 4 consecutive items. |
There was a problem hiding this comment.
Suggested change
| //! where each thread owns 4 consecutive items. | |
| //! where each thread owns 2 consecutive items. |
| // warp #2 input: {2, 3, 4, 5, ..., 33} | ||
| // warp #4 input: {3, 4, 5, 6, ..., 34} | ||
|
|
||
| // Collectively compute the block-wide inclusive prefix max scan |
There was a problem hiding this comment.
Suggested change
| // Collectively compute the block-wide inclusive prefix max scan | |
| // Collectively compute warp-wide inclusive prefix max scans |
| // warp #2 input: {1, 1, 1, 1, ..., 1} | ||
| // warp #4 input: {1, 1, 1, 1, ..., 1} | ||
|
|
||
| // Collectively compute the block-wide inclusive prefix max scan |
There was a problem hiding this comment.
Suggested change
| // Collectively compute the block-wide inclusive prefix max scan | |
| // Collectively compute warp-wide inclusive prefix max scans |
| __global__ void InclusiveScanKernelAggr(int* output, int* d_warp_aggregate) | ||
| { | ||
| // Specialize WarpScan for type int | ||
| typedef cub::WarpScan<int> warp_scan_t; |
There was a problem hiding this comment.
Suggested change
| typedef cub::WarpScan<int> warp_scan_t; | |
| using warp_scan_t = cub::WarpScan<int>; |
|
|
||
| cuda::std::inclusive_scan(start, end, start, sum_op<int>{}, init_val); | ||
|
|
||
| expected_aggr.push_back(expected[i * 32 + 31] - init_val); // warp aggregate doed not take |
There was a problem hiding this comment.
Suggested change
| expected_aggr.push_back(expected[i * 32 + 31] - init_val); // warp aggregate doed not take | |
| expected_aggr.push_back(expected[i * 32 + 31] - init_val); // warp aggregate does not take |
There was a problem hiding this comment.
question: is this always equivalent to:
Suggested change
| expected_aggr.push_back(expected[i * 32 + 31] - init_val); // warp aggregate doed not take | |
| expected_aggr.push_back(32); // warp aggregate does not take |
| }; | ||
| // input: {[0, -1], [2, -3],[4, -5], ... [126, -127]} | ||
|
|
||
| // Collectively compute the block-wide inclusive prefix max scan |
There was a problem hiding this comment.
question: I was under impression that it's "prefix sum" or "scan". I never heart of "prefix scan". Is this term used in the wild or is it a typo?
🟨 CI finished in 5h 44m: Pass: 99%/249 | Total: 2d 07h | Avg: 13m 20s | Max: 2h 01m | Hits: 91%/246866
|
| 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: 249)
| # | Runner |
|---|---|
| 178 | linux-amd64-cpu16 |
| 40 | linux-amd64-gpu-v100-latest-1 |
| 16 | linux-arm64-cpu16 |
| 15 | windows-amd64-cpu16 |
🟩 CI finished in 19h 49m: Pass: 100%/249 | Total: 2d 07h | Avg: 13m 28s | Max: 2h 01m | Hits: 91%/248572
|
| 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: 249)
| # | Runner |
|---|---|
| 178 | linux-amd64-cpu16 |
| 40 | linux-amd64-gpu-v100-latest-1 |
| 16 | linux-arm64-cpu16 |
| 15 | windows-amd64-cpu16 |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
Doesn't fix but works towards #693.
Adds warp_scan with initial value support on warp and block level. Tests added.
todo: