[CUB] `DeviceReduce` with device resident problem size by NaderAlAwar · Pull Request #9722 · NVIDIA/cccl

NaderAlAwar · 2026-07-07T01:04:32Z

Description

No SASS changes to any existing benchmarks.

We expect the device-resident problem size to be slightly slower: the host cannot read the problem size, so it must assume it is large, launching a worst-case grid and always taking the two-pass path. Compared to the host-provided num_items, the results show an overhead of at most ~6 us, which is most noticeable at the smallest problem sizes.

Additionally, on the run-to-run path, when the problem-size element is 64-bit we observe larger overheads: the work partitioning that the host normally computes once from the immediate value must instead be recomputed by each block on device (where 64-bit integer division is expensive), and carrying the 64-bit problem size widens the kernel's register footprint relative to the 32-bit instantiation (40 vs. 32 registers on SM90), which can reduce occupancy. The not-guaranteed (atomic) kernel absorbs the deferred machinery with no register increase and shows no 64-bit overhead, as its table below confirms. Users whose problem sizes fit in 32 bits should prefer a 32-bit element for the deferred source.

Run-to-run reduce

## [0] NVIDIA H200

|  T{ct}  |  OffsetT{ct}  |  Elements{io}  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |      Diff |   %Diff |  Status  |
|---------|---------------|----------------|------------|-------------|------------|-------------|-----------|---------|----------|
|   I8    |      I32      |      2^4       |   3.820 us |       2.85% |   9.532 us |       3.84% |  5.712 us | 149.55% | 🔴 SLOW  |
|   I8    |      I32      |      2^8       |   3.717 us |       3.67% |   9.740 us |       3.84% |  6.024 us | 162.07% | 🔴 SLOW  |
|   I8    |      I32      |      2^12      |   3.894 us |       3.36% |   9.895 us |       4.49% |  6.000 us | 154.08% | 🔴 SLOW  |
|   I8    |      I32      |      2^16      |   6.738 us |       5.67% |   9.793 us |       4.09% |  3.056 us |  45.35% | 🔴 SLOW  |
|   I8    |      I32      |      2^20      |   7.095 us |       5.71% |   9.618 us |       4.34% |  2.524 us |  35.57% | 🔴 SLOW  |
|   I8    |      I32      |      2^24      |  12.864 us |       3.41% |  13.598 us |       3.03% |  0.734 us |   5.71% | 🔴 SLOW  |
|   I8    |      I32      |      2^28      |  78.761 us |       1.48% |  79.451 us |       1.45% |  0.690 us |   0.88% | 🔵 SAME  |
|   I8    |      I64      |      2^4       |   3.782 us |       3.75% |   9.100 us |       4.53% |  5.317 us | 140.58% | 🔴 SLOW  |
|   I8    |      I64      |      2^8       |   3.849 us |       2.83% |   9.171 us |       5.07% |  5.322 us | 138.27% | 🔴 SLOW  |
|   I8    |      I64      |      2^12      |   3.895 us |       3.48% |   9.108 us |       4.80% |  5.213 us | 133.83% | 🔴 SLOW  |
|   I8    |      I64      |      2^16      |   6.740 us |       5.76% |   8.996 us |       3.94% |  2.256 us |  33.48% | 🔴 SLOW  |
|   I8    |      I64      |      2^20      |   7.163 us |       5.61% |   9.017 us |       4.38% |  1.854 us |  25.88% | 🔴 SLOW  |
|   I8    |      I64      |      2^24      |  12.987 us |       3.13% |  15.735 us |       2.74% |  2.748 us |  21.16% | 🔴 SLOW  |
|   I8    |      I64      |      2^28      |  78.785 us |       0.91% |  83.755 us |       0.81% |  4.971 us |   6.31% | 🔴 SLOW  |
|   I16   |      I32      |      2^4       |   3.762 us |       3.35% |   9.784 us |       4.43% |  6.022 us | 160.08% | 🔴 SLOW  |
|   I16   |      I32      |      2^8       |   3.716 us |       3.79% |   9.760 us |       4.37% |  6.043 us | 162.61% | 🔴 SLOW  |
|   I16   |      I32      |      2^12      |   4.033 us |       3.60% |   9.905 us |       3.60% |  5.872 us | 145.58% | 🔴 SLOW  |
|   I16   |      I32      |      2^16      |   6.830 us |       5.82% |   9.801 us |       3.70% |  2.971 us |  43.49% | 🔴 SLOW  |
|   I16   |      I32      |      2^20      |   7.657 us |       5.68% |   9.657 us |       4.55% |  2.000 us |  26.12% | 🔴 SLOW  |
|   I16   |      I32      |      2^24      |  17.722 us |       2.81% |  18.332 us |       2.50% |  0.610 us |   3.44% | 🔴 SLOW  |
|   I16   |      I32      |      2^28      | 135.091 us |       1.15% | 135.734 us |       1.10% |  0.644 us |   0.48% | 🔵 SAME  |
|   I16   |      I64      |      2^4       |   3.838 us |       3.81% |   9.072 us |       4.76% |  5.235 us | 136.41% | 🔴 SLOW  |
|   I16   |      I64      |      2^8       |   3.941 us |       3.18% |   9.187 us |       4.92% |  5.246 us | 133.13% | 🔴 SLOW  |
|   I16   |      I64      |      2^12      |   4.051 us |       3.95% |   9.113 us |       3.95% |  5.062 us | 124.95% | 🔴 SLOW  |
|   I16   |      I64      |      2^16      |   6.810 us |       6.09% |   8.968 us |       5.30% |  2.158 us |  31.70% | 🔴 SLOW  |
|   I16   |      I64      |      2^20      |   7.628 us |       5.33% |   8.915 us |       4.55% |  1.286 us |  16.86% | 🔴 SLOW  |
|   I16   |      I64      |      2^24      |  17.378 us |       2.74% |  19.275 us |       2.28% |  1.897 us |  10.92% | 🔴 SLOW  |
|   I16   |      I64      |      2^28      | 135.099 us |       1.28% | 137.044 us |       1.22% |  1.945 us |   1.44% | 🔴 SLOW  |
|   I32   |      I32      |      2^4       |   3.632 us |       3.62% |   9.465 us |       3.94% |  5.833 us | 160.61% | 🔴 SLOW  |
|   I32   |      I32      |      2^8       |   3.550 us |       3.78% |   9.428 us |       4.14% |  5.879 us | 165.62% | 🔴 SLOW  |
|   I32   |      I32      |      2^12      |   3.920 us |       3.84% |   9.570 us |       4.22% |  5.649 us | 144.11% | 🔴 SLOW  |
|   I32   |      I32      |      2^16      |   6.686 us |       5.38% |   9.453 us |       3.82% |  2.767 us |  41.38% | 🔴 SLOW  |
|   I32   |      I32      |      2^20      |   8.223 us |       4.85% |   9.475 us |       4.03% |  1.252 us |  15.23% | 🔴 SLOW  |
|   I32   |      I32      |      2^24      |  27.523 us |       2.29% |  28.162 us |       2.28% |  0.639 us |   2.32% | 🔴 SLOW  |
|   I32   |      I32      |      2^28      | 252.882 us |       0.82% | 253.720 us |       0.83% |  0.838 us |   0.33% | 🔵 SAME  |
|   I32   |      I64      |      2^4       |   3.640 us |       3.83% |   8.830 us |       4.54% |  5.189 us | 142.56% | 🔴 SLOW  |
|   I32   |      I64      |      2^8       |   3.578 us |       3.81% |   8.901 us |       3.82% |  5.323 us | 148.78% | 🔴 SLOW  |
|   I32   |      I64      |      2^12      |   3.989 us |       3.62% |   8.752 us |       4.39% |  4.763 us | 119.39% | 🔴 SLOW  |
|   I32   |      I64      |      2^16      |   6.851 us |       6.12% |   8.920 us |       3.81% |  2.070 us |  30.21% | 🔴 SLOW  |
|   I32   |      I64      |      2^20      |   8.493 us |       4.86% |   8.943 us |       4.54% |  0.449 us |   5.29% | 🔴 SLOW  |
|   I32   |      I64      |      2^24      |  27.841 us |       2.32% |  28.962 us |       2.26% |  1.122 us |   4.03% | 🔴 SLOW  |
|   I32   |      I64      |      2^28      | 253.023 us |       0.44% | 253.842 us |       0.39% |  0.819 us |   0.32% | 🔵 SAME  |
|   I64   |      I32      |      2^4       |   3.709 us |       3.67% |   9.768 us |       3.88% |  6.059 us | 163.37% | 🔴 SLOW  |
|   I64   |      I32      |      2^8       |   3.876 us |       3.48% |   9.822 us |       3.45% |  5.945 us | 153.38% | 🔴 SLOW  |
|   I64   |      I32      |      2^12      |   6.909 us |       5.92% |   9.756 us |       4.47% |  2.847 us |  41.21% | 🔴 SLOW  |
|   I64   |      I32      |      2^16      |   7.057 us |       5.02% |   9.574 us |       4.02% |  2.517 us |  35.67% | 🔴 SLOW  |
|   I64   |      I32      |      2^20      |  10.121 us |       4.97% |  10.522 us |       4.68% |  0.401 us |   3.96% | 🔵 SAME  |
|   I64   |      I32      |      2^24      |  46.186 us |       1.64% |  46.703 us |       1.55% |  0.516 us |   1.12% | 🔵 SAME  |
|   I64   |      I32      |      2^28      | 493.509 us |       0.55% | 492.765 us |       0.55% | -0.743 us |  -0.15% | 🔵 SAME  |
|   I64   |      I64      |      2^4       |   3.869 us |       3.07% |   9.977 us |       4.73% |  6.108 us | 157.86% | 🔴 SLOW  |
|   I64   |      I64      |      2^8       |   3.928 us |       3.54% |   9.914 us |       3.71% |  5.986 us | 152.40% | 🔴 SLOW  |
|   I64   |      I64      |      2^12      |   6.811 us |       5.47% |   9.851 us |       3.73% |  3.040 us |  44.64% | 🔴 SLOW  |
|   I64   |      I64      |      2^16      |   7.186 us |       6.35% |   9.801 us |       3.20% |  2.614 us |  36.38% | 🔴 SLOW  |
|   I64   |      I64      |      2^20      |  10.062 us |       5.00% |  10.757 us |       3.86% |  0.695 us |   6.91% | 🔴 SLOW  |
|   I64   |      I64      |      2^24      |  46.321 us |       1.63% |  46.991 us |       1.74% |  0.669 us |   1.44% | 🔵 SAME  |
|   I64   |      I64      |      2^28      | 493.137 us |       0.54% | 493.278 us |       0.56% |  0.141 us |   0.03% | 🔵 SAME  |
|  I128   |      I32      |      2^4       |   4.081 us |       3.00% |   9.380 us |       4.66% |  5.298 us | 129.82% | 🔴 SLOW  |
|  I128   |      I32      |      2^8       |   4.126 us |       3.39% |   9.321 us |       4.26% |  5.196 us | 125.93% | 🔴 SLOW  |
|  I128   |      I32      |      2^12      |   7.495 us |       4.60% |   9.100 us |       3.90% |  1.606 us |  21.42% | 🔴 SLOW  |
|  I128   |      I32      |      2^16      |   7.915 us |       5.17% |   9.221 us |       4.86% |  1.306 us |  16.50% | 🔴 SLOW  |
|  I128   |      I32      |      2^20      |  16.397 us |       2.90% |  16.860 us |       2.49% |  0.463 us |   2.82% | 🔴 SLOW  |
|  I128   |      I32      |      2^24      |  91.614 us |       1.14% |  93.772 us |       1.30% |  2.158 us |   2.36% | 🔴 SLOW  |
|  I128   |      I32      |      2^28      |   1.010 ms |       0.27% |   1.011 ms |       0.27% |  0.625 us |   0.06% | 🔵 SAME  |
|  I128   |      I64      |      2^4       |   4.079 us |       3.85% |   9.073 us |       4.34% |  4.994 us | 122.43% | 🔴 SLOW  |
|  I128   |      I64      |      2^8       |   4.275 us |       3.22% |   9.022 us |       4.12% |  4.748 us | 111.07% | 🔴 SLOW  |
|  I128   |      I64      |      2^12      |   7.467 us |       5.28% |   8.796 us |       4.85% |  1.329 us |  17.80% | 🔴 SLOW  |
|  I128   |      I64      |      2^16      |   7.931 us |       4.57% |   8.880 us |       3.67% |  0.949 us |  11.96% | 🔴 SLOW  |
|  I128   |      I64      |      2^20      |  16.218 us |       3.03% |  16.915 us |       3.13% |  0.698 us |   4.30% | 🔴 SLOW  |
|  I128   |      I64      |      2^24      |  87.591 us |       1.01% |  89.105 us |       1.08% |  1.514 us |   1.73% | 🔴 SLOW  |
|  I128   |      I64      |      2^28      | 997.151 us |       0.47% | 997.273 us |       0.49% |  0.122 us |   0.01% | 🔵 SAME  |
|   F32   |      I32      |      2^4       |   3.652 us |       4.01% |   9.783 us |       3.72% |  6.131 us | 167.86% | 🔴 SLOW  |
|   F32   |      I32      |      2^8       |   3.672 us |       3.49% |   9.813 us |       3.77% |  6.141 us | 167.23% | 🔴 SLOW  |
|   F32   |      I32      |      2^12      |   4.035 us |       3.19% |   9.853 us |       3.46% |  5.818 us | 144.21% | 🔴 SLOW  |
|   F32   |      I32      |      2^16      |   6.778 us |       5.04% |   9.684 us |       3.77% |  2.906 us |  42.88% | 🔴 SLOW  |
|   F32   |      I32      |      2^20      |   8.580 us |       5.00% |   9.749 us |       3.77% |  1.169 us |  13.62% | 🔴 SLOW  |
|   F32   |      I32      |      2^24      |  27.687 us |       2.25% |  28.313 us |       2.31% |  0.626 us |   2.26% | 🔴 SLOW  |
|   F32   |      I32      |      2^28      | 253.394 us |       0.87% | 254.394 us |       0.83% |  1.000 us |   0.39% | 🔵 SAME  |
|   F32   |      I64      |      2^4       |   3.616 us |       3.67% |   9.101 us |       4.24% |  5.485 us | 151.66% | 🔴 SLOW  |
|   F32   |      I64      |      2^8       |   3.665 us |       3.86% |   9.123 us |       4.75% |  5.457 us | 148.89% | 🔴 SLOW  |
|   F32   |      I64      |      2^12      |   4.002 us |       3.12% |   9.069 us |       3.73% |  5.067 us | 126.62% | 🔴 SLOW  |
|   F32   |      I64      |      2^16      |   6.794 us |       5.00% |   8.873 us |       4.59% |  2.079 us |  30.61% | 🔴 SLOW  |
|   F32   |      I64      |      2^20      |   8.419 us |       4.17% |   9.246 us |       4.33% |  0.827 us |   9.82% | 🔴 SLOW  |
|   F32   |      I64      |      2^24      |  27.975 us |       2.32% |  29.490 us |       2.23% |  1.515 us |   5.41% | 🔴 SLOW  |
|   F32   |      I64      |      2^28      | 253.099 us |       0.85% | 253.870 us |       0.85% |  0.770 us |   0.30% | 🔵 SAME  |
|   F64   |      I32      |      2^4       |   3.688 us |       3.23% |   9.837 us |       3.58% |  6.149 us | 166.75% | 🔴 SLOW  |
|   F64   |      I32      |      2^8       |   3.993 us |       4.23% |   9.835 us |       4.15% |  5.843 us | 146.33% | 🔴 SLOW  |
|   F64   |      I32      |      2^12      |   6.810 us |       6.39% |   9.906 us |       3.95% |  3.096 us |  45.46% | 🔴 SLOW  |
|   F64   |      I32      |      2^16      |   7.162 us |       4.72% |   9.704 us |       3.19% |  2.542 us |  35.49% | 🔴 SLOW  |
|   F64   |      I32      |      2^20      |   9.900 us |       4.41% |  10.700 us |       4.38% |  0.799 us |   8.07% | 🔴 SLOW  |
|   F64   |      I32      |      2^24      |  46.532 us |       1.67% |  47.010 us |       1.62% |  0.478 us |   1.03% | 🔵 SAME  |
|   F64   |      I32      |      2^28      | 493.946 us |       0.52% | 493.737 us |       0.56% | -0.209 us |  -0.04% | 🔵 SAME  |
|   F64   |      I64      |      2^4       |   3.985 us |       2.75% |  10.257 us |       4.39% |  6.271 us | 157.36% | 🔴 SLOW  |
|   F64   |      I64      |      2^8       |   4.093 us |       3.25% |  10.248 us |       3.56% |  6.154 us | 150.34% | 🔴 SLOW  |
|   F64   |      I64      |      2^12      |   7.142 us |       5.89% |  10.260 us |       4.32% |  3.118 us |  43.66% | 🔴 SLOW  |
|   F64   |      I64      |      2^16      |   7.412 us |       6.12% |  10.039 us |       3.22% |  2.627 us |  35.44% | 🔴 SLOW  |
|   F64   |      I64      |      2^20      |  10.542 us |       4.29% |  10.994 us |       4.37% |  0.453 us |   4.29% | 🔴 SLOW  |
|   F64   |      I64      |      2^24      |  46.493 us |       1.75% |  47.105 us |       1.63% |  0.613 us |   1.32% | 🔵 SAME  |
|   F64   |      I64      |      2^28      | 493.359 us |       0.31% | 487.053 us |       0.22% | -6.306 us |  -1.28% | 🟢 FAST  |
|   C32   |      I32      |      2^4       |   3.922 us |       3.02% |  10.122 us |       4.14% |  6.200 us | 158.06% | 🔴 SLOW  |
|   C32   |      I32      |      2^8       |   3.839 us |       3.73% |  10.075 us |       4.13% |  6.235 us | 162.41% | 🔴 SLOW  |
|   C32   |      I32      |      2^12      |   7.273 us |       5.91% |  10.156 us |       4.41% |  2.883 us |  39.64% | 🔴 SLOW  |
|   C32   |      I32      |      2^16      |   7.331 us |       5.78% |   9.893 us |       3.68% |  2.562 us |  34.94% | 🔴 SLOW  |
|   C32   |      I32      |      2^20      |  10.309 us |       4.54% |  10.983 us |       4.33% |  0.674 us |   6.54% | 🔴 SLOW  |
|   C32   |      I32      |      2^24      |  47.283 us |       1.55% |  47.671 us |       1.52% |  0.387 us |   0.82% | 🔵 SAME  |
|   C32   |      I32      |      2^28      | 487.881 us |       0.51% | 489.613 us |       0.56% |  1.732 us |   0.36% | 🔵 SAME  |
|   C32   |      I64      |      2^4       |   4.004 us |       3.48% |  10.164 us |       3.86% |  6.160 us | 153.83% | 🔴 SLOW  |
|   C32   |      I64      |      2^8       |   3.879 us |       4.34% |  10.260 us |       4.04% |  6.381 us | 164.48% | 🔴 SLOW  |
|   C32   |      I64      |      2^12      |   7.099 us |       5.80% |  10.176 us |       4.55% |  3.077 us |  43.34% | 🔴 SLOW  |
|   C32   |      I64      |      2^16      |   7.316 us |       4.75% |  10.025 us |       3.71% |  2.709 us |  37.03% | 🔴 SLOW  |
|   C32   |      I64      |      2^20      |  10.478 us |       3.69% |  10.954 us |       4.13% |  0.477 us |   4.55% | 🔴 SLOW  |
|   C32   |      I64      |      2^24      |  47.548 us |       1.53% |  47.938 us |       1.51% |  0.390 us |   0.82% | 🔵 SAME  |
|   C32   |      I64      |      2^28      | 489.389 us |       0.53% | 489.802 us |       0.54% |  0.413 us |   0.08% | 🔵 SAME  |

