Blackwell cluster-level work-stealing#3908
Closed
gonzalobg wants to merge 10 commits intoNVIDIA:mainfrom
Closed
Conversation
Contributor
Contributor
Author
|
pre-commit.ci autofix |
Contributor
|
pre-commit.ci autofix |
Contributor
|
/ok to test |
Contributor
🟨 CI finished in 1h 08m: Pass: 75%/158 | Total: 22h 42m | Avg: 8m 37s | Max: 40m 15s | Hits: 95%/151073
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | libcu++ |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | libcu++ |
| +/- | CUB |
| +/- | Thrust |
| +/- | CUDA Experimental |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 158)
| # | Runner |
|---|---|
| 111 | linux-amd64-cpu16 |
| 15 | windows-amd64-cpu16 |
| 10 | linux-arm64-cpu16 |
| 8 | linux-amd64-gpu-rtx2080-latest-1 |
| 6 | linux-amd64-gpu-rtxa6000-latest-1 |
| 5 | linux-amd64-gpu-h100-latest-1 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
miscco
requested changes
Feb 24, 2025
| #include <cuda/std/__utility/unreachable.h> | ||
| #include <cuda/std/cstdint> | ||
|
|
||
| #include <cooperative_groups.h> |
Contributor
There was a problem hiding this comment.
cooperative groups does not work with nvc++ currently so this needs to be guarded and we need to ensure that we might build without it
Collaborator
|
@gonzalobg are we still interested in carrying this forward? |
Contributor
Author
|
Yes, but i don't know when I can get to it. |
This file contains hidden or 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
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.
Description
This PR extends the work-stealing APIs with
for_each_canceled_clusterto enable cluster-level work stealing (see work-stealing tracking issue: #3870).It improves the example to show prologue and epilogue reuse, and moves the example to the examples directory so that it can be tested.
We should probably explore whether we actually need two APIs, or whether
for_each_canceled_blockcould internally detect a cluster and just internally switch to do cluster-level work stealing. Exploring this is tracked in #3870.Checklist