Optimize `cuda::minimum/maximum` for `float`, `double`, `half`, `nv_bfloat16`, `__float128` #5034

fbusato · 2025-06-18T22:27:51Z

Description

cuda::minimum/maximum is implemented as a < b ? a : b. On the other hand, it produces non-optimal code for device code.
This PR adds the specializations for float, double, __half, __nv_bfloat16, __float128.

I was tempted to add support for vector types but the semantic is ambiguous, e.g.
a.x < b.x ? a : (a.y < b.y ? a : b) or
T{a.x < b.x ? a.x, a.y < b.y ? a.y : b.y}

…_bfloat16/__float128

libcudacxx/test/libcudacxx/cuda/functional/maximum.pass.cpp

Co-authored-by: Bernhard Manfred Gruber <[email protected]>

github-actions · 2025-06-19T02:54:25Z

🟨 CI finished in 2h 25m: Pass: 99%/183 | Total: 2d 03h | Avg: 16m 48s | Max: 1h 40m | Hits: 83%/297220

🟨 libcudacxx: Pass: 97%/45 | Total: 12h 01m | Avg: 16m 02s | Max: 43m 50s | Hits: 73%/134079

🔍 cpu: amd64 🔍
  🔍 amd64              Pass:  97%/43  | Total: 11h 42m | Avg: 16m 19s | Max: 43m 50s | Hits:  73%/127138
  🟩 arm64              Pass: 100%/2   | Total: 19m 40s | Avg:  9m 50s | Max: 15m 12s | Hits:  88%/6941  
🔍 ctk: 12.9 🔍
  🟩 12.0               Pass: 100%/5   | Total: 53m 46s | Avg: 10m 45s | Max: 31m 24s | Hits:  84%/16988 
  🔍 12.9               Pass:  97%/40  | Total: 11h 07m | Avg: 16m 41s | Max: 43m 50s | Hits:  72%/117091
🔍 cudacxx: nvcc12.9 🔍
  🟩 ClangCUDA19        Pass: 100%/2   | Total: 50m 31s | Avg: 25m 15s | Max: 26m 38s | Hits:  26%/6905  
  🟩 nvcc12.0           Pass: 100%/5   | Total: 53m 46s | Avg: 10m 45s | Max: 31m 24s | Hits:  84%/16988 
  🔍 nvcc12.9           Pass:  97%/38  | Total: 10h 17m | Avg: 16m 14s | Max: 43m 50s | Hits:  75%/110186
🔍 cudacxx_family: nvcc 🔍
  🟩 ClangCUDA          Pass: 100%/2   | Total: 50m 31s | Avg: 25m 15s | Max: 26m 38s | Hits:  26%/6905  
  🔍 nvcc               Pass:  97%/43  | Total: 11h 11m | Avg: 15m 36s | Max: 43m 50s | Hits:  76%/127174
🔍 cxx: Clang19 🔍
  🟩 Clang14            Pass: 100%/4   | Total: 46m 43s | Avg: 11m 40s | Max: 28m 05s | Hits:  79%/13766 
  🟩 Clang15            Pass: 100%/2   | Total: 36m 39s | Avg: 18m 19s | Max: 31m 37s | Hits:  64%/6901  
  🟩 Clang16            Pass: 100%/2   | Total: 26m 35s | Avg: 13m 17s | Max: 15m 48s | Hits:  85%/6901  
  🟩 Clang17            Pass: 100%/2   | Total:  9m 59s | Avg:  4m 59s | Max:  5m 04s | Hits:  99%/6901  
  🟩 Clang18            Pass: 100%/2   | Total: 20m 53s | Avg: 10m 26s | Max: 16m 24s | Hits:  88%/6901  
  🔍 Clang19            Pass:  83%/6   | Total:  1h 21m | Avg: 13m 31s | Max: 26m 38s | Hits:  68%/17276 
  🟩 GCC7               Pass: 100%/2   | Total: 32m 48s | Avg: 16m 24s | Max: 28m 46s | Hits:  64%/6837  
  🟩 GCC8               Pass: 100%/1   | Total:  4m 20s | Avg:  4m 20s | Max:  4m 20s | Hits:  99%/3429  
  🟩 GCC9               Pass: 100%/2   | Total: 34m 36s | Avg: 17m 18s | Max: 30m 06s | Hits:  64%/6849  
  🟩 GCC10              Pass: 100%/2   | Total: 32m 19s | Avg: 16m 09s | Max: 17m 17s | Hits:  77%/6903  
  🟩 GCC11              Pass: 100%/2   | Total: 36m 46s | Avg: 18m 23s | Max: 30m 18s | Hits:  61%/6899  
  🟩 GCC12              Pass: 100%/2   | Total: 37m 56s | Avg: 18m 58s | Max: 32m 54s | Hits:  64%/6903  
  🟩 GCC13              Pass: 100%/10  | Total:  1h 53m | Avg: 11m 21s | Max: 21m 58s | Hits:  93%/17522 
  🟩 MSVC14.29          Pass: 100%/2   | Total:  1h 00m | Avg: 30m 24s | Max: 31m 24s | Hits:  63%/6575  
  🟩 MSVC14.43          Pass: 100%/2   | Total: 59m 48s | Avg: 29m 54s | Max: 31m 07s | Hits:  63%/6627  
  🟩 NVHPC25.5          Pass: 100%/2   | Total:  1h 26m | Avg: 43m 24s | Max: 43m 50s | Hits:  27%/6889  
