[CUB] Add `cub::detail::ThreadReducePartial` by pauleonix · Pull Request #5324 · NVIDIA/cccl

pauleonix · 2025-07-21T19:24:46Z

Description

Add ThreadReducePartial which in contrast to the existing cub::ThreadReduce takes a valid_items (runtime-) parameter. This new function guarantees that the reduction operator is not called with invalid items. To be usable within *PartialTile methods of BlockScan, it also needs to guarantee to avoid commutation which should not be a problem for a single thread.

closes #5192

@elstehle an me decided to not guarantee any specific output in the valid_items == 0 case as the neutral/identity value is not known. In practice this case will for now return the first element of the input array, which is invalid. This avoids adding complexity and hopefully avoids further performance pessimization.

Checklist

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

fbusato · 2025-07-21T20:01:26Z

@elstehle an me decided to not guarantee any specific output in the valid_items == 0 case as the neutral/identity value is not known. In practice this case will for now return the first element of the input array, which is invalid. This avoids adding complexity and hopefully avoids further performance pessimization.

this looks undefined behavior. Would make sense to add an assertion to check wrong usage in debug mode?

pauleonix · 2025-07-21T23:55:53Z

@fbusato Do you mean

assert(valid_items > 0);

That would not let us use it where we want it in BlockScan. Some threads will not have any valid items. The point is that their output is not written to global memory either way. If we were planning to make these APIs public, I would be more concerned about allowing this UB. The argument is that guaranteeing that e.g. T{} is returned for valid_items < 1 can still cause bugs when the operator is multiplies or rather a custom version of multiplies that has some side effect and can't be handled by sth like identity_v. Or is the plan to let users specialize it for their custom operators?

fbusato · 2025-07-22T00:00:15Z

I meant _CCCL_ASSERT(valid_items >= 0) but I understand that there are situations where valid_items < 0 is considered valid

pauleonix · 2025-07-22T00:05:23Z

But valid_items == 0 would already trigger the issue.

I should have included it in OP, but yes we also decided to allow valid_items > num_items and valid_items < 0 being equivalent to the clamp(valid_items, 0, num_items) because mandating explicit clamping would only make code more complicated.

pauleonix · 2025-07-28T13:32:36Z

Rebased so reviewers don't have to look through the test helper file from #5300 again (cub/test/thread_reduce/catch2_test_thread_reduce_helper.cuh).

github-actions · 2025-07-30T14:24:27Z

🟩 CI finished in 1h 53m: Pass: 100%/162 | Total: 3d 20h | Avg: 34m 10s | Max: 1h 46m | Hits: 66%/152679

🟩 cub: Pass: 100%/50 | Total: 2d 03h | Avg: 1h 01m | Max: 1h 46m | Hits: 41%/51974

🟩 cpu
  🟩 amd64              Pass: 100%/48  | Total:  2d 01h | Avg:  1h 01m | Max:  1h 46m | Hits:  41%/49442 
  🟩 arm64              Pass: 100%/2   | Total:  2h 05m | Avg:  1h 02m | Max:  1h 06m | Hits:  41%/2532  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total:  6h 03m | Avg:  1h 12m | Max:  1h 33m | Hits:  41%/6226  
  🟩 12.9               Pass: 100%/45  | Total:  1d 21h | Avg:  1h 00m | Max:  1h 46m | Hits:  41%/45748 