Not-guaranteed reduce

## [0] NVIDIA H200

|  T{ct}  |  OffsetT{ct}  |  Elements{io}  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |      Diff |   %Diff |  Status  |
|---------|---------------|----------------|------------|-------------|------------|-------------|-----------|---------|----------|
|   I32   |      I32      |      2^4       |   5.011 us |       3.16% |   7.366 us |       1.84% |  2.355 us |  46.99% | 🔴 SLOW  |
|   I32   |      I32      |      2^8       |   4.855 us |       3.05% |   7.576 us |       1.80% |  2.721 us |  56.05% | 🔴 SLOW  |
|   I32   |      I32      |      2^12      |   5.294 us |       3.23% |   7.580 us |       1.77% |  2.286 us |  43.19% | 🔴 SLOW  |
|   I32   |      I32      |      2^16      |   5.354 us |       3.30% |   7.543 us |       1.83% |  2.189 us |  40.90% | 🔴 SLOW  |
|   I32   |      I32      |      2^20      |   7.126 us |       2.57% |   7.753 us |       1.86% |  0.627 us |   8.79% | 🔴 SLOW  |
|   I32   |      I32      |      2^24      |  25.729 us |       2.04% |  25.983 us |       2.15% |  0.254 us |   0.99% | 🔵 SAME  |
|   I32   |      I32      |      2^28      | 249.298 us |       0.85% | 249.723 us |       0.79% |  0.425 us |   0.17% | 🔵 SAME  |
|   I32   |      I64      |      2^4       |   4.987 us |       2.88% |   6.968 us |       2.29% |  1.980 us |  39.71% | 🔴 SLOW  |
|   I32   |      I64      |      2^8       |   5.161 us |       3.25% |   6.917 us |       2.27% |  1.757 us |  34.04% | 🔴 SLOW  |
|   I32   |      I64      |      2^12      |   5.335 us |       3.31% |   6.885 us |       1.96% |  1.551 us |  29.07% | 🔴 SLOW  |
|   I32   |      I64      |      2^16      |   5.440 us |       2.93% |   6.926 us |       2.19% |  1.486 us |  27.31% | 🔴 SLOW  |
|   I32   |      I64      |      2^20      |   7.143 us |       2.94% |   7.413 us |       2.36% |  0.270 us |   3.78% | 🔴 SLOW  |
|   I32   |      I64      |      2^24      |  25.848 us |       1.97% |  26.099 us |       2.16% |  0.252 us |   0.97% | 🔵 SAME  |
|   I32   |      I64      |      2^28      | 249.348 us |       0.37% | 248.801 us |       0.38% | -0.547 us |  -0.22% | 🔵 SAME  |
|   I64   |      I32      |      2^4       |   5.031 us |       3.32% |   7.706 us |       1.82% |  2.675 us |  53.17% | 🔴 SLOW  |
|   I64   |      I32      |      2^8       |   5.085 us |       3.26% |   7.685 us |       1.79% |  2.600 us |  51.13% | 🔴 SLOW  |
|   I64   |      I32      |      2^12      |   5.431 us |       2.80% |   7.678 us |       1.78% |  2.247 us |  41.37% | 🔴 SLOW  |
|   I64   |      I32      |      2^16      |   5.644 us |       2.23% |   7.668 us |       1.79% |  2.024 us |  35.86% | 🔴 SLOW  |
|   I64   |      I32      |      2^20      |   8.503 us |       2.68% |   8.746 us |       2.57% |  0.243 us |   2.86% | 🔴 SLOW  |
|   I64   |      I32      |      2^24      |  43.003 us |       1.61% |  43.459 us |       1.54% |  0.456 us |   1.06% | 🔵 SAME  |
|   I64   |      I32      |      2^28      | 488.318 us |       0.51% | 488.198 us |       0.54% | -0.120 us |  -0.02% | 🔵 SAME  |
|   I64   |      I64      |      2^4       |   5.054 us |       3.34% |   7.755 us |       1.77% |  2.701 us |  53.44% | 🔴 SLOW  |
|   I64   |      I64      |      2^8       |   5.096 us |       3.12% |   7.744 us |       1.83% |  2.648 us |  51.96% | 🔴 SLOW  |
|   I64   |      I64      |      2^12      |   5.420 us |       2.92% |   7.702 us |       1.61% |  2.282 us |  42.11% | 🔴 SLOW  |
|   I64   |      I64      |      2^16      |   5.647 us |       2.52% |   7.695 us |       1.90% |  2.048 us |  36.26% | 🔴 SLOW  |
|   I64   |      I64      |      2^20      |   8.416 us |       2.52% |   8.576 us |       2.84% |  0.160 us |   1.90% | 🔵 SAME  |
|   I64   |      I64      |      2^24      |  42.954 us |       1.64% |  43.495 us |       1.57% |  0.542 us |   1.26% | 🔵 SAME  |
|   I64   |      I64      |      2^28      | 488.282 us |       0.54% | 482.597 us |       0.52% | -5.685 us |  -1.16% | 🟢 FAST  |
|   F32   |      I32      |      2^4       |   4.952 us |       2.93% |   7.611 us |       1.46% |  2.660 us |  53.72% | 🔴 SLOW  |
|   F32   |      I32      |      2^8       |   5.203 us |       3.02% |   7.597 us |       2.15% |  2.394 us |  46.02% | 🔴 SLOW  |
|   F32   |      I32      |      2^12      |   5.371 us |       3.04% |   7.499 us |       1.91% |  2.128 us |  39.62% | 🔴 SLOW  |
|   F32   |      I32      |      2^16      |   5.420 us |       2.86% |   7.558 us |       1.89% |  2.138 us |  39.45% | 🔴 SLOW  |
|   F32   |      I32      |      2^20      |   7.212 us |       2.48% |   7.661 us |       2.23% |  0.449 us |   6.23% | 🔴 SLOW  |
|   F32   |      I32      |      2^24      |  25.858 us |       2.13% |  26.046 us |       2.08% |  0.188 us |   0.73% | 🔵 SAME  |
|   F32   |      I32      |      2^28      | 249.570 us |       0.87% | 249.801 us |       0.87% |  0.232 us |   0.09% | 🔵 SAME  |
|   F32   |      I64      |      2^4       |   5.258 us |       3.02% |   7.528 us |       1.59% |  2.270 us |  43.18% | 🔴 SLOW  |
|   F32   |      I64      |      2^8       |   5.022 us |       3.08% |   7.657 us |       1.63% |  2.635 us |  52.47% | 🔴 SLOW  |
|   F32   |      I64      |      2^12      |   5.413 us |       2.94% |   7.549 us |       2.11% |  2.136 us |  39.47% | 🔴 SLOW  |
|   F32   |      I64      |      2^16      |   5.484 us |       3.06% |   7.641 us |       1.58% |  2.158 us |  39.35% | 🔴 SLOW  |
|   F32   |      I64      |      2^20      |   7.204 us |       2.79% |   7.763 us |       2.23% |  0.559 us |   7.77% | 🔴 SLOW  |
|   F32   |      I64      |      2^24      |  25.980 us |       2.03% |  26.179 us |       2.09% |  0.200 us |   0.77% | 🔵 SAME  |
|   F32   |      I64      |      2^28      | 249.918 us |       0.38% | 249.166 us |       0.36% | -0.753 us |  -0.30% | 🔵 SAME  |
|   F64   |      I32      |      2^4       |   5.286 us |       2.86% |   7.510 us |       1.83% |  2.224 us |  42.07% | 🔴 SLOW  |
|   F64   |      I32      |      2^8       |   5.284 us |       2.74% |   7.537 us |       1.94% |  2.253 us |  42.63% | 🔴 SLOW  |
|   F64   |      I32      |      2^12      |   5.403 us |       2.72% |   7.610 us |       1.53% |  2.206 us |  40.83% | 🔴 SLOW  |
|   F64   |      I32      |      2^16      |   5.674 us |       2.67% |   7.565 us |       1.63% |  1.891 us |  33.33% | 🔴 SLOW  |
|   F64   |      I32      |      2^20      |   8.474 us |       2.46% |   8.673 us |       2.66% |  0.199 us |   2.35% | 🔵 SAME  |
|   F64   |      I32      |      2^24      |  43.066 us |       1.59% |  43.354 us |       1.58% |  0.288 us |   0.67% | 🔵 SAME  |
|   F64   |      I32      |      2^28      | 488.468 us |       0.56% | 487.513 us |       0.51% | -0.956 us |  -0.20% | 🔵 SAME  |
|   F64   |      I64      |      2^4       |   5.143 us |       2.69% |   7.626 us |       1.87% |  2.483 us |  48.28% | 🔴 SLOW  |
|   F64   |      I64      |      2^8       |   5.075 us |       3.32% |   7.620 us |       1.84% |  2.544 us |  50.13% | 🔴 SLOW  |
|   F64   |      I64      |      2^12      |   5.393 us |       2.98% |   7.618 us |       1.72% |  2.225 us |  41.25% | 🔴 SLOW  |
|   F64   |      I64      |      2^16      |   5.617 us |       2.71% |   7.561 us |       1.64% |  1.944 us |  34.61% | 🔴 SLOW  |
|   F64   |      I64      |      2^20      |   8.552 us |       2.63% |   8.710 us |       2.65% |  0.158 us |   1.85% | 🔵 SAME  |
|   F64   |      I64      |      2^24      |  43.210 us |       1.62% |  43.686 us |       1.60% |  0.476 us |   1.10% | 🔵 SAME  |
|   F64   |      I64      |      2^28      | 488.758 us |       0.54% | 483.060 us |       0.50% | -5.698 us |  -1.17% | 🟢 FAST  |