🔍 cxx_family: Clang 🔍
  🔍 Clang              Pass:  94%/18  | Total:  3h 42m | Avg: 12m 20s | Max: 31m 37s | Hits:  78%/58646 
  🟩 GCC                Pass: 100%/21  | Total:  4h 52m | Avg: 13m 55s | Max: 32m 54s | Hits:  77%/55342 
  🟩 MSVC               Pass: 100%/4   | Total:  2h 00m | Avg: 30m 09s | Max: 31m 24s | Hits:  63%/13202 
  🟩 NVHPC              Pass: 100%/2   | Total:  1h 26m | Avg: 43m 24s | Max: 43m 50s | Hits:  27%/6889  
🔍 gpu: rtx2080 🔍
  🟩 h100               Pass: 100%/2   | Total: 22m 33s | Avg: 11m 16s | Max: 18m 05s | Hits:  99%/3554  
  🔍 rtx2080            Pass:  97%/43  | Total: 11h 39m | Avg: 16m 15s | Max: 43m 50s | Hits:  73%/130525
🔍 jobs: Test 🔍
  🟩 Build              Pass: 100%/39  | Total: 10h 37m | Avg: 16m 20s | Max: 43m 50s | Hits:  73%/134039
  🟩 NVRTC              Pass: 100%/2   | Total: 41m 56s | Avg: 20m 58s | Max: 21m 58s | Hits:  90%/40    
  🔍 Test               Pass:  66%/3   | Total: 40m 41s | Avg: 13m 33s | Max: 18m 05s
  🟩 VerifyCodegen      Pass: 100%/1   | Total:  1m 53s | Avg:  1m 53s | Max:  1m 53s
🔍 std: 20 🔍
  🟩 17                 Pass: 100%/22  | Total:  5h 44m | Avg: 15m 38s | Max: 42m 58s | Hits:  75%/71587 
  🔍 20                 Pass:  95%/22  | Total:  6h 15m | Avg: 17m 04s | Max: 43m 50s | Hits:  71%/62492 
🟩 sm
  🟩 75                 Pass: 100%/2   | Total: 41m 56s | Avg: 20m 58s | Max: 21m 58s | Hits:  90%/40    
  🟩 90                 Pass: 100%/2   | Total: 22m 33s | Avg: 11m 16s | Max: 18m 05s | Hits:  99%/3554  
  🟩 90;90a;100         Pass: 100%/1   | Total:  5m 05s | Avg:  5m 05s | Max:  5m 05s | Hits:  99%/3554

🟩 cub: Pass: 100%/47 | Total: 20h 55m | Avg: 26m 42s | Max: 1h 30m | Hits: 90%/57682

🟩 cpu
  🟩 amd64              Pass: 100%/45  | Total: 20h 29m | Avg: 27m 18s | Max:  1h 30m | Hits:  89%/55172 
  🟩 arm64              Pass: 100%/2   | Total: 26m 33s | Avg: 13m 16s | Max: 20m 27s | Hits:  99%/2510  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total:  2h 16m | Avg: 27m 18s | Max:  1h 14m | Hits:  87%/6091  
  🟩 12.9               Pass: 100%/42  | Total: 18h 39m | Avg: 26m 38s | Max:  1h 30m | Hits:  90%/51591 
🟩 cudacxx
  🟩 ClangCUDA19        Pass: 100%/2   | Total: 10m 34s | Avg:  5m 17s | Max:  5m 30s | Hits:  99%/2161  
  🟩 nvcc12.0           Pass: 100%/5   | Total:  2h 16m | Avg: 27m 18s | Max:  1h 14m | Hits:  87%/6091  
  🟩 nvcc12.9           Pass: 100%/40  | Total: 18h 28m | Avg: 27m 42s | Max:  1h 30m | Hits:  90%/49430 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total: 10m 34s | Avg:  5m 17s | Max:  5m 30s | Hits:  99%/2161  
  🟩 nvcc               Pass: 100%/45  | Total: 20h 45m | Avg: 27m 40s | Max:  1h 30m | Hits:  89%/55521 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total:  1h 01m | Avg: 15m 18s | Max: 22m 01s | Hits:  95%/5022  
  🟩 Clang15            Pass: 100%/2   | Total: 24m 12s | Avg: 12m 06s | Max: 15m 14s | Hits:  99%/2507  
  🟩 Clang16            Pass: 100%/2   | Total: 24m 33s | Avg: 12m 16s | Max: 15m 23s | Hits:  98%/2507  
  🟩 Clang17            Pass: 100%/2   | Total: 28m 31s | Avg: 14m 15s | Max: 14m 16s | Hits:  99%/2507  
  🟩 Clang18            Pass: 100%/2   | Total: 20m 34s | Avg: 10m 17s | Max: 13m 33s | Hits:  99%/2507  
  🟩 Clang19            Pass: 100%/7   | Total:  1h 41m | Avg: 14m 31s | Max: 33m 24s | Hits:  99%/8430  
  🟩 GCC7               Pass: 100%/2   | Total: 39m 14s | Avg: 19m 37s | Max: 21m 08s | Hits:  94%/2510  
  🟩 GCC8               Pass: 100%/1   | Total: 16m 57s | Avg: 16m 57s | Max: 16m 57s | Hits:  95%/1255  
  🟩 GCC9               Pass: 100%/2   | Total: 39m 18s | Avg: 19m 39s | Max: 21m 52s | Hits:  94%/2510  
  🟩 GCC10              Pass: 100%/2   | Total: 37m 14s | Avg: 18m 37s | Max: 18m 52s | Hits:  91%/2511  
  🟩 GCC11              Pass: 100%/2   | Total: 33m 59s | Avg: 16m 59s | Max: 17m 15s | Hits:  97%/2507  
  🟩 GCC12              Pass: 100%/2   | Total: 38m 44s | Avg: 19m 22s | Max: 19m 43s | Hits:  91%/2507  
  🟩 GCC13              Pass: 100%/11  | Total:  5h 16m | Avg: 28m 46s | Max: 49m 41s | Hits:  96%/13813 
  🟩 MSVC14.29          Pass: 100%/2   | Total:  2h 31m | Avg:  1h 15m | Max:  1h 17m | Hits:  41%/2140  
  🟩 MSVC14.43          Pass: 100%/2   | Total:  2h 51m | Avg:  1h 25m | Max:  1h 30m | Hits:  39%/2140  
  🟩 NVHPC25.5          Pass: 100%/2   | Total:  2h 30m | Avg:  1h 15m | Max:  1h 15m | Hits:  36%/2309  
