[libcudacxx] Stable abstraction for Blackwell work-stealing (PTX try_cancel) by gonzalobg · Pull Request #3671 · NVIDIA/cccl

gonzalobg · 2025-02-04T14:24:49Z

Description

This PR provides a stable low-level abstraction for Blackwell work-stealing features (PTX try_cancel) in libcu++. It's targeting CUDA 13.0.

The main goals of this low-level abstraction are to:

deliver a reasonable amount of functionality with reasonable performance,
in a way that enables us to deliver low-level bugfixes, performance tuning, and more finer grained control in the future,
with an API that is as easy to use as possible (hiding the gnarly synchronization),
while still hitting the CUDA 13.0 deadline.

This PR only covers the thread-block level work-stealing. Once it lands, a follow-up PR will add cluster level work-stealing in a 1:1 fashion.

The following features have made MVP:

Stable API.
Portable to all supported compute capabilities.
Supports all C++ Standard versions supported by libcu++.

There are many extensions to this feature that are worth evaluating but not made it into the MVP. These are tracked here: #3870 .

Checklist

I am familiar with the Contributing Guidelines.
New or existing tests cover these changes.
The documentation is up to date with these changes.

copy-pr-bot · 2025-02-04T14:24:53Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

libcudacxx/include/cuda/try_cancel

Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>

gonzalobg · 2025-02-07T09:53:50Z

pre-commit.ci autofix

docs/libcudacxx/extended_api/work_stealing.rst

libcudacxx/include/cuda/try_cancel

libcudacxx/test/libcudacxx/cuda/try_cancel/try_cancel.pass.cpp

libcudacxx/include/cuda/try_cancel

bernhardmgruber · 2025-02-17T19:33:48Z

We discussed the PR on VC today. @gonzalobg will push a few fixes, @miscco will polish, and I shall review.

libcudacxx/test/libcudacxx/cuda/try_cancel/try_cancel.pass.cpp

miscco · 2025-02-19T12:13:09Z

libcudacxx/test/libcudacxx/cuda/for_each_canceled/for_each_canceled.pass.cpp

+    if (auto e = cudaDeviceSynchronize(); e != cudaSuccess)
+    {
+      std::cerr << "ERROR: synchronize failed" << std::endl;
+      return false;
+    }


assert is fine

docs/libcudacxx/extended_api/work_stealing.rst

miscco · 2025-02-19T13:03:52Z

/ok to test

miscco · 2025-02-19T13:07:34Z

/ok to test

miscco · 2025-02-19T13:10:52Z

/ok to test

libcudacxx/test/libcudacxx/cuda/for_each_canceled/for_each_canceled.pass.cpp

libcudacxx/include/cuda/for_each_canceled

miscco · 2025-02-19T14:38:29Z

/ok to test

miscco · 2025-02-19T14:43:38Z

/ok to test

miscco · 2025-02-19T14:58:21Z

/ok to test

miscco · 2025-02-19T15:07:40Z

/ok to test

github-actions · 2025-02-19T16:37:00Z

🟩 CI finished in 1h 36m: Pass: 100%/158 | Total: 3d 01h | Avg: 27m 50s | Max: 1h 24m | Hits: 76%/247898

🟩 cub: Pass: 100%/45 | Total: 1d 17h | Avg: 55m 37s | Max: 1h 24m | Hits: 47%/53221

🟩 cpu
  🟩 amd64              Pass: 100%/43  | Total:  1d 15h | Avg: 55m 20s | Max:  1h 24m | Hits:  48%/50803 
  🟩 arm64              Pass: 100%/2   | Total:  2h 03m | Avg:  1h 01m | Max:  1h 02m | Hits:  39%/2418  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total:  5h 13m | Avg:  1h 02m | Max:  1h 05m | Hits:  34%/5879  
  🟩 12.5               Pass: 100%/2   | Total:  2h 16m | Avg:  1h 08m | Max:  1h 08m | Hits:  36%/2236  
  🟩 12.8               Pass: 100%/38  | Total:  1d 10h | Avg: 54m 01s | Max:  1h 24m | Hits:  50%/45106 
🟩 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total:  2h 01m | Avg:  1h 00m | Max:  1h 01m | Hits:  40%/2090  
  🟩 nvcc12.0           Pass: 100%/5   | Total:  5h 13m | Avg:  1h 02m | Max:  1h 05m | Hits:  34%/5879  
  🟩 nvcc12.5           Pass: 100%/2   | Total:  2h 16m | Avg:  1h 08m | Max:  1h 08m | Hits:  36%/2236  
  🟩 nvcc12.8           Pass: 100%/36  | Total:  1d 08h | Avg: 53m 38s | Max:  1h 24m | Hits:  50%/43016 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total:  2h 01m | Avg:  1h 00m | Max:  1h 01m | Hits:  40%/2090  
  🟩 nvcc               Pass: 100%/43  | Total:  1d 15h | Avg: 55m 22s | Max:  1h 24m | Hits:  48%/51131 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total:  4h 10m | Avg:  1h 02m | Max:  1h 03m | Hits:  39%/4844  
  🟩 Clang15            Pass: 100%/2   | Total:  2h 05m | Avg:  1h 02m | Max:  1h 04m | Hits:  39%/2418  
  🟩 Clang16            Pass: 100%/2   | Total:  2h 02m | Avg:  1h 01m | Max:  1h 01m | Hits:  39%/2418  
  🟩 Clang17            Pass: 100%/2   | Total:  2h 10m | Avg:  1h 05m | Max:  1h 06m | Hits:  39%/2418  
  🟩 Clang18            Pass: 100%/7   | Total:  5h 50m | Avg: 50m 03s | Max:  1h 03m | Hits:  57%/8135  
  🟩 GCC7               Pass: 100%/2   | Total:  2h 03m | Avg:  1h 01m | Max:  1h 02m | Hits:  38%/2422  
  🟩 GCC8               Pass: 100%/1   | Total: 56m 53s | Avg: 56m 53s | Max: 56m 53s | Hits:  38%/1211  
  🟩 GCC9               Pass: 100%/2   | Total:  2h 03m | Avg:  1h 01m | Max:  1h 02m | Hits:  38%/2422  
  🟩 GCC10              Pass: 100%/2   | Total:  2h 04m | Avg:  1h 02m | Max:  1h 04m | Hits:  38%/2422  
  🟩 GCC11              Pass: 100%/2   | Total:  2h 07m | Avg:  1h 03m | Max:  1h 09m | Hits:  38%/2418  
  🟩 GCC12              Pass: 100%/2   | Total:  1h 58m | Avg: 59m 28s | Max:  1h 00m | Hits:  38%/2418  
  🟩 GCC13              Pass: 100%/11  | Total:  6h 49m | Avg: 37m 15s | Max:  1h 07m | Hits:  72%/13299 
  🟩 MSVC14.29          Pass: 100%/2   | Total:  2h 23m | Avg:  1h 11m | Max:  1h 18m | Hits:  12%/2070  
  🟩 MSVC14.42          Pass: 100%/2   | Total:  2h 40m | Avg:  1h 20m | Max:  1h 24m | Hits:  12%/2070  
  🟩 NVHPC24.7          Pass: 100%/2   | Total:  2h 16m | Avg:  1h 08m | Max:  1h 08m | Hits:  36%/2236  
🟩 cxx_family
  🟩 Clang              Pass: 100%/17  | Total: 16h 18m | Avg: 57m 34s | Max:  1h 06m | Hits:  46%/20233 
  🟩 GCC                Pass: 100%/22  | Total: 18h 04m | Avg: 49m 16s | Max:  1h 09m | Hits:  55%/26612 
  🟩 MSVC               Pass: 100%/4   | Total:  5h 03m | Avg:  1h 15m | Max:  1h 24m | Hits:  12%/4140  
  🟩 NVHPC              Pass: 100%/2   | Total:  2h 16m | Avg:  1h 08m | Max:  1h 08m | Hits:  36%/2236  
🟩 gpu
  🟩 h100               Pass: 100%/3   | Total:  1h 17m | Avg: 25m 59s | Max: 29m 16s | Hits:  79%/3627  
  🟩 rtx2080            Pass: 100%/34  | Total:  1d 12h | Avg:  1h 04m | Max:  1h 24m | Hits:  36%/39922 
  🟩 rtxa6000           Pass: 100%/8   | Total:  4h 07m | Avg: 30m 57s | Max:  1h 01m | Hits:  84%/9672  