Checklist

New or existing tests cover these changes.
The documentation is up to date with these changes.

coderabbitai · 2026-07-07T01:16:22Z

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

@coderabbitai resume to resume automatic reviews.
@coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

▶️ Resume reviews
🔍 Trigger review

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 44665a28-fe79-4bf4-85dc-aa808a92133c

📥 Commits

Reviewing files that changed from the base of the PR and between cd93de7 and 8851b0f.

📒 Files selected for processing (1)

cub/cub/device/dispatch/dispatch_reduce.cuh

🚧 Files skipped from review as they are similar to previous changes (1)

cub/cub/device/dispatch/dispatch_reduce.cuh

📝 Walkthrough

Summary by CodeRabbit

New Features
- DeviceReduce now supports deferred, device-resident num_items across Reduce/Sum/Min/Max/TransformReduce, including stream-ordered reading and CUDA Graph usage.
- Added NVBench benchmarks for deferred reduce execution, including nondeterministic coverage.
Bug Fixes
- Improved deferred problem-size handling across stable and non-stable reduction paths, including correct dispatch/kernel behavior.
- Deferred problem sizes are rejected when using gpu_to_gpu determinism.
Documentation
- Added a “Deferred problem sizes” section covering requirements and limitations.
Tests
- Added deferred-count DeviceReduce test suite (including determinism modes, coordination, and Graph replay).
- Updated kernel-selection instantiations for environment-related reduce tests.