🟩 cxx_family
  🟩 Clang              Pass: 100%/19  | Total:  4h 20m | Avg: 13m 43s | Max: 33m 24s | Hits:  98%/23480 
  🟩 GCC                Pass: 100%/22  | Total:  8h 41m | Avg: 23m 43s | Max: 49m 41s | Hits:  95%/27613 
  🟩 MSVC               Pass: 100%/4   | Total:  5h 22m | Avg:  1h 20m | Max:  1h 30m | Hits:  40%/4280  
  🟩 NVHPC              Pass: 100%/2   | Total:  2h 30m | Avg:  1h 15m | Max:  1h 15m | Hits:  36%/2309  
🟩 gpu
  🟩 h100               Pass: 100%/3   | Total:  1h 00m | Avg: 20m 05s | Max: 29m 21s | Hits:  99%/3768  
  🟩 rtx2080            Pass: 100%/36  | Total: 15h 57m | Avg: 26m 36s | Max:  1h 30m | Hits:  87%/43872 
  🟩 rtxa6000           Pass: 100%/8   | Total:  3h 57m | Avg: 29m 42s | Max: 39m 20s | Hits:  99%/10042 
🟩 jobs
  🟩 Build              Pass: 100%/39  | Total: 16h 29m | Avg: 25m 21s | Max:  1h 30m | Hits:  88%/47638 
  🟩 DeviceLaunch       Pass: 100%/1   | Total: 38m 10s | Avg: 38m 10s | Max: 38m 10s | Hits:  99%/1256  
  🟩 GraphCapture       Pass: 100%/1   | Total: 34m 27s | Avg: 34m 27s | Max: 34m 27s | Hits:  99%/1256  
  🟩 HostLaunch         Pass: 100%/3   | Total:  1h 42m | Avg: 34m 01s | Max: 39m 20s | Hits:  99%/3766  
  🟩 TestGPU            Pass: 100%/3   | Total:  1h 31m | Avg: 30m 36s | Max: 36m 20s | Hits:  99%/3766  
🟩 sm
  🟩 90                 Pass: 100%/3   | Total:  1h 00m | Avg: 20m 05s | Max: 29m 21s | Hits:  99%/3768  
  🟩 90;90a;100         Pass: 100%/1   | Total: 19m 09s | Avg: 19m 09s | Max: 19m 09s | Hits:  94%/1256  
🟩 std
  🟩 17                 Pass: 100%/21  | Total: 10h 12m | Avg: 29m 08s | Max:  1h 20m | Hits:  85%/25508 
  🟩 20                 Pass: 100%/26  | Total: 10h 43m | Avg: 24m 45s | Max:  1h 30m | Hits:  94%/32174

🟩 thrust: Pass: 100%/47 | Total: 13h 33m | Avg: 17m 18s | Max: 1h 40m | Hits: 92%/89895

🟩 cmake_options
  🟩 -DTHRUST_DISPATCH_TYPE=Force32bit Pass: 100%/2   | Total: 21m 44s | Avg: 10m 52s | Max: 14m 04s | Hits:  99%/3828  
🟩 cpu
  🟩 amd64              Pass: 100%/45  | Total: 13h 21m | Avg: 17m 48s | Max:  1h 40m | Hits:  92%/86068 
  🟩 arm64              Pass: 100%/2   | Total: 11m 47s | Avg:  5m 53s | Max:  6m 39s | Hits:  99%/3827  
🟩 ctk
  🟩 12.0               Pass: 100%/5   | Total:  1h 27m | Avg: 17m 25s | Max:  1h 03m | Hits:  89%/9560  
  🟩 12.9               Pass: 100%/42  | Total: 12h 06m | Avg: 17m 17s | Max:  1h 40m | Hits:  92%/80335 