🟩 cudacxx
  🟩 ClangCUDA19        Pass: 100%/2   | Total:  1h 14m | Avg: 37m 15s | Max: 38m 03s | Hits:  43%/2181  
  🟩 nvcc12.0           Pass: 100%/5   | Total:  6h 03m | Avg:  1h 12m | Max:  1h 33m | Hits:  41%/6226  
  🟩 nvcc12.9           Pass: 100%/43  | Total:  1d 19h | Avg:  1h 01m | Max:  1h 46m | Hits:  41%/43567 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total:  1h 14m | Avg: 37m 15s | Max: 38m 03s | Hits:  43%/2181  
  🟩 nvcc               Pass: 100%/48  | Total:  2d 02h | Avg:  1h 02m | Max:  1h 46m | Hits:  41%/49793 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total:  4h 21m | Avg:  1h 05m | Max:  1h 08m | Hits:  41%/5066  
  🟩 Clang15            Pass: 100%/2   | Total:  2h 05m | Avg:  1h 02m | Max:  1h 03m | Hits:  41%/2529  
  🟩 Clang16            Pass: 100%/2   | Total:  2h 10m | Avg:  1h 05m | Max:  1h 06m | Hits:  41%/2529  
  🟩 Clang17            Pass: 100%/2   | Total:  2h 04m | Avg:  1h 02m | Max:  1h 03m | Hits:  41%/2529  
  🟩 Clang18            Pass: 100%/2   | Total:  2h 04m | Avg:  1h 02m | Max:  1h 03m | Hits:  41%/2529  
  🟩 Clang19            Pass: 100%/7   | Total:  5h 04m | Avg: 43m 33s | Max:  1h 02m | Hits:  42%/5975  
  🟩 GCC7               Pass: 100%/2   | Total:  2h 17m | Avg:  1h 08m | Max:  1h 11m | Hits:  41%/2532  
  🟩 GCC8               Pass: 100%/1   | Total:  1h 06m | Avg:  1h 06m | Max:  1h 06m | Hits:  41%/1266  
  🟩 GCC9               Pass: 100%/2   | Total:  2h 30m | Avg:  1h 15m | Max:  1h 20m | Hits:  41%/2532  
  🟩 GCC10              Pass: 100%/2   | Total:  2h 29m | Avg:  1h 14m | Max:  1h 18m | Hits:  41%/2533  
  🟩 GCC11              Pass: 100%/2   | Total:  2h 12m | Avg:  1h 06m | Max:  1h 07m | Hits:  41%/2529  
  🟩 GCC12              Pass: 100%/2   | Total:  2h 15m | Avg:  1h 07m | Max:  1h 08m | Hits:  41%/2529  
  🟩 GCC13              Pass: 100%/12  | Total:  8h 15m | Avg: 41m 17s | Max:  1h 16m | Hits:  41%/7599  
  🟩 MSVC14.29          Pass: 100%/2   | Total:  3h 16m | Avg:  1h 38m | Max:  1h 43m | Hits:  39%/2322  
  🟩 MSVC14.43          Pass: 100%/4   | Total:  6h 12m | Avg:  1h 33m | Max:  1h 46m | Hits:  39%/4644  
  🟩 NVHPC25.5          Pass: 100%/2   | Total:  2h 47m | Avg:  1h 23m | Max:  1h 25m | Hits:  39%/2331  
🟩 cxx_family
  🟩 Clang              Pass: 100%/19  | Total: 17h 51m | Avg: 56m 22s | Max:  1h 08m | Hits:  41%/21157 
  🟩 GCC                Pass: 100%/23  | Total: 21h 07m | Avg: 55m 06s | Max:  1h 20m | Hits:  41%/21520 
  🟩 MSVC               Pass: 100%/6   | Total:  9h 29m | Avg:  1h 34m | Max:  1h 46m | Hits:  39%/6966  
  🟩 NVHPC              Pass: 100%/2   | Total:  2h 47m | Avg:  1h 23m | Max:  1h 25m | Hits:  39%/2331  
🟩 gpu
  🟩 h100               Pass: 100%/3   | Total:  1h 31m | Avg: 30m 36s | Max: 33m 28s | Hits:  41%/1267  
  🟩 rtx2080            Pass: 100%/39  | Total:  1d 21h | Avg:  1h 09m | Max:  1h 46m | Hits:  41%/48175 
  🟩 rtxa6000           Pass: 100%/8   | Total:  4h 29m | Avg: 33m 41s | Max:  1h 11m | Hits:  41%/2532  
🟩 jobs
  🟩 Build              Pass: 100%/42  | Total:  1d 23h | Avg:  1h 08m | Max:  1h 46m | Hits:  41%/51974 
  🟩 DeviceLaunch       Pass: 100%/1   | Total: 23m 01s | Avg: 23m 01s | Max: 23m 01s
  🟩 GraphCapture       Pass: 100%/1   | Total: 16m 06s | Avg: 16m 06s | Max: 16m 06s
  🟩 HostLaunch         Pass: 100%/3   | Total:  1h 18m | Avg: 26m 08s | Max: 29m 28s
  🟩 TestGPU            Pass: 100%/3   | Total:  1h 18m | Avg: 26m 06s | Max: 28m 52s
🟩 sm
  🟩 90                 Pass: 100%/3   | Total:  1h 31m | Avg: 30m 36s | Max: 33m 28s | Hits:  41%/1267  
  🟩 90;90a             Pass: 100%/2   | Total:  2h 15m | Avg:  1h 07m | Max:  1h 24m | Hits:  40%/2428  
  🟩 100;120            Pass: 100%/2   | Total:  2h 09m | Avg:  1h 04m | Max:  1h 20m | Hits:  40%/2428  