Walkthrough

Adds deferred device-resident num_items support for cub::DeviceReduce, threads it through dispatch and kernels, updates JIT symbol naming, and adds benchmarks and tests for deferred reduce paths.

Changes

Deferred num_items for DeviceReduce

Layer / File(s)	Summary
Deferred parameter helpers `cub/cub/detail/deferred_parameter.cuh`	Adds internal helpers for normalizing deferred parameters and resolving normalized values to a target type.
DeviceReduce public API `cub/cub/device/device_reduce.cuh`	Adds deferred `num_items` handling to DeviceReduce docs and dispatch call sites for `Reduce`, `Sum`, `Min`, `Max`, and `TransformReduce`.
Dispatch normalization and kernel wiring `cub/cub/device/dispatch/dispatch_reduce.cuh`	Adds normalized `num_items` helpers and updates dispatch paths, grid sizing, and kernel source wiring for deferred sizing.
Kernel-level deferred sizing `cub/cub/device/dispatch/kernels/kernel_reduce.cuh`	Updates `DeviceReduceKernel` to accept normalized problem sizes and adds a deferred single-tile kernel variant.
JIT kernel name update `c/parallel/src/reduce.cu`	Updates the `DeviceReduceKernel` symbol template and arguments used for JIT compilation and linking.
Deferred reduce benchmarks `cub/benchmarks/bench/reduce/deferred_sum.cu`, `cub/benchmarks/bench/reduce/deferred_nondeterministic.cu`	Adds nvbench benchmarks for deferred `DeviceReduce::Reduce` under deterministic and not-guaranteed execution paths.
Deferred num_items tests `cub/test/catch2_test_device_reduce_deferred.cu`, `cub/test/catch2_test_device_reduce_env.cu`	Adds deferred-num_items coverage across direct calls, selection pipelines, stream/event handoff, CUDA graph replay, environment entry points, and kernel-selection instantiations.