🟩 cudacxx
  🟩 ClangCUDA19        Pass: 100%/2   | Total: 11m 11s | Avg:  5m 35s | Max:  5m 39s | Hits: 100%/3826  
  🟩 nvcc12.0           Pass: 100%/5   | Total:  1h 27m | Avg: 17m 25s | Max:  1h 03m | Hits:  89%/9560  
  🟩 nvcc12.9           Pass: 100%/40  | Total: 11h 55m | Avg: 17m 52s | Max:  1h 40m | Hits:  92%/76509 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total: 11m 11s | Avg:  5m 35s | Max:  5m 39s | Hits: 100%/3826  
  🟩 nvcc               Pass: 100%/45  | Total: 13h 22m | Avg: 17m 49s | Max:  1h 40m | Hits:  92%/86069 
🟩 cxx
  🟩 Clang14            Pass: 100%/4   | Total: 21m 53s | Avg:  5m 28s | Max:  5m 57s | Hits: 100%/7652  
  🟩 Clang15            Pass: 100%/2   | Total: 11m 26s | Avg:  5m 43s | Max:  5m 48s | Hits: 100%/3826  
  🟩 Clang16            Pass: 100%/2   | Total: 12m 07s | Avg:  6m 03s | Max:  6m 12s | Hits: 100%/3826  
  🟩 Clang17            Pass: 100%/2   | Total: 11m 55s | Avg:  5m 57s | Max:  6m 13s | Hits: 100%/3826  
  🟩 Clang18            Pass: 100%/2   | Total: 11m 21s | Avg:  5m 40s | Max:  5m 50s | Hits: 100%/3826  
  🟩 Clang19            Pass: 100%/7   | Total: 48m 27s | Avg:  6m 55s | Max: 11m 56s | Hits:  99%/13391 
  🟩 GCC7               Pass: 100%/2   | Total: 13m 14s | Avg:  6m 37s | Max:  6m 40s | Hits:  99%/3828  
  🟩 GCC8               Pass: 100%/1   | Total: 13m 43s | Avg: 13m 43s | Max: 13m 43s | Hits:  89%/1914  
  🟩 GCC9               Pass: 100%/2   | Total: 14m 22s | Avg:  7m 11s | Max:  7m 27s | Hits:  99%/3828  
  🟩 GCC10              Pass: 100%/2   | Total: 14m 05s | Avg:  7m 02s | Max:  7m 06s | Hits:  99%/3828  
  🟩 GCC11              Pass: 100%/2   | Total: 14m 41s | Avg:  7m 20s | Max:  7m 22s | Hits:  99%/3828  
  🟩 GCC12              Pass: 100%/2   | Total: 14m 40s | Avg:  7m 20s | Max:  7m 22s | Hits:  99%/3828  
  🟩 GCC13              Pass: 100%/10  | Total:  1h 58m | Avg: 11m 48s | Max: 32m 53s | Hits:  98%/19140 
  🟩 MSVC14.29          Pass: 100%/2   | Total:  2h 06m | Avg:  1h 03m | Max:  1h 03m | Hits:  58%/3812  
  🟩 MSVC14.43          Pass: 100%/3   | Total:  2h 54m | Avg: 58m 11s | Max:  1h 11m | Hits:  62%/5718  
  🟩 NVHPC25.5          Pass: 100%/2   | Total:  3h 12m | Avg:  1h 36m | Max:  1h 40m | Hits:  32%/3824  
🟩 cxx_family
  🟩 Clang              Pass: 100%/19  | Total:  1h 57m | Avg:  6m 09s | Max: 11m 56s | Hits:  99%/36347 
  🟩 GCC                Pass: 100%/21  | Total:  3h 22m | Avg:  9m 39s | Max: 32m 53s | Hits:  98%/40194 
  🟩 MSVC               Pass: 100%/5   | Total:  5h 01m | Avg:  1h 00m | Max:  1h 11m | Hits:  60%/9530  
  🟩 NVHPC              Pass: 100%/2   | Total:  3h 12m | Avg:  1h 36m | Max:  1h 40m | Hits:  32%/3824  
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 18m 39s | Avg:  9m 19s | Max: 13m 16s | Hits:  99%/3828  
  🟩 rtx2080            Pass: 100%/35  | Total: 10h 10m | Avg: 17m 26s | Max:  1h 40m | Hits:  91%/66946 
  🟩 rtx4090            Pass: 100%/10  | Total:  3h 04m | Avg: 18m 24s | Max:  1h 11m | Hits:  94%/19121 
🟩 jobs
  🟩 Build              Pass: 100%/40  | Total: 11h 49m | Avg: 17m 43s | Max:  1h 40m | Hits:  91%/76507 
  🟩 TestCPU            Pass: 100%/3   | Total: 51m 26s | Avg: 17m 08s | Max: 33m 54s | Hits:  99%/5733  
  🟩 TestGPU            Pass: 100%/4   | Total: 52m 48s | Avg: 13m 12s | Max: 14m 04s | Hits:  99%/7655  
🟩 sm
  🟩 90                 Pass: 100%/2   | Total: 18m 39s | Avg:  9m 19s | Max: 13m 16s | Hits:  99%/3828  
  🟩 90;90a;100         Pass: 100%/1   | Total: 32m 53s | Avg: 32m 53s | Max: 32m 53s | Hits:  87%/1914  