🟩 std
  🟩 17                 Pass: 100%/21  | Total:  1d 01h | Avg:  1h 12m | Max:  1h 43m | Hits:  41%/25978 
  🟩 20                 Pass: 100%/29  | Total:  1d 02h | Avg: 53m 53s | Max:  1h 46m | Hits:  41%/25996

🟩 thrust: Pass: 100%/50 | Total: 1d 08h | Avg: 38m 52s | Max: 1h 16m | Hits: 77%/84139

🟩 cmake_options
  🟩 -DTHRUST_DISPATCH_TYPE=Force32bit Pass: 100%/2   | Total: 40m 47s | Avg: 20m 23s | Max: 35m 24s | Hits:  78%/1914  
🟩 cpu
  🟩 amd64              Pass: 100%/48  | Total:  1d 07h | Avg: 38m 53s | Max:  1h 16m | Hits:  76%/80312 
  🟩 arm64              Pass: 100%/2   | Total:  1h 16m | Avg: 38m 09s | Max: 41m 50s | Hits:  78%/3827  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total:  3h 24m | Avg: 40m 57s | Max:  1h 05m | Hits:  76%/9560  
  🟩 12.9               Pass: 100%/45  | Total:  1d 04h | Avg: 38m 38s | Max:  1h 16m | Hits:  77%/74579 
🟩 cudacxx
  🟩 ClangCUDA19        Pass: 100%/2   | Total: 56m 55s | Avg: 28m 27s | Max: 28m 49s | Hits:  78%/3826  
  🟩 nvcc12.0           Pass: 100%/5   | Total:  3h 24m | Avg: 40m 57s | Max:  1h 05m | Hits:  76%/9560  
  🟩 nvcc12.9           Pass: 100%/43  | Total:  1d 04h | Avg: 39m 06s | Max:  1h 16m | Hits:  77%/70753 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total: 56m 55s | Avg: 28m 27s | Max: 28m 49s | Hits:  78%/3826  
  🟩 nvcc               Pass: 100%/48  | Total:  1d 07h | Avg: 39m 18s | Max:  1h 16m | Hits:  76%/80313 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total:  2h 23m | Avg: 35m 45s | Max: 41m 29s | Hits:  78%/7652  
  🟩 Clang15            Pass: 100%/2   | Total:  1h 25m | Avg: 42m 54s | Max: 45m 53s | Hits:  78%/3826  
  🟩 Clang16            Pass: 100%/2   | Total:  1h 13m | Avg: 36m 54s | Max: 37m 26s | Hits:  78%/3826  
  🟩 Clang17            Pass: 100%/2   | Total:  1h 16m | Avg: 38m 10s | Max: 39m 05s | Hits:  78%/3826  
  🟩 Clang18            Pass: 100%/2   | Total:  1h 24m | Avg: 42m 01s | Max: 45m 16s | Hits:  78%/3826  
  🟩 Clang19            Pass: 100%/7   | Total:  2h 54m | Avg: 24m 57s | Max: 38m 19s | Hits:  78%/9565  
  🟩 GCC7               Pass: 100%/2   | Total:  1h 17m | Avg: 38m 58s | Max: 42m 35s | Hits:  78%/3828  
  🟩 GCC8               Pass: 100%/1   | Total: 39m 03s | Avg: 39m 03s | Max: 39m 03s | Hits:  78%/1914  
  🟩 GCC9               Pass: 100%/2   | Total:  1h 19m | Avg: 39m 34s | Max: 41m 46s | Hits:  78%/3828  
  🟩 GCC10              Pass: 100%/2   | Total:  1h 25m | Avg: 42m 34s | Max: 44m 46s | Hits:  78%/3828  
  🟩 GCC11              Pass: 100%/2   | Total:  1h 27m | Avg: 43m 40s | Max: 47m 12s | Hits:  78%/3828  
  🟩 GCC12              Pass: 100%/2   | Total:  1h 29m | Avg: 44m 33s | Max: 45m 51s | Hits:  78%/3828  
  🟩 GCC13              Pass: 100%/11  | Total:  4h 26m | Avg: 24m 16s | Max: 41m 50s | Hits:  78%/13398 
  🟩 MSVC14.29          Pass: 100%/2   | Total:  2h 16m | Avg:  1h 08m | Max:  1h 10m | Hits:  67%/3812  
  🟩 MSVC14.43          Pass: 100%/5   | Total:  4h 58m | Avg: 59m 43s | Max:  1h 16m | Hits:  74%/9530  
  🟩 NVHPC25.5          Pass: 100%/2   | Total:  2h 26m | Avg:  1h 13m | Max:  1h 15m | Hits:  67%/3824  