🟩 jobs
  🟩 Build              Pass: 100%/37  | Total:  1d 14h | Avg:  1h 02m | Max:  1h 24m | Hits:  36%/43549 
  🟩 DeviceLaunch       Pass: 100%/1   | Total: 21m 07s | Avg: 21m 07s | Max: 21m 07s | Hits:  99%/1209  
  🟩 GraphCapture       Pass: 100%/1   | Total: 16m 07s | Avg: 16m 07s | Max: 16m 07s | Hits:  99%/1209  
  🟩 HostLaunch         Pass: 100%/3   | Total:  1h 13m | Avg: 24m 27s | Max: 25m 51s | Hits:  99%/3627  
  🟩 TestGPU            Pass: 100%/3   | Total:  1h 03m | Avg: 21m 07s | Max: 22m 51s | Hits:  99%/3627  
🟩 sm
  🟩 90                 Pass: 100%/3   | Total:  1h 17m | Avg: 25m 59s | Max: 29m 16s | Hits:  79%/3627  
  🟩 90;90a;100         Pass: 100%/1   | Total:  1h 07m | Avg:  1h 07m | Max:  1h 07m | Hits:  38%/1209  
🟩 std
  🟩 17                 Pass: 100%/20  | Total: 21h 05m | Avg:  1h 03m | Max:  1h 18m | Hits:  35%/23419 
  🟩 20                 Pass: 100%/25  | Total: 20h 37m | Avg: 49m 31s | Max:  1h 24m | Hits:  57%/29802

🟩 thrust: Pass: 100%/45 | Total: 22h 31m | Avg: 30m 02s | Max: 1h 07m | Hits: 72%/80136

🟩 cmake_options
  🟩 -DTHRUST_DISPATCH_TYPE=Force32bit Pass: 100%/2   | Total: 36m 41s | Avg: 18m 20s | Max: 25m 39s | Hits:  88%/3564  
🟩 cpu
  🟩 amd64              Pass: 100%/43  | Total: 21h 38m | Avg: 30m 12s | Max:  1h 07m | Hits:  72%/76573 
  🟩 arm64              Pass: 100%/2   | Total: 52m 54s | Avg: 26m 27s | Max: 27m 04s | Hits:  76%/3563  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total:  2h 52m | Avg: 34m 25s | Max: 55m 28s | Hits:  65%/8901  
  🟩 12.5               Pass: 100%/2   | Total:  2h 11m | Avg:  1h 05m | Max:  1h 07m | Hits:  27%/3562  
  🟩 12.8               Pass: 100%/38  | Total: 17h 28m | Avg: 27m 34s | Max:  1h 01m | Hits:  76%/67673 
🟩 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total: 45m 22s | Avg: 22m 41s | Max: 23m 16s | Hits:  76%/3562  
  🟩 nvcc12.0           Pass: 100%/5   | Total:  2h 52m | Avg: 34m 25s | Max: 55m 28s | Hits:  65%/8901  
  🟩 nvcc12.5           Pass: 100%/2   | Total:  2h 11m | Avg:  1h 05m | Max:  1h 07m | Hits:  27%/3562  
  🟩 nvcc12.8           Pass: 100%/36  | Total: 16h 42m | Avg: 27m 51s | Max:  1h 01m | Hits:  76%/64111 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total: 45m 22s | Avg: 22m 41s | Max: 23m 16s | Hits:  76%/3562  
  🟩 nvcc               Pass: 100%/43  | Total: 21h 46m | Avg: 30m 23s | Max:  1h 07m | Hits:  72%/76574 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total:  1h 53m | Avg: 28m 23s | Max: 29m 42s | Hits:  76%/7124  
  🟩 Clang15            Pass: 100%/2   | Total: 53m 44s | Avg: 26m 52s | Max: 27m 33s | Hits:  76%/3562  
  🟩 Clang16            Pass: 100%/2   | Total: 57m 55s | Avg: 28m 57s | Max: 30m 00s | Hits:  76%/3562  
  🟩 Clang17            Pass: 100%/2   | Total:  1h 00m | Avg: 30m 24s | Max: 30m 31s | Hits:  76%/3562  
  🟩 Clang18            Pass: 100%/7   | Total:  2h 19m | Avg: 19m 56s | Max: 26m 39s | Hits:  83%/12467 
  🟩 GCC7               Pass: 100%/2   | Total: 56m 41s | Avg: 28m 20s | Max: 28m 44s | Hits:  76%/3564  
  🟩 GCC8               Pass: 100%/1   | Total: 28m 05s | Avg: 28m 05s | Max: 28m 05s | Hits:  76%/1782  
  🟩 GCC9               Pass: 100%/2   | Total:  1h 00m | Avg: 30m 16s | Max: 31m 17s | Hits:  76%/3564  
  🟩 GCC10              Pass: 100%/2   | Total: 59m 47s | Avg: 29m 53s | Max: 31m 41s | Hits:  76%/3564  
  🟩 GCC11              Pass: 100%/2   | Total:  1h 02m | Avg: 31m 03s | Max: 31m 22s | Hits:  76%/3564  
  🟩 GCC12              Pass: 100%/2   | Total: 58m 06s | Avg: 29m 03s | Max: 29m 28s | Hits:  76%/3564  
  🟩 GCC13              Pass: 100%/10  | Total:  3h 29m | Avg: 20m 57s | Max: 33m 10s | Hits:  85%/17820 
  🟩 MSVC14.29          Pass: 100%/2   | Total:  1h 51m | Avg: 55m 43s | Max: 55m 58s | Hits:  35%/3550  
  🟩 MSVC14.42          Pass: 100%/3   | Total:  2h 28m | Avg: 49m 28s | Max:  1h 01m | Hits:  37%/5325  
  🟩 NVHPC24.7          Pass: 100%/2   | Total:  2h 11m | Avg:  1h 05m | Max:  1h 07m | Hits:  27%/3562  
🟩 cxx_family
  🟩 Clang              Pass: 100%/17  | Total:  7h 05m | Avg: 25m 02s | Max: 30m 31s | Hits:  79%/30277 
  🟩 GCC                Pass: 100%/21  | Total:  8h 54m | Avg: 25m 28s | Max: 33m 10s | Hits:  80%/37422 
  🟩 MSVC               Pass: 100%/5   | Total:  4h 19m | Avg: 51m 58s | Max:  1h 01m | Hits:  36%/8875  
  🟩 NVHPC              Pass: 100%/2   | Total:  2h 11m | Avg:  1h 05m | Max:  1h 07m | Hits:  27%/3562  
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 29m 40s | Avg: 14m 50s | Max: 18m 12s | Hits:  88%/3564  
  🟩 rtx2080            Pass: 100%/33  | Total: 18h 20m | Avg: 33m 20s | Max:  1h 07m | Hits:  69%/58769 
  🟩 rtx4090            Pass: 100%/10  | Total:  3h 41m | Avg: 22m 10s | Max:  1h 01m | Hits:  82%/17803 
🟩 jobs
  🟩 Build              Pass: 100%/38  | Total: 21h 02m | Avg: 33m 14s | Max:  1h 07m | Hits:  68%/67671 
  🟩 TestCPU            Pass: 100%/3   | Total: 45m 30s | Avg: 15m 10s | Max: 30m 05s | Hits:  90%/5338  
  🟩 TestGPU            Pass: 100%/4   | Total: 43m 27s | Avg: 10m 51s | Max: 11m 28s | Hits:  99%/7127  
🟩 sm
  🟩 90                 Pass: 100%/2   | Total: 29m 40s | Avg: 14m 50s | Max: 18m 12s | Hits:  88%/3564  
  🟩 90;90a;100         Pass: 100%/1   | Total: 33m 10s | Avg: 33m 10s | Max: 33m 10s | Hits:  76%/1782  
🟩 std
  🟩 17                 Pass: 100%/20  | Total: 11h 34m | Avg: 34m 43s | Max:  1h 07m | Hits:  67%/35611 
  🟩 20                 Pass: 100%/23  | Total: 10h 20m | Avg: 26m 59s | Max:  1h 03m | Hits:  76%/40961

🟩 libcudacxx: Pass: 100%/43 | Total: 6h 09m | Avg: 8m 35s | Max: 25m 41s | Hits: 92%/102981