Assessment against linked issues

Objective	Addressed	Explanation
Implement device-resident problem size support for run-to-run and not-guaranteed determinism [`#9561`]	✅
Provide performance numbers/benchmark results [`#9561`]	❌	Benchmarks were added, but no performance numbers or results are included.

Out-of-scope changes

Code Change	Explanation
`DeviceReduce::Sum`, `Min`, `Max`, and `TransformReduce` now accept deferred `num_items` (`cub/cub/device/device_reduce.cuh:775-786, 879-899, 1526-1545, 2131-2142`)	The issue asks for proof-of-concept support in run-to-run and not-guaranteed reduce paths, not expansion to these additional entry points.
Deferred-size documentation and CUDA Graph/stream lifetime requirements were added (`cub/cub/device/device_reduce.cuh:126-146`)	The issue requests implementation and performance numbers, not documentation changes.
Selection-pipeline and graph replay tests for deferred counts (`cub/test/catch2_test_device_reduce_deferred.cu:115-699`)	These cover broader deferred-count flows beyond the stated reduce determinism proof of concept.

Suggested reviewers: fbusato, bernhardmgruber

Warning

Review ran into problems

🔥 Problems

Git: Failed to clone repository. Please run the @coderabbitai full review command to re-trigger a full review. If the issue persists, set path_filters to include or exclude specific files.