🟩 cxx_family
  🟩 Clang              Pass: 100%/19  | Total: 10h 37m | Avg: 33m 34s | Max: 45m 53s | Hits:  78%/32521 
  🟩 GCC                Pass: 100%/22  | Total: 12h 04m | Avg: 32m 56s | Max: 47m 12s | Hits:  78%/34452 
  🟩 MSVC               Pass: 100%/7   | Total:  7h 14m | Avg:  1h 02m | Max:  1h 16m | Hits:  72%/13342 
  🟩 NVHPC              Pass: 100%/2   | Total:  2h 26m | Avg:  1h 13m | Max:  1h 15m | Hits:  67%/3824  
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 30m 03s | Avg: 15m 01s | Max: 21m 10s | Hits:  78%/1914  
  🟩 rtx2080            Pass: 100%/38  | Total:  1d 03h | Avg: 43m 59s | Max:  1h 16m | Hits:  76%/72672 
  🟩 rtx4090            Pass: 100%/10  | Total:  4h 01m | Avg: 24m 08s | Max:  1h 12m | Hits:  80%/9553  
🟩 jobs
  🟩 Build              Pass: 100%/43  | Total:  1d 07h | Avg: 43m 38s | Max:  1h 16m | Hits:  76%/82233 
  🟩 TestCPU            Pass: 100%/3   | Total: 40m 58s | Avg: 13m 39s | Max: 33m 00s | Hits:  99%/1906  
  🟩 TestGPU            Pass: 100%/4   | Total: 26m 03s | Avg:  6m 30s | Max:  8m 53s
🟩 sm
  🟩 90                 Pass: 100%/2   | Total: 30m 03s | Avg: 15m 01s | Max: 21m 10s | Hits:  78%/1914  
  🟩 90;90a             Pass: 100%/2   | Total:  1h 30m | Avg: 45m 04s | Max: 59m 54s | Hits:  73%/3820  
  🟩 100;120            Pass: 100%/2   | Total:  1h 28m | Avg: 44m 27s | Max: 56m 24s | Hits:  73%/3820  
🟩 std
  🟩 17                 Pass: 100%/21  | Total: 15h 50m | Avg: 45m 14s | Max:  1h 16m | Hits:  76%/40160 
  🟩 20                 Pass: 100%/27  | Total: 15h 52m | Avg: 35m 16s | Max:  1h 15m | Hits:  77%/42065

🟩 cudax: Pass: 100%/28 | Total: 3h 37m | Avg: 7m 46s | Max: 15m 39s | Hits: 90%/15906

🟩 cpu
  🟩 amd64              Pass: 100%/24  | Total:  3h 14m | Avg:  8m 06s | Max: 15m 39s | Hits:  90%/13462 
  🟩 arm64              Pass: 100%/4   | Total: 22m 59s | Avg:  5m 44s | Max:  6m 35s | Hits:  90%/2444  
🟩 ctk
  🟩 12.0               Pass: 100%/3   | Total: 25m 50s | Avg:  8m 36s | Max: 15m 39s | Hits:  88%/1531  
  🟩 12.9               Pass: 100%/25  | Total:  3h 11m | Avg:  7m 39s | Max: 14m 18s | Hits:  90%/14375 
🟩 cudacxx
  🟩 nvcc12.0           Pass: 100%/3   | Total: 25m 50s | Avg:  8m 36s | Max: 15m 39s | Hits:  88%/1531  
  🟩 nvcc12.9           Pass: 100%/25  | Total:  3h 11m | Avg:  7m 39s | Max: 14m 18s | Hits:  90%/14375 
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/28  | Total:  3h 37m | Avg:  7m 46s | Max: 15m 39s | Hits:  90%/15906 
🟩 cxx
  🟩 Clang14            Pass: 100%/2   | Total: 11m 17s | Avg:  5m 38s | Max:  6m 34s | Hits:  90%/1224  
  🟩 Clang15            Pass: 100%/1   | Total:  5m 48s | Avg:  5m 48s | Max:  5m 48s | Hits:  90%/611   
  🟩 Clang16            Pass: 100%/1   | Total:  5m 57s | Avg:  5m 57s | Max:  5m 57s | Hits:  90%/611   
  🟩 Clang17            Pass: 100%/1   | Total:  6m 12s | Avg:  6m 12s | Max:  6m 12s | Hits:  90%/611   
  🟩 Clang18            Pass: 100%/1   | Total:  5m 53s | Avg:  5m 53s | Max:  5m 53s | Hits:  90%/611   
  🟩 Clang19            Pass: 100%/4   | Total: 25m 10s | Avg:  6m 17s | Max:  9m 17s | Hits:  92%/2444  
  🟩 GCC10              Pass: 100%/2   | Total: 11m 56s | Avg:  5m 58s | Max:  6m 28s | Hits:  90%/1224  
  🟩 GCC11              Pass: 100%/1   | Total:  6m 32s | Avg:  6m 32s | Max:  6m 32s | Hits:  90%/611   
  🟩 GCC12              Pass: 100%/1   | Total:  6m 23s | Avg:  6m 23s | Max:  6m 23s | Hits:  90%/611   
  🟩 GCC13              Pass: 100%/8   | Total: 55m 11s | Avg:  6m 53s | Max:  9m 41s | Hits:  92%/4888  
  🟩 MSVC14.39          Pass: 100%/1   | Total: 15m 39s | Avg: 15m 39s | Max: 15m 39s | Hits:  78%/309   
  🟩 MSVC14.43          Pass: 100%/3   | Total: 40m 28s | Avg: 13m 29s | Max: 14m 18s | Hits:  78%/933   
  🟩 NVHPC25.5          Pass: 100%/2   | Total: 21m 03s | Avg: 10m 31s | Max: 10m 48s | Hits:  88%/1218  