🟩 cpu
  🟩 amd64              Pass: 100%/41  | Total:  6h 01m | Avg:  8m 49s | Max: 25m 41s | Hits:  91%/97332 
  🟩 arm64              Pass: 100%/2   | Total:  7m 54s | Avg:  3m 57s | Max:  4m 02s | Hits:  98%/5649  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total: 39m 12s | Avg:  7m 50s | Max: 23m 23s | Hits:  89%/13662 
  🟩 12.5               Pass: 100%/2   | Total: 19m 34s | Avg:  9m 47s | Max:  9m 58s | Hits:  98%/5594  
  🟩 12.8               Pass: 100%/36  | Total:  5h 10m | Avg:  8m 38s | Max: 25m 41s | Hits:  92%/83725 
🟩 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total: 43m 38s | Avg: 21m 49s | Max: 23m 57s | Hits:  26%/5612  
  🟩 nvcc12.0           Pass: 100%/5   | Total: 39m 12s | Avg:  7m 50s | Max: 23m 23s | Hits:  89%/13662 
  🟩 nvcc12.5           Pass: 100%/2   | Total: 19m 34s | Avg:  9m 47s | Max:  9m 58s | Hits:  98%/5594  
  🟩 nvcc12.8           Pass: 100%/34  | Total:  4h 27m | Avg:  7m 51s | Max: 25m 41s | Hits:  97%/78113 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total: 43m 38s | Avg: 21m 49s | Max: 23m 57s | Hits:  26%/5612  
  🟩 nvcc               Pass: 100%/41  | Total:  5h 26m | Avg:  7m 57s | Max: 25m 41s | Hits:  96%/97369 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total: 18m 07s | Avg:  4m 31s | Max:  4m 48s | Hits:  98%/11192 
  🟩 Clang15            Pass: 100%/2   | Total:  9m 09s | Avg:  4m 34s | Max:  4m 44s | Hits:  99%/5606  
  🟩 Clang16            Pass: 100%/2   | Total:  9m 05s | Avg:  4m 32s | Max:  4m 50s | Hits:  99%/5606  
  🟩 Clang17            Pass: 100%/2   | Total:  9m 06s | Avg:  4m 33s | Max:  4m 42s | Hits:  99%/5606  
  🟩 Clang18            Pass: 100%/6   | Total:  1h 14m | Avg: 12m 22s | Max: 23m 57s | Hits:  63%/14042 
  🟩 GCC7               Pass: 100%/2   | Total:  7m 38s | Avg:  3m 49s | Max:  4m 03s | Hits:  99%/5544  
  🟩 GCC8               Pass: 100%/1   | Total:  4m 26s | Avg:  4m 26s | Max:  4m 26s | Hits:  99%/2782  
  🟩 GCC9               Pass: 100%/2   | Total:  8m 00s | Avg:  4m 00s | Max:  4m 18s | Hits:  99%/5556  
  🟩 GCC10              Pass: 100%/2   | Total:  8m 24s | Avg:  4m 12s | Max:  4m 13s | Hits:  98%/5612  
  🟩 GCC11              Pass: 100%/2   | Total:  8m 36s | Avg:  4m 18s | Max:  4m 27s | Hits:  98%/5608  
  🟩 GCC12              Pass: 100%/2   | Total:  8m 42s | Avg:  4m 21s | Max:  4m 28s | Hits:  98%/5608  
  🟩 GCC13              Pass: 100%/10  | Total:  1h 27m | Avg:  8m 45s | Max: 16m 28s | Hits:  95%/14301 
  🟩 MSVC14.29          Pass: 100%/2   | Total: 49m 04s | Avg: 24m 32s | Max: 25m 41s | Hits:  74%/5082  
  🟩 MSVC14.42          Pass: 100%/2   | Total: 47m 56s | Avg: 23m 58s | Max: 24m 18s | Hits:  98%/5242  
  🟩 NVHPC24.7          Pass: 100%/2   | Total: 19m 34s | Avg:  9m 47s | Max:  9m 58s | Hits:  98%/5594  
🟩 cxx_family
  🟩 Clang              Pass: 100%/16  | Total:  1h 59m | Avg:  7m 29s | Max: 23m 57s | Hits:  87%/42052 
  🟩 GCC                Pass: 100%/21  | Total:  2h 13m | Avg:  6m 21s | Max: 16m 28s | Hits:  97%/45011 
  🟩 MSVC               Pass: 100%/4   | Total:  1h 37m | Avg: 24m 15s | Max: 25m 41s | Hits:  86%/10324 
  🟩 NVHPC              Pass: 100%/2   | Total: 19m 34s | Avg:  9m 47s | Max:  9m 58s | Hits:  98%/5594  
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 17m 58s | Avg:  8m 59s | Max: 13m 27s | Hits:  98%/2914  
  🟩 rtx2080            Pass: 100%/41  | Total:  5h 51m | Avg:  8m 34s | Max: 25m 41s | Hits:  92%/100067
🟩 jobs
  🟩 Build              Pass: 100%/37  | Total:  5h 02m | Avg:  8m 10s | Max: 25m 41s | Hits:  92%/102941
  🟩 NVRTC              Pass: 100%/2   | Total: 32m 54s | Avg: 16m 27s | Max: 16m 28s | Hits:  90%/40    
  🟩 Test               Pass: 100%/3   | Total: 31m 40s | Avg: 10m 33s | Max: 13m 27s
  🟩 VerifyCodegen      Pass: 100%/1   | Total:  2m 28s | Avg:  2m 28s | Max:  2m 28s
🟩 sm
  🟩 75                 Pass: 100%/2   | Total: 32m 54s | Avg: 16m 27s | Max: 16m 28s | Hits:  90%/40    
  🟩 90                 Pass: 100%/2   | Total: 17m 58s | Avg:  8m 59s | Max: 13m 27s | Hits:  98%/2914  
  🟩 90;90a;100         Pass: 100%/1   | Total:  4m 47s | Avg:  4m 47s | Max:  4m 47s | Hits:  98%/2914  
🟩 std
  🟩 17                 Pass: 100%/21  | Total:  3h 18m | Avg:  9m 28s | Max: 25m 41s | Hits:  90%/54907 
  🟩 20                 Pass: 100%/21  | Total:  2h 48m | Avg:  8m 01s | Max: 24m 18s | Hits:  94%/48074

🟩 cudax: Pass: 100%/22 | Total: 2h 11m | Avg: 5m 58s | Max: 13m 50s | Hits: 96%/11264

🟩 cpu
  🟩 amd64              Pass: 100%/18  | Total:  1h 56m | Avg:  6m 28s | Max: 13m 50s | Hits:  96%/9036  
  🟩 arm64              Pass: 100%/4   | Total: 15m 08s | Avg:  3m 47s | Max:  3m 58s | Hits:  98%/2228  
🟩 ctk
  🟩 12.0               Pass: 100%/1   | Total:  9m 34s | Avg:  9m 34s | Max:  9m 34s | Hits:  61%/262   
  🟩 12.5               Pass: 100%/2   | Total: 14m 36s | Avg:  7m 18s | Max:  7m 21s | Hits:  95%/710   
  🟩 12.8               Pass: 100%/19  | Total:  1h 47m | Avg:  5m 39s | Max: 13m 50s | Hits:  97%/10292 