_{Comment @coderabbitai help to get the list of available commands.}

coderabbitai

Actionable comments posted: 1

🧹 Nitpick comments (1)

cub/cub/detail/deferred_parameter.cuh (1)
50-50: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: resolve_parameter is device-only and cannot throw, but is missing noexcept while its siblings normalize_parameter and resolve_num_items have it. Mark it noexcept.
-  [[nodiscard]] _CCCL_DEVICE _CCCL_FORCEINLINE TargetT resolve_parameter(ParameterT parameter)
+  [[nodiscard]] _CCCL_DEVICE _CCCL_FORCEINLINE TargetT resolve_parameter(ParameterT parameter) noexcept
As per coding guidelines: "Functions that do not throw exceptions must be marked noexcept."

Source: Coding guidelines

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 305953b0-8b67-4c17-943a-5463e090b0b1

📥 Commits

Reviewing files that changed from the base of the PR and between 48ff8d1 and 9a963e9.

📒 Files selected for processing (9)

c/parallel/src/reduce.cu
cub/benchmarks/bench/reduce/deferred_nondeterministic.cu
cub/benchmarks/bench/reduce/deferred_sum.cu
cub/cub/detail/deferred_parameter.cuh
cub/cub/device/device_reduce.cuh
cub/cub/device/dispatch/dispatch_reduce.cuh
cub/cub/device/dispatch/kernels/kernel_reduce.cuh
cub/test/catch2_test_device_reduce_deferred.cu
cub/test/catch2_test_device_reduce_env.cu