🟩 cxx_family
  🟩 Clang              Pass: 100%/10  | Total:  1h 00m | Avg:  6m 01s | Max:  9m 17s | Hits:  91%/6112  
  🟩 GCC                Pass: 100%/12  | Total:  1h 20m | Avg:  6m 40s | Max:  9m 41s | Hits:  91%/7334  
  🟩 MSVC               Pass: 100%/4   | Total: 56m 07s | Avg: 14m 01s | Max: 15m 39s | Hits:  78%/1242  
  🟩 NVHPC              Pass: 100%/2   | Total: 21m 03s | Avg: 10m 31s | Max: 10m 48s | Hits:  88%/1218  
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 13m 26s | Avg:  6m 43s | Max:  8m 47s | Hits:  94%/1222  
  🟩 rtx2080            Pass: 100%/26  | Total:  3h 24m | Avg:  7m 50s | Max: 15m 39s | Hits:  89%/14684 
🟩 jobs
  🟩 Build              Pass: 100%/25  | Total:  3h 09m | Avg:  7m 35s | Max: 15m 39s | Hits:  89%/14073 
  🟩 Test               Pass: 100%/3   | Total: 27m 45s | Avg:  9m 15s | Max:  9m 41s | Hits:  99%/1833  
🟩 sm
  🟩 90                 Pass: 100%/2   | Total: 13m 26s | Avg:  6m 43s | Max:  8m 47s | Hits:  94%/1222  
  🟩 90;90a             Pass: 100%/2   | Total: 18m 41s | Avg:  9m 20s | Max: 12m 45s | Hits:  86%/922   
  🟩 100;120            Pass: 100%/2   | Total: 19m 38s | Avg:  9m 49s | Max: 13m 25s | Hits:  86%/922   
🟩 std
  🟩 17                 Pass: 100%/3   | Total: 22m 09s | Avg:  7m 23s | Max: 10m 48s | Hits:  89%/1831  
  🟩 20                 Pass: 100%/25  | Total:  3h 15m | Avg:  7m 48s | Max: 15m 39s | Hits:  90%/14075

🟩 python: Pass: 100%/22 | Total: 3h 35m | Avg: 9m 47s | Max: 18m 31s

🟩 cpu
  🟩 amd64              Pass: 100%/22  | Total:  3h 35m | Avg:  9m 47s | Max: 18m 31s
🟩 ctk
  🟩 12.5               Pass: 100%/6   | Total: 42m 33s | Avg:  7m 05s | Max: 13m 18s
  🟩 12.8               Pass: 100%/2   | Total: 36m 17s | Avg: 18m 08s | Max: 18m 31s
  🟩 12.9               Pass: 100%/14  | Total:  2h 16m | Avg:  9m 45s | Max: 18m 28s
🟩 cudacxx
  🟩 nvcc12.5           Pass: 100%/6   | Total: 42m 33s | Avg:  7m 05s | Max: 13m 18s
  🟩 nvcc12.8           Pass: 100%/2   | Total: 36m 17s | Avg: 18m 08s | Max: 18m 31s
  🟩 nvcc12.9           Pass: 100%/14  | Total:  2h 16m | Avg:  9m 45s | Max: 18m 28s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/22  | Total:  3h 35m | Avg:  9m 47s | Max: 18m 31s
🟩 cxx
  🟩 GCC13              Pass: 100%/22  | Total:  3h 35m | Avg:  9m 47s | Max: 18m 31s