🟩 std
  🟩 17                 Pass: 100%/21  | Total:  6h 53m | Avg: 19m 42s | Max:  1h 40m | Hits:  89%/40160 
  🟩 20                 Pass: 100%/24  | Total:  6h 17m | Avg: 15m 44s | Max:  1h 31m | Hits:  94%/45907

🟩 cudax: Pass: 100%/26 | Total: 2h 26m | Avg: 5m 37s | Max: 15m 26s | Hits: 98%/15234

🟩 cpu
  🟩 amd64              Pass: 100%/22  | Total:  2h 13m | Avg:  6m 03s | Max: 15m 26s | Hits:  97%/12798 
  🟩 arm64              Pass: 100%/4   | Total: 13m 04s | Avg:  3m 16s | Max:  3m 30s | Hits:  99%/2436  
🟩 ctk
  🟩 12.0               Pass: 100%/3   | Total: 21m 53s | Avg:  7m 17s | Max: 15m 26s | Hits:  96%/1526  
  🟩 12.9               Pass: 100%/23  | Total:  2h 04m | Avg:  5m 24s | Max: 13m 21s | Hits:  98%/13708 
🟩 cudacxx
  🟩 nvcc12.0           Pass: 100%/3   | Total: 21m 53s | Avg:  7m 17s | Max: 15m 26s | Hits:  96%/1526  
  🟩 nvcc12.9           Pass: 100%/23  | Total:  2h 04m | Avg:  5m 24s | Max: 13m 21s | Hits:  98%/13708 
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/26  | Total:  2h 26m | Avg:  5m 37s | Max: 15m 26s | Hits:  98%/15234 
🟩 cxx
  🟩 Clang14            Pass: 100%/2   | Total:  6m 18s | Avg:  3m 09s | Max:  3m 18s | Hits: 100%/1220  
  🟩 Clang15            Pass: 100%/1   | Total:  3m 25s | Avg:  3m 25s | Max:  3m 25s | Hits: 100%/609   
  🟩 Clang16            Pass: 100%/1   | Total:  3m 34s | Avg:  3m 34s | Max:  3m 34s | Hits: 100%/609   
  🟩 Clang17            Pass: 100%/1   | Total:  3m 25s | Avg:  3m 25s | Max:  3m 25s | Hits: 100%/609   
  🟩 Clang18            Pass: 100%/1   | Total:  3m 13s | Avg:  3m 13s | Max:  3m 13s | Hits: 100%/609   
  🟩 Clang19            Pass: 100%/4   | Total: 17m 23s | Avg:  4m 20s | Max:  8m 00s | Hits: 100%/2436  
  🟩 GCC10              Pass: 100%/2   | Total:  7m 26s | Avg:  3m 43s | Max:  3m 59s | Hits:  99%/1220  
  🟩 GCC11              Pass: 100%/1   | Total:  3m 54s | Avg:  3m 54s | Max:  3m 54s | Hits:  99%/609   
  🟩 GCC12              Pass: 100%/1   | Total:  3m 59s | Avg:  3m 59s | Max:  3m 59s | Hits:  99%/609   
  🟩 GCC13              Pass: 100%/8   | Total: 43m 08s | Avg:  5m 23s | Max: 12m 14s | Hits:  99%/4872  
  🟩 MSVC14.39          Pass: 100%/1   | Total: 15m 26s | Avg: 15m 26s | Max: 15m 26s | Hits:  81%/308   
  🟩 MSVC14.43          Pass: 100%/1   | Total: 13m 21s | Avg: 13m 21s | Max: 13m 21s | Hits:  81%/310   
  🟩 NVHPC25.5          Pass: 100%/2   | Total: 21m 54s | Avg: 10m 57s | Max: 11m 30s | Hits:  87%/1214  
🟩 cxx_family
  🟩 Clang              Pass: 100%/10  | Total: 37m 18s | Avg:  3m 43s | Max:  8m 00s | Hits: 100%/6092  
  🟩 GCC                Pass: 100%/12  | Total: 58m 27s | Avg:  4m 52s | Max: 12m 14s | Hits:  99%/7310  
  🟩 MSVC               Pass: 100%/2   | Total: 28m 47s | Avg: 14m 23s | Max: 15m 26s | Hits:  81%/618   
  🟩 NVHPC              Pass: 100%/2   | Total: 21m 54s | Avg: 10m 57s | Max: 11m 30s | Hits:  87%/1214  
🟩 gpu
  🟩 h100               Pass: 100%/2   | Total: 13m 11s | Avg:  6m 35s | Max: 10m 07s | Hits:  99%/1218  
  🟩 rtx2080            Pass: 100%/24  | Total:  2h 13m | Avg:  5m 33s | Max: 15m 26s | Hits:  98%/14016 
🟩 jobs
  🟩 Build              Pass: 100%/23  | Total:  1h 56m | Avg:  5m 02s | Max: 15m 26s | Hits:  97%/13407 
  🟩 Test               Pass: 100%/3   | Total: 30m 21s | Avg: 10m 07s | Max: 12m 14s | Hits:  99%/1827  
🟩 sm
  🟩 90                 Pass: 100%/3   | Total: 16m 37s | Avg:  5m 32s | Max: 10m 07s | Hits:  99%/1827  
  🟩 90a                Pass: 100%/1   | Total:  3m 14s | Avg:  3m 14s | Max:  3m 14s | Hits:  99%/609   