🟩 cudacxx
  🟩 nvcc12.0           Pass: 100%/1   | Total:  9m 34s | Avg:  9m 34s | Max:  9m 34s | Hits:  61%/262   
  🟩 nvcc12.5           Pass: 100%/2   | Total: 14m 36s | Avg:  7m 18s | Max:  7m 21s | Hits:  95%/710   
  🟩 nvcc12.8           Pass: 100%/19  | Total:  1h 47m | Avg:  5m 39s | Max: 13m 50s | Hits:  97%/10292 
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/22  | Total:  2h 11m | Avg:  5m 58s | Max: 13m 50s | Hits:  96%/11264 
🟩 cxx
  🟩 Clang14            Pass: 100%/1   | Total:  4m 05s | Avg:  4m 05s | Max:  4m 05s | Hits:  98%/559   
  🟩 Clang15            Pass: 100%/1   | Total:  4m 20s | Avg:  4m 20s | Max:  4m 20s | Hits:  98%/557   
  🟩 Clang16            Pass: 100%/1   | Total:  4m 07s | Avg:  4m 07s | Max:  4m 07s | Hits:  98%/557   
  🟩 Clang17            Pass: 100%/1   | Total:  4m 20s | Avg:  4m 20s | Max:  4m 20s | Hits:  98%/557   
  🟩 Clang18            Pass: 100%/4   | Total: 23m 59s | Avg:  5m 59s | Max: 12m 35s | Hits:  98%/2228  
  🟩 GCC10              Pass: 100%/1   | Total:  4m 01s | Avg:  4m 01s | Max:  4m 01s | Hits:  98%/559   
  🟩 GCC11              Pass: 100%/1   | Total:  4m 11s | Avg:  4m 11s | Max:  4m 11s | Hits:  98%/557   
  🟩 GCC12              Pass: 100%/2   | Total: 16m 26s | Avg:  8m 13s | Max: 12m 10s | Hits:  98%/1114  
  🟩 GCC13              Pass: 100%/6   | Total: 32m 38s | Avg:  5m 26s | Max: 13m 50s | Hits:  98%/3342  
  🟩 MSVC14.39          Pass: 100%/1   | Total:  9m 34s | Avg:  9m 34s | Max:  9m 34s | Hits:  61%/262   
  🟩 MSVC14.42          Pass: 100%/1   | Total:  9m 18s | Avg:  9m 18s | Max:  9m 18s | Hits:  61%/262   
  🟩 NVHPC24.7          Pass: 100%/2   | Total: 14m 36s | Avg:  7m 18s | Max:  7m 21s | Hits:  95%/710   
🟩 cxx_family
  🟩 Clang              Pass: 100%/8   | Total: 40m 51s | Avg:  5m 06s | Max: 12m 35s | Hits:  98%/4458  
  🟩 GCC                Pass: 100%/10  | Total: 57m 16s | Avg:  5m 43s | Max: 13m 50s | Hits:  98%/5572  
  🟩 MSVC               Pass: 100%/2   | Total: 18m 52s | Avg:  9m 26s | Max:  9m 34s | Hits:  61%/524   
  🟩 NVHPC              Pass: 100%/2   | Total: 14m 36s | Avg:  7m 18s | Max:  7m 21s | Hits:  95%/710   
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 17m 35s | Avg:  8m 47s | Max: 13m 50s | Hits:  98%/1114  
  🟩 rtx2080            Pass: 100%/20  | Total:  1h 54m | Avg:  5m 42s | Max: 12m 35s | Hits:  96%/10150 
🟩 jobs
  🟩 Build              Pass: 100%/19  | Total:  1h 33m | Avg:  4m 53s | Max:  9m 34s | Hits:  96%/9593  
  🟩 Test               Pass: 100%/3   | Total: 38m 35s | Avg: 12m 51s | Max: 13m 50s | Hits:  99%/1671  
🟩 sm
  🟩 90                 Pass: 100%/3   | Total: 21m 15s | Avg:  7m 05s | Max: 13m 50s | Hits:  98%/1671  
  🟩 90a                Pass: 100%/1   | Total:  3m 31s | Avg:  3m 31s | Max:  3m 31s | Hits:  98%/557   
🟩 std
  🟩 17                 Pass: 100%/4   | Total: 18m 33s | Avg:  4m 38s | Max:  7m 15s | Hits:  97%/2026  
  🟩 20                 Pass: 100%/18  | Total:  1h 53m | Avg:  6m 16s | Max: 13m 50s | Hits:  96%/9238

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 13m 28s | Avg: 6m 44s | Max: 10m 58s | Hits: 97%/296

🟩 cpu
  🟩 amd64              Pass: 100%/2   | Total: 13m 28s | Avg:  6m 44s | Max: 10m 58s | Hits:  97%/296   
🟩 ctk
  🟩 12.8               Pass: 100%/2   | Total: 13m 28s | Avg:  6m 44s | Max: 10m 58s | Hits:  97%/296   
🟩 cudacxx
  🟩 nvcc12.8           Pass: 100%/2   | Total: 13m 28s | Avg:  6m 44s | Max: 10m 58s | Hits:  97%/296   
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/2   | Total: 13m 28s | Avg:  6m 44s | Max: 10m 58s | Hits:  97%/296   
🟩 cxx
  🟩 GCC13              Pass: 100%/2   | Total: 13m 28s | Avg:  6m 44s | Max: 10m 58s | Hits:  97%/296   
🟩 cxx_family
  🟩 GCC                Pass: 100%/2   | Total: 13m 28s | Avg:  6m 44s | Max: 10m 58s | Hits:  97%/296   
🟩 gpu
  🟩 rtx2080            Pass: 100%/2   | Total: 13m 28s | Avg:  6m 44s | Max: 10m 58s | Hits:  97%/296   
🟩 jobs
  🟩 Build              Pass: 100%/1   | Total:  2m 30s | Avg:  2m 30s | Max:  2m 30s | Hits:  97%/148   
  🟩 Test               Pass: 100%/1   | Total: 10m 58s | Avg: 10m 58s | Max: 10m 58s | Hits:  98%/148

🟩 python: Pass: 100%/1 | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s

🟩 cpu
  🟩 amd64              Pass: 100%/1   | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s
🟩 ctk
  🟩 12.8               Pass: 100%/1   | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s
🟩 cudacxx
  🟩 nvcc12.8           Pass: 100%/1   | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/1   | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s
🟩 cxx
  🟩 GCC13              Pass: 100%/1   | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s
🟩 cxx_family
  🟩 GCC                Pass: 100%/1   | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s
🟩 gpu
  🟩 rtx2080            Pass: 100%/1   | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s
🟩 jobs
  🟩 Test               Pass: 100%/1   | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s

👃 Inspect Changes

Modifications in project?

	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 · 2025-02-20T09:19:41Z

/ok to test

miscco · 2025-02-20T09:52:33Z

/ok to test

github-actions · 2025-02-20T10:59:43Z

🟩 CI finished in 1h 04m: Pass: 100%/158 | Total: 23h 42m | Avg: 9m 00s | Max: 31m 09s | Hits: 94%/247972

🟩 cub: Pass: 100%/45 | Total: 8h 16m | Avg: 11m 02s | Max: 31m 09s | Hits: 93%/53221

🟩 cpu
  🟩 amd64              Pass: 100%/43  | Total:  8h 05m | Avg: 11m 17s | Max: 31m 09s | Hits:  92%/50803 
  🟩 arm64              Pass: 100%/2   | Total: 11m 00s | Avg:  5m 30s | Max:  5m 47s | Hits:  99%/2418  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total: 50m 13s | Avg: 10m 02s | Max: 27m 53s | Hits:  85%/5879  
  🟩 12.5               Pass: 100%/2   | Total: 19m 22s | Avg:  9m 41s | Max:  9m 51s | Hits:  98%/2236  
  🟩 12.8               Pass: 100%/38  | Total:  7h 07m | Avg: 11m 14s | Max: 31m 09s | Hits:  94%/45106 