🟩 cxx_family
  🟩 GCC                Pass: 100%/22  | Total:  3h 35m | Avg:  9m 47s | Max: 18m 31s
🟩 gpu
  🟩 h100               Pass: 100%/4   | Total: 36m 03s | Avg:  9m 00s | Max: 16m 12s
  🟩 l4                 Pass: 100%/18  | Total:  2h 59m | Avg:  9m 57s | Max: 18m 31s
🟩 jobs
  🟩 Build cuda.cccl    Pass: 100%/2   | Total: 19m 02s | Avg:  9m 31s | Max:  9m 39s
  🟩 Test cuda.cccl.cooperative Pass: 100%/5   | Total:  1h 05m | Avg: 13m 05s | Max: 13m 44s
  🟩 Test cuda.cccl.examples Pass: 100%/5   | Total: 22m 05s | Avg:  4m 25s | Max:  4m 44s
  🟩 Test cuda.cccl.headers Pass: 100%/5   | Total: 19m 29s | Avg:  3m 53s | Max:  4m 11s
  🟩 Test cuda.cccl.parallel Pass: 100%/5   | Total:  1h 29m | Avg: 17m 51s | Max: 18m 31s
🟩 py_version
  🟩 3.10               Pass: 100%/9   | Total:  1h 30m | Avg: 10m 02s | Max: 18m 31s
  🟩 3.13               Pass: 100%/13  | Total:  2h 05m | Avg:  9m 37s | Max: 18m 28s

🟩 cccl_c_parallel: Pass: 100%/4 | Total: 53m 32s | Avg: 13m 23s | Max: 20m 43s | Hits: 97%/660

🟩 cpu
  🟩 amd64              Pass: 100%/4   | Total: 53m 32s | Avg: 13m 23s | Max: 20m 43s | Hits:  97%/660   
🟩 ctk
  🟩 12.9               Pass: 100%/4   | Total: 53m 32s | Avg: 13m 23s | Max: 20m 43s | Hits:  97%/660   
🟩 cudacxx
  🟩 nvcc12.9           Pass: 100%/4   | Total: 53m 32s | Avg: 13m 23s | Max: 20m 43s | Hits:  97%/660   
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/4   | Total: 53m 32s | Avg: 13m 23s | Max: 20m 43s | Hits:  97%/660   
🟩 cxx
  🟩 GCC13              Pass: 100%/4   | Total: 53m 32s | Avg: 13m 23s | Max: 20m 43s | Hits:  97%/660   
🟩 cxx_family
  🟩 GCC                Pass: 100%/4   | Total: 53m 32s | Avg: 13m 23s | Max: 20m 43s | Hits:  97%/660   
🟩 gpu
  🟩 h100               Pass: 100%/1   | Total: 20m 43s | Avg: 20m 43s | Max: 20m 43s | Hits:  98%/165   
  🟩 l4                 Pass: 100%/1   | Total: 16m 14s | Avg: 16m 14s | Max: 16m 14s | Hits:  98%/165   
  🟩 rtx2080            Pass: 100%/2   | Total: 16m 35s | Avg:  8m 17s | Max: 13m 34s | Hits:  96%/330   
🟩 jobs
  🟩 Build              Pass: 100%/1   | Total:  3m 01s | Avg:  3m 01s | Max:  3m 01s | Hits:  93%/165   
  🟩 Test               Pass: 100%/3   | Total: 50m 31s | Avg: 16m 50s | Max: 20m 43s | Hits:  98%/495

🟩 packaging: Pass: 100%/4 | Total: 14m 45s | Avg: 3m 41s | Max: 4m 15s

🟩 cpu
  🟩 amd64              Pass: 100%/4   | Total: 14m 45s | Avg:  3m 41s | Max:  4m 15s
🟩 ctk
  🟩 12.0               Pass: 100%/2   | Total:  7m 24s | Avg:  3m 42s | Max:  4m 15s
  🟩 12.9               Pass: 100%/2   | Total:  7m 21s | Avg:  3m 40s | Max:  3m 42s
🟩 cudacxx
  🟩 nvcc12.0           Pass: 100%/2   | Total:  7m 24s | Avg:  3m 42s | Max:  4m 15s
  🟩 nvcc12.9           Pass: 100%/2   | Total:  7m 21s | Avg:  3m 40s | Max:  3m 42s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/4   | Total: 14m 45s | Avg:  3m 41s | Max:  4m 15s
