Port thrust::minmax_element to CUB

[bernhardmgruber]

Replaces #4970

• Merge before: Port thrust::min|max_element to CUB #8291
• Merge before: Add cub::DeviceReduce::ArgMin[Last]Max #10731

Benchmark looks great:

# minmax_element

## [0] NVIDIA B200

|  T{ct}  |  Elements  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|---------|------------|------------|-------------|------------|-------------|-------------|---------|----------|
|   I8    |    2^16    |  31.204 us |       2.78% |  24.298 us |       8.28% |   -6.905 us | -22.13% |  🟢 FAST  |
|   I8    |    2^20    |  38.606 us |       2.44% |  33.895 us |       4.91% |   -4.711 us | -12.20% |  🟢 FAST  |
|   I8    |    2^24    |  66.596 us |       0.99% |  58.992 us |       1.79% |   -7.604 us | -11.42% |  🟢 FAST  |
|   I8    |    2^28    | 459.885 us |       0.14% | 303.848 us |       0.77% | -156.037 us | -33.93% |  🟢 FAST  |
|   I16   |    2^16    |  30.834 us |       4.60% |  23.595 us |       3.94% |   -7.240 us | -23.48% |  🟢 FAST  |
|   I16   |    2^20    |  38.357 us |       2.62% |  31.118 us |      10.41% |   -7.239 us | -18.87% |  🟢 FAST  |
|   I16   |    2^24    |  67.187 us |       1.34% |  47.518 us |       2.60% |  -19.669 us | -29.27% |  🟢 FAST  |
|   I16   |    2^28    | 462.640 us |       0.14% | 230.710 us |       0.97% | -231.930 us | -50.13% |  🟢 FAST  |
|   I32   |    2^16    |  30.828 us |       3.09% |  23.626 us |       5.67% |   -7.203 us | -23.36% |  🟢 FAST  |
|   I32   |    2^20    |  37.615 us |       3.02% |  25.445 us |       3.69% |  -12.170 us | -32.35% |  🟢 FAST  |
|   I32   |    2^24    |  65.216 us |       0.97% |  47.517 us |       2.21% |  -17.699 us | -27.14% |  🟢 FAST  |
|   I32   |    2^28    | 463.344 us |       0.20% | 220.140 us |       0.62% | -243.204 us | -52.49% |  🟢 FAST  |
|   I64   |    2^16    |  30.227 us |       2.00% |  23.534 us |       4.79% |   -6.692 us | -22.14% |  🟢 FAST  |
|   I64   |    2^20    |  36.643 us |       2.12% |  34.074 us |       3.32% |   -2.568 us |  -7.01% |  🟢 FAST  |
|   I64   |    2^24    |  77.376 us |       1.38% |  70.959 us |       3.67% |   -6.417 us |  -8.29% |  🟢 FAST  |
|   I64   |    2^28    | 602.920 us |       0.13% | 383.571 us |       0.83% | -219.349 us | -36.38% |  🟢 FAST  |
|  I128   |    2^16    |  36.408 us |       2.60% |  31.618 us |       9.98% |   -4.789 us | -13.15% |  🟢 FAST  |
|  I128   |    2^20    |  46.258 us |       1.74% |  47.672 us |       2.64% |    1.414 us |   3.06% |  🔴 SLOW  |
|  I128   |    2^24    | 152.925 us |       0.42% | 109.189 us |       0.90% |  -43.736 us | -28.60% |  🟢 FAST  |
|  I128   |    2^28    |   1.715 ms |       0.05% | 760.571 us |       0.20% | -954.178 us | -55.65% |  🟢 FAST  |
|   F32   |    2^16    |  30.554 us |       2.75% |  23.443 us |       3.77% |   -7.111 us | -23.27% |  🟢 FAST  |
|   F32   |    2^20    |  36.941 us |       2.50% |  25.588 us |       3.34% |  -11.353 us | -30.73% |  🟢 FAST  |
|   F32   |    2^24    |  65.409 us |       1.35% |  47.415 us |       1.74% |  -17.995 us | -27.51% |  🟢 FAST  |
|   F32   |    2^28    | 464.065 us |       0.13% | 220.033 us |       0.62% | -244.032 us | -52.59% |  🟢 FAST  |
|   F64   |    2^16    |  30.649 us |       2.49% |  23.927 us |       7.17% |   -6.721 us | -21.93% |  🟢 FAST  |
|   F64   |    2^20    |  36.484 us |       1.94% |  34.401 us |       3.91% |   -2.083 us |  -5.71% |  🟢 FAST  |
|   F64   |    2^24    |  75.188 us |       0.90% |  63.183 us |       2.02% |  -12.005 us | -15.97% |  🟢 FAST  |
|   F64   |    2^28    | 594.403 us |       0.14% | 373.279 us |       1.03% | -221.124 us | -37.20% |  🟢 FAST  |

Fixes: #1626
Fixes: #699
Fixes: #762

[copy-pr-bot]

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 2h 39m: Pass: 100%/112 | Total: 4d 03h | Max: 2h 19m | Hits: 60%/179180

See results here.

[bernhardmgruber]

Important: this is a silent breaking change in behavior, but I am not sure how big of a problem it is in practice and whether we should defer this to a major release.

Alternatively: we could extend cub::DeviceReduce::ArgMinMax with an option to return the first equivalent maximum instead of the last and use this API in Thrust for now.

[bernhardmgruber]

Yes, we should use ArgMinLastMax until CCCL 4.0