🟩 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total: 10m 04s | Avg:  5m 02s | Max:  5m 05s | Hits: 100%/2090  
  🟩 nvcc12.0           Pass: 100%/5   | Total: 50m 13s | Avg: 10m 02s | Max: 27m 53s | Hits:  85%/5879  
  🟩 nvcc12.5           Pass: 100%/2   | Total: 19m 22s | Avg:  9m 41s | Max:  9m 51s | Hits:  98%/2236  
  🟩 nvcc12.8           Pass: 100%/36  | Total:  6h 57m | Avg: 11m 35s | Max: 31m 09s | Hits:  93%/43016 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total: 10m 04s | Avg:  5m 02s | Max:  5m 05s | Hits: 100%/2090  
  🟩 nvcc               Pass: 100%/43  | Total:  8h 06m | Avg: 11m 19s | Max: 31m 09s | Hits:  92%/51131 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total: 22m 23s | Avg:  5m 35s | Max:  5m 55s | Hits: 100%/4844  
  🟩 Clang15            Pass: 100%/2   | Total: 12m 55s | Avg:  6m 27s | Max:  6m 36s | Hits: 100%/2418  
  🟩 Clang16            Pass: 100%/2   | Total: 12m 21s | Avg:  6m 10s | Max:  6m 29s | Hits: 100%/2418  
  🟩 Clang17            Pass: 100%/2   | Total: 12m 01s | Avg:  6m 00s | Max:  6m 03s | Hits: 100%/2418  
  🟩 Clang18            Pass: 100%/7   | Total:  1h 12m | Avg: 10m 22s | Max: 23m 59s | Hits: 100%/8135  
  🟩 GCC7               Pass: 100%/2   | Total: 12m 04s | Avg:  6m 02s | Max:  6m 12s | Hits:  99%/2422  
  🟩 GCC8               Pass: 100%/1   | Total:  6m 29s | Avg:  6m 29s | Max:  6m 29s | Hits:  99%/1211  
  🟩 GCC9               Pass: 100%/2   | Total: 11m 44s | Avg:  5m 52s | Max:  5m 58s | Hits:  99%/2422  
  🟩 GCC10              Pass: 100%/2   | Total: 13m 34s | Avg:  6m 47s | Max:  6m 50s | Hits:  99%/2422  
  🟩 GCC11              Pass: 100%/2   | Total: 13m 10s | Avg:  6m 35s | Max:  6m 46s | Hits:  99%/2418  
  🟩 GCC12              Pass: 100%/2   | Total: 13m 07s | Avg:  6m 33s | Max:  6m 42s | Hits:  99%/2418  
  🟩 GCC13              Pass: 100%/11  | Total:  2h 40m | Avg: 14m 37s | Max: 25m 15s | Hits:  99%/13299 
  🟩 MSVC14.29          Pass: 100%/2   | Total: 55m 05s | Avg: 27m 32s | Max: 27m 53s | Hits:  16%/2070  
  🟩 MSVC14.42          Pass: 100%/2   | Total: 59m 02s | Avg: 29m 31s | Max: 31m 09s | Hits:  16%/2070  
  🟩 NVHPC24.7          Pass: 100%/2   | Total: 19m 22s | Avg:  9m 41s | Max:  9m 51s | Hits:  98%/2236  
🟩 cxx_family
  🟩 Clang              Pass: 100%/17  | Total:  2h 12m | Avg:  7m 46s | Max: 23m 59s | Hits: 100%/20233 
  🟩 GCC                Pass: 100%/22  | Total:  3h 51m | Avg: 10m 30s | Max: 25m 15s | Hits:  99%/26612 
  🟩 MSVC               Pass: 100%/4   | Total:  1h 54m | Avg: 28m 31s | Max: 31m 09s | Hits:  16%/4140  
  🟩 NVHPC              Pass: 100%/2   | Total: 19m 22s | Avg:  9m 41s | Max:  9m 51s | Hits:  98%/2236  
🟩 gpu
  🟩 h100               Pass: 100%/3   | Total: 53m 20s | Avg: 17m 46s | Max: 25m 15s | Hits:  99%/3627  
  🟩 rtx2080            Pass: 100%/34  | Total:  5h 03m | Avg:  8m 55s | Max: 31m 09s | Hits:  91%/39922 
  🟩 rtxa6000           Pass: 100%/8   | Total:  2h 19m | Avg: 17m 28s | Max: 24m 08s | Hits:  99%/9672  
🟩 jobs
  🟩 Build              Pass: 100%/37  | Total:  5h 21m | Avg:  8m 41s | Max: 31m 09s | Hits:  91%/43549 
  🟩 DeviceLaunch       Pass: 100%/1   | Total: 20m 51s | Avg: 20m 51s | Max: 20m 51s | Hits:  99%/1209  
  🟩 GraphCapture       Pass: 100%/1   | Total: 16m 08s | Avg: 16m 08s | Max: 16m 08s | Hits:  99%/1209  
  🟩 HostLaunch         Pass: 100%/3   | Total:  1h 13m | Avg: 24m 27s | Max: 25m 15s | Hits:  99%/3627  
  🟩 TestGPU            Pass: 100%/3   | Total:  1h 04m | Avg: 21m 35s | Max: 23m 16s | Hits:  99%/3627  
🟩 sm
  🟩 90                 Pass: 100%/3   | Total: 53m 20s | Avg: 17m 46s | Max: 25m 15s | Hits:  99%/3627  
  🟩 90;90a;100         Pass: 100%/1   | Total:  6m 54s | Avg:  6m 54s | Max:  6m 54s | Hits:  99%/1209  
🟩 std
  🟩 17                 Pass: 100%/20  | Total:  3h 10m | Avg:  9m 30s | Max: 27m 53s | Hits:  88%/23419 
  🟩 20                 Pass: 100%/25  | Total:  5h 06m | Avg: 12m 16s | Max: 31m 09s | Hits:  96%/29802

🟩 thrust: Pass: 100%/45 | Total: 6h 25m | Avg: 8m 34s | Max: 29m 22s | Hits: 96%/80136

🟩 cmake_options
  🟩 -DTHRUST_DISPATCH_TYPE=Force32bit Pass: 100%/2   | Total: 16m 50s | Avg:  8m 25s | Max: 11m 01s | Hits:  99%/3564  
🟩 cpu
  🟩 amd64              Pass: 100%/43  | Total:  6h 16m | Avg:  8m 44s | Max: 29m 22s | Hits:  96%/76573 
  🟩 arm64              Pass: 100%/2   | Total:  9m 29s | Avg:  4m 44s | Max:  5m 04s | Hits:  99%/3563  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total: 43m 23s | Avg:  8m 40s | Max: 24m 02s | Hits:  94%/8901  
  🟩 12.5               Pass: 100%/2   | Total: 26m 49s | Avg: 13m 24s | Max: 13m 34s | Hits:  99%/3562  
  🟩 12.8               Pass: 100%/38  | Total:  5h 15m | Avg:  8m 18s | Max: 29m 22s | Hits:  96%/67673 
🟩 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total: 10m 10s | Avg:  5m 05s | Max:  5m 11s | Hits: 100%/3562  
  🟩 nvcc12.0           Pass: 100%/5   | Total: 43m 23s | Avg:  8m 40s | Max: 24m 02s | Hits:  94%/8901  
  🟩 nvcc12.5           Pass: 100%/2   | Total: 26m 49s | Avg: 13m 24s | Max: 13m 34s | Hits:  99%/3562  
  🟩 nvcc12.8           Pass: 100%/36  | Total:  5h 05m | Avg:  8m 28s | Max: 29m 22s | Hits:  96%/64111 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total: 10m 10s | Avg:  5m 05s | Max:  5m 11s | Hits: 100%/3562  
  🟩 nvcc               Pass: 100%/43  | Total:  6h 15m | Avg:  8m 44s | Max: 29m 22s | Hits:  96%/76574 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total: 20m 20s | Avg:  5m 05s | Max:  5m 45s | Hits: 100%/7124  
  🟩 Clang15            Pass: 100%/2   | Total: 10m 50s | Avg:  5m 25s | Max:  5m 35s | Hits: 100%/3562  
  🟩 Clang16            Pass: 100%/2   | Total: 10m 50s | Avg:  5m 25s | Max:  5m 27s | Hits: 100%/3562  
  🟩 Clang17            Pass: 100%/2   | Total: 11m 29s | Avg:  5m 44s | Max:  5m 49s | Hits: 100%/3562  
  🟩 Clang18            Pass: 100%/7   | Total: 43m 05s | Avg:  6m 09s | Max: 10m 20s | Hits:  99%/12467 
  🟩 GCC7               Pass: 100%/2   | Total: 10m 12s | Avg:  5m 06s | Max:  5m 17s | Hits:  99%/3564  
  🟩 GCC8               Pass: 100%/1   | Total:  5m 18s | Avg:  5m 18s | Max:  5m 18s | Hits:  99%/1782  
  🟩 GCC9               Pass: 100%/2   | Total: 10m 33s | Avg:  5m 16s | Max:  5m 33s | Hits:  99%/3564  
  🟩 GCC10              Pass: 100%/2   | Total: 10m 55s | Avg:  5m 27s | Max:  5m 37s | Hits:  99%/3564  
  🟩 GCC11              Pass: 100%/2   | Total: 11m 31s | Avg:  5m 45s | Max:  5m 55s | Hits:  99%/3564  
  🟩 GCC12              Pass: 100%/2   | Total: 11m 40s | Avg:  5m 50s | Max:  5m 56s | Hits:  99%/3564  
  🟩 GCC13              Pass: 100%/10  | Total:  1h 16m | Avg:  7m 39s | Max: 12m 02s | Hits:  99%/17820 
  🟩 MSVC14.29          Pass: 100%/2   | Total: 47m 19s | Avg: 23m 39s | Max: 24m 02s | Hits:  70%/3550  
  🟩 MSVC14.42          Pass: 100%/3   | Total:  1h 18m | Avg: 26m 06s | Max: 29m 22s | Hits:  70%/5325  
  🟩 NVHPC24.7          Pass: 100%/2   | Total: 26m 49s | Avg: 13m 24s | Max: 13m 34s | Hits:  99%/3562  