🟩 cxx
  🟩 Clang14            Pass: 100%/1   | Total:  3m 09s | Avg:  3m 09s | Max:  3m 09s
  🟩 Clang19            Pass: 100%/1   | Total:  3m 42s | Avg:  3m 42s | Max:  3m 42s
  🟩 GCC12              Pass: 100%/1   | Total:  4m 15s | Avg:  4m 15s | Max:  4m 15s
  🟩 GCC13              Pass: 100%/1   | Total:  3m 39s | Avg:  3m 39s | Max:  3m 39s
🟩 cxx_family
  🟩 Clang              Pass: 100%/2   | Total:  6m 51s | Avg:  3m 25s | Max:  3m 42s
  🟩 GCC                Pass: 100%/2   | Total:  7m 54s | Avg:  3m 57s | Max:  4m 15s
🟩 gpu
  🟩 rtx2080            Pass: 100%/4   | Total: 14m 45s | Avg:  3m 41s | Max:  4m 15s
🟩 jobs
  🟩 Test               Pass: 100%/4   | Total: 14m 45s | Avg:  3m 41s | Max:  4m 15s

🟩 stdpar: Pass: 100%/4 | Total: 16m 02s | Avg: 4m 00s | Max: 4m 04s

🟩 cpu
  🟩 amd64              Pass: 100%/2   | Total:  8m 05s | Avg:  4m 02s | Max:  4m 04s
  🟩 arm64              Pass: 100%/2   | Total:  7m 57s | Avg:  3m 58s | Max:  4m 01s
🟩 ctk
  🟩 12.9               Pass: 100%/4   | Total: 16m 02s | Avg:  4m 00s | Max:  4m 04s
🟩 cudacxx
  🟩 nvcc12.9           Pass: 100%/4   | Total: 16m 02s | Avg:  4m 00s | Max:  4m 04s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/4   | Total: 16m 02s | Avg:  4m 00s | Max:  4m 04s
🟩 cxx
  🟩 NVHPC25.5          Pass: 100%/4   | Total: 16m 02s | Avg:  4m 00s | Max:  4m 04s
🟩 cxx_family
  🟩 NVHPC              Pass: 100%/4   | Total: 16m 02s | Avg:  4m 00s | Max:  4m 04s
🟩 gpu
  🟩 rtx2080            Pass: 100%/4   | Total: 16m 02s | Avg:  4m 00s | Max:  4m 04s
🟩 jobs
  🟩 Build              Pass: 100%/4   | Total: 16m 02s | Avg:  4m 00s | Max:  4m 04s
🟩 std
  🟩 17                 Pass: 100%/2   | Total:  8m 00s | Avg:  4m 00s | Max:  4m 04s
  🟩 20                 Pass: 100%/2   | Total:  8m 02s | Avg:  4m 01s | Max:  4m 01s

👃 Inspect Changes

Modifications in project?

	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: 162)

#	Runner
93	`linux-amd64-cpu16`
17	`linux-amd64-gpu-l4-latest-1`
17	`windows-amd64-cpu16`
10	`linux-arm64-cpu16`
9	`linux-amd64-gpu-h100-latest-1`
7	`linux-amd64-gpu-rtx2080-latest-1`
6	`linux-amd64-gpu-rtxa6000-latest-1`
3	`linux-amd64-gpu-rtx4090-latest-1`

* Add ThreadReducePartial and tests

pauleonix requested a review from a team as a code owner July 21, 2025 19:24

pauleonix requested a review from alliepiper July 21, 2025 19:24

github-project-automation Bot added this to CCCL Jul 21, 2025

github-project-automation Bot moved this to Todo in CCCL Jul 21, 2025

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jul 21, 2025

pauleonix requested a review from elstehle July 21, 2025 19:26

pauleonix self-assigned this Jul 21, 2025

pauleonix changed the title ~~[CUB} Add cub::detail::ThreadReducePartial~~ [CUB] Add cub::detail::ThreadReducePartial Jul 21, 2025

This comment was marked as resolved.

Sign in to view

This comment was marked as outdated.

Sign in to view

pauleonix commented Jul 22, 2025

View reviewed changes

Comment thread cub/test/thread_reduce/catch2_test_thread_reduce_partial.cu Resolved

pauleonix commented Jul 22, 2025

View reviewed changes

Comment thread cub/test/thread_reduce/catch2_test_thread_reduce_partial.cu Resolved

pauleonix commented Jul 22, 2025

View reviewed changes

Comment thread cub/test/thread_reduce/catch2_test_thread_reduce_partial.cu Resolved

pauleonix commented Jul 23, 2025

View reviewed changes

Comment thread cub/test/thread_reduce/catch2_test_thread_reduce_partial.cu Resolved