🟩 std
  🟩 17                 Pass: 100%/4   | Total: 21m 30s | Avg:  5m 22s | Max: 11m 30s | Hits:  96%/2434  
  🟩 20                 Pass: 100%/22  | Total:  2h 04m | Avg:  5m 40s | Max: 15m 26s | Hits:  98%/12800

🟩 python: Pass: 100%/8 | Total: 1h 35m | Avg: 11m 58s | Max: 18m 19s

🟩 cpu
  🟩 amd64              Pass: 100%/8   | Total:  1h 35m | Avg: 11m 58s | Max: 18m 19s
🟩 ctk
  🟩 12.9               Pass: 100%/8   | Total:  1h 35m | Avg: 11m 58s | Max: 18m 19s
🟩 cudacxx
  🟩 nvcc12.9           Pass: 100%/8   | Total:  1h 35m | Avg: 11m 58s | Max: 18m 19s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/8   | Total:  1h 35m | Avg: 11m 58s | Max: 18m 19s
🟩 cxx
  🟩 GCC13              Pass: 100%/8   | Total:  1h 35m | Avg: 11m 58s | Max: 18m 19s
🟩 cxx_family
  🟩 GCC                Pass: 100%/8   | Total:  1h 35m | Avg: 11m 58s | Max: 18m 19s
🟩 gpu
  🟩 rtxa6000           Pass: 100%/8   | Total:  1h 35m | Avg: 11m 58s | Max: 18m 19s
🟩 jobs
  🟩 Build cuda.cccl    Pass: 100%/2   | Total: 20m 01s | Avg: 10m 00s | Max: 10m 21s
  🟩 Test cuda.cccl.cooperative Pass: 100%/2   | Total: 35m 39s | Avg: 17m 49s | Max: 18m 19s
  🟩 Test cuda.cccl.headers Pass: 100%/2   | Total:  7m 31s | Avg:  3m 45s | Max:  3m 50s
  🟩 Test cuda.cccl.parallel Pass: 100%/2   | Total: 32m 39s | Avg: 16m 19s | Max: 16m 39s
🟩 py_version
  🟩 3.10               Pass: 100%/4   | Total: 47m 31s | Avg: 11m 52s | Max: 17m 20s
  🟩 3.13               Pass: 100%/4   | Total: 48m 19s | Avg: 12m 04s | Max: 18m 19s

🟩 packaging: Pass: 100%/4 | Total: 11m 36s | Avg: 2m 54s | Max: 3m 02s

🟩 cpu
  🟩 amd64              Pass: 100%/4   | Total: 11m 36s | Avg:  2m 54s | Max:  3m 02s
🟩 ctk
  🟩 12.0               Pass: 100%/2   | Total:  5m 53s | Avg:  2m 56s | Max:  3m 02s
  🟩 12.9               Pass: 100%/2   | Total:  5m 43s | Avg:  2m 51s | Max:  2m 57s
🟩 cudacxx
  🟩 nvcc12.0           Pass: 100%/2   | Total:  5m 53s | Avg:  2m 56s | Max:  3m 02s
  🟩 nvcc12.9           Pass: 100%/2   | Total:  5m 43s | Avg:  2m 51s | Max:  2m 57s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/4   | Total: 11m 36s | Avg:  2m 54s | Max:  3m 02s
🟩 cxx
  🟩 Clang14            Pass: 100%/1   | Total:  3m 02s | Avg:  3m 02s | Max:  3m 02s
  🟩 Clang19            Pass: 100%/1   | Total:  2m 57s | Avg:  2m 57s | Max:  2m 57s
  🟩 GCC12              Pass: 100%/1   | Total:  2m 51s | Avg:  2m 51s | Max:  2m 51s
  🟩 GCC13              Pass: 100%/1   | Total:  2m 46s | Avg:  2m 46s | Max:  2m 46s
🟩 cxx_family
  🟩 Clang              Pass: 100%/2   | Total:  5m 59s | Avg:  2m 59s | Max:  3m 02s
  🟩 GCC                Pass: 100%/2   | Total:  5m 37s | Avg:  2m 48s | Max:  2m 51s
🟩 gpu
  🟩 rtx2080            Pass: 100%/4   | Total: 11m 36s | Avg:  2m 54s | Max:  3m 02s
🟩 jobs
  🟩 Test               Pass: 100%/4   | Total: 11m 36s | Avg:  2m 54s | Max:  3m 02s

🟩 stdpar: Pass: 100%/4 | Total: 18m 18s | Avg: 4m 34s | Max: 4m 51s

🟩 cpu
  🟩 amd64              Pass: 100%/2   | Total:  9m 36s | Avg:  4m 48s | Max:  4m 51s
  🟩 arm64              Pass: 100%/2   | Total:  8m 42s | Avg:  4m 21s | Max:  4m 29s
🟩 ctk
  🟩 12.9               Pass: 100%/4   | Total: 18m 18s | Avg:  4m 34s | Max:  4m 51s
🟩 cudacxx
  🟩 nvcc12.9           Pass: 100%/4   | Total: 18m 18s | Avg:  4m 34s | Max:  4m 51s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/4   | Total: 18m 18s | Avg:  4m 34s | Max:  4m 51s