Jacobfaib

Misc stuff LGTM

github-actions · 2026-07-09T17:41:03Z

🥳 CI Workflow Results

🟩 Finished in 4h 28m: Pass: 100%/291 | Total: 12d 05h | Max: 2h 34m | Hits: 20%/972783

See results here.

NaderAlAwar added 5 commits July 2, 2026 16:44

Initial deferred num items implementation

736d434

Add tests and benchmarks

3596659

Make benchmarks consistent and update tests

ecd0399

Fix sass differences

2decb22

pre-commit

9a963e9

NaderAlAwar requested review from a team as code owners July 7, 2026 01:04

NaderAlAwar requested a review from oleksandr-pavlyk July 7, 2026 01:04

github-project-automation Bot added this to CCCL Jul 7, 2026

NaderAlAwar requested a review from fbusato July 7, 2026 01:04

github-project-automation Bot moved this to Todo in CCCL Jul 7, 2026

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

coderabbitai Bot reviewed Jul 7, 2026

View reviewed changes

Comment thread cub/benchmarks/bench/reduce/deferred_sum.cu Outdated

NaderAlAwar added 4 commits July 6, 2026 20:33

Merge branch 'main' into device-reduce-deferred-num-items

af8fedc

Address comments

663ec1b

Fix CI error