pauleonix marked this pull request as draft July 24, 2025 14:52

This comment was marked as outdated.

Sign in to view

cccl-authenticator-app Bot moved this from In Review to In Progress in CCCL Jul 24, 2025

pauleonix marked this pull request as ready for review July 28, 2025 00:53

cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Jul 28, 2025

This comment was marked as outdated.

Sign in to view

pauleonix added 3 commits July 28, 2025 15:26

Add ThreadReducePartial and tests

79fbb86

FIx MSVC warnings

f61345b

Implement Feedback from ThreadScan*Partial

570da2f

pauleonix force-pushed the thread_reduce_partial branch from 5b6c3e6 to 570da2f Compare July 28, 2025 13:30

elstehle approved these changes Jul 28, 2025

View reviewed changes

Comment thread cub/cub/thread/thread_reduce.cuh Resolved

Comment thread cub/test/thread_reduce/catch2_test_thread_reduce_partial.cu Resolved

Comment thread cub/test/thread_reduce/catch2_test_thread_reduce_partial.cu Resolved

This comment was marked as outdated.

Sign in to view

fbusato requested changes Jul 28, 2025

View reviewed changes

Comment thread cub/cub/thread/thread_reduce.cuh Resolved

Comment thread cub/cub/thread/thread_reduce.cuh Resolved

Comment thread cub/cub/thread/thread_reduce.cuh Resolved

Comment thread cub/test/thread_reduce/catch2_test_thread_reduce_partial.cu Resolved

github-project-automation Bot moved this from In Review to In Progress in CCCL Jul 28, 2025

Address some feedback

7a92078

This comment was marked as resolved.

Sign in to view

fbusato requested changes Jul 29, 2025

View reviewed changes

Comment thread cub/cub/thread/thread_reduce.cuh Resolved

Implement more feedback

2e9c287

pauleonix requested a review from fbusato July 29, 2025 23:51

This comment was marked as outdated.

Sign in to view

Replace _CUDA_VSTD according to new guidelines

adcaace

fbusato approved these changes Jul 30, 2025

View reviewed changes

github-project-automation Bot moved this from In Progress to In Review in CCCL Jul 30, 2025

pauleonix merged commit 95b7a4c into NVIDIA:main Jul 30, 2025

github-project-automation Bot moved this from In Review to Done in CCCL Jul 30, 2025

pauleonix deleted the thread_reduce_partial branch July 30, 2025 17:34

davebayer pushed a commit to davebayer/cccl that referenced this pull request Sep 23, 2025

[CUB] Add cub::detail::ThreadReducePartial (NVIDIA#5324)

d4d1cca

* Add ThreadReducePartial and tests

pauleonix mentioned this pull request Feb 16, 2026

[CUB] Implement internal PartialTile member functions for BlockScan #5433

Closed

3 tasks done

Conversation

pauleonix commented Jul 21, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Checklist

Uh oh!

fbusato commented Jul 21, 2025

Uh oh!

This comment was marked as resolved.

pauleonix commented Jul 21, 2025

Uh oh!

fbusato commented Jul 22, 2025

Uh oh!

pauleonix commented Jul 22, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

This comment was marked as outdated.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment was marked as outdated.

This comment was marked as outdated.

pauleonix commented Jul 28, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment was marked as outdated.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment was marked as resolved.

Uh oh!

This comment was marked as outdated.

github-actions Bot commented Jul 30, 2025

🟩 cub: Pass: 100%/50 | Total: 2d 03h | Avg: 1h 01m | Max: 1h 46m | Hits: 41%/51974

🟩 thrust: Pass: 100%/50 | Total: 1d 08h | Avg: 38m 52s | Max: 1h 16m | Hits: 77%/84139

🟩 cudax: Pass: 100%/28 | Total: 3h 37m | Avg: 7m 46s | Max: 15m 39s | Hits: 90%/15906

🟩 python: Pass: 100%/22 | Total: 3h 35m | Avg: 9m 47s | Max: 18m 31s

🟩 cccl_c_parallel: Pass: 100%/4 | Total: 53m 32s | Avg: 13m 23s | Max: 20m 43s | Hits: 97%/660

🟩 packaging: Pass: 100%/4 | Total: 14m 45s | Avg: 3m 41s | Max: 4m 15s

🟩 stdpar: Pass: 100%/4 | Total: 16m 02s | Avg: 4m 00s | Max: 4m 04s

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 162)

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

pauleonix commented Jul 21, 2025 •

edited

Loading

pauleonix commented Jul 22, 2025 •

edited

Loading

pauleonix commented Jul 28, 2025 •

edited

Loading