🟩 cxx_family
  🟩 Clang              Pass: 100%/17  | Total:  1h 36m | Avg:  5m 40s | Max: 10m 20s | Hits:  99%/30277 
  🟩 GCC                Pass: 100%/21  | Total:  2h 16m | Avg:  6m 30s | Max: 12m 02s | Hits:  99%/37422 
  🟩 MSVC               Pass: 100%/5   | Total:  2h 05m | Avg: 25m 07s | Max: 29m 22s | Hits:  70%/8875  
  🟩 NVHPC              Pass: 100%/2   | Total: 26m 49s | Avg: 13m 24s | Max: 13m 34s | Hits:  99%/3562  
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 17m 03s | Avg:  8m 31s | Max: 12m 02s | Hits:  99%/3564  
  🟩 rtx2080            Pass: 100%/33  | Total:  4h 10m | Avg:  7m 34s | Max: 24m 43s | Hits:  97%/58769 
  🟩 rtx4090            Pass: 100%/10  | Total:  1h 58m | Avg: 11m 51s | Max: 29m 22s | Hits:  94%/17803 
🟩 jobs
  🟩 Build              Pass: 100%/38  | Total:  4h 56m | Avg:  7m 48s | Max: 24m 43s | Hits:  96%/67671 
  🟩 TestCPU            Pass: 100%/3   | Total: 44m 23s | Avg: 14m 47s | Max: 29m 22s | Hits:  90%/5338  
  🟩 TestGPU            Pass: 100%/4   | Total: 44m 41s | Avg: 11m 10s | Max: 12m 02s | Hits:  99%/7127  
🟩 sm
  🟩 90                 Pass: 100%/2   | Total: 17m 03s | Avg:  8m 31s | Max: 12m 02s | Hits:  99%/3564  
  🟩 90;90a;100         Pass: 100%/1   | Total:  6m 20s | Avg:  6m 20s | Max:  6m 20s | Hits:  99%/1782  
🟩 std
  🟩 17                 Pass: 100%/20  | Total:  2h 52m | Avg:  8m 36s | Max: 24m 43s | Hits:  95%/35611 
  🟩 20                 Pass: 100%/23  | Total:  3h 16m | Avg:  8m 33s | Max: 29m 22s | Hits:  97%/40961

🟩 libcudacxx: Pass: 100%/43 | Total: 6h 18m | Avg: 8m 48s | Max: 26m 26s | Hits: 92%/103055

🟩 cpu
  🟩 amd64              Pass: 100%/41  | Total:  6h 11m | Avg:  9m 03s | Max: 26m 26s | Hits:  91%/97402 
  🟩 arm64              Pass: 100%/2   | Total:  7m 22s | Avg:  3m 41s | Max:  3m 45s | Hits:  99%/5653  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total: 41m 07s | Avg:  8m 13s | Max: 26m 26s | Hits:  99%/13672 
  🟩 12.5               Pass: 100%/2   | Total: 18m 05s | Avg:  9m 02s | Max:  9m 06s | Hits:  98%/5598  
  🟩 12.8               Pass: 100%/36  | Total:  5h 19m | Avg:  8m 52s | Max: 25m 59s | Hits:  90%/83785 
🟩 cudacxx
  🟩 ClangCUDA18        Pass: 100%/2   | Total: 43m 50s | Avg: 21m 55s | Max: 22m 41s | Hits:  26%/5616  
  🟩 nvcc12.0           Pass: 100%/5   | Total: 41m 07s | Avg:  8m 13s | Max: 26m 26s | Hits:  99%/13672 
  🟩 nvcc12.5           Pass: 100%/2   | Total: 18m 05s | Avg:  9m 02s | Max:  9m 06s | Hits:  98%/5598  
  🟩 nvcc12.8           Pass: 100%/34  | Total:  4h 35m | Avg:  8m 06s | Max: 25m 59s | Hits:  95%/78169 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total: 43m 50s | Avg: 21m 55s | Max: 22m 41s | Hits:  26%/5616  
  🟩 nvcc               Pass: 100%/41  | Total:  5h 34m | Avg:  8m 10s | Max: 26m 26s | Hits:  96%/97439 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total: 16m 28s | Avg:  4m 07s | Max:  4m 35s | Hits:  99%/11200 
  🟩 Clang15            Pass: 100%/2   | Total:  9m 09s | Avg:  4m 34s | Max:  4m 41s | Hits:  99%/5610  
  🟩 Clang16            Pass: 100%/2   | Total:  9m 14s | Avg:  4m 37s | Max:  4m 51s | Hits:  99%/5610  
  🟩 Clang17            Pass: 100%/2   | Total:  9m 29s | Avg:  4m 44s | Max:  4m 48s | Hits:  99%/5610  
  🟩 Clang18            Pass: 100%/6   | Total:  1h 05m | Avg: 10m 50s | Max: 22m 41s | Hits:  70%/14052 
  🟩 GCC7               Pass: 100%/2   | Total:  7m 31s | Avg:  3m 45s | Max:  3m 57s | Hits:  99%/5548  
  🟩 GCC8               Pass: 100%/1   | Total:  3m 43s | Avg:  3m 43s | Max:  3m 43s | Hits:  99%/2784  
  🟩 GCC9               Pass: 100%/2   | Total:  7m 33s | Avg:  3m 46s | Max:  4m 06s | Hits:  99%/5560  
  🟩 GCC10              Pass: 100%/2   | Total: 31m 38s | Avg: 15m 49s | Max: 21m 25s | Hits:  55%/5616  
  🟩 GCC11              Pass: 100%/2   | Total: 13m 21s | Avg:  6m 40s | Max:  9m 07s | Hits:  88%/5612  
  🟩 GCC12              Pass: 100%/2   | Total:  8m 10s | Avg:  4m 05s | Max:  4m 08s | Hits:  99%/5612  
  🟩 GCC13              Pass: 100%/10  | Total:  1h 17m | Avg:  7m 45s | Max: 16m 52s | Hits:  99%/14311 
  🟩 MSVC14.29          Pass: 100%/2   | Total: 52m 25s | Avg: 26m 12s | Max: 26m 26s | Hits:  99%/5086  
  🟩 MSVC14.42          Pass: 100%/2   | Total: 49m 18s | Avg: 24m 39s | Max: 24m 44s | Hits:  98%/5246  
  🟩 NVHPC24.7          Pass: 100%/2   | Total: 18m 05s | Avg:  9m 02s | Max:  9m 06s | Hits:  98%/5598  
🟩 cxx_family
  🟩 Clang              Pass: 100%/16  | Total:  1h 49m | Avg:  6m 50s | Max: 22m 41s | Hits:  89%/42082 
  🟩 GCC                Pass: 100%/21  | Total:  2h 29m | Avg:  7m 07s | Max: 21m 25s | Hits:  92%/45043 
  🟩 MSVC               Pass: 100%/4   | Total:  1h 41m | Avg: 25m 25s | Max: 26m 26s | Hits:  98%/10332 
  🟩 NVHPC              Pass: 100%/2   | Total: 18m 05s | Avg:  9m 02s | Max:  9m 06s | Hits:  98%/5598  
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 17m 32s | Avg:  8m 46s | Max: 13m 39s | Hits:  99%/2916  
  🟩 rtx2080            Pass: 100%/41  | Total:  6h 01m | Avg:  8m 48s | Max: 26m 26s | Hits:  92%/100139
🟩 jobs
  🟩 Build              Pass: 100%/37  | Total:  5h 13m | Avg:  8m 28s | Max: 26m 26s | Hits:  92%/103015
  🟩 NVRTC              Pass: 100%/2   | Total: 31m 50s | Avg: 15m 55s | Max: 16m 52s | Hits:  90%/40    
  🟩 Test               Pass: 100%/3   | Total: 31m 22s | Avg: 10m 27s | Max: 13m 39s
  🟩 VerifyCodegen      Pass: 100%/1   | Total:  2m 08s | Avg:  2m 08s | Max:  2m 08s