b3b6097

Fix CI

af33be0

This comment has been minimized.

Sign in to view

NaderAlAwar added 2 commits July 7, 2026 08:02

Merge branch 'main' into device-reduce-deferred-num-items

5c2ef8d

MSVC fix

0be8893

Jacobfaib reviewed Jul 7, 2026

View reviewed changes

Comment thread cub/cub/device/dispatch/kernels/kernel_reduce.cuh Outdated

Comment thread cub/cub/device/dispatch/kernels/kernel_reduce.cuh Outdated

Comment thread cub/cub/device/dispatch/dispatch_reduce.cuh Outdated

Comment thread cub/cub/device/dispatch/dispatch_reduce.cuh Outdated

This comment has been minimized.

Sign in to view

Address comments

057a85b

This comment has been minimized.

Sign in to view

Merge branch 'main' into device-reduce-deferred-num-items

123c142

NaderAlAwar mentioned this pull request Jul 7, 2026

[CUB] GPU-to-GPU DeviceReduce with device resident problem size #9740

Open

2 tasks

This comment has been minimized.

Sign in to view

NaderAlAwar requested a review from Jacobfaib July 8, 2026 13:01

Jacobfaib approved these changes Jul 8, 2026

View reviewed changes

Comment thread cub/cub/detail/deferred_parameter.cuh Outdated

fbusato reviewed Jul 8, 2026

View reviewed changes

address reviews

f824032

NaderAlAwar requested a review from fbusato July 8, 2026 21:36

fix CI

4b19824

fbusato approved these changes Jul 8, 2026

View reviewed changes

Comment thread cub/cub/detail/deferred_parameter.cuh Outdated

address reviews

cd93de7

NaderAlAwar enabled auto-merge (squash) July 8, 2026 21:59

This comment has been minimized.

Sign in to view

bernhardmgruber reviewed Jul 9, 2026

View reviewed changes

Comment thread cub/cub/device/dispatch/dispatch_reduce.cuh

address review

8851b0f

NaderAlAwar requested a review from bernhardmgruber July 9, 2026 13:06

shwina approved these changes Jul 9, 2026

View reviewed changes

This comment has been minimized.

Sign in to view

NaderAlAwar merged commit d47ca8e into NVIDIA:main Jul 9, 2026
613 of 621 checks passed

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

This was referenced Jul 10, 2026

Reduce duplication in reduce kernels if extracting common code to device function does not hurt performance #9789

Open

[EPIC]: Support device-resident problem sizes in CUB device-level algorithms #8789

Open

Uh oh!

Conversation

NaderAlAwar commented Jul 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Checklist

Uh oh!

coderabbitai Bot commented Jul 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Reviews paused

Summary by CodeRabbit

Walkthrough

Changes

Assessment against linked issues

Review ran into problems

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

This comment has been minimized.

This comment has been minimized.

Jacobfaib left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

This comment has been minimized.

github-actions Bot commented Jul 9, 2026

🥳 CI Workflow Results

🟩 Finished in 4h 28m: Pass: 100%/291 | Total: 12d 05h | Max: 2h 34m | Hits: 20%/972783

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

NaderAlAwar commented Jul 7, 2026 •

edited

Loading

coderabbitai Bot commented Jul 7, 2026 •

edited

Loading