🟩 cxx
  🟩 NVHPC25.5          Pass: 100%/4   | Total: 18m 18s | Avg:  4m 34s | Max:  4m 51s
🟩 cxx_family
  🟩 NVHPC              Pass: 100%/4   | Total: 18m 18s | Avg:  4m 34s | Max:  4m 51s
🟩 gpu
  🟩 rtx2080            Pass: 100%/4   | Total: 18m 18s | Avg:  4m 34s | Max:  4m 51s
🟩 jobs
  🟩 Build              Pass: 100%/4   | Total: 18m 18s | Avg:  4m 34s | Max:  4m 51s
🟩 std
  🟩 17                 Pass: 100%/2   | Total:  9m 14s | Avg:  4m 37s | Max:  4m 45s
  🟩 20                 Pass: 100%/2   | Total:  9m 04s | Avg:  4m 32s | Max:  4m 51s

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 12m 15s | Avg: 6m 07s | Max: 10m 11s | Hits: 98%/330

🟩 cpu
  🟩 amd64              Pass: 100%/2   | Total: 12m 15s | Avg:  6m 07s | Max: 10m 11s | Hits:  98%/330   
🟩 ctk
  🟩 12.9               Pass: 100%/2   | Total: 12m 15s | Avg:  6m 07s | Max: 10m 11s | Hits:  98%/330   
🟩 cudacxx
  🟩 nvcc12.9           Pass: 100%/2   | Total: 12m 15s | Avg:  6m 07s | Max: 10m 11s | Hits:  98%/330   
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/2   | Total: 12m 15s | Avg:  6m 07s | Max: 10m 11s | Hits:  98%/330   
🟩 cxx
  🟩 GCC13              Pass: 100%/2   | Total: 12m 15s | Avg:  6m 07s | Max: 10m 11s | Hits:  98%/330   
🟩 cxx_family
  🟩 GCC                Pass: 100%/2   | Total: 12m 15s | Avg:  6m 07s | Max: 10m 11s | Hits:  98%/330   
🟩 gpu
  🟩 rtx2080            Pass: 100%/2   | Total: 12m 15s | Avg:  6m 07s | Max: 10m 11s | Hits:  98%/330   
🟩 jobs
  🟩 Build              Pass: 100%/1   | Total:  2m 04s | Avg:  2m 04s | Max:  2m 04s | Hits:  98%/165   
  🟩 Test               Pass: 100%/1   | Total: 10m 11s | Avg: 10m 11s | Max: 10m 11s | Hits:  98%/165

👃 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: 183)

#	Runner
125	`linux-amd64-cpu16`
15	`windows-amd64-cpu16`
12	`linux-arm64-cpu16`
12	`linux-amd64-gpu-rtxa6000-latest-1`
11	`linux-amd64-gpu-rtx2080-latest-1`
5	`linux-amd64-gpu-h100-latest-1`
3	`linux-amd64-gpu-rtx4090-latest-1`

miscco · 2025-06-23T06:46:33Z

libcudacxx/include/cuda/__functional/maximum.h

+#if _CCCL_HAS_NVBF16()
+#  include <cuda_bf16.h>
+#endif
+#if _CCCL_HAS_NVFP16()
+#  include <cuda_fp16.h>
+#endif
+#if _CCCL_HAS_FLOAT128()
+#  include <crt/device_fp128_functions.h>
+#endif


We should not include those headers, but the generic one

We should also move the fp128 one there

Suggested change

#if _CCCL_HAS_NVBF16()

# include <cuda_bf16.h>

#endif

#if _CCCL_HAS_NVFP16()

# include <cuda_fp16.h>

#endif

#if _CCCL_HAS_FLOAT128()

# include <crt/device_fp128_functions.h>

#endif

# include <cuda/std/__floating_point/nvfp_types.h>

#if _CCCL_HAS_FLOAT128()

# include <crt/device_fp128_functions.h>

#endif

miscco · 2025-06-23T06:49:05Z

libcudacxx/test/libcudacxx/cuda/functional/maximum.pass.cpp

@@ -32,12 +32,37 @@ __host__ __device__ constexpr bool test()
         test<int>(1, 0, 1) && //
         test<int>(0, 0, 0) && //
         test<int>(-1, 1, 1) && //
-         test<char>('a', 'b', 'b');
+         test<char>('a', 'b', 'b') && //


Please refactor this function to just call test<meow>() individually and then return true

bernhardmgruber · 2025-06-24T08:20:44Z

libcudacxx/include/cuda/__functional/maximum.h