🟩 sm
  🟩 75                 Pass: 100%/2   | Total: 31m 50s | Avg: 15m 55s | Max: 16m 52s | Hits:  90%/40    
  🟩 90                 Pass: 100%/2   | Total: 17m 32s | Avg:  8m 46s | Max: 13m 39s | Hits:  99%/2916  
  🟩 90;90a;100         Pass: 100%/1   | Total:  4m 31s | Avg:  4m 31s | Max:  4m 31s | Hits:  99%/2916  
🟩 std
  🟩 17                 Pass: 100%/21  | Total:  3h 26m | Avg:  9m 49s | Max: 26m 26s | Hits:  91%/54947 
  🟩 20                 Pass: 100%/21  | Total:  2h 50m | Avg:  8m 06s | Max: 24m 34s | Hits:  93%/48108

🟩 cudax: Pass: 100%/22 | Total: 1h 56m | Avg: 5m 18s | Max: 13m 33s | Hits: 97%/11264

🟩 cpu
  🟩 amd64              Pass: 100%/18  | Total:  1h 45m | Avg:  5m 51s | Max: 13m 33s | Hits:  97%/9036  
  🟩 arm64              Pass: 100%/4   | Total: 11m 11s | Avg:  2m 47s | Max:  2m 53s | Hits:  99%/2228  
🟩 ctk
  🟩 12.0               Pass: 100%/1   | Total: 10m 44s | Avg: 10m 44s | Max: 10m 44s | Hits:  61%/262   
  🟩 12.5               Pass: 100%/2   | Total: 10m 39s | Avg:  5m 19s | Max:  5m 26s | Hits:  96%/710   
  🟩 12.8               Pass: 100%/19  | Total:  1h 35m | Avg:  5m 00s | Max: 13m 33s | Hits:  98%/10292 
🟩 cudacxx
  🟩 nvcc12.0           Pass: 100%/1   | Total: 10m 44s | Avg: 10m 44s | Max: 10m 44s | Hits:  61%/262   
  🟩 nvcc12.5           Pass: 100%/2   | Total: 10m 39s | Avg:  5m 19s | Max:  5m 26s | Hits:  96%/710   
  🟩 nvcc12.8           Pass: 100%/19  | Total:  1h 35m | Avg:  5m 00s | Max: 13m 33s | Hits:  98%/10292 
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/22  | Total:  1h 56m | Avg:  5m 18s | Max: 13m 33s | Hits:  97%/11264 
🟩 cxx
  🟩 Clang14            Pass: 100%/1   | Total:  3m 17s | Avg:  3m 17s | Max:  3m 17s | Hits: 100%/559   
  🟩 Clang15            Pass: 100%/1   | Total:  3m 15s | Avg:  3m 15s | Max:  3m 15s | Hits: 100%/557   
  🟩 Clang16            Pass: 100%/1   | Total:  3m 15s | Avg:  3m 15s | Max:  3m 15s | Hits: 100%/557   
  🟩 Clang17            Pass: 100%/1   | Total:  3m 37s | Avg:  3m 37s | Max:  3m 37s | Hits: 100%/557   
  🟩 Clang18            Pass: 100%/4   | Total: 21m 10s | Avg:  5m 17s | Max: 12m 24s | Hits: 100%/2228  
  🟩 GCC10              Pass: 100%/1   | Total:  3m 11s | Avg:  3m 11s | Max:  3m 11s | Hits:  99%/559   
  🟩 GCC11              Pass: 100%/1   | Total:  3m 14s | Avg:  3m 14s | Max:  3m 14s | Hits:  99%/557   
  🟩 GCC12              Pass: 100%/2   | Total: 17m 02s | Avg:  8m 31s | Max: 13m 31s | Hits:  99%/1114  
  🟩 GCC13              Pass: 100%/6   | Total: 28m 24s | Avg:  4m 44s | Max: 13m 33s | Hits:  99%/3342  
  🟩 MSVC14.39          Pass: 100%/1   | Total: 10m 44s | Avg: 10m 44s | Max: 10m 44s | Hits:  61%/262   
  🟩 MSVC14.42          Pass: 100%/1   | Total:  8m 53s | Avg:  8m 53s | Max:  8m 53s | Hits:  61%/262   
  🟩 NVHPC24.7          Pass: 100%/2   | Total: 10m 39s | Avg:  5m 19s | Max:  5m 26s | Hits:  96%/710   
🟩 cxx_family
  🟩 Clang              Pass: 100%/8   | Total: 34m 34s | Avg:  4m 19s | Max: 12m 24s | Hits: 100%/4458  
  🟩 GCC                Pass: 100%/10  | Total: 51m 51s | Avg:  5m 11s | Max: 13m 33s | Hits:  99%/5572  
  🟩 MSVC               Pass: 100%/2   | Total: 19m 37s | Avg:  9m 48s | Max: 10m 44s | Hits:  61%/524   
  🟩 NVHPC              Pass: 100%/2   | Total: 10m 39s | Avg:  5m 19s | Max:  5m 26s | Hits:  96%/710   
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 16m 29s | Avg:  8m 14s | Max: 13m 33s | Hits:  99%/1114  
  🟩 rtx2080            Pass: 100%/20  | Total:  1h 40m | Avg:  5m 00s | Max: 13m 31s | Hits:  97%/10150 
🟩 jobs
  🟩 Build              Pass: 100%/19  | Total:  1h 17m | Avg:  4m 03s | Max: 10m 44s | Hits:  97%/9593  
  🟩 Test               Pass: 100%/3   | Total: 39m 28s | Avg: 13m 09s | Max: 13m 33s | Hits:  99%/1671  
🟩 sm
  🟩 90                 Pass: 100%/3   | Total: 19m 31s | Avg:  6m 30s | Max: 13m 33s | Hits:  99%/1671  
  🟩 90a                Pass: 100%/1   | Total:  3m 11s | Avg:  3m 11s | Max:  3m 11s | Hits:  99%/557   
🟩 std
  🟩 17                 Pass: 100%/4   | Total: 13m 58s | Avg:  3m 29s | Max:  5m 26s | Hits:  99%/2026  
  🟩 20                 Pass: 100%/18  | Total:  1h 42m | Avg:  5m 42s | Max: 13m 33s | Hits:  97%/9238

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 13m 18s | Avg: 6m 39s | Max: 11m 01s | Hits: 98%/296

🟩 cpu
  🟩 amd64              Pass: 100%/2   | Total: 13m 18s | Avg:  6m 39s | Max: 11m 01s | Hits:  98%/296   
🟩 ctk
  🟩 12.8               Pass: 100%/2   | Total: 13m 18s | Avg:  6m 39s | Max: 11m 01s | Hits:  98%/296   
🟩 cudacxx
  🟩 nvcc12.8           Pass: 100%/2   | Total: 13m 18s | Avg:  6m 39s | Max: 11m 01s | Hits:  98%/296   
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/2   | Total: 13m 18s | Avg:  6m 39s | Max: 11m 01s | Hits:  98%/296   
🟩 cxx
  🟩 GCC13              Pass: 100%/2   | Total: 13m 18s | Avg:  6m 39s | Max: 11m 01s | Hits:  98%/296   
🟩 cxx_family
  🟩 GCC                Pass: 100%/2   | Total: 13m 18s | Avg:  6m 39s | Max: 11m 01s | Hits:  98%/296   
🟩 gpu
  🟩 rtx2080            Pass: 100%/2   | Total: 13m 18s | Avg:  6m 39s | Max: 11m 01s | Hits:  98%/296   
🟩 jobs
  🟩 Build              Pass: 100%/1   | Total:  2m 17s | Avg:  2m 17s | Max:  2m 17s | Hits:  98%/148   
  🟩 Test               Pass: 100%/1   | Total: 11m 01s | Avg: 11m 01s | Max: 11m 01s | Hits:  98%/148

🟩 python: Pass: 100%/1 | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s

🟩 cpu
  🟩 amd64              Pass: 100%/1   | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s
🟩 ctk
  🟩 12.8               Pass: 100%/1   | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s
🟩 cudacxx
  🟩 nvcc12.8           Pass: 100%/1   | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/1   | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s
🟩 cxx
  🟩 GCC13              Pass: 100%/1   | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s
🟩 cxx_family
  🟩 GCC                Pass: 100%/1   | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s
🟩 gpu
  🟩 rtx2080            Pass: 100%/1   | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s
🟩 jobs
  🟩 Test               Pass: 100%/1   | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s

👃 Inspect Changes

Modifications in project?

	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`

…cancel) (NVIDIA#3671) * [libcudacxx] Experimental try_cancel exposure * Update documentation * Update try_cancel_blocks ABI Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com> * Update ABI of __cluster_get_dim Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com> * Use if target and provide SW fallback * Use simple license * Fix if guard * Guard for C++20 or newer * Simplify API * Add tests * Clarify C++20 support in docs * Test main function should only run in host * Rename to for_each_cancelled_block and extend docs * Fix typo in docs * [pre-commit.ci] auto code formatting * Support C++17, move to different file, improve docs * Fix two typos * Free memory in doc example * Fix typos and add suggestions * [pre-commit.ci] auto code formatting * remove dangling requires clauses * More comments; initial arrive can be relaxed * cancelled 2 cancelled for consistency with PTX * Add missing invocable include * Add missing __detail namespace closing brace * [pre-commit.ci] auto code formatting * Stabilize API for CTK * [pre-commit.ci] auto code formatting * Update docs * Update docs Co-authored-by: Bernhard Manfred Gruber <bernhardmgruber@gmail.com> * Enable tests in C++17 Co-authored-by: Bernhard Manfred Gruber <bernhardmgruber@gmail.com> * Update test Co-authored-by: Bernhard Manfred Gruber <bernhardmgruber@gmail.com> * Update test Co-authored-by: Bernhard Manfred Gruber <bernhardmgruber@gmail.com> * Run test on all silicon * Fix Bernhard suggestions * Improve docs clarity * This needs a cuda compiler * [pre-commit.ci] auto code formatting * Use int as the size type * Use `assert` in tests * Use functions from cuda::std * Reduce includes to necessary ones * Use proper license * Drop unnecessary `__detail` namespace Everything that is `__ugly` is already internal * Cleanup the test a bit more * Drop unsupported dialects * Move to `<cuda/functional>` * clusterlaunchcontrol.try_cancel requires PTX 8.7 * Add missing include * Drop superfluous header guard * Use `NV_DISPATCH_TARGET` because that is more future proof * document requirement on PTX ISA 8.7 * Add check back ^^ * Fix namespace * Add pre PTX ISA 8.7 fallback and use invoke to support function pointers * Move to `<cuda/work_stealing>` * move test file --------- Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com> Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> Co-authored-by: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>

[libcudacxx] Experimental try_cancel exposure

803e781

gonzalobg requested review from a team as code owners February 4, 2025 14:24

gonzalobg requested a review from griwes February 4, 2025 14:24

gonzalobg assigned bernhardmgruber, miscco and ahendriksen Feb 4, 2025

Update documentation

3740f0f

miscco requested changes Feb 4, 2025

View reviewed changes

gonzalobg and others added 12 commits February 4, 2025 17:01

Update try_cancel_blocks ABI

09c48f8

Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>

Update ABI of __cluster_get_dim

6d15626

Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>

Use if target and provide SW fallback

4b9f6a3

Use simple license

57dcd42

Fix if guard

aba388b

Guard for C++20 or newer

220cf29

Simplify API

d2baa62

Add tests

27a6a52

Clarify C++20 support in docs

e1f090b

Test main function should only run in host

7e99c76

Rename to for_each_cancelled_block and extend docs

0b62b46

Fix typo in docs

0202b52

[pre-commit.ci] auto code formatting

a9b46bc

bernhardmgruber self-requested a review February 17, 2025 15:42

bernhardmgruber reviewed Feb 17, 2025

View reviewed changes

libcudacxx/include/cuda/try_cancel Outdated Show resolved Hide resolved

libcudacxx/include/cuda/try_cancel Outdated Show resolved Hide resolved

gonzalobg added 2 commits February 18, 2025 06:11

Support C++17, move to different file, improve docs

dd6c971

Fix two typos

cbd4885

Add missing include

0272ffe

miscco reviewed Feb 19, 2025

View reviewed changes

miscco added 2 commits February 19, 2025 14:07

Drop superfluous header guard

7ea70fb

Use NV_DISPATCH_TARGET because that is more future proof

284d84b

document requirement on PTX ISA 8.7

8829e90

gonzalobg commented Feb 19, 2025

View reviewed changes

libcudacxx/test/libcudacxx/cuda/for_each_canceled/for_each_canceled.pass.cpp Show resolved Hide resolved

gonzalobg commented Feb 19, 2025

View reviewed changes

libcudacxx/include/cuda/for_each_canceled Outdated Show resolved Hide resolved

Add check back ^^

0d1991e

Fix namespace

da83903

Add pre PTX ISA 8.7 fallback and use invoke to support function pointers

7068da5

Move to <cuda/work_stealing>

91f3e4d

miscco approved these changes Feb 20, 2025

View reviewed changes

gonzalobg changed the title ~~[libcudacxx] Experimental try_cancel exposure~~ [libcudacxx] Stable abstraction for Blackwell work-stealing (PTX try_cancel) Feb 20, 2025

gonzalobg mentioned this pull request Feb 20, 2025

[FEA]: work-stealing improvements meta-issue #3870

Open

12 tasks

move test file

990bb85

miscco merged commit 180e6c0 into NVIDIA:main Feb 20, 2025
170 of 173 checks passed

bernhardmgruber mentioned this pull request Feb 20, 2025

[EPIC] Use work stealing in all relevant CUB algorithms #3871

Open

18 tasks

Conversation

gonzalobg commented Feb 4, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Checklist

Uh oh!

copy-pr-bot bot commented Feb 4, 2025

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

gonzalobg commented Feb 7, 2025

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

bernhardmgruber commented Feb 17, 2025

Uh oh!

Uh oh!

miscco Feb 19, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

miscco commented Feb 19, 2025

Uh oh!

miscco commented Feb 19, 2025

Uh oh!

miscco commented Feb 19, 2025

Uh oh!

Uh oh!

Uh oh!

miscco commented Feb 19, 2025

Uh oh!

miscco commented Feb 19, 2025

Uh oh!

miscco commented Feb 19, 2025

Uh oh!

miscco commented Feb 19, 2025

Uh oh!

github-actions bot commented Feb 19, 2025

🟩 cub: Pass: 100%/45 | Total: 1d 17h | Avg: 55m 37s | Max: 1h 24m | Hits: 47%/53221

🟩 thrust: Pass: 100%/45 | Total: 22h 31m | Avg: 30m 02s | Max: 1h 07m | Hits: 72%/80136

🟩 libcudacxx: Pass: 100%/43 | Total: 6h 09m | Avg: 8m 35s | Max: 25m 41s | Hits: 92%/102981

🟩 cudax: Pass: 100%/22 | Total: 2h 11m | Avg: 5m 58s | Max: 13m 50s | Hits: 96%/11264

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 13m 28s | Avg: 6m 44s | Max: 10m 58s | Hits: 97%/296

🟩 python: Pass: 100%/1 | Total: 30m 15s | Avg: 30m 15s | Max: 30m 15s

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 158)

Uh oh!

miscco commented Feb 20, 2025

Uh oh!

miscco commented Feb 20, 2025

Uh oh!

github-actions bot commented Feb 20, 2025

🟩 cub: Pass: 100%/45 | Total: 8h 16m | Avg: 11m 02s | Max: 31m 09s | Hits: 93%/53221

🟩 thrust: Pass: 100%/45 | Total: 6h 25m | Avg: 8m 34s | Max: 29m 22s | Hits: 96%/80136

🟩 libcudacxx: Pass: 100%/43 | Total: 6h 18m | Avg: 8m 48s | Max: 26m 26s | Hits: 92%/103055

🟩 cudax: Pass: 100%/22 | Total: 1h 56m | Avg: 5m 18s | Max: 13m 33s | Hits: 97%/11264

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 13m 18s | Avg: 6m 39s | Max: 11m 01s | Hits: 98%/296

🟩 python: Pass: 100%/1 | Total: 30m 50s | Avg: 30m 50s | Max: 30m 50s

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

gonzalobg commented Feb 4, 2025 •

edited

Loading