+    {
+      if constexpr (_CUDA_VSTD::is_same_v<_Tp, float>)
+      {
+        NV_IF_TARGET(NV_IS_DEVICE, (return ::fmaxf(__lhs, __rhs);))


Important: Isn't this a behavior change for NaNs?

Before maximum{}(NaN, 1) == NaN
After: maximum{}(NaN, 1) == 1

The documentation for fmaxf says:

Returns the larger of two floating-point arguments, treating NaNs as missing data (between a NaN and a numeric value, the numeric value is chosen).

Is this what we want? And is this consistent with any fast path in CUB that detects that the operator is maximum and then applies a custom logic, side-stepping calling this actual operator?

add cuda::minimum/maximum specialization for float/double/__half/__nv…

316cd29

…_bfloat16/__float128

fbusato self-assigned this Jun 18, 2025

fbusato requested a review from a team as a code owner June 18, 2025 22:27

fbusato requested a review from griwes June 18, 2025 22:27

fbusato added this to CCCL Jun 18, 2025

github-project-automation bot moved this to Todo in CCCL Jun 18, 2025

fbusato added the 3.1.0 Targeted for 3.1 release label Jun 18, 2025

fbusato moved this from Todo to In Review in CCCL Jun 18, 2025

bernhardmgruber reviewed Jun 18, 2025

View reviewed changes

libcudacxx/test/libcudacxx/cuda/functional/maximum.pass.cpp Outdated Show resolved Hide resolved

bernhardmgruber requested review from miscco and bernhardmgruber and removed request for bernhardmgruber June 18, 2025 22:47

fbusato and others added 5 commits June 18, 2025 16:37

protect half/bfloat with macro

69c26e1

Update libcudacxx/test/libcudacxx/cuda/functional/maximum.pass.cpp

99706e7

Co-authored-by: Bernhard Manfred Gruber <[email protected]>

enable tests for CTK >= 12.2

e2b3dea

protect float128 functions with macro

77565fd

comment consistency

0c440cc

miscco requested changes Jun 23, 2025

View reviewed changes

github-project-automation bot moved this from In Review to In Progress in CCCL Jun 23, 2025

bernhardmgruber reviewed Jun 24, 2025

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Optimize `cuda::minimum/maximum` for `float`, `double`, `half`, `nv_bfloat16`, `__float128` #5034

Optimize `cuda::minimum/maximum` for `float`, `double`, `half`, `nv_bfloat16`, `__float128` #5034

Uh oh!

fbusato commented Jun 18, 2025

Uh oh!

Uh oh!

github-actions bot commented Jun 19, 2025

🟨 libcudacxx: Pass: 97%/45 | Total: 12h 01m | Avg: 16m 02s | Max: 43m 50s | Hits: 73%/134079

🟩 cub: Pass: 100%/47 | Total: 20h 55m | Avg: 26m 42s | Max: 1h 30m | Hits: 90%/57682

🟩 thrust: Pass: 100%/47 | Total: 13h 33m | Avg: 17m 18s | Max: 1h 40m | Hits: 92%/89895

🟩 cudax: Pass: 100%/26 | Total: 2h 26m | Avg: 5m 37s | Max: 15m 26s | Hits: 98%/15234

🟩 python: Pass: 100%/8 | Total: 1h 35m | Avg: 11m 58s | Max: 18m 19s

🟩 packaging: Pass: 100%/4 | Total: 11m 36s | Avg: 2m 54s | Max: 3m 02s

🟩 stdpar: Pass: 100%/4 | Total: 18m 18s | Avg: 4m 34s | Max: 4m 51s

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 12m 15s | Avg: 6m 07s | Max: 10m 11s | Hits: 98%/330

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 183)

Uh oh!

miscco Jun 23, 2025

Uh oh!

miscco Jun 23, 2025

Uh oh!

bernhardmgruber Jun 24, 2025

Uh oh!

Uh oh!

Optimize cuda::minimum/maximum for float, double, __half, __nv_bfloat16, __float128 #5034

Are you sure you want to change the base?

Optimize cuda::minimum/maximum for float, double, __half, __nv_bfloat16, __float128 #5034

Uh oh!

Conversation

fbusato commented Jun 18, 2025

Description

Uh oh!

Uh oh!

github-actions bot commented Jun 19, 2025

🟨 libcudacxx: Pass: 97%/45 | Total: 12h 01m | Avg: 16m 02s | Max: 43m 50s | Hits: 73%/134079

🟩 cub: Pass: 100%/47 | Total: 20h 55m | Avg: 26m 42s | Max: 1h 30m | Hits: 90%/57682

🟩 thrust: Pass: 100%/47 | Total: 13h 33m | Avg: 17m 18s | Max: 1h 40m | Hits: 92%/89895

🟩 cudax: Pass: 100%/26 | Total: 2h 26m | Avg: 5m 37s | Max: 15m 26s | Hits: 98%/15234

🟩 python: Pass: 100%/8 | Total: 1h 35m | Avg: 11m 58s | Max: 18m 19s

🟩 packaging: Pass: 100%/4 | Total: 11m 36s | Avg: 2m 54s | Max: 3m 02s

🟩 stdpar: Pass: 100%/4 | Total: 18m 18s | Avg: 4m 34s | Max: 4m 51s

🟩 cccl_c_parallel: Pass: 100%/2 | Total: 12m 15s | Avg: 6m 07s | Max: 10m 11s | Hits: 98%/330

👃 Inspect Changes

Modifications in project?

Modifications in project or dependencies?

🏃‍ Runner counts (total jobs: 183)

Uh oh!

miscco Jun 23, 2025

Choose a reason for hiding this comment

Uh oh!

miscco Jun 23, 2025

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Jun 24, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Optimize `cuda::minimum/maximum` for `float`, `double`, `half`, `nv_bfloat16`, `__float128` #5034

Optimize `cuda::minimum/maximum` for `float`, `double`, `half`, `nv_bfloat16`, `__float128` #5034