Improvements to: Use BlockLoadToShared in DeviceMerge #6077 #6460

bernhardmgruber · 2025-11-03T19:22:34Z

If BlockLoadToShared is disabled, this PR has no SASS differences to #6077 on sm90.

Perf on Blackwell is amazing!

Perf and tunings on Ampere look ok. I am willing to take the few regressions for improvements elsewhere.

Perf and tunings on Hopper look okish. There are some regressions mixed with improvements where I am not entirely convinced we should take them. Waiting for reviewer feedback.

Summary of all regressions:

pairs on A100
|     I8     |     I32      |      I32      |      2^16      |     1     |  19.891 us |       2.62% |  20.622 us |       2.34% |    0.731 us |   3.68% |   SLOW   |
|     I8     |     I32      |      I32      |      2^16      |   0.201   |  19.256 us |       2.63% |  20.072 us |       2.60% |    0.816 us |   4.24% |   SLOW   |
|     I8     |     I32      |      I64      |      2^16      |     1     |  19.897 us |       2.62% |  20.575 us |       2.21% |    0.678 us |   3.41% |   SLOW   |
|     I8     |     I32      |      I64      |      2^16      |   0.201   |  19.537 us |       2.42% |  20.251 us |       2.58% |    0.715 us |   3.66% |   SLOW   |
|     I8     |     I64      |      I32      |      2^16      |     1     |  18.769 us |       2.63% |  19.808 us |       2.52% |    1.039 us |   5.53% |   SLOW   |
|     I8     |     I64      |      I32      |      2^16      |   0.201   |  18.638 us |       2.55% |  19.789 us |       2.61% |    1.152 us |   6.18% |   SLOW   |
|     I8     |     I64      |      I64      |      2^16      |     1     |  19.354 us |       2.72% |  20.215 us |       2.56% |    0.861 us |   4.45% |   SLOW   |
|     I8     |     I64      |      I64      |      2^16      |   0.201   |  19.059 us |       2.92% |  20.212 us |       2.50% |    1.153 us |   6.05% |   SLOW   |
|    I32     |      I8      |      I32      |      2^16      |     1     |  21.357 us |       2.49% |  22.041 us |       2.64% |    0.683 us |   3.20% |   SLOW   |
|    I32     |      I8      |      I32      |      2^16      |   0.201   |  20.507 us |       2.27% |  21.337 us |       2.48% |    0.830 us |   4.05% |   SLOW   |
|    I32     |      I8      |      I64      |      2^16      |     1     |  21.703 us |       2.30% |  22.307 us |       2.23% |    0.604 us |   2.78% |   SLOW   |
|    I32     |      I8      |      I64      |      2^16      |   0.201   |  21.039 us |       3.38% |  21.941 us |       2.71% |    0.902 us |   4.29% |   SLOW   |
|    I32     |     I16      |      I32      |      2^16      |     1     |  20.828 us |       2.59% |  21.562 us |       2.38% |    0.734 us |   3.52% |   SLOW   |
|    I32     |     I16      |      I32      |      2^20      |     1     |  30.818 us |       2.27% |  31.578 us |       2.52% |    0.760 us |   2.47% |   SLOW   |
|    I64     |     I32      |      I64      |      2^20      |   0.201   |  39.912 us |       1.48% |  40.762 us |       1.61% |    0.850 us |   2.13% |   SLOW   |
|    F32     |      I8      |      I32      |      2^16      |   0.201   |  21.271 us |       3.19% |  22.294 us |       3.74% |    1.023 us |   4.81% |   SLOW   |
|    F32     |      I8      |      I64      |      2^16      |     1     |  21.998 us |       2.66% |  23.316 us |       3.47% |    1.318 us |   5.99% |   SLOW   |
|    F32     |     I16      |      I32      |      2^16      |     1     |  21.222 us |       3.15% |  21.993 us |       2.91% |    0.771 us |   3.63% |   SLOW   |
|    F32     |     I16      |      I64      |      2^20      |     1     |  31.649 us |       2.40% |  32.601 us |       1.35% |    0.953 us |   3.01% |   SLOW   |

keys on H200
|    I128    |      I32      |      2^20      |   0.201   |  29.638 us |       2.30% |  30.614 us |       2.07% |    0.976 us |   3.29% |   SLOW   |
|    I128    |      I32      |      2^24      |   0.201   | 165.223 us |       0.92% | 168.570 us |       0.85% |    3.347 us |   2.03% |   SLOW   |
|    I128    |      I64      |      2^20      |   0.201   |  29.976 us |       2.25% |  30.916 us |       1.87% |    0.940 us |   3.14% |   SLOW   |

pairs on H200
|    I32     |     I64      |      I64      |      2^16      |   0.201   |  16.660 us |       3.56% |  17.505 us |       2.77% |    0.845 us |   5.07% |   SLOW   |
|    F32     |     I64      |      I64      |      2^16      |     1     |  17.377 us |       3.09% |  17.898 us |       2.79% |    0.521 us |   3.00% |   SLOW   |
|    F32     |     I64      |      I64      |      2^16      |   0.201   |  16.541 us |       2.94% |  17.036 us |       3.51% |    0.495 us |   2.99% |   SLOW   |
|    F64     |     I32      |      I32      |      2^16      |   0.201   |  17.472 us |       3.48% |  18.169 us |       2.85% |    0.697 us |   3.99% |   SLOW   |

keys on B200
|    F64     |      I64      |      2^16      |     1     |  18.428 us |       1.26% |  19.081 us |       5.82% |    0.653 us |   3.54% |   SLOW   |

pairs on B200
|     I8     |     I64      |      I32      |      2^16      |     1     |  17.509 us |       5.81% |  18.284 us |       2.76% |    0.775 us |   4.43% |   SLOW   |
|     I8     |     I64      |      I64      |      2^16      |     1     |  17.377 us |       5.83% |  18.099 us |       3.95% |    0.722 us |   4.15% |   SLOW   |
|    I32     |     I32      |      I64      |      2^16      |   0.201   |  18.238 us |       1.59% |  18.613 us |       4.67% |    0.375 us |   2.06% |   SLOW   |
|    I64     |     I16      |      I32      |      2^16      |     1     |  19.996 us |       4.36% |  20.467 us |       0.22% |    0.471 us |   2.36% |   SLOW   |
|    F32     |     I32      |      I32      |      2^20      |     1     |  25.994 us |       3.60% |  26.611 us |       0.18% |    0.617 us |   2.37% |   SLOW   |
|    F64     |     I16      |      I32      |      2^16      |     1     |  20.017 us |       5.01% |  20.463 us |       0.54% |    0.446 us |   2.23% |   SLOW   |
|    F64     |     I32      |      I32      |      2^16      |     1     |  19.828 us |       4.92% |  20.511 us |       1.56% |    0.683 us |   3.45% |   SLOW   |
|    F64     |     I32      |      I64      |      2^16      |   0.201   |  18.273 us |       1.68% |  18.777 us |       5.10% |    0.505 us |   2.76% |   SLOW   |

copy-pr-bot · 2025-11-03T19:22:37Z

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.

bernhardmgruber · 2025-11-04T10:45:30Z

Edit: I later found a regression in #6077 that caused a higher register consumption and so occupancy dropped. We would have to reevaluate the content below (not important for now though):

I had a commit that inlined the translate_indices lambda into the loop where the items are gathered from shared memory, so the registers are not touched twice. That would no longer overlap the index translation with the bulk copy. This was faster in some cases, but slower in others. I think we could consider customizing this.

pairs on H200, translating indices while gathering and vs. separate translation during bulk copy, 2% threshold:

## [0] NVIDIA H200

|  KeyT{ct}  |  ValueT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|------------|--------------|---------------|----------------|-----------|------------|-------------|------------|-------------|-------------|---------|----------|
|     I8     |      I8      |      I32      |      2^28      |     1     | 860.779 us |       0.14% | 881.686 us |       0.12% |   20.907 us |   2.43% |   SLOW   |
|     I8     |     I16      |      I32      |      2^24      |     1     |  74.358 us |       0.74% |  75.882 us |       0.73% |    1.524 us |   2.05% |   SLOW   |
|     I8     |     I16      |      I32      |      2^28      |     1     | 889.974 us |       0.13% | 914.533 us |       0.14% |   24.559 us |   2.76% |   SLOW   |
|     I8     |     I16      |      I32      |      2^24      |   0.201   |  73.406 us |       0.82% |  74.965 us |       0.75% |    1.559 us |   2.12% |   SLOW   |
|     I8     |     I16      |      I32      |      2^28      |   0.201   | 869.890 us |       0.11% | 896.038 us |       0.11% |   26.147 us |   3.01% |   SLOW   |
|     I8     |     I16      |      I64      |      2^24      |     1     |  75.109 us |       0.74% |  76.808 us |       0.75% |    1.699 us |   2.26% |   SLOW   |
|     I8     |     I16      |      I64      |      2^28      |     1     | 898.696 us |       0.16% | 926.100 us |       0.17% |   27.404 us |   3.05% |   SLOW   |
|     I8     |     I16      |      I64      |      2^24      |   0.201   |  73.909 us |       0.78% |  75.649 us |       0.71% |    1.740 us |   2.35% |   SLOW   |
|     I8     |     I16      |      I64      |      2^28      |   0.201   | 880.266 us |       0.12% | 911.406 us |       0.13% |   31.140 us |   3.54% |   SLOW   |
|    I16     |     I16      |      I32      |      2^28      |     1     |   1.024 ms |       0.10% |   1.048 ms |       0.10% |   24.643 us |   2.41% |   SLOW   |
|    I16     |     I16      |      I32      |      2^24      |   0.201   |  80.755 us |       0.72% |  82.551 us |       0.71% |    1.795 us |   2.22% |   SLOW   |
|    I16     |     I16      |      I32      |      2^28      |   0.201   | 972.396 us |       0.37% | 998.223 us |       0.34% |   25.827 us |   2.66% |   SLOW   |
|    I16     |     I16      |      I64      |      2^28      |     1     |   1.035 ms |       0.13% |   1.060 ms |       0.11% |   24.693 us |   2.38% |   SLOW   |
|    I16     |     I16      |      I64      |      2^28      |   0.201   | 990.923 us |       0.45% |   1.014 ms |       0.38% |   22.791 us |   2.30% |   SLOW   |
|    I32     |      I8      |      I32      |      2^24      |     1     |  77.941 us |       1.05% |  90.412 us |       1.09% |   12.471 us |  16.00% |   SLOW   |
|    I32     |      I8      |      I32      |      2^28      |     1     | 899.882 us |       0.31% |   1.158 ms |       0.69% |  258.088 us |  28.68% |   SLOW   |
|    I32     |      I8      |      I32      |      2^24      |   0.201   |  73.969 us |       4.24% |  84.392 us |       1.26% |   10.423 us |  14.09% |   SLOW   |
|    I32     |      I8      |      I32      |      2^28      |   0.201   | 842.219 us |       0.37% |   1.079 ms |       0.53% |  237.096 us |  28.15% |   SLOW   |
|    I32     |      I8      |      I64      |      2^24      |     1     |  94.693 us |       1.47% |  79.804 us |       0.79% |  -14.890 us | -15.72% |   FAST   |
|    I32     |      I8      |      I64      |      2^28      |     1     |   1.229 ms |       0.52% | 922.911 us |       0.30% | -305.596 us | -24.88% |   FAST   |
|    I32     |      I8      |      I64      |      2^24      |   0.201   |  86.803 us |       1.34% |  75.024 us |       0.89% |  -11.779 us | -13.57% |   FAST   |
|    I32     |      I8      |      I64      |      2^28      |   0.201   |   1.095 ms |       0.48% | 860.327 us |       0.36% | -234.388 us | -21.41% |   FAST   |
|    I128    |      I8      |      I64      |      2^16      |   0.201   |  17.472 us |       2.89% |  17.924 us |       2.53% |    0.452 us |   2.58% |   SLOW   |
|    F32     |      I8      |      I32      |      2^24      |     1     |  77.802 us |       0.83% |  90.072 us |       1.15% |   12.270 us |  15.77% |   SLOW   |
|    F32     |      I8      |      I32      |      2^28      |     1     | 900.433 us |       0.30% |   1.159 ms |       0.67% |  258.293 us |  28.69% |   SLOW   |
|    F32     |      I8      |      I32      |      2^24      |   0.201   |  73.650 us |       0.99% |  85.415 us |       1.26% |   11.766 us |  15.98% |   SLOW   |
|    F32     |      I8      |      I32      |      2^28      |   0.201   | 842.622 us |       0.37% |   1.079 ms |       0.50% |  235.979 us |  28.01% |   SLOW   |
|    F32     |      I8      |      I64      |      2^24      |     1     |  94.450 us |       1.33% |  79.389 us |       0.85% |  -15.062 us | -15.95% |   FAST   |
|    F32     |      I8      |      I64      |      2^28      |     1     |   1.227 ms |       0.51% | 922.699 us |       0.29% | -303.833 us | -24.77% |   FAST   |
|    F32     |      I8      |      I64      |      2^24      |   0.201   |  86.526 us |       1.22% |  75.365 us |       0.92% |  -11.160 us | -12.90% |   FAST   |
|    F32     |      I8      |      I64      |      2^28      |   0.201   |   1.095 ms |       0.47% | 860.671 us |       0.35% | -234.038 us | -21.38% |   FAST   |
|    F64     |      I8      |      I64      |      2^16      |     1     |  17.288 us |       3.05% |  17.736 us |       2.85% |    0.448 us |   2.59% |   SAME   |

So it seems for I8 value and I64 offset, we should translate indices while gathering.

copy-pr-bot · 2025-11-05T01:14:48Z

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

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

bernhardmgruber · 2025-11-05T10:59:54Z

pairs on H200, #6077 vs this PR, 2% threshold (NOT the final perf after merging to main, look below)::

## [0] NVIDIA H200 NVL

|  KeyT{ct}  |  ValueT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|------------|--------------|---------------|----------------|-----------|------------|-------------|------------|-------------|-------------|---------|----------|
|    I32     |     I64      |      I32      |      2^16      |     1     |  17.235 us |       2.92% |  17.651 us |       2.34% |    0.416 us |   2.41% |   SLOW   |
|    I32     |     I64      |      I64      |      2^16      |   0.201   |  16.908 us |       3.34% |  17.505 us |       2.77% |    0.597 us |   3.53% |   SLOW   |
|    I64     |     I16      |      I32      |      2^16      |     1     |  18.723 us |       2.84% |  18.274 us |       3.29% |   -0.449 us |  -2.40% |   SAME   |
|    I64     |     I64      |      I32      |      2^24      |   0.201   | 174.576 us |       0.63% | 166.990 us |       0.69% |   -7.586 us |  -4.35% |   FAST   |
|    I64     |     I64      |      I32      |      2^28      |   0.201   |   2.516 ms |       0.56% |   2.402 ms |       0.59% | -113.495 us |  -4.51% |   FAST   |
|    I64     |     I64      |      I64      |      2^20      |   0.201   |  30.033 us |       2.28% |  29.219 us |       2.04% |   -0.814 us |  -2.71% |   FAST   |
|    I64     |     I64      |      I64      |      2^24      |   0.201   | 175.519 us |       0.63% | 167.780 us |       0.67% |   -7.739 us |  -4.41% |   FAST   |
|    I64     |     I64      |      I64      |      2^28      |   0.201   |   2.524 ms |       0.56% |   2.409 ms |       0.61% | -115.931 us |  -4.59% |   FAST   |
|    I128    |      I8      |      I64      |      2^16      |     1     |  20.365 us |       2.72% |  19.897 us |       3.05% |   -0.468 us |  -2.30% |   SAME   |
|    I128    |      I8      |      I64      |      2^20      |   0.201   |  33.627 us |       1.86% |  32.566 us |       2.16% |   -1.061 us |  -3.16% |   FAST   |
|    I128    |      I8      |      I64      |      2^24      |   0.201   | 206.912 us |       0.66% | 198.574 us |       0.68% |   -8.337 us |  -4.03% |   FAST   |
|    I128    |      I8      |      I64      |      2^28      |   0.201   |   2.990 ms |       0.48% |   2.870 ms |       0.59% | -120.169 us |  -4.02% |   FAST   |
|    I128    |     I16      |      I32      |      2^20      |   0.201   |  33.973 us |       1.85% |  32.770 us |       2.23% |   -1.204 us |  -3.54% |   FAST   |
|    I128    |     I16      |      I32      |      2^24      |   0.201   | 213.634 us |       0.61% | 203.837 us |       0.62% |   -9.797 us |  -4.59% |   FAST   |
|    I128    |     I16      |      I32      |      2^28      |   0.201   |   3.109 ms |       0.47% |   2.953 ms |       0.62% | -155.956 us |  -5.02% |   FAST   |
|    I128    |     I16      |      I64      |      2^20      |   0.201   |  33.868 us |       2.13% |  33.016 us |       1.97% |   -0.852 us |  -2.52% |   FAST   |
|    I128    |     I16      |      I64      |      2^24      |   0.201   | 213.796 us |       0.64% | 206.954 us |       0.60% |   -6.842 us |  -3.20% |   FAST   |
|    I128    |     I16      |      I64      |      2^28      |   0.201   |   3.101 ms |       0.51% |   3.005 ms |       0.60% |  -95.195 us |  -3.07% |   FAST   |
|    I128    |     I32      |      I32      |      2^20      |   0.201   |  34.591 us |       2.00% |  33.274 us |       2.22% |   -1.317 us |  -3.81% |   FAST   |
|    I128    |     I32      |      I32      |      2^24      |   0.201   | 221.274 us |       0.60% | 213.361 us |       0.62% |   -7.913 us |  -3.58% |   FAST   |
|    I128    |     I32      |      I32      |      2^28      |   0.201   |   3.225 ms |       0.57% |   3.113 ms |       0.66% | -111.218 us |  -3.45% |   FAST   |
|    I128    |     I32      |      I64      |      2^16      |   0.201   |  19.455 us |       2.20% |  19.044 us |       2.22% |   -0.411 us |  -2.11% |   SAME   |
|    I128    |     I32      |      I64      |      2^20      |   0.201   |  34.398 us |       2.03% |  33.572 us |       1.94% |   -0.825 us |  -2.40% |   FAST   |
|    I128    |     I32      |      I64      |      2^24      |   0.201   | 222.366 us |       0.61% | 215.657 us |       0.59% |   -6.710 us |  -3.02% |   FAST   |
|    I128    |     I32      |      I64      |      2^28      |   0.201   |   3.240 ms |       0.57% |   3.149 ms |       0.60% |  -90.502 us |  -2.79% |   FAST   |
|    F32     |      I8      |      I32      |      2^16      |   0.201   |  16.953 us |       3.08% |  17.367 us |       2.76% |    0.413 us |   2.44% |   SAME   |
|    F32     |     I64      |      I32      |      2^16      |   0.201   |  17.041 us |       2.62% |  16.632 us |       3.19% |   -0.409 us |  -2.40% |   SAME   |
|    F64     |      I8      |      I64      |      2^16      |     1     |  18.492 us |       2.80% |  18.101 us |       2.84% |   -0.391 us |  -2.11% |   SAME   |
|    F64     |     I64      |      I32      |      2^20      |   0.201   |  29.848 us |       2.10% |  29.054 us |       2.14% |   -0.794 us |  -2.66% |   FAST   |
|    F64     |     I64      |      I32      |      2^24      |   0.201   | 174.619 us |       0.66% | 166.494 us |       0.67% |   -8.125 us |  -4.65% |   FAST   |
|    F64     |     I64      |      I32      |      2^28      |   0.201   |   2.515 ms |       0.60% |   2.398 ms |       0.57% | -116.653 us |  -4.64% |   FAST   |
|    F64     |     I64      |      I64      |      2^16      |     1     |  18.905 us |       2.86% |  18.338 us |       3.08% |   -0.567 us |  -3.00% |   FAST   |
|    F64     |     I64      |      I64      |      2^16      |   0.201   |  17.755 us |       2.50% |  17.352 us |       2.56% |   -0.403 us |  -2.27% |   SAME   |
|    F64     |     I64      |      I64      |      2^20      |   0.201   |  29.871 us |       2.39% |  29.084 us |       2.11% |   -0.786 us |  -2.63% |   FAST   |
|    F64     |     I64      |      I64      |      2^24      |   0.201   | 175.236 us |       0.62% | 167.228 us |       0.68% |   -8.008 us |  -4.57% |   FAST   |
|    F64     |     I64      |      I64      |      2^28      |   0.201   |   2.521 ms |       0.59% |   2.405 ms |       0.59% | -115.896 us |  -4.60% |   FAST   |

pairs on H200, main vs this PR, 2% threshold:

## [0] NVIDIA H200 NVL

|  KeyT{ct}  |  ValueT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|------------|--------------|---------------|----------------|-----------|------------|-------------|------------|-------------|-------------|---------|----------|
|     I8     |      I8      |      I32      |      2^16      |     1     |  17.335 us |       3.09% |  16.589 us |       3.14% |   -0.746 us |  -4.30% |   FAST   |
|     I8     |      I8      |      I32      |      2^20      |     1     |  28.025 us |       1.92% |  24.168 us |       2.28% |   -3.858 us | -13.76% |   FAST   |
|     I8     |      I8      |      I32      |      2^24      |     1     | 112.987 us |       0.41% |  80.148 us |       0.65% |  -32.839 us | -29.06% |   FAST   |
|     I8     |      I8      |      I32      |      2^28      |     1     |   1.457 ms |       0.14% | 941.832 us |       0.24% | -514.982 us | -35.35% |   FAST   |
|     I8     |      I8      |      I32      |      2^16      |   0.201   |  17.025 us |       2.65% |  16.363 us |       2.73% |   -0.662 us |  -3.89% |   FAST   |
|     I8     |      I8      |      I32      |      2^20      |   0.201   |  26.615 us |       2.11% |  23.430 us |       2.33% |   -3.185 us | -11.97% |   FAST   |
|     I8     |      I8      |      I32      |      2^24      |   0.201   | 111.511 us |       0.47% |  78.698 us |       0.62% |  -32.813 us | -29.43% |   FAST   |
|     I8     |      I8      |      I32      |      2^28      |   0.201   |   1.452 ms |       0.15% | 927.566 us |       0.20% | -524.006 us | -36.10% |   FAST   |
|     I8     |      I8      |      I64      |      2^16      |     1     |  17.508 us |       2.76% |  16.745 us |       2.72% |   -0.763 us |  -4.36% |   FAST   |
|     I8     |      I8      |      I64      |      2^20      |     1     |  28.308 us |       2.08% |  24.191 us |       2.39% |   -4.117 us | -14.54% |   FAST   |
|     I8     |      I8      |      I64      |      2^24      |     1     | 114.841 us |       0.43% |  79.956 us |       0.63% |  -34.885 us | -30.38% |   FAST   |
|     I8     |      I8      |      I64      |      2^28      |     1     |   1.492 ms |       0.15% | 927.009 us |       0.19% | -564.798 us | -37.86% |   FAST   |
|     I8     |      I8      |      I64      |      2^16      |   0.201   |  17.105 us |       2.71% |  16.407 us |       2.86% |   -0.698 us |  -4.08% |   FAST   |
|     I8     |      I8      |      I64      |      2^20      |   0.201   |  26.730 us |       1.88% |  23.401 us |       1.93% |   -3.329 us | -12.45% |   FAST   |
|     I8     |      I8      |      I64      |      2^24      |   0.201   | 113.230 us |       0.61% |  78.120 us |       0.81% |  -35.110 us | -31.01% |   FAST   |
|     I8     |      I8      |      I64      |      2^28      |   0.201   |   1.487 ms |       0.14% | 920.487 us |       0.20% | -566.044 us | -38.08% |   FAST   |
|     I8     |     I16      |      I32      |      2^16      |     1     |  17.280 us |       3.26% |  16.297 us |       3.36% |   -0.983 us |  -5.69% |   FAST   |
|     I8     |     I16      |      I32      |      2^20      |     1     |  28.296 us |       1.99% |  23.876 us |       2.29% |   -4.420 us | -15.62% |   FAST   |
|     I8     |     I16      |      I32      |      2^24      |     1     | 116.143 us |       0.44% |  80.754 us |       0.70% |  -35.389 us | -30.47% |   FAST   |
|     I8     |     I16      |      I32      |      2^28      |     1     |   1.497 ms |       0.17% | 971.736 us |       0.24% | -525.282 us | -35.09% |   FAST   |
|     I8     |     I16      |      I32      |      2^16      |   0.201   |  16.786 us |       2.53% |  15.968 us |       2.71% |   -0.818 us |  -4.87% |   FAST   |
|     I8     |     I16      |      I32      |      2^20      |   0.201   |  26.769 us |       1.89% |  23.264 us |       2.10% |   -3.505 us | -13.09% |   FAST   |
|     I8     |     I16      |      I32      |      2^24      |   0.201   | 114.516 us |       0.66% |  79.530 us |       0.85% |  -34.986 us | -30.55% |   FAST   |
|     I8     |     I16      |      I32      |      2^28      |   0.201   |   1.490 ms |       0.19% | 951.763 us |       0.20% | -538.557 us | -36.14% |   FAST   |
|     I8     |     I16      |      I64      |      2^16      |     1     |  17.305 us |       3.10% |  16.433 us |       3.21% |   -0.872 us |  -5.04% |   FAST   |
|     I8     |     I16      |      I64      |      2^20      |     1     |  28.325 us |       1.89% |  23.965 us |       2.34% |   -4.360 us | -15.39% |   FAST   |
|     I8     |     I16      |      I64      |      2^24      |     1     | 116.659 us |       0.41% |  81.379 us |       0.64% |  -35.280 us | -30.24% |   FAST   |
|     I8     |     I16      |      I64      |      2^28      |     1     |   1.514 ms |       0.18% | 979.021 us |       0.21% | -534.608 us | -35.32% |   FAST   |
|     I8     |     I16      |      I64      |      2^16      |   0.201   |  16.821 us |       2.66% |  16.090 us |       2.90% |   -0.731 us |  -4.35% |   FAST   |
|     I8     |     I16      |      I64      |      2^20      |   0.201   |  26.844 us |       2.11% |  22.992 us |       2.44% |   -3.852 us | -14.35% |   FAST   |
|     I8     |     I16      |      I64      |      2^24      |   0.201   | 115.129 us |       0.54% |  80.130 us |       0.78% |  -35.000 us | -30.40% |   FAST   |
|     I8     |     I16      |      I64      |      2^28      |   0.201   |   1.507 ms |       0.20% | 961.206 us |       0.20% | -546.193 us | -36.23% |   FAST   |
|     I8     |     I32      |      I32      |      2^20      |     1     |  27.272 us |       1.91% |  23.339 us |       2.47% |   -3.933 us | -14.42% |   FAST   |
|     I8     |     I32      |      I32      |      2^24      |     1     | 129.403 us |       0.54% |  81.387 us |       0.85% |  -48.016 us | -37.11% |   FAST   |
|     I8     |     I32      |      I32      |      2^28      |     1     |   1.760 ms |       0.23% | 966.360 us |       0.26% | -793.962 us | -45.10% |   FAST   |
|     I8     |     I32      |      I32      |      2^20      |   0.201   |  25.893 us |       2.02% |  22.308 us |       2.03% |   -3.585 us | -13.85% |   FAST   |
|     I8     |     I32      |      I32      |      2^24      |   0.201   | 128.437 us |       0.86% |  80.533 us |       1.03% |  -47.904 us | -37.30% |   FAST   |
|     I8     |     I32      |      I32      |      2^28      |   0.201   |   1.742 ms |       0.27% | 960.113 us |       0.29% | -781.478 us | -44.87% |   FAST   |
|     I8     |     I32      |      I64      |      2^24      |     1     |  94.690 us |       0.68% |  82.880 us |       0.78% |  -11.810 us | -12.47% |   FAST   |
|     I8     |     I32      |      I64      |      2^28      |     1     |   1.201 ms |       0.34% | 988.747 us |       0.31% | -212.540 us | -17.69% |   FAST   |
|     I8     |     I32      |      I64      |      2^20      |   0.201   |  23.378 us |       2.52% |  22.537 us |       2.29% |   -0.841 us |  -3.60% |   FAST   |
|     I8     |     I32      |      I64      |      2^24      |   0.201   |  92.456 us |       0.91% |  82.073 us |       0.93% |  -10.384 us | -11.23% |   FAST   |
|     I8     |     I32      |      I64      |      2^28      |   0.201   |   1.208 ms |       0.41% | 992.727 us |       0.39% | -215.519 us | -17.84% |   FAST   |
|     I8     |     I64      |      I32      |      2^20      |     1     |  26.556 us |       2.06% |  24.733 us |       2.09% |   -1.823 us |  -6.87% |   FAST   |
|     I8     |     I64      |      I32      |      2^24      |     1     | 133.506 us |       0.94% | 108.338 us |       0.58% |  -25.167 us | -18.85% |   FAST   |
|     I8     |     I64      |      I32      |      2^28      |     1     |   1.914 ms |       0.19% |   1.423 ms |       0.33% | -491.088 us | -25.66% |   FAST   |
|     I8     |     I64      |      I32      |      2^20      |   0.201   |  25.180 us |       2.36% |  23.704 us |       2.66% |   -1.477 us |  -5.86% |   FAST   |
|     I8     |     I64      |      I32      |      2^24      |   0.201   | 134.733 us |       1.16% | 107.604 us |       0.48% |  -27.129 us | -20.14% |   FAST   |
|     I8     |     I64      |      I32      |      2^28      |   0.201   |   1.898 ms |       0.24% |   1.408 ms |       0.38% | -490.214 us | -25.83% |   FAST   |
|     I8     |     I64      |      I64      |      2^20      |     1     |  26.876 us |       1.85% |  25.143 us |       2.12% |   -1.733 us |  -6.45% |   FAST   |
|     I8     |     I64      |      I64      |      2^24      |     1     | 135.134 us |       1.03% | 109.286 us |       0.52% |  -25.849 us | -19.13% |   FAST   |
|     I8     |     I64      |      I64      |      2^28      |     1     |   1.923 ms |       0.19% |   1.430 ms |       0.33% | -492.272 us | -25.60% |   FAST   |
|     I8     |     I64      |      I64      |      2^20      |   0.201   |  25.378 us |       2.44% |  24.145 us |       2.51% |   -1.233 us |  -4.86% |   FAST   |
|     I8     |     I64      |      I64      |      2^24      |   0.201   | 135.358 us |       1.22% | 108.813 us |       0.66% |  -26.546 us | -19.61% |   FAST   |
|     I8     |     I64      |      I64      |      2^28      |   0.201   |   1.905 ms |       0.24% |   1.417 ms |       0.38% | -487.669 us | -25.60% |   FAST   |
|    I16     |      I8      |      I32      |      2^16      |     1     |  18.486 us |       2.57% |  17.401 us |       2.70% |   -1.085 us |  -5.87% |   FAST   |
|    I16     |      I8      |      I32      |      2^20      |     1     |  30.239 us |       1.71% |  25.921 us |       1.99% |   -4.318 us | -14.28% |   FAST   |
|    I16     |      I8      |      I32      |      2^24      |     1     | 132.267 us |       0.66% |  95.136 us |       0.73% |  -37.131 us | -28.07% |   FAST   |
|    I16     |      I8      |      I32      |      2^28      |     1     |   1.716 ms |       0.16% |   1.112 ms |       0.31% | -603.818 us | -35.19% |   FAST   |
|    I16     |      I8      |      I32      |      2^16      |   0.201   |  17.537 us |       3.38% |  17.118 us |       3.11% |   -0.419 us |  -2.39% |   SAME   |
|    I16     |      I8      |      I32      |      2^20      |   0.201   |  28.118 us |       1.86% |  24.814 us |       1.94% |   -3.304 us | -11.75% |   FAST   |
|    I16     |      I8      |      I32      |      2^24      |   0.201   | 122.310 us |       0.81% |  86.901 us |       1.15% |  -35.409 us | -28.95% |   FAST   |
|    I16     |      I8      |      I32      |      2^28      |   0.201   |   1.578 ms |       0.18% |   1.026 ms |       0.57% | -551.820 us | -34.97% |   FAST   |
|    I16     |      I8      |      I64      |      2^16      |     1     |  18.810 us |       3.04% |  17.726 us |       3.32% |   -1.083 us |  -5.76% |   FAST   |
|    I16     |      I8      |      I64      |      2^20      |     1     |  30.886 us |       1.51% |  25.935 us |       1.96% |   -4.951 us | -16.03% |   FAST   |
|    I16     |      I8      |      I64      |      2^24      |     1     | 134.077 us |       0.68% |  95.580 us |       0.75% |  -38.498 us | -28.71% |   FAST   |
|    I16     |      I8      |      I64      |      2^28      |     1     |   1.746 ms |       0.17% |   1.115 ms |       0.34% | -630.277 us | -36.11% |   FAST   |
|    I16     |      I8      |      I64      |      2^16      |   0.201   |  18.049 us |       3.40% |  17.250 us |       3.21% |   -0.799 us |  -4.43% |   FAST   |
|    I16     |      I8      |      I64      |      2^20      |   0.201   |  28.632 us |       1.91% |  24.811 us |       2.00% |   -3.821 us | -13.35% |   FAST   |
|    I16     |      I8      |      I64      |      2^24      |   0.201   | 125.032 us |       0.81% |  87.719 us |       0.93% |  -37.313 us | -29.84% |   FAST   |
|    I16     |      I8      |      I64      |      2^28      |   0.201   |   1.614 ms |       0.23% |   1.037 ms |       0.44% | -577.369 us | -35.77% |   FAST   |
|    I16     |     I16      |      I32      |      2^16      |     1     |  18.106 us |       2.80% |  17.158 us |       3.09% |   -0.947 us |  -5.23% |   FAST   |
|    I16     |     I16      |      I32      |      2^20      |     1     |  30.343 us |       1.58% |  25.768 us |       1.89% |   -4.575 us | -15.08% |   FAST   |
|    I16     |     I16      |      I32      |      2^24      |     1     | 137.896 us |       0.74% |  95.469 us |       0.87% |  -42.427 us | -30.77% |   FAST   |
|    I16     |     I16      |      I32      |      2^28      |     1     |   1.777 ms |       0.16% |   1.125 ms |       0.21% | -651.849 us | -36.69% |   FAST   |
|    I16     |     I16      |      I32      |      2^16      |   0.201   |  17.671 us |       3.31% |  16.945 us |       3.47% |   -0.726 us |  -4.11% |   FAST   |
|    I16     |     I16      |      I32      |      2^20      |   0.201   |  28.254 us |       1.86% |  24.647 us |       2.23% |   -3.607 us | -12.77% |   FAST   |
|    I16     |     I16      |      I32      |      2^24      |   0.201   | 126.904 us |       0.70% |  87.810 us |       0.66% |  -39.093 us | -30.81% |   FAST   |
|    I16     |     I16      |      I32      |      2^28      |   0.201   |   1.644 ms |       0.24% |   1.067 ms |       0.29% | -577.001 us | -35.10% |   FAST   |
|    I16     |     I16      |      I64      |      2^16      |     1     |  18.647 us |       2.58% |  17.732 us |       2.84% |   -0.915 us |  -4.91% |   FAST   |
|    I16     |     I16      |      I64      |      2^20      |     1     |  30.902 us |       1.54% |  25.929 us |       1.90% |   -4.973 us | -16.09% |   FAST   |
|    I16     |     I16      |      I64      |      2^24      |     1     | 139.079 us |       0.75% |  95.658 us |       0.80% |  -43.421 us | -31.22% |   FAST   |
|    I16     |     I16      |      I64      |      2^28      |     1     |   1.809 ms |       0.20% |   1.135 ms |       0.19% | -673.778 us | -37.25% |   FAST   |
|    I16     |     I16      |      I64      |      2^16      |   0.201   |  18.070 us |       3.40% |  17.307 us |       3.42% |   -0.762 us |  -4.22% |   FAST   |
|    I16     |     I16      |      I64      |      2^20      |   0.201   |  28.597 us |       2.06% |  24.872 us |       2.27% |   -3.724 us | -13.02% |   FAST   |
|    I16     |     I16      |      I64      |      2^24      |   0.201   | 128.716 us |       0.71% |  88.206 us |       0.67% |  -40.510 us | -31.47% |   FAST   |
|    I16     |     I16      |      I64      |      2^28      |   0.201   |   1.683 ms |       0.30% |   1.085 ms |       0.43% | -597.386 us | -35.51% |   FAST   |
|    I16     |     I32      |      I32      |      2^16      |     1     |  17.071 us |       2.96% |  16.549 us |       2.95% |   -0.521 us |  -3.05% |   FAST   |
|    I16     |     I32      |      I32      |      2^24      |     1     | 120.327 us |       0.97% |  93.413 us |       0.54% |  -26.913 us | -22.37% |   FAST   |
|    I16     |     I32      |      I32      |      2^28      |     1     |   1.587 ms |       0.33% |   1.125 ms |       0.34% | -462.428 us | -29.14% |   FAST   |
|    I16     |     I32      |      I32      |      2^24      |   0.201   | 108.139 us |       1.08% |  89.299 us |       0.92% |  -18.840 us | -17.42% |   FAST   |
|    I16     |     I32      |      I32      |      2^28      |   0.201   |   1.364 ms |       0.31% |   1.072 ms |       0.35% | -292.896 us | -21.47% |   FAST   |
|    I16     |     I32      |      I64      |      2^16      |     1     |  17.369 us |       3.43% |  16.896 us |       3.31% |   -0.473 us |  -2.73% |   SAME   |
|    I16     |     I32      |      I64      |      2^24      |     1     | 121.303 us |       1.01% |  94.328 us |       0.72% |  -26.976 us | -22.24% |   FAST   |
|    I16     |     I32      |      I64      |      2^28      |     1     |   1.610 ms |       0.20% |   1.132 ms |       0.35% | -478.129 us | -29.69% |   FAST   |
|    I16     |     I32      |      I64      |      2^24      |   0.201   | 108.606 us |       1.16% |  90.094 us |       0.80% |  -18.512 us | -17.05% |   FAST   |
|    I16     |     I32      |      I64      |      2^28      |   0.201   |   1.379 ms |       0.41% |   1.084 ms |       0.34% | -295.035 us | -21.40% |   FAST   |
|    I16     |     I64      |      I32      |      2^20      |     1     |  28.482 us |       1.90% |  26.352 us |       1.99% |   -2.131 us |  -7.48% |   FAST   |
|    I16     |     I64      |      I32      |      2^24      |     1     | 161.184 us |       0.64% | 124.895 us |       0.73% |  -36.289 us | -22.51% |   FAST   |
|    I16     |     I64      |      I32      |      2^28      |     1     |   2.258 ms |       0.19% |   1.634 ms |       0.24% | -623.929 us | -27.63% |   FAST   |
|    I16     |     I64      |      I32      |      2^20      |   0.201   |  26.835 us |       1.91% |  25.386 us |       2.14% |   -1.449 us |  -5.40% |   FAST   |
|    I16     |     I64      |      I32      |      2^24      |   0.201   | 147.507 us |       0.96% | 117.472 us |       0.64% |  -30.035 us | -20.36% |   FAST   |
|    I16     |     I64      |      I32      |      2^28      |   0.201   |   1.989 ms |       0.23% |   1.547 ms |       0.40% | -441.698 us | -22.21% |   FAST   |
|    I16     |     I64      |      I64      |      2^20      |     1     |  28.597 us |       2.06% |  26.569 us |       2.05% |   -2.029 us |  -7.09% |   FAST   |
|    I16     |     I64      |      I64      |      2^24      |     1     | 163.409 us |       0.69% | 125.744 us |       0.70% |  -37.664 us | -23.05% |   FAST   |
|    I16     |     I64      |      I64      |      2^28      |     1     |   2.284 ms |       0.16% |   1.644 ms |       0.26% | -639.274 us | -27.99% |   FAST   |
|    I16     |     I64      |      I64      |      2^20      |   0.201   |  27.055 us |       1.96% |  25.600 us |       2.20% |   -1.456 us |  -5.38% |   FAST   |
|    I16     |     I64      |      I64      |      2^24      |   0.201   | 148.450 us |       0.82% | 118.352 us |       0.69% |  -30.097 us | -20.27% |   FAST   |
|    I16     |     I64      |      I64      |      2^28      |   0.201   |   1.995 ms |       0.23% |   1.551 ms |       0.41% | -444.430 us | -22.27% |   FAST   |
|    I32     |      I8      |      I32      |      2^16      |     1     |  18.168 us |       2.75% |  17.743 us |       2.94% |   -0.425 us |  -2.34% |   SAME   |
|    I32     |      I8      |      I32      |      2^20      |     1     |  29.333 us |       1.54% |  25.483 us |       1.96% |   -3.850 us | -13.12% |   FAST   |
|    I32     |      I8      |      I32      |      2^24      |     1     | 141.894 us |       0.90% |  84.897 us |       0.87% |  -56.997 us | -40.17% |   FAST   |
|    I32     |      I8      |      I32      |      2^28      |     1     |   1.896 ms |       0.38% | 974.290 us |       0.37% | -921.935 us | -48.62% |   FAST   |
|    I32     |      I8      |      I32      |      2^16      |   0.201   |  17.525 us |       3.53% |  17.130 us |       3.57% |   -0.396 us |  -2.26% |   SAME   |
|    I32     |      I8      |      I32      |      2^20      |   0.201   |  27.531 us |       2.21% |  24.463 us |       2.41% |   -3.068 us | -11.14% |   FAST   |
|    I32     |      I8      |      I32      |      2^24      |   0.201   | 132.578 us |       0.99% |  80.057 us |       1.02% |  -52.522 us | -39.62% |   FAST   |
|    I32     |      I8      |      I32      |      2^28      |   0.201   |   1.724 ms |       0.37% | 909.130 us |       0.46% | -814.624 us | -47.26% |   FAST   |
|    I32     |      I8      |      I64      |      2^16      |     1     |  18.302 us |       2.77% |  17.806 us |       2.81% |   -0.496 us |  -2.71% |   SAME   |
|    I32     |      I8      |      I64      |      2^20      |     1     |  29.532 us |       1.69% |  26.019 us |       2.25% |   -3.513 us | -11.90% |   FAST   |
|    I32     |      I8      |      I64      |      2^24      |     1     | 145.677 us |       1.05% | 104.237 us |       1.34% |  -41.441 us | -28.45% |   FAST   |
|    I32     |      I8      |      I64      |      2^28      |     1     |   1.971 ms |       0.39% |   1.346 ms |       0.68% | -624.805 us | -31.70% |   FAST   |
|    I32     |      I8      |      I64      |      2^20      |   0.201   |  27.732 us |       2.01% |  24.609 us |       2.22% |   -3.124 us | -11.26% |   FAST   |
|    I32     |      I8      |      I64      |      2^24      |   0.201   | 133.076 us |       0.99% |  94.949 us |       1.30% |  -38.128 us | -28.65% |   FAST   |
|    I32     |      I8      |      I64      |      2^28      |   0.201   |   1.774 ms |       0.39% |   1.201 ms |       0.63% | -573.019 us | -32.31% |   FAST   |
|    I32     |     I16      |      I32      |      2^16      |     1     |  17.982 us |       2.83% |  17.352 us |       3.15% |   -0.630 us |  -3.50% |   FAST   |
|    I32     |     I16      |      I32      |      2^20      |     1     |  25.685 us |       2.33% |  25.004 us |       2.00% |   -0.681 us |  -2.65% |   FAST   |
|    I32     |     I16      |      I32      |      2^24      |     1     | 113.586 us |       0.69% |  92.400 us |       0.83% |  -21.186 us | -18.65% |   FAST   |
|    I32     |     I16      |      I32      |      2^28      |     1     |   1.473 ms |       0.32% |   1.111 ms |       0.41% | -362.665 us | -24.62% |   FAST   |
|    I32     |     I16      |      I32      |      2^24      |   0.201   | 104.046 us |       1.08% |  85.719 us |       1.03% |  -18.327 us | -17.61% |   FAST   |
|    I32     |     I16      |      I32      |      2^28      |   0.201   |   1.338 ms |       0.38% |   1.022 ms |       0.51% | -316.455 us | -23.64% |   FAST   |
|    I32     |     I16      |      I64      |      2^16      |     1     |  18.018 us |       3.09% |  17.513 us |       3.13% |   -0.505 us |  -2.80% |   SAME   |
|    I32     |     I16      |      I64      |      2^20      |     1     |  25.934 us |       2.04% |  25.335 us |       2.10% |   -0.600 us |  -2.31% |   FAST   |
|    I32     |     I16      |      I64      |      2^24      |     1     | 117.864 us |       0.97% |  93.509 us |       0.68% |  -24.355 us | -20.66% |   FAST   |
|    I32     |     I16      |      I64      |      2^28      |     1     |   1.560 ms |       0.47% |   1.125 ms |       0.38% | -435.218 us | -27.90% |   FAST   |
|    I32     |     I16      |      I64      |      2^24      |   0.201   | 104.975 us |       1.22% |  86.828 us |       0.92% |  -18.147 us | -17.29% |   FAST   |
|    I32     |     I16      |      I64      |      2^28      |   0.201   |   1.338 ms |       0.46% |   1.029 ms |       0.53% | -309.439 us | -23.12% |   FAST   |
|    I32     |     I32      |      I32      |      2^16      |     1     |  17.354 us |       2.68% |  16.979 us |       3.24% |   -0.375 us |  -2.16% |   SAME   |
|    I32     |     I32      |      I32      |      2^20      |     1     |  27.723 us |       2.25% |  25.520 us |       2.19% |   -2.203 us |  -7.95% |   FAST   |
|    I32     |     I32      |      I32      |      2^24      |     1     | 133.188 us |       0.74% | 102.544 us |       0.56% |  -30.644 us | -23.01% |   FAST   |
|    I32     |     I32      |      I32      |      2^28      |     1     |   1.813 ms |       0.23% |   1.302 ms |       0.32% | -510.371 us | -28.16% |   FAST   |
|    I32     |     I32      |      I32      |      2^20      |   0.201   |  25.890 us |       2.23% |  24.402 us |       2.14% |   -1.488 us |  -5.75% |   FAST   |
|    I32     |     I32      |      I32      |      2^24      |   0.201   | 121.399 us |       1.02% |  97.597 us |       0.70% |  -23.802 us | -19.61% |   FAST   |
|    I32     |     I32      |      I32      |      2^28      |   0.201   |   1.649 ms |       0.30% |   1.245 ms |       0.56% | -404.278 us | -24.51% |   FAST   |
|    I32     |     I32      |      I64      |      2^20      |     1     |  27.972 us |       1.88% |  25.933 us |       2.21% |   -2.039 us |  -7.29% |   FAST   |
|    I32     |     I32      |      I64      |      2^24      |     1     | 135.083 us |       0.67% | 103.265 us |       0.58% |  -31.818 us | -23.55% |   FAST   |
|    I32     |     I32      |      I64      |      2^28      |     1     |   1.853 ms |       0.28% |   1.310 ms |       0.37% | -543.487 us | -29.33% |   FAST   |
|    I32     |     I32      |      I64      |      2^20      |   0.201   |  26.057 us |       2.14% |  24.412 us |       2.31% |   -1.645 us |  -6.31% |   FAST   |
|    I32     |     I32      |      I64      |      2^24      |   0.201   | 121.896 us |       1.05% |  98.003 us |       0.68% |  -23.892 us | -19.60% |   FAST   |
|    I32     |     I32      |      I64      |      2^28      |   0.201   |   1.651 ms |       0.31% |   1.250 ms |       0.57% | -401.028 us | -24.29% |   FAST   |
|    I32     |     I64      |      I32      |      2^16      |     1     |  17.286 us |       2.84% |  17.651 us |       2.34% |    0.364 us |   2.11% |   SAME   |
|    I32     |     I64      |      I32      |      2^20      |     1     |  28.629 us |       1.78% |  26.994 us |       1.80% |   -1.635 us |  -5.71% |   FAST   |
|    I32     |     I64      |      I32      |      2^24      |     1     | 148.699 us |       0.57% | 138.496 us |       0.64% |  -10.202 us |  -6.86% |   FAST   |
|    I32     |     I64      |      I32      |      2^28      |     1     |   2.054 ms |       0.28% |   1.904 ms |       0.27% | -150.440 us |  -7.32% |   FAST   |
|    I32     |     I64      |      I32      |      2^16      |   0.201   |  16.457 us |       3.33% |  16.988 us |       3.27% |    0.530 us |   3.22% |   SAME   |
|    I32     |     I64      |      I32      |      2^20      |   0.201   |  27.057 us |       1.96% |  26.251 us |       2.14% |   -0.806 us |  -2.98% |   FAST   |
|    I32     |     I64      |      I32      |      2^28      |   0.201   |   1.880 ms |       0.38% |   1.842 ms |       0.41% |  -38.425 us |  -2.04% |   FAST   |
|    I32     |     I64      |      I64      |      2^16      |     1     |  17.432 us |       3.32% |  17.918 us |       2.84% |    0.486 us |   2.79% |   SAME   |
|    I32     |     I64      |      I64      |      2^20      |     1     |  28.947 us |       1.69% |  27.493 us |       2.14% |   -1.454 us |  -5.02% |   FAST   |
|    I32     |     I64      |      I64      |      2^24      |     1     | 149.672 us |       0.52% | 139.520 us |       0.66% |  -10.152 us |  -6.78% |   FAST   |
|    I32     |     I64      |      I64      |      2^28      |     1     |   2.060 ms |       0.29% |   1.914 ms |       0.29% | -145.195 us |  -7.05% |   FAST   |
|    I32     |     I64      |      I64      |      2^16      |   0.201   |  16.660 us |       3.56% |  17.505 us |       2.77% |    0.845 us |   5.07% |   SLOW   |
|    I32     |     I64      |      I64      |      2^20      |   0.201   |  27.671 us |       2.26% |  26.625 us |       1.80% |   -1.047 us |  -3.78% |   FAST   |
|    I32     |     I64      |      I64      |      2^28      |   0.201   |   1.887 ms |       0.38% |   1.847 ms |       0.39% |  -40.337 us |  -2.14% |   FAST   |
|    I64     |      I8      |      I32      |      2^16      |     1     |  18.376 us |       2.77% |  18.818 us |       3.01% |    0.442 us |   2.40% |   SAME   |
|    I64     |      I8      |      I32      |      2^20      |     1     |  30.392 us |       1.88% |  28.807 us |       2.00% |   -1.585 us |  -5.22% |   FAST   |
|    I64     |      I8      |      I32      |      2^24      |     1     | 153.943 us |       0.67% | 126.312 us |       0.78% |  -27.631 us | -17.95% |   FAST   |
|    I64     |      I8      |      I32      |      2^28      |     1     |   2.091 ms |       0.38% |   1.633 ms |       0.36% | -457.639 us | -21.88% |   FAST   |
|    I64     |      I8      |      I32      |      2^16      |   0.201   |  17.734 us |       2.78% |  18.200 us |       2.80% |    0.466 us |   2.63% |   SAME   |
|    I64     |      I8      |      I32      |      2^20      |   0.201   |  27.844 us |       2.26% |  26.434 us |       2.08% |   -1.409 us |  -5.06% |   FAST   |
|    I64     |      I8      |      I32      |      2^24      |   0.201   | 139.428 us |       1.05% | 114.008 us |       0.83% |  -25.420 us | -18.23% |   FAST   |
|    I64     |      I8      |      I32      |      2^28      |   0.201   |   1.910 ms |       0.37% |   1.461 ms |       0.67% | -448.523 us | -23.48% |   FAST   |
|    I64     |      I8      |      I64      |      2^20      |     1     |  30.316 us |       1.76% |  28.629 us |       1.86% |   -1.687 us |  -5.57% |   FAST   |
|    I64     |      I8      |      I64      |      2^24      |     1     | 156.016 us |       0.72% | 126.880 us |       0.76% |  -29.135 us | -18.67% |   FAST   |
|    I64     |      I8      |      I64      |      2^28      |     1     |   2.112 ms |       0.36% |   1.642 ms |       0.32% | -469.639 us | -22.24% |   FAST   |
|    I64     |      I8      |      I64      |      2^20      |   0.201   |  27.906 us |       2.00% |  26.523 us |       1.94% |   -1.383 us |  -4.96% |   FAST   |
|    I64     |      I8      |      I64      |      2^24      |   0.201   | 141.052 us |       1.02% | 114.924 us |       0.88% |  -26.129 us | -18.52% |   FAST   |
|    I64     |      I8      |      I64      |      2^28      |   0.201   |   1.943 ms |       0.33% |   1.467 ms |       0.66% | -476.306 us | -24.51% |   FAST   |
|    I64     |     I16      |      I32      |      2^16      |     1     |  18.663 us |       2.93% |  18.274 us |       3.29% |   -0.389 us |  -2.08% |   SAME   |
|    I64     |     I16      |      I32      |      2^20      |     1     |  30.741 us |       1.79% |  28.916 us |       2.03% |   -1.825 us |  -5.94% |   FAST   |
|    I64     |     I16      |      I32      |      2^24      |     1     | 159.162 us |       0.64% | 129.052 us |       0.74% |  -30.109 us | -18.92% |   FAST   |
|    I64     |     I16      |      I32      |      2^28      |     1     |   2.171 ms |       0.41% |   1.675 ms |       0.38% | -496.488 us | -22.87% |   FAST   |
|    I64     |     I16      |      I32      |      2^20      |   0.201   |  28.306 us |       2.50% |  26.632 us |       2.12% |   -1.674 us |  -5.92% |   FAST   |
|    I64     |     I16      |      I32      |      2^24      |   0.201   | 145.686 us |       1.04% | 119.428 us |       0.74% |  -26.258 us | -18.02% |   FAST   |
|    I64     |     I16      |      I32      |      2^28      |   0.201   |   2.041 ms |       0.33% |   1.553 ms |       0.72% | -487.777 us | -23.90% |   FAST   |
|    I64     |     I16      |      I64      |      2^20      |     1     |  30.778 us |       1.62% |  29.069 us |       1.78% |   -1.709 us |  -5.55% |   FAST   |
|    I64     |     I16      |      I64      |      2^24      |     1     | 160.858 us |       0.72% | 138.038 us |       0.77% |  -22.820 us | -14.19% |   FAST   |
|    I64     |     I16      |      I64      |      2^28      |     1     |   2.204 ms |       0.40% |   1.825 ms |       0.40% | -378.754 us | -17.18% |   FAST   |
|    I64     |     I16      |      I64      |      2^20      |   0.201   |  28.567 us |       2.32% |  26.755 us |       1.93% |   -1.812 us |  -6.34% |   FAST   |
|    I64     |     I16      |      I64      |      2^24      |   0.201   | 146.692 us |       1.19% | 126.698 us |       0.67% |  -19.993 us | -13.63% |   FAST   |
|    I64     |     I16      |      I64      |      2^28      |   0.201   |   2.065 ms |       0.34% |   1.676 ms |       0.55% | -388.376 us | -18.81% |   FAST   |
|    I64     |     I32      |      I32      |      2^20      |     1     |  30.725 us |       1.69% |  29.740 us |       2.11% |   -0.984 us |  -3.20% |   FAST   |
|    I64     |     I32      |      I32      |      2^24      |     1     | 155.030 us |       0.66% | 142.693 us |       0.72% |  -12.337 us |  -7.96% |   FAST   |
|    I64     |     I32      |      I32      |      2^28      |     1     |   2.109 ms |       0.44% |   1.908 ms |       0.43% | -200.805 us |  -9.52% |   FAST   |
|    I64     |     I32      |      I32      |      2^20      |   0.201   |  27.816 us |       2.08% |  26.738 us |       2.38% |   -1.078 us |  -3.88% |   FAST   |
|    I64     |     I32      |      I32      |      2^28      |   0.201   |   1.878 ms |       0.61% |   1.835 ms |       0.69% |  -43.232 us |  -2.30% |   FAST   |
|    I64     |     I32      |      I64      |      2^20      |     1     |  30.731 us |       2.26% |  29.676 us |       2.17% |   -1.055 us |  -3.43% |   FAST   |
|    I64     |     I32      |      I64      |      2^24      |     1     | 155.268 us |       0.64% | 142.991 us |       0.70% |  -12.277 us |  -7.91% |   FAST   |
|    I64     |     I32      |      I64      |      2^28      |     1     |   2.121 ms |       0.41% |   1.919 ms |       0.44% | -202.275 us |  -9.54% |   FAST   |
|    I64     |     I32      |      I64      |      2^20      |   0.201   |  28.107 us |       1.87% |  27.005 us |       1.66% |   -1.102 us |  -3.92% |   FAST   |
|    I64     |     I32      |      I64      |      2^28      |   0.201   |   1.886 ms |       0.56% |   1.838 ms |       0.70% |  -48.215 us |  -2.56% |   FAST   |
|    I128    |      I8      |      I32      |      2^20      |     1     |  38.764 us |       1.56% |  37.850 us |       1.61% |   -0.914 us |  -2.36% |   FAST   |
|    I128    |      I8      |      I32      |      2^24      |     1     | 245.186 us |       0.56% | 232.224 us |       0.70% |  -12.962 us |  -5.29% |   FAST   |
|    I128    |      I8      |      I32      |      2^28      |     1     |   3.546 ms |       0.28% |   3.321 ms |       0.29% | -225.112 us |  -6.35% |   FAST   |
|    I128    |      I8      |      I32      |      2^24      |   0.201   | 194.739 us |       0.69% | 188.944 us |       0.74% |   -5.796 us |  -2.98% |   FAST   |
|    I128    |      I8      |      I32      |      2^28      |   0.201   |   2.814 ms |       0.65% |   2.692 ms |       0.69% | -122.826 us |  -4.36% |   FAST   |
|    I128    |     I64      |      I32      |      2^20      |     1     |  40.652 us |       1.60% |  39.142 us |       1.65% |   -1.511 us |  -3.72% |   FAST   |
|    I128    |     I64      |      I32      |      2^24      |     1     | 304.515 us |       0.48% | 279.126 us |       0.45% |  -25.390 us |  -8.34% |   FAST   |
|    I128    |     I64      |      I32      |      2^28      |     1     |   4.563 ms |       0.43% |   4.158 ms |       0.38% | -405.289 us |  -8.88% |   FAST   |
|    I128    |     I64      |      I32      |      2^24      |   0.201   | 270.542 us |       0.60% | 253.529 us |       0.58% |  -17.013 us |  -6.29% |   FAST   |
|    I128    |     I64      |      I32      |      2^28      |   0.201   |   4.078 ms |       0.51% |   3.790 ms |       0.62% | -287.480 us |  -7.05% |   FAST   |
|    I128    |     I64      |      I64      |      2^20      |     1     |  41.331 us |       2.03% |  39.522 us |       1.71% |   -1.809 us |  -4.38% |   FAST   |
|    I128    |     I64      |      I64      |      2^24      |     1     | 307.520 us |       0.41% | 277.689 us |       0.45% |  -29.831 us |  -9.70% |   FAST   |
|    I128    |     I64      |      I64      |      2^28      |     1     |   4.608 ms |       0.39% |   4.132 ms |       0.43% | -476.698 us | -10.34% |   FAST   |
|    I128    |     I64      |      I64      |      2^24      |   0.201   | 272.195 us |       0.57% | 253.328 us |       0.60% |  -18.867 us |  -6.93% |   FAST   |
|    I128    |     I64      |      I64      |      2^28      |   0.201   |   4.096 ms |       0.49% |   3.780 ms |       0.65% | -315.578 us |  -7.70% |   FAST   |
|    F32     |      I8      |      I32      |      2^16      |     1     |  18.158 us |       2.87% |  17.664 us |       2.67% |   -0.494 us |  -2.72% |   FAST   |
|    F32     |      I8      |      I32      |      2^20      |     1     |  29.145 us |       1.75% |  25.469 us |       2.28% |   -3.676 us | -12.61% |   FAST   |
|    F32     |      I8      |      I32      |      2^24      |     1     | 142.093 us |       0.97% |  84.759 us |       0.71% |  -57.334 us | -40.35% |   FAST   |
|    F32     |      I8      |      I32      |      2^28      |     1     |   1.895 ms |       0.37% | 974.197 us |       0.39% | -920.782 us | -48.59% |   FAST   |
|    F32     |      I8      |      I32      |      2^20      |   0.201   |  27.607 us |       2.22% |  24.340 us |       2.16% |   -3.267 us | -11.83% |   FAST   |
|    F32     |      I8      |      I32      |      2^24      |   0.201   | 131.329 us |       0.83% |  79.951 us |       1.12% |  -51.378 us | -39.12% |   FAST   |
|    F32     |      I8      |      I32      |      2^28      |   0.201   |   1.722 ms |       0.37% | 909.399 us |       0.45% | -813.019 us | -47.20% |   FAST   |
|    F32     |      I8      |      I64      |      2^20      |     1     |  29.281 us |       1.69% |  26.171 us |       2.28% |   -3.110 us | -10.62% |   FAST   |
|    F32     |      I8      |      I64      |      2^24      |     1     | 146.060 us |       1.09% | 103.680 us |       1.52% |  -42.380 us | -29.02% |   FAST   |
|    F32     |      I8      |      I64      |      2^28      |     1     |   1.973 ms |       0.39% |   1.340 ms |       0.68% | -632.768 us | -32.07% |   FAST   |
|    F32     |      I8      |      I64      |      2^16      |   0.201   |  17.639 us |       3.20% |  16.993 us |       2.68% |   -0.646 us |  -3.66% |   FAST   |
|    F32     |      I8      |      I64      |      2^20      |   0.201   |  27.649 us |       2.12% |  24.654 us |       2.02% |   -2.996 us | -10.83% |   FAST   |
|    F32     |      I8      |      I64      |      2^24      |   0.201   | 133.790 us |       0.96% |  96.409 us |       1.13% |  -37.381 us | -27.94% |   FAST   |
|    F32     |      I8      |      I64      |      2^28      |   0.201   |   1.775 ms |       0.37% |   1.196 ms |       0.61% | -579.063 us | -32.62% |   FAST   |
|    F32     |     I16      |      I32      |      2^16      |     1     |  18.196 us |       2.97% |  17.463 us |       3.19% |   -0.733 us |  -4.03% |   FAST   |
|    F32     |     I16      |      I32      |      2^20      |     1     |  25.891 us |       2.33% |  25.332 us |       1.98% |   -0.558 us |  -2.16% |   FAST   |
|    F32     |     I16      |      I32      |      2^24      |     1     | 113.430 us |       0.70% |  92.665 us |       0.74% |  -20.765 us | -18.31% |   FAST   |
|    F32     |     I16      |      I32      |      2^28      |     1     |   1.474 ms |       0.29% |   1.110 ms |       0.42% | -363.294 us | -24.65% |   FAST   |
|    F32     |     I16      |      I32      |      2^16      |   0.201   |  17.070 us |       2.96% |  16.716 us |       3.57% |   -0.353 us |  -2.07% |   SAME   |
|    F32     |     I16      |      I32      |      2^24      |   0.201   | 104.793 us |       1.25% |  86.513 us |       0.90% |  -18.280 us | -17.44% |   FAST   |
|    F32     |     I16      |      I32      |      2^28      |   0.201   |   1.339 ms |       0.40% |   1.021 ms |       0.54% | -317.214 us | -23.70% |   FAST   |
|    F32     |     I16      |      I64      |      2^16      |     1     |  18.330 us |       3.13% |  17.394 us |       2.76% |   -0.936 us |  -5.11% |   FAST   |
|    F32     |     I16      |      I64      |      2^20      |     1     |  26.189 us |       2.01% |  25.583 us |       2.08% |   -0.606 us |  -2.31% |   FAST   |
|    F32     |     I16      |      I64      |      2^24      |     1     | 116.833 us |       0.83% |  93.500 us |       0.88% |  -23.333 us | -19.97% |   FAST   |
|    F32     |     I16      |      I64      |      2^28      |     1     |   1.561 ms |       0.47% |   1.125 ms |       0.39% | -435.398 us | -27.90% |   FAST   |
|    F32     |     I16      |      I64      |      2^24      |   0.201   | 106.065 us |       1.29% |  87.061 us |       0.96% |  -19.004 us | -17.92% |   FAST   |
|    F32     |     I16      |      I64      |      2^28      |   0.201   |   1.338 ms |       0.46% |   1.029 ms |       0.50% | -309.978 us | -23.16% |   FAST   |
|    F32     |     I32      |      I32      |      2^20      |     1     |  28.073 us |       2.23% |  25.865 us |       2.30% |   -2.208 us |  -7.86% |   FAST   |
|    F32     |     I32      |      I32      |      2^24      |     1     | 133.354 us |       0.68% | 102.744 us |       0.59% |  -30.610 us | -22.95% |   FAST   |
|    F32     |     I32      |      I32      |      2^28      |     1     |   1.815 ms |       0.27% |   1.304 ms |       0.39% | -511.045 us | -28.16% |   FAST   |
|    F32     |     I32      |      I32      |      2^20      |   0.201   |  25.891 us |       2.09% |  24.257 us |       2.42% |   -1.634 us |  -6.31% |   FAST   |
|    F32     |     I32      |      I32      |      2^24      |   0.201   | 121.760 us |       0.94% |  97.554 us |       0.69% |  -24.205 us | -19.88% |   FAST   |
|    F32     |     I32      |      I32      |      2^28      |   0.201   |   1.650 ms |       0.33% |   1.245 ms |       0.55% | -404.666 us | -24.52% |   FAST   |
|    F32     |     I32      |      I64      |      2^20      |     1     |  28.041 us |       2.06% |  25.899 us |       2.02% |   -2.142 us |  -7.64% |   FAST   |
|    F32     |     I32      |      I64      |      2^24      |     1     | 135.335 us |       0.69% | 102.934 us |       0.56% |  -32.401 us | -23.94% |   FAST   |
|    F32     |     I32      |      I64      |      2^28      |     1     |   1.857 ms |       0.28% |   1.311 ms |       0.35% | -545.935 us | -29.41% |   FAST   |
|    F32     |     I32      |      I64      |      2^16      |   0.201   |  16.547 us |       2.88% |  16.893 us |       3.11% |    0.347 us |   2.10% |   SAME   |
|    F32     |     I32      |      I64      |      2^20      |   0.201   |  26.225 us |       2.21% |  24.507 us |       2.18% |   -1.718 us |  -6.55% |   FAST   |
|    F32     |     I32      |      I64      |      2^24      |   0.201   | 120.750 us |       1.18% |  98.155 us |       0.69% |  -22.595 us | -18.71% |   FAST   |
|    F32     |     I32      |      I64      |      2^28      |   0.201   |   1.650 ms |       0.33% |   1.249 ms |       0.56% | -401.525 us | -24.33% |   FAST   |
|    F32     |     I64      |      I32      |      2^16      |     1     |  17.172 us |       2.79% |  17.526 us |       3.07% |    0.354 us |   2.06% |   SAME   |
|    F32     |     I64      |      I32      |      2^20      |     1     |  28.463 us |       1.98% |  27.108 us |       2.07% |   -1.355 us |  -4.76% |   FAST   |
|    F32     |     I64      |      I32      |      2^24      |     1     | 148.967 us |       0.58% | 138.248 us |       0.68% |  -10.720 us |  -7.20% |   FAST   |
|    F32     |     I64      |      I32      |      2^28      |     1     |   2.054 ms |       0.28% |   1.904 ms |       0.29% | -149.714 us |  -7.29% |   FAST   |
|    F32     |     I64      |      I32      |      2^20      |   0.201   |  27.195 us |       1.99% |  26.065 us |       2.11% |   -1.131 us |  -4.16% |   FAST   |
|    F32     |     I64      |      I32      |      2^28      |   0.201   |   1.881 ms |       0.37% |   1.843 ms |       0.40% |  -37.881 us |  -2.01% |   FAST   |
|    F32     |     I64      |      I64      |      2^16      |     1     |  17.377 us |       3.09% |  17.898 us |       2.79% |    0.521 us |   3.00% |   SLOW   |
|    F32     |     I64      |      I64      |      2^20      |     1     |  28.733 us |       2.08% |  27.538 us |       2.01% |   -1.195 us |  -4.16% |   FAST   |
|    F32     |     I64      |      I64      |      2^24      |     1     | 149.023 us |       0.53% | 139.269 us |       0.66% |   -9.755 us |  -6.55% |   FAST   |
|    F32     |     I64      |      I64      |      2^28      |     1     |   2.059 ms |       0.30% |   1.914 ms |       0.29% | -144.580 us |  -7.02% |   FAST   |
|    F32     |     I64      |      I64      |      2^16      |   0.201   |  16.541 us |       2.94% |  17.036 us |       3.51% |    0.495 us |   2.99% |   SLOW   |
|    F32     |     I64      |      I64      |      2^20      |   0.201   |  27.293 us |       2.23% |  26.741 us |       2.01% |   -0.552 us |  -2.02% |   FAST   |
|    F32     |     I64      |      I64      |      2^28      |   0.201   |   1.887 ms |       0.36% |   1.848 ms |       0.41% |  -39.181 us |  -2.08% |   FAST   |
|    F64     |      I8      |      I32      |      2^20      |     1     |  29.865 us |       1.77% |  28.333 us |       1.91% |   -1.532 us |  -5.13% |   FAST   |
|    F64     |      I8      |      I32      |      2^24      |     1     | 153.523 us |       0.85% | 125.498 us |       0.82% |  -28.024 us | -18.25% |   FAST   |
|    F64     |      I8      |      I32      |      2^28      |     1     |   2.077 ms |       0.39% |   1.621 ms |       0.39% | -455.827 us | -21.94% |   FAST   |
|    F64     |      I8      |      I32      |      2^20      |   0.201   |  27.144 us |       2.00% |  26.086 us |       2.26% |   -1.058 us |  -3.90% |   FAST   |
|    F64     |      I8      |      I32      |      2^24      |   0.201   | 137.781 us |       1.03% | 113.549 us |       0.83% |  -24.232 us | -17.59% |   FAST   |
|    F64     |      I8      |      I32      |      2^28      |   0.201   |   1.886 ms |       0.39% |   1.456 ms |       0.60% | -429.257 us | -22.76% |   FAST   |
|    F64     |      I8      |      I64      |      2^20      |     1     |  30.168 us |       1.92% |  28.483 us |       1.80% |   -1.685 us |  -5.58% |   FAST   |
|    F64     |      I8      |      I64      |      2^24      |     1     | 155.092 us |       0.91% | 125.827 us |       0.81% |  -29.266 us | -18.87% |   FAST   |
|    F64     |      I8      |      I64      |      2^28      |     1     |   2.103 ms |       0.44% |   1.636 ms |       0.34% | -467.211 us | -22.21% |   FAST   |
|    F64     |      I8      |      I64      |      2^20      |   0.201   |  27.983 us |       2.35% |  26.262 us |       2.12% |   -1.721 us |  -6.15% |   FAST   |
|    F64     |      I8      |      I64      |      2^24      |   0.201   | 140.980 us |       1.15% | 114.524 us |       0.80% |  -26.456 us | -18.77% |   FAST   |
|    F64     |      I8      |      I64      |      2^28      |   0.201   |   1.962 ms |       0.41% |   1.463 ms |       0.58% | -499.360 us | -25.45% |   FAST   |
|    F64     |     I16      |      I32      |      2^20      |     1     |  30.600 us |       2.30% |  28.795 us |       1.91% |   -1.805 us |  -5.90% |   FAST   |
|    F64     |     I16      |      I32      |      2^24      |     1     | 158.709 us |       1.06% | 128.190 us |       0.81% |  -30.520 us | -19.23% |   FAST   |
|    F64     |     I16      |      I32      |      2^28      |     1     |   2.173 ms |       0.51% |   1.672 ms |       0.78% | -501.201 us | -23.06% |   FAST   |
|    F64     |     I16      |      I32      |      2^20      |   0.201   |  28.281 us |       2.13% |  26.636 us |       2.31% |   -1.645 us |  -5.82% |   FAST   |
|    F64     |     I16      |      I32      |      2^24      |   0.201   | 141.018 us |       1.20% | 119.174 us |       0.78% |  -21.844 us | -15.49% |   FAST   |
|    F64     |     I16      |      I32      |      2^28      |   0.201   |   1.950 ms |       0.38% |   1.552 ms |       0.65% | -397.531 us | -20.39% |   FAST   |
|    F64     |     I16      |      I64      |      2^20      |     1     |  30.581 us |       1.95% |  28.994 us |       1.94% |   -1.587 us |  -5.19% |   FAST   |
|    F64     |     I16      |      I64      |      2^24      |     1     | 160.808 us |       1.32% | 138.092 us |       0.74% |  -22.716 us | -14.13% |   FAST   |
|    F64     |     I16      |      I64      |      2^28      |     1     |   2.217 ms |       0.55% |   1.825 ms |       0.43% | -392.485 us | -17.70% |   FAST   |
|    F64     |     I16      |      I64      |      2^20      |   0.201   |  28.338 us |       2.16% |  26.743 us |       2.01% |   -1.595 us |  -5.63% |   FAST   |
|    F64     |     I16      |      I64      |      2^24      |   0.201   | 142.362 us |       1.16% | 126.431 us |       0.79% |  -15.930 us | -11.19% |   FAST   |
|    F64     |     I16      |      I64      |      2^28      |   0.201   |   1.973 ms |       0.38% |   1.671 ms |       0.59% | -301.756 us | -15.30% |   FAST   |
|    F64     |     I32      |      I32      |      2^20      |     1     |  31.824 us |       1.78% |  29.674 us |       1.72% |   -2.150 us |  -6.75% |   FAST   |
|    F64     |     I32      |      I32      |      2^24      |     1     | 177.075 us |       0.77% | 141.938 us |       0.73% |  -35.137 us | -19.84% |   FAST   |
|    F64     |     I32      |      I32      |      2^28      |     1     |   2.474 ms |       0.37% |   1.914 ms |       0.76% | -560.316 us | -22.65% |   FAST   |
|    F64     |     I32      |      I32      |      2^16      |   0.201   |  17.472 us |       3.48% |  18.169 us |       2.85% |    0.697 us |   3.99% |   SLOW   |
|    F64     |     I32      |      I32      |      2^20      |   0.201   |  28.537 us |       2.34% |  26.970 us |       2.02% |   -1.566 us |  -5.49% |   FAST   |
|    F64     |     I32      |      I32      |      2^24      |   0.201   | 161.618 us |       1.11% | 135.459 us |       0.73% |  -26.159 us | -16.19% |   FAST   |
|    F64     |     I32      |      I32      |      2^28      |   0.201   |   2.297 ms |       0.33% |   1.834 ms |       0.75% | -462.957 us | -20.15% |   FAST   |
|    F64     |     I32      |      I64      |      2^20      |     1     |  31.909 us |       1.63% |  29.977 us |       1.82% |   -1.933 us |  -6.06% |   FAST   |
|    F64     |     I32      |      I64      |      2^24      |     1     | 178.282 us |       0.77% | 142.907 us |       0.70% |  -35.375 us | -19.84% |   FAST   |
|    F64     |     I32      |      I64      |      2^28      |     1     |   2.500 ms |       0.40% |   1.926 ms |       0.77% | -573.414 us | -22.94% |   FAST   |
|    F64     |     I32      |      I64      |      2^20      |   0.201   |  28.822 us |       2.40% |  26.935 us |       2.35% |   -1.887 us |  -6.55% |   FAST   |
|    F64     |     I32      |      I64      |      2^24      |   0.201   | 162.828 us |       1.18% | 135.747 us |       0.71% |  -27.081 us | -16.63% |   FAST   |
|    F64     |     I32      |      I64      |      2^28      |   0.201   |   2.320 ms |       0.33% |   1.837 ms |       0.67% | -483.441 us | -20.83% |   FAST   |

bernhardmgruber · 2025-11-05T11:04:10Z

Keys WITHOUT tuning (this is NOT proposed by this PR)

keys on H200, #6077 vs this PR, 2% threshold:

(same perf)

keys on H200, main vs this PR, 2% threshold:

|  KeyT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|------------|---------------|----------------|-----------|------------|-------------|------------|-------------|-------------|---------|----------|
|     I8     |      I32      |      2^24      |     1     |  53.553 us |       0.93% |  49.400 us |       0.94% |   -4.153 us |  -7.75% |   FAST   |
|     I8     |      I32      |      2^28      |     1     | 524.000 us |       0.22% | 456.980 us |       0.19% |  -67.020 us | -12.79% |   FAST   |
|     I8     |      I32      |      2^24      |   0.201   |  51.930 us |       1.00% |  47.767 us |       1.15% |   -4.163 us |  -8.02% |   FAST   |
|     I8     |      I32      |      2^28      |   0.201   | 518.773 us |       0.22% | 451.721 us |       0.20% |  -67.052 us | -12.93% |   FAST   |
|     I8     |      I64      |      2^24      |     1     |  53.475 us |       1.01% |  49.901 us |       1.12% |   -3.575 us |  -6.68% |   FAST   |
|     I8     |      I64      |      2^28      |     1     | 513.522 us |       0.21% | 459.907 us |       0.23% |  -53.615 us | -10.44% |   FAST   |
|     I8     |      I64      |      2^24      |   0.201   |  51.747 us |       1.09% |  47.712 us |       1.14% |   -4.035 us |  -7.80% |   FAST   |
|     I8     |      I64      |      2^28      |   0.201   | 508.798 us |       0.21% | 454.459 us |       0.20% |  -54.340 us | -10.68% |   FAST   |
|    I16     |      I32      |      2^24      |     1     |  66.515 us |       0.80% |  59.263 us |       0.96% |   -7.252 us | -10.90% |   FAST   |
|    I16     |      I32      |      2^28      |     1     | 739.462 us |       0.22% | 614.673 us |       0.16% | -124.789 us | -16.88% |   FAST   |
|    I16     |      I32      |      2^24      |   0.201   |  59.502 us |       0.95% |  53.354 us |       1.11% |   -6.147 us | -10.33% |   FAST   |
|    I16     |      I32      |      2^28      |   0.201   | 621.409 us |       0.19% | 523.866 us |       0.25% |  -97.544 us | -15.70% |   FAST   |
|    I16     |      I64      |      2^16      |     1     |  15.626 us |       2.87% |  15.146 us |       2.83% |   -0.480 us |  -3.07% |   FAST   |
|    I16     |      I64      |      2^20      |     1     |  22.882 us |       1.96% |  22.075 us |       2.13% |   -0.807 us |  -3.53% |   FAST   |
|    I16     |      I64      |      2^24      |     1     |  65.482 us |       0.84% |  58.395 us |       1.05% |   -7.087 us | -10.82% |   FAST   |
|    I16     |      I64      |      2^28      |     1     | 705.779 us |       0.19% | 599.146 us |       0.17% | -106.633 us | -15.11% |   FAST   |
|    I16     |      I64      |      2^16      |   0.201   |  15.232 us |       3.24% |  14.774 us |       3.34% |   -0.458 us |  -3.01% |   SAME   |
|    I16     |      I64      |      2^20      |   0.201   |  21.770 us |       2.63% |  21.022 us |       2.68% |   -0.748 us |  -3.44% |   FAST   |
|    I16     |      I64      |      2^24      |   0.201   |  58.975 us |       0.98% |  52.286 us |       1.08% |   -6.689 us | -11.34% |   FAST   |
|    I16     |      I64      |      2^28      |   0.201   | 606.815 us |       0.20% | 510.236 us |       0.23% |  -96.580 us | -15.92% |   FAST   |
|    I32     |      I32      |      2^24      |     1     |  64.982 us |       0.99% |  61.051 us |       1.26% |   -3.931 us |  -6.05% |   FAST   |
|    I32     |      I32      |      2^28      |     1     | 696.443 us |       0.25% | 613.061 us |       0.62% |  -83.383 us | -11.97% |   FAST   |
|    I32     |      I64      |      2^24      |     1     |  65.744 us |       1.10% |  61.443 us |       1.37% |   -4.301 us |  -6.54% |   FAST   |
|    I32     |      I64      |      2^28      |     1     | 701.632 us |       0.25% | 615.338 us |       0.65% |  -86.294 us | -12.30% |   FAST   |
|    I64     |      I32      |      2^24      |     1     | 109.395 us |       0.58% | 103.137 us |       0.83% |   -6.258 us |  -5.72% |   FAST   |
|    I64     |      I32      |      2^28      |     1     |   1.392 ms |       0.18% |   1.269 ms |       0.39% | -122.551 us |  -8.80% |   FAST   |
|    I64     |      I32      |      2^24      |   0.201   |  91.789 us |       0.90% |  96.803 us |       0.83% |    5.014 us |   5.46% |   SLOW   |
|    I64     |      I32      |      2^28      |   0.201   |   1.154 ms |       1.11% |   1.218 ms |       0.99% |   63.979 us |   5.54% |   SLOW   |
|    I64     |      I64      |      2^24      |     1     | 110.577 us |       0.71% | 105.699 us |       0.80% |   -4.879 us |  -4.41% |   FAST   |
|    I64     |      I64      |      2^28      |     1     |   1.395 ms |       0.17% |   1.293 ms |       0.31% | -102.579 us |  -7.35% |   FAST   |
|    I64     |      I64      |      2^24      |   0.201   |  92.541 us |       0.80% |  97.500 us |       0.79% |    4.959 us |   5.36% |   SLOW   |
|    I64     |      I64      |      2^28      |   0.201   |   1.155 ms |       1.09% |   1.219 ms |       1.00% |   63.986 us |   5.54% |   SLOW   |
|    I128    |      I32      |      2^24      |     1     | 221.340 us |       0.41% | 205.435 us |       0.51% |  -15.905 us |  -7.19% |   FAST   |
|    I128    |      I32      |      2^28      |     1     |   3.158 ms |       0.18% |   2.880 ms |       0.23% | -278.283 us |  -8.81% |   FAST   |
|    I128    |      I32      |      2^20      |   0.201   |  29.638 us |       2.30% |  30.650 us |       2.02% |    1.012 us |   3.41% |   SLOW   |
|    I128    |      I32      |      2^24      |   0.201   | 165.223 us |       0.92% | 168.584 us |       0.88% |    3.361 us |   2.03% |   SLOW   |
|    I128    |      I64      |      2^24      |     1     | 222.850 us |       0.45% | 204.981 us |       0.52% |  -17.870 us |  -8.02% |   FAST   |
|    I128    |      I64      |      2^28      |     1     |   3.176 ms |       0.17% |   2.871 ms |       0.25% | -305.465 us |  -9.62% |   FAST   |
|    I128    |      I64      |      2^20      |   0.201   |  29.976 us |       2.25% |  30.941 us |       1.95% |    0.965 us |   3.22% |   SLOW   |
|    F32     |      I32      |      2^24      |     1     |  65.320 us |       0.98% |  60.948 us |       1.26% |   -4.372 us |  -6.69% |   FAST   |
|    F32     |      I32      |      2^28      |     1     | 695.382 us |       0.27% | 612.969 us |       0.65% |  -82.413 us | -11.85% |   FAST   |
|    F32     |      I64      |      2^24      |     1     |  65.409 us |       1.09% |  61.347 us |       1.32% |   -4.061 us |  -6.21% |   FAST   |
|    F32     |      I64      |      2^28      |     1     | 701.967 us |       0.26% | 614.687 us |       0.56% |  -87.280 us | -12.43% |   FAST   |
|    F64     |      I32      |      2^24      |     1     | 108.908 us |       0.66% | 104.196 us |       0.83% |   -4.713 us |  -4.33% |   FAST   |
|    F64     |      I32      |      2^28      |     1     |   1.384 ms |       0.19% |   1.281 ms |       0.35% | -103.316 us |  -7.46% |   FAST   |
|    F64     |      I32      |      2^24      |   0.201   |  91.880 us |       0.86% |  97.066 us |       0.83% |    5.186 us |   5.64% |   SLOW   |
|    F64     |      I32      |      2^28      |   0.201   |   1.153 ms |       1.09% |   1.219 ms |       0.97% |   65.510 us |   5.68% |   SLOW   |
|    F64     |      I64      |      2^28      |     1     |   1.392 ms |       0.16% |   1.353 ms |       0.39% |  -39.924 us |  -2.87% |   FAST   |
|    F64     |      I64      |      2^24      |   0.201   |  92.442 us |       0.82% |  99.081 us |       0.80% |    6.639 us |   7.18% |   SLOW   |
|    F64     |      I64      |      2^28      |   0.201   |   1.154 ms |       1.12% |   1.239 ms |       0.86% |   85.358 us |   7.40% |   SLOW   |

bernhardmgruber · 2025-11-05T12:06:35Z

Keys WITH tuning:

keys on H200, #6077 vs this PR, 2% threshold (NOT the final perf after merging to main, look below):

|  KeyT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |       Diff |   %Diff |  Status  |
|------------|---------------|----------------|-----------|------------|-------------|------------|-------------|------------|---------|----------|
|    I64     |      I32      |      2^24      |     1     | 103.747 us |       0.82% | 109.350 us |       0.60% |   5.603 us |   5.40% |   SLOW   |
|    I64     |      I32      |      2^28      |     1     |   1.269 ms |       0.41% |   1.392 ms |       0.17% | 122.890 us |   9.68% |   SLOW   |
|    I64     |      I32      |      2^24      |   0.201   |  96.832 us |       0.80% |  91.797 us |       0.93% |  -5.035 us |  -5.20% |   FAST   |
|    I64     |      I32      |      2^28      |   0.201   |   1.218 ms |       0.97% |   1.154 ms |       1.08% | -64.224 us |  -5.27% |   FAST   |
|    I64     |      I64      |      2^24      |     1     | 105.563 us |       0.81% | 110.189 us |       0.59% |   4.626 us |   4.38% |   SLOW   |
|    I64     |      I64      |      2^28      |     1     |   1.292 ms |       0.30% |   1.395 ms |       0.18% | 103.047 us |   7.97% |   SLOW   |
|    I64     |      I64      |      2^24      |   0.201   |  97.482 us |       0.82% |  92.565 us |       0.80% |  -4.917 us |  -5.04% |   FAST   |
|    I64     |      I64      |      2^28      |   0.201   |   1.219 ms |       0.91% |   1.156 ms |       1.11% | -63.486 us |  -5.21% |   FAST   |
|    F64     |      I32      |      2^24      |     1     | 104.030 us |       0.82% | 109.583 us |       0.84% |   5.553 us |   5.34% |   SLOW   |
|    F64     |      I32      |      2^28      |     1     |   1.280 ms |       0.35% |   1.385 ms |       0.18% | 104.544 us |   8.17% |   SLOW   |
|    F64     |      I32      |      2^24      |   0.201   |  96.977 us |       0.80% |  91.956 us |       0.84% |  -5.021 us |  -5.18% |   FAST   |
|    F64     |      I32      |      2^28      |   0.201   |   1.218 ms |       0.94% |   1.153 ms |       1.10% | -65.123 us |  -5.34% |   FAST   |
|    F64     |      I64      |      2^28      |     1     |   1.352 ms |       0.43% |   1.393 ms |       0.16% |  40.510 us |   3.00% |   SLOW   |
|    F64     |      I64      |      2^24      |   0.201   |  98.996 us |       0.80% |  92.621 us |       0.78% |  -6.375 us |  -6.44% |   FAST   |
|    F64     |      I64      |      2^28      |   0.201   |   1.239 ms |       0.88% |   1.154 ms |       1.12% | -85.082 us |  -6.87% |   FAST   |

keys on H200, main vs this PR, 2% threshold:

|  KeyT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|------------|---------------|----------------|-----------|------------|-------------|------------|-------------|-------------|---------|----------|
|     I8     |      I32      |      2^24      |     1     |  53.553 us |       0.93% |  49.545 us |       1.04% |   -4.008 us |  -7.48% |   FAST   |
|     I8     |      I32      |      2^28      |     1     | 524.000 us |       0.22% | 456.986 us |       0.21% |  -67.014 us | -12.79% |   FAST   |
|     I8     |      I32      |      2^24      |   0.201   |  51.930 us |       1.00% |  47.660 us |       1.05% |   -4.270 us |  -8.22% |   FAST   |
|     I8     |      I32      |      2^28      |   0.201   | 518.773 us |       0.22% | 451.773 us |       0.20% |  -67.000 us | -12.92% |   FAST   |
|     I8     |      I64      |      2^24      |     1     |  53.475 us |       1.01% |  49.824 us |       1.14% |   -3.651 us |  -6.83% |   FAST   |
|     I8     |      I64      |      2^28      |     1     | 513.522 us |       0.21% | 460.018 us |       0.20% |  -53.504 us | -10.42% |   FAST   |
|     I8     |      I64      |      2^24      |   0.201   |  51.747 us |       1.09% |  47.664 us |       1.17% |   -4.083 us |  -7.89% |   FAST   |
|     I8     |      I64      |      2^28      |   0.201   | 508.798 us |       0.21% | 454.412 us |       0.20% |  -54.387 us | -10.69% |   FAST   |
|    I16     |      I32      |      2^24      |     1     |  66.515 us |       0.80% |  59.333 us |       0.92% |   -7.182 us | -10.80% |   FAST   |
|    I16     |      I32      |      2^28      |     1     | 739.462 us |       0.22% | 614.621 us |       0.17% | -124.841 us | -16.88% |   FAST   |
|    I16     |      I32      |      2^24      |   0.201   |  59.502 us |       0.95% |  53.252 us |       1.08% |   -6.250 us | -10.50% |   FAST   |
|    I16     |      I32      |      2^28      |   0.201   | 621.409 us |       0.19% | 523.954 us |       0.25% |  -97.456 us | -15.68% |   FAST   |
|    I16     |      I64      |      2^16      |     1     |  15.626 us |       2.87% |  15.117 us |       2.79% |   -0.509 us |  -3.26% |   FAST   |
|    I16     |      I64      |      2^20      |     1     |  22.882 us |       1.96% |  22.079 us |       2.05% |   -0.803 us |  -3.51% |   FAST   |
|    I16     |      I64      |      2^24      |     1     |  65.482 us |       0.84% |  58.191 us |       1.02% |   -7.291 us | -11.13% |   FAST   |
|    I16     |      I64      |      2^28      |     1     | 705.779 us |       0.19% | 599.145 us |       0.18% | -106.634 us | -15.11% |   FAST   |
|    I16     |      I64      |      2^16      |   0.201   |  15.232 us |       3.24% |  14.765 us |       3.53% |   -0.467 us |  -3.07% |   SAME   |
|    I16     |      I64      |      2^20      |   0.201   |  21.770 us |       2.63% |  21.028 us |       2.72% |   -0.743 us |  -3.41% |   FAST   |
|    I16     |      I64      |      2^24      |   0.201   |  58.975 us |       0.98% |  52.248 us |       1.03% |   -6.728 us | -11.41% |   FAST   |
|    I16     |      I64      |      2^28      |   0.201   | 606.815 us |       0.20% | 510.240 us |       0.23% |  -96.575 us | -15.92% |   FAST   |
|    I32     |      I32      |      2^24      |     1     |  64.982 us |       0.99% |  60.799 us |       1.14% |   -4.183 us |  -6.44% |   FAST   |
|    I32     |      I32      |      2^28      |     1     | 696.443 us |       0.25% | 613.023 us |       0.62% |  -83.420 us | -11.98% |   FAST   |
|    I32     |      I64      |      2^24      |     1     |  65.744 us |       1.10% |  61.541 us |       1.35% |   -4.203 us |  -6.39% |   FAST   |
|    I32     |      I64      |      2^28      |     1     | 701.632 us |       0.25% | 615.223 us |       0.61% |  -86.408 us | -12.32% |   FAST   |
|    I128    |      I32      |      2^24      |     1     | 221.340 us |       0.41% | 204.987 us |       0.56% |  -16.354 us |  -7.39% |   FAST   |
|    I128    |      I32      |      2^28      |     1     |   3.158 ms |       0.18% |   2.880 ms |       0.23% | -278.312 us |  -8.81% |   FAST   |
|    I128    |      I32      |      2^20      |   0.201   |  29.638 us |       2.30% |  30.614 us |       2.07% |    0.976 us |   3.29% |   SLOW   |
|    I128    |      I32      |      2^24      |   0.201   | 165.223 us |       0.92% | 168.570 us |       0.85% |    3.347 us |   2.03% |   SLOW   |
|    I128    |      I64      |      2^24      |     1     | 222.850 us |       0.45% | 204.981 us |       0.55% |  -17.869 us |  -8.02% |   FAST   |
|    I128    |      I64      |      2^28      |     1     |   3.176 ms |       0.17% |   2.870 ms |       0.24% | -305.749 us |  -9.63% |   FAST   |
|    I128    |      I64      |      2^20      |   0.201   |  29.976 us |       2.25% |  30.916 us |       1.87% |    0.940 us |   3.14% |   SLOW   |
|    F32     |      I32      |      2^24      |     1     |  65.320 us |       0.98% |  60.886 us |       1.24% |   -4.434 us |  -6.79% |   FAST   |
|    F32     |      I32      |      2^28      |     1     | 695.382 us |       0.27% | 612.822 us |       0.61% |  -82.560 us | -11.87% |   FAST   |
|    F32     |      I64      |      2^24      |     1     |  65.409 us |       1.09% |  61.374 us |       1.34% |   -4.035 us |  -6.17% |   FAST   |
|    F32     |      I64      |      2^28      |     1     | 701.967 us |       0.26% | 614.665 us |       0.55% |  -87.302 us | -12.44% |   FAST   |

We can see that disabling BlockLoadToShared for sizeof(KeyT) == 8 regresses over #6077, but avoids regression against main, which is the goal.

bernhardmgruber · 2025-11-05T12:23:32Z

On B200, keys, main vs this PR, 2% threshold:

|  KeyT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|------------|---------------|----------------|-----------|------------|-------------|------------|-------------|-------------|---------|----------|
|     I8     |      I32      |      2^24      |     1     |  49.111 us |       0.30% |  46.986 us |       0.49% |   -2.125 us |  -4.33% |   FAST   |
|     I8     |      I32      |      2^28      |     1     | 426.097 us |       0.12% | 393.235 us |       0.11% |  -32.863 us |  -7.71% |   FAST   |
|     I8     |      I32      |      2^24      |   0.201   |  47.064 us |       0.38% |  45.173 us |       1.43% |   -1.891 us |  -4.02% |   FAST   |
|     I8     |      I32      |      2^28      |   0.201   | 423.558 us |       0.17% | 389.140 us |       0.14% |  -34.418 us |  -8.13% |   FAST   |
|     I8     |      I64      |      2^24      |     1     |  49.130 us |       0.45% |  46.983 us |       0.35% |   -2.148 us |  -4.37% |   FAST   |
|     I8     |      I64      |      2^28      |     1     | 427.941 us |       0.16% | 382.031 us |       0.32% |  -45.910 us | -10.73% |   FAST   |
|     I8     |      I64      |      2^24      |   0.201   |  46.972 us |       0.42% |  44.850 us |       0.20% |   -2.122 us |  -4.52% |   FAST   |
|     I8     |      I64      |      2^28      |   0.201   | 424.527 us |       0.22% | 378.856 us |       0.05% |  -45.671 us | -10.76% |   FAST   |
|    I16     |      I32      |      2^24      |     1     |  59.403 us |       0.60% |  55.963 us |       1.73% |   -3.441 us |  -5.79% |   FAST   |
|    I16     |      I32      |      2^28      |     1     | 582.168 us |       0.20% | 534.424 us |       0.14% |  -47.745 us |  -8.20% |   FAST   |
|    I16     |      I32      |      2^24      |   0.201   |  52.741 us |       2.84% |  50.120 us |       1.98% |   -2.621 us |  -4.97% |   FAST   |
|    I16     |      I32      |      2^28      |   0.201   | 493.831 us |       0.24% | 457.971 us |       0.29% |  -35.860 us |  -7.26% |   FAST   |
|    I16     |      I64      |      2^24      |   0.201   |  52.420 us |       1.82% |  50.507 us |       1.68% |   -1.913 us |  -3.65% |   FAST   |
|    I16     |      I64      |      2^28      |   0.201   | 473.467 us |       0.26% | 457.576 us |       0.24% |  -15.892 us |  -3.36% |   FAST   |
|    I32     |      I32      |      2^24      |     1     |  59.445 us |       0.62% |  55.333 us |       0.64% |   -4.113 us |  -6.92% |   FAST   |
|    I32     |      I32      |      2^28      |     1     | 575.404 us |       0.18% | 499.993 us |       0.16% |  -75.411 us | -13.11% |   FAST   |
|    I32     |      I32      |      2^24      |   0.201   |  52.719 us |       1.67% |  50.629 us |       1.74% |   -2.090 us |  -3.96% |   FAST   |
|    I32     |      I32      |      2^28      |   0.201   | 456.562 us |       0.11% | 425.839 us |       0.10% |  -30.723 us |  -6.73% |   FAST   |
|    I32     |      I64      |      2^24      |     1     |  59.473 us |       1.01% |  55.385 us |       1.12% |   -4.088 us |  -6.87% |   FAST   |
|    I32     |      I64      |      2^28      |     1     | 575.088 us |       0.25% | 498.816 us |       0.26% |  -76.272 us | -13.26% |   FAST   |
|    I32     |      I64      |      2^24      |   0.201   |  52.759 us |       1.47% |  49.685 us |       2.01% |   -3.074 us |  -5.83% |   FAST   |
|    I32     |      I64      |      2^28      |   0.201   | 450.063 us |       0.19% | 422.680 us |       0.24% |  -27.383 us |  -6.08% |   FAST   |
|    I64     |      I32      |      2^24      |     1     |  97.328 us |       1.05% |  90.103 us |       0.20% |   -7.225 us |  -7.42% |   FAST   |
|    I64     |      I32      |      2^28      |     1     |   1.189 ms |       0.10% |   1.056 ms |       0.12% | -132.981 us | -11.18% |   FAST   |
|    I64     |      I32      |      2^24      |   0.201   |  73.428 us |       0.91% |  71.413 us |       0.86% |   -2.015 us |  -2.74% |   FAST   |
|    I64     |      I32      |      2^28      |   0.201   | 780.502 us |       0.15% | 759.669 us |       0.16% |  -20.834 us |  -2.67% |   FAST   |
|    I64     |      I64      |      2^24      |     1     |  98.066 us |       0.60% |  90.192 us |       0.63% |   -7.875 us |  -8.03% |   FAST   |
|    I64     |      I64      |      2^28      |     1     |   1.197 ms |       0.11% |   1.072 ms |       0.13% | -125.076 us | -10.45% |   FAST   |
|    I64     |      I64      |      2^24      |   0.201   |  73.152 us |       1.24% |  71.436 us |       0.51% |   -1.716 us |  -2.35% |   FAST   |
|    I128    |      I32      |      2^24      |     1     | 204.807 us |       0.10% | 178.176 us |       0.00% |  -26.631 us | -13.00% |   ????   |
|    I128    |      I32      |      2^28      |     1     |   2.930 ms |       0.05% |   2.488 ms |       0.06% | -441.426 us | -15.07% |   FAST   |
|    I128    |      I32      |      2^24      |   0.201   | 141.322 us |       0.28% | 121.007 us |       0.57% |  -20.315 us | -14.38% |   FAST   |
|    I128    |      I32      |      2^28      |   0.201   |   1.923 ms |       0.10% |   1.600 ms |       0.10% | -322.221 us | -16.76% |   FAST   |
|    I128    |      I64      |      2^24      |     1     | 203.990 us |       0.55% | 176.532 us |       0.52% |  -27.457 us | -13.46% |   FAST   |
|    I128    |      I64      |      2^28      |     1     |   2.934 ms |       0.05% |   2.467 ms |       0.06% | -466.352 us | -15.90% |   FAST   |
|    I128    |      I64      |      2^24      |   0.201   | 131.328 us |       0.68% | 122.418 us |       0.86% |   -8.910 us |  -6.78% |   FAST   |
|    I128    |      I64      |      2^28      |   0.201   |   1.770 ms |       0.08% |   1.609 ms |       0.09% | -161.688 us |  -9.13% |   FAST   |
|    F32     |      I32      |      2^24      |     1     |  59.384 us |       0.19% |  55.359 us |       0.72% |   -4.025 us |  -6.78% |   FAST   |
|    F32     |      I32      |      2^28      |     1     | 575.434 us |       0.24% | 500.952 us |       0.20% |  -74.482 us | -12.94% |   FAST   |
|    F32     |      I32      |      2^24      |   0.201   |  52.794 us |       1.56% |  50.607 us |       1.79% |   -2.186 us |  -4.14% |   FAST   |
|    F32     |      I32      |      2^28      |   0.201   | 456.597 us |       0.08% | 425.786 us |       0.13% |  -30.811 us |  -6.75% |   FAST   |
|    F32     |      I64      |      2^24      |     1     |  59.593 us |       1.14% |  55.636 us |       1.48% |   -3.957 us |  -6.64% |   FAST   |
|    F32     |      I64      |      2^28      |     1     | 576.243 us |       0.22% | 499.657 us |       0.27% |  -76.586 us | -13.29% |   FAST   |
|    F32     |      I64      |      2^24      |   0.201   |  52.789 us |       1.42% |  49.640 us |       2.08% |   -3.149 us |  -5.97% |   FAST   |
|    F32     |      I64      |      2^28      |   0.201   | 449.673 us |       0.23% | 423.498 us |       0.19% |  -26.175 us |  -5.82% |   FAST   |
|    F64     |      I32      |      2^24      |     1     |  97.957 us |       0.79% |  90.088 us |       0.19% |   -7.868 us |  -8.03% |   FAST   |
|    F64     |      I32      |      2^28      |     1     |   1.191 ms |       0.10% |   1.057 ms |       0.14% | -134.383 us | -11.28% |   FAST   |
|    F64     |      I32      |      2^24      |   0.201   |  73.393 us |       0.97% |  70.381 us |       1.63% |   -3.012 us |  -4.10% |   FAST   |
|    F64     |      I32      |      2^28      |   0.201   | 787.483 us |       1.26% | 757.011 us |       0.92% |  -30.472 us |  -3.87% |   FAST   |
|    F64     |      I64      |      2^16      |     1     |  18.428 us |       1.26% |  19.081 us |       5.82% |    0.653 us |   3.54% |   SLOW   |
|    F64     |      I64      |      2^24      |     1     |  98.275 us |       0.16% |  91.068 us |       1.16% |   -7.207 us |  -7.33% |   FAST   |
|    F64     |      I64      |      2^28      |     1     |   1.201 ms |       0.09% |   1.066 ms |       0.10% | -135.156 us | -11.25% |   FAST   |
|    F64     |      I64      |      2^24      |   0.201   |  73.589 us |       0.14% |  71.481 us |       0.12% |   -2.109 us |  -2.87% |   FAST   |
|    F64     |      I64      |      2^28      |   0.201   | 788.293 us |       0.16% | 763.760 us |       0.15% |  -24.533 us |  -3.11% |   FAST   |

On B200, pairs, main vs this PR, 2% threshold:

|  KeyT{ct}  |  ValueT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|------------|--------------|---------------|----------------|-----------|------------|-------------|------------|-------------|-------------|---------|----------|
|     I8     |      I8      |      I32      |      2^16      |     1     |  18.903 us |       5.34% |  18.232 us |       1.72% |   -0.671 us |  -3.55% |   FAST   |
|     I8     |      I8      |      I32      |      2^20      |     1     |  23.972 us |       3.71% |  22.513 us |       1.41% |   -1.459 us |  -6.08% |   FAST   |
|     I8     |      I8      |      I32      |      2^24      |     1     | 100.262 us |       0.12% |  73.697 us |       0.26% |  -26.565 us | -26.50% |   FAST   |
|     I8     |      I8      |      I32      |      2^28      |     1     |   1.252 ms |       0.07% | 786.443 us |       0.05% | -465.344 us | -37.17% |   FAST   |
|     I8     |      I8      |      I32      |      2^24      |   0.201   |  98.238 us |       0.19% |  71.621 us |       0.13% |  -26.616 us | -27.09% |   FAST   |
|     I8     |      I8      |      I32      |      2^28      |   0.201   |   1.247 ms |       0.06% | 780.144 us |       0.07% | -467.352 us | -37.46% |   FAST   |
|     I8     |      I8      |      I64      |      2^16      |     1     |  19.109 us |       5.19% |  18.523 us |       2.78% |   -0.586 us |  -3.07% |   FAST   |
|     I8     |      I8      |      I64      |      2^20      |     1     |  24.130 us |       3.23% |  22.705 us |       2.83% |   -1.425 us |  -5.91% |   FAST   |
|     I8     |      I8      |      I64      |      2^24      |     1     | 102.396 us |       0.41% |  96.249 us |       0.29% |   -6.147 us |  -6.00% |   FAST   |
|     I8     |      I8      |      I64      |      2^28      |     1     |   1.316 ms |       0.08% |   1.178 ms |       0.06% | -137.973 us | -10.49% |   FAST   |
|     I8     |      I8      |      I64      |      2^24      |   0.201   | 101.215 us |       1.01% |  94.043 us |       0.10% |   -7.172 us |  -7.09% |   FAST   |
|     I8     |      I8      |      I64      |      2^28      |   0.201   |   1.312 ms |       0.08% |   1.170 ms |       0.06% | -141.894 us | -10.82% |   FAST   |
|     I8     |     I16      |      I32      |      2^16      |     1     |  19.621 us |       5.63% |  18.419 us |       3.01% |   -1.203 us |  -6.13% |   FAST   |
|     I8     |     I16      |      I32      |      2^20      |     1     |  24.524 us |       0.98% |  22.980 us |       3.72% |   -1.545 us |  -6.30% |   FAST   |
|     I8     |     I16      |      I32      |      2^24      |     1     | 102.601 us |       0.62% |  73.711 us |       0.07% |  -28.889 us | -28.16% |   FAST   |
|     I8     |     I16      |      I32      |      2^28      |     1     |   1.281 ms |       0.08% | 792.607 us |       0.08% | -488.082 us | -38.11% |   FAST   |
|     I8     |     I16      |      I32      |      2^24      |   0.201   | 101.082 us |       1.00% |  71.623 us |       0.17% |  -29.459 us | -29.14% |   FAST   |
|     I8     |     I16      |      I32      |      2^28      |   0.201   |   1.276 ms |       0.05% | 788.674 us |       0.10% | -487.121 us | -38.18% |   FAST   |
|     I8     |     I16      |      I64      |      2^16      |     1     |  19.896 us |       4.40% |  18.419 us |       2.18% |   -1.477 us |  -7.42% |   FAST   |
|     I8     |     I16      |      I64      |      2^20      |     1     |  24.483 us |       0.49% |  22.621 us |       2.47% |   -1.862 us |  -7.61% |   FAST   |
|     I8     |     I16      |      I64      |      2^24      |     1     | 106.409 us |       0.62% |  73.682 us |       0.36% |  -32.727 us | -30.76% |   FAST   |
|     I8     |     I16      |      I64      |      2^28      |     1     |   1.347 ms |       0.07% | 796.764 us |       0.09% | -550.506 us | -40.86% |   FAST   |
|     I8     |     I16      |      I64      |      2^16      |   0.201   |  18.727 us |       4.61% |  18.258 us |       0.53% |   -0.470 us |  -2.51% |   FAST   |
|     I8     |     I16      |      I64      |      2^20      |   0.201   |  23.495 us |       4.32% |  22.359 us |       0.40% |   -1.136 us |  -4.84% |   FAST   |
|     I8     |     I16      |      I64      |      2^24      |   0.201   | 104.260 us |       0.21% |  71.622 us |       0.63% |  -32.638 us | -31.30% |   FAST   |
|     I8     |     I16      |      I64      |      2^28      |   0.201   |   1.344 ms |       0.06% | 794.496 us |       0.08% | -549.304 us | -40.88% |   FAST   |
|     I8     |     I32      |      I32      |      2^20      |     1     |  25.995 us |       3.59% |  24.571 us |       0.71% |   -1.424 us |  -5.48% |   FAST   |
|     I8     |     I32      |      I32      |      2^24      |     1     |  88.599 us |       1.10% |  81.264 us |       1.25% |   -7.335 us |  -8.28% |   FAST   |
|     I8     |     I32      |      I32      |      2^28      |     1     |   1.048 ms |       0.16% | 919.589 us |       0.11% | -128.454 us | -12.26% |   FAST   |
|     I8     |     I32      |      I32      |      2^24      |   0.201   |  87.960 us |       0.31% |  79.822 us |       0.11% |   -8.138 us |  -9.25% |   FAST   |
|     I8     |     I32      |      I32      |      2^28      |   0.201   |   1.041 ms |       0.14% | 913.807 us |       0.11% | -127.462 us | -12.24% |   FAST   |
|     I8     |     I32      |      I64      |      2^20      |     1     |  25.051 us |       3.63% |  24.517 us |       0.56% |   -0.534 us |  -2.13% |   FAST   |
|     I8     |     I32      |      I64      |      2^24      |     1     |  88.139 us |       0.60% |  82.126 us |       0.82% |   -6.013 us |  -6.82% |   FAST   |
|     I8     |     I32      |      I64      |      2^28      |     1     |   1.055 ms |       0.09% | 923.405 us |       0.10% | -131.744 us | -12.49% |   FAST   |
|     I8     |     I32      |      I64      |      2^24      |   0.201   |  87.295 us |       1.09% |  79.721 us |       0.11% |   -7.574 us |  -8.68% |   FAST   |
|     I8     |     I32      |      I64      |      2^28      |   0.201   |   1.051 ms |       0.09% | 917.790 us |       0.11% | -133.157 us | -12.67% |   FAST   |
|     I8     |     I64      |      I32      |      2^16      |     1     |  17.509 us |       5.81% |  18.284 us |       2.76% |    0.775 us |   4.43% |   SLOW   |
|     I8     |     I64      |      I32      |      2^24      |     1     |  97.297 us |       1.16% |  87.801 us |       0.84% |   -9.496 us |  -9.76% |   FAST   |
|     I8     |     I64      |      I32      |      2^28      |     1     |   1.211 ms |       0.13% |   1.042 ms |       0.14% | -169.132 us | -13.97% |   FAST   |
|     I8     |     I64      |      I32      |      2^24      |   0.201   |  96.182 us |       0.19% |  86.100 us |       0.58% |  -10.082 us | -10.48% |   FAST   |
|     I8     |     I64      |      I32      |      2^28      |   0.201   |   1.195 ms |       0.14% |   1.026 ms |       0.13% | -169.637 us | -14.19% |   FAST   |
|     I8     |     I64      |      I64      |      2^16      |     1     |  17.377 us |       5.83% |  18.099 us |       3.95% |    0.722 us |   4.15% |   SLOW   |
|     I8     |     I64      |      I64      |      2^24      |     1     |  97.793 us |       1.40% |  88.070 us |       1.16% |   -9.724 us |  -9.94% |   FAST   |
|     I8     |     I64      |      I64      |      2^28      |     1     |   1.218 ms |       0.15% |   1.042 ms |       0.13% | -176.155 us | -14.46% |   FAST   |
|     I8     |     I64      |      I64      |      2^24      |   0.201   |  96.230 us |       0.59% |  85.887 us |       0.40% |  -10.343 us | -10.75% |   FAST   |
|     I8     |     I64      |      I64      |      2^28      |   0.201   |   1.203 ms |       0.14% |   1.025 ms |       0.14% | -178.629 us | -14.84% |   FAST   |
|    I16     |      I8      |      I32      |      2^16      |     1     |  20.390 us |       1.89% |  18.718 us |       3.83% |   -1.672 us |  -8.20% |   FAST   |
|    I16     |      I8      |      I32      |      2^20      |     1     |  25.191 us |       3.74% |  24.456 us |       2.91% |   -0.735 us |  -2.92% |   FAST   |
|    I16     |      I8      |      I32      |      2^24      |     1     | 122.854 us |       0.10% | 113.267 us |       0.84% |   -9.587 us |  -7.80% |   FAST   |
|    I16     |      I8      |      I32      |      2^28      |     1     |   1.548 ms |       0.06% |   1.423 ms |       0.08% | -125.501 us |  -8.11% |   FAST   |
|    I16     |      I8      |      I32      |      2^16      |   0.201   |  19.726 us |       4.88% |  18.406 us |       0.33% |   -1.320 us |  -6.69% |   FAST   |
|    I16     |      I8      |      I32      |      2^20      |   0.201   |  24.593 us |       1.63% |  22.852 us |       3.63% |   -1.741 us |  -7.08% |   FAST   |
|    I16     |      I8      |      I32      |      2^24      |   0.201   | 109.083 us |       0.83% | 102.744 us |       1.49% |   -6.339 us |  -5.81% |   FAST   |
|    I16     |      I8      |      I32      |      2^28      |   0.201   |   1.383 ms |       0.09% |   1.276 ms |       0.11% | -107.035 us |  -7.74% |   FAST   |
|    I16     |      I8      |      I64      |      2^16      |     1     |  20.398 us |       0.57% |  18.746 us |       4.19% |   -1.652 us |  -8.10% |   FAST   |
|    I16     |      I8      |      I64      |      2^20      |     1     |  25.617 us |       3.92% |  24.586 us |       5.71% |   -1.031 us |  -4.02% |   FAST   |
|    I16     |      I8      |      I64      |      2^24      |     1     | 122.815 us |       0.08% |  88.022 us |       0.16% |  -34.793 us | -28.33% |   FAST   |
|    I16     |      I8      |      I64      |      2^28      |     1     |   1.550 ms |       0.06% |   1.008 ms |       0.09% | -542.272 us | -34.98% |   FAST   |
|    I16     |      I8      |      I64      |      2^16      |   0.201   |  19.628 us |       5.16% |  18.293 us |       0.93% |   -1.335 us |  -6.80% |   FAST   |
|    I16     |      I8      |      I64      |      2^20      |   0.201   |  24.678 us |       2.88% |  22.938 us |       4.19% |   -1.740 us |  -7.05% |   FAST   |
|    I16     |      I8      |      I64      |      2^24      |   0.201   | 113.560 us |       1.11% |  79.919 us |       0.91% |  -33.641 us | -29.62% |   FAST   |
|    I16     |      I8      |      I64      |      2^28      |   0.201   |   1.435 ms |       0.09% | 896.053 us |       0.15% | -538.660 us | -37.54% |   FAST   |
|    I16     |     I16      |      I32      |      2^16      |     1     |  20.218 us |       3.30% |  18.434 us |       1.02% |   -1.784 us |  -8.82% |   FAST   |
|    I16     |     I16      |      I32      |      2^24      |     1     | 126.021 us |       0.82% |  88.064 us |       0.00% |  -37.957 us | -30.12% |   ????   |
|    I16     |     I16      |      I32      |      2^28      |     1     |   1.579 ms |       0.07% | 995.085 us |       0.10% | -583.972 us | -36.98% |   FAST   |
|    I16     |     I16      |      I32      |      2^16      |   0.201   |  19.085 us |       5.55% |  18.509 us |       2.37% |   -0.577 us |  -3.02% |   FAST   |
|    I16     |     I16      |      I32      |      2^20      |   0.201   |  24.576 us |       0.00% |  23.615 us |       4.61% |   -0.961 us |  -3.91% |   ????   |
|    I16     |     I16      |      I32      |      2^24      |   0.201   | 115.642 us |       0.88% |  80.489 us |       1.55% |  -35.153 us | -30.40% |   FAST   |
|    I16     |     I16      |      I32      |      2^28      |   0.201   |   1.463 ms |       0.08% | 906.525 us |       0.15% | -556.440 us | -38.04% |   FAST   |
|    I16     |     I16      |      I64      |      2^16      |     1     |  19.875 us |       5.15% |  18.439 us |       1.60% |   -1.435 us |  -7.22% |   FAST   |
|    I16     |     I16      |      I64      |      2^24      |     1     | 124.791 us |       0.27% |  88.835 us |       1.13% |  -35.956 us | -28.81% |   FAST   |
|    I16     |     I16      |      I64      |      2^28      |     1     |   1.558 ms |       0.07% |   1.003 ms |       0.12% | -555.487 us | -35.65% |   FAST   |
|    I16     |     I16      |      I64      |      2^16      |   0.201   |  19.500 us |       7.60% |  18.733 us |       4.47% |   -0.768 us |  -3.94% |   SAME   |
|    I16     |     I16      |      I64      |      2^20      |   0.201   |  25.008 us |       3.72% |  23.496 us |       5.57% |   -1.512 us |  -6.05% |   FAST   |
|    I16     |     I16      |      I64      |      2^24      |   0.201   | 114.508 us |       1.17% |  78.892 us |       1.39% |  -35.616 us | -31.10% |   FAST   |
|    I16     |     I16      |      I64      |      2^28      |   0.201   |   1.444 ms |       0.10% | 885.622 us |       0.15% | -558.821 us | -38.69% |   FAST   |
|    I16     |     I32      |      I32      |      2^24      |     1     | 102.567 us |       0.57% |  81.920 us |       0.25% |  -20.647 us | -20.13% |   FAST   |
|    I16     |     I32      |      I32      |      2^28      |     1     |   1.238 ms |       0.11% | 914.044 us |       0.13% | -324.183 us | -26.18% |   FAST   |
|    I16     |     I32      |      I32      |      2^16      |   0.201   |  18.390 us |       1.17% |  17.736 us |       5.60% |   -0.654 us |  -3.56% |   FAST   |
|    I16     |     I32      |      I32      |      2^20      |   0.201   |  24.779 us |       2.65% |  23.498 us |       5.33% |   -1.281 us |  -5.17% |   FAST   |
|    I16     |     I32      |      I32      |      2^24      |   0.201   |  96.444 us |       1.21% |  75.339 us |       1.55% |  -21.105 us | -21.88% |   FAST   |
|    I16     |     I32      |      I32      |      2^28      |   0.201   |   1.149 ms |       0.11% | 817.758 us |       0.14% | -331.622 us | -28.85% |   FAST   |
|    I16     |     I32      |      I64      |      2^24      |     1     | 100.317 us |       0.25% |  81.879 us |       0.09% |  -18.439 us | -18.38% |   FAST   |
|    I16     |     I32      |      I64      |      2^28      |     1     |   1.203 ms |       0.10% | 913.203 us |       0.17% | -289.591 us | -24.08% |   FAST   |
|    I16     |     I32      |      I64      |      2^16      |   0.201   |  18.849 us |       4.89% |  18.136 us |       7.80% |   -0.714 us |  -3.79% |   SAME   |
|    I16     |     I32      |      I64      |      2^24      |   0.201   |  94.118 us |       0.47% |  76.094 us |       1.22% |  -18.024 us | -19.15% |   FAST   |
|    I16     |     I32      |      I64      |      2^28      |   0.201   |   1.117 ms |       0.10% | 819.392 us |       0.14% | -297.367 us | -26.63% |   FAST   |
|    I16     |     I64      |      I32      |      2^24      |     1     | 114.071 us |       0.82% | 106.494 us |       0.11% |   -7.576 us |  -6.64% |   FAST   |
|    I16     |     I64      |      I32      |      2^28      |     1     |   1.433 ms |       0.11% |   1.321 ms |       0.07% | -112.051 us |  -7.82% |   FAST   |
|    I16     |     I64      |      I32      |      2^24      |   0.201   | 104.215 us |       1.11% |  95.289 us |       1.24% |   -8.926 us |  -8.57% |   FAST   |
|    I16     |     I64      |      I32      |      2^28      |   0.201   |   1.299 ms |       0.16% |   1.136 ms |       0.13% | -162.936 us | -12.54% |   FAST   |
|    I16     |     I64      |      I64      |      2^20      |     1     |  29.436 us |       3.51% |  28.633 us |       0.26% |   -0.803 us |  -2.73% |   FAST   |
|    I16     |     I64      |      I64      |      2^24      |     1     | 129.352 us |       0.63% | 107.868 us |       0.91% |  -21.483 us | -16.61% |   FAST   |
|    I16     |     I64      |      I64      |      2^28      |     1     |   1.713 ms |       0.08% |   1.337 ms |       0.11% | -375.777 us | -21.94% |   FAST   |
|    I16     |     I64      |      I64      |      2^16      |   0.201   |  17.495 us |       7.87% |  18.540 us |       9.01% |    1.045 us |   5.97% |   SAME   |
|    I16     |     I64      |      I64      |      2^20      |   0.201   |  27.898 us |       3.21% |  25.575 us |       3.85% |   -2.324 us |  -8.33% |   FAST   |
|    I16     |     I64      |      I64      |      2^24      |   0.201   | 119.381 us |       0.93% |  96.056 us |       0.36% |  -23.325 us | -19.54% |   FAST   |
|    I16     |     I64      |      I64      |      2^28      |   0.201   |   1.554 ms |       0.12% |   1.146 ms |       0.12% | -407.987 us | -26.25% |   FAST   |
|    I32     |      I8      |      I32      |      2^20      |     1     |  28.070 us |       3.30% |  26.611 us |       0.18% |   -1.459 us |  -5.20% |   FAST   |
|    I32     |      I8      |      I32      |      2^24      |     1     |  94.819 us |       1.13% |  86.644 us |       1.28% |   -8.175 us |  -8.62% |   FAST   |
|    I32     |      I8      |      I32      |      2^28      |     1     |   1.123 ms |       0.10% |   1.016 ms |       0.09% | -107.236 us |  -9.55% |   FAST   |
|    I32     |      I8      |      I32      |      2^16      |   0.201   |  19.389 us |       5.19% |  18.936 us |       8.50% |   -0.452 us |  -2.33% |   SAME   |
|    I32     |      I8      |      I32      |      2^24      |   0.201   |  86.323 us |       0.93% |  78.117 us |       0.97% |   -8.206 us |  -9.51% |   FAST   |
|    I32     |      I8      |      I32      |      2^28      |   0.201   |   1.002 ms |       0.13% | 878.617 us |       0.14% | -123.649 us | -12.34% |   FAST   |
|    I32     |      I8      |      I64      |      2^20      |     1     |  27.789 us |       3.58% |  26.586 us |       0.80% |   -1.203 us |  -4.33% |   FAST   |
|    I32     |      I8      |      I64      |      2^24      |     1     |  95.074 us |       1.14% |  90.181 us |       1.55% |   -4.893 us |  -5.15% |   FAST   |
|    I32     |      I8      |      I64      |      2^28      |     1     |   1.121 ms |       0.10% |   1.057 ms |       0.13% |  -64.083 us |  -5.72% |   FAST   |
|    I32     |      I8      |      I64      |      2^24      |   0.201   |  86.363 us |       1.15% |  80.938 us |       1.30% |   -5.425 us |  -6.28% |   FAST   |
|    I32     |      I8      |      I64      |      2^28      |   0.201   |   1.004 ms |       0.13% | 911.841 us |       0.18% |  -92.651 us |  -9.22% |   FAST   |
|    I32     |     I16      |      I32      |      2^24      |     1     | 100.638 us |       0.73% |  80.114 us |       0.84% |  -20.523 us | -20.39% |   FAST   |
|    I32     |     I16      |      I32      |      2^28      |     1     |   1.231 ms |       0.09% | 903.631 us |       0.12% | -327.222 us | -26.58% |   FAST   |
|    I32     |     I16      |      I32      |      2^16      |   0.201   |  19.836 us |       4.49% |  19.394 us |       7.14% |   -0.442 us |  -2.23% |   SAME   |
|    I32     |     I16      |      I32      |      2^20      |   0.201   |  26.484 us |       1.18% |  25.235 us |       4.02% |   -1.249 us |  -4.72% |   FAST   |
|    I32     |     I16      |      I32      |      2^24      |   0.201   |  91.848 us |       0.72% |  71.705 us |       0.45% |  -20.143 us | -21.93% |   FAST   |
|    I32     |     I16      |      I32      |      2^28      |   0.201   |   1.089 ms |       0.10% | 776.356 us |       0.17% | -312.858 us | -28.72% |   FAST   |
|    I32     |     I16      |      I64      |      2^24      |     1     | 100.819 us |       1.51% |  80.817 us |       1.30% |  -20.002 us | -19.84% |   FAST   |
|    I32     |     I16      |      I64      |      2^28      |     1     |   1.224 ms |       0.10% | 909.583 us |       0.11% | -314.338 us | -25.68% |   FAST   |
|    I32     |     I16      |      I64      |      2^20      |   0.201   |  26.128 us |       2.64% |  24.496 us |       1.62% |   -1.632 us |  -6.25% |   FAST   |
|    I32     |     I16      |      I64      |      2^24      |   0.201   |  90.169 us |       0.79% |  71.872 us |       1.17% |  -18.297 us | -20.29% |   FAST   |
|    I32     |     I16      |      I64      |      2^28      |   0.201   |   1.085 ms |       0.11% | 780.531 us |       0.17% | -304.255 us | -28.05% |   FAST   |
|    I32     |     I32      |      I32      |      2^24      |     1     |  95.318 us |       1.20% |  92.916 us |       1.07% |   -2.402 us |  -2.52% |   FAST   |
|    I32     |     I32      |      I32      |      2^28      |     1     |   1.166 ms |       0.10% |   1.129 ms |       0.11% |  -37.569 us |  -3.22% |   FAST   |
|    I32     |     I32      |      I32      |      2^28      |   0.201   | 993.212 us |       0.15% | 970.701 us |       0.13% |  -22.511 us |  -2.27% |   FAST   |
|    I32     |     I32      |      I64      |      2^24      |     1     |  96.297 us |       1.19% |  87.835 us |       1.10% |   -8.462 us |  -8.79% |   FAST   |
|    I32     |     I32      |      I64      |      2^28      |     1     |   1.169 ms |       0.09% |   1.044 ms |       0.08% | -124.538 us | -10.66% |   FAST   |
|    I32     |     I32      |      I64      |      2^16      |   0.201   |  18.238 us |       1.59% |  18.613 us |       4.67% |    0.375 us |   2.06% |   SLOW   |
|    I32     |     I32      |      I64      |      2^20      |   0.201   |  25.028 us |       3.81% |  24.428 us |       0.65% |   -0.600 us |  -2.40% |   FAST   |
|    I32     |     I32      |      I64      |      2^24      |   0.201   |  84.993 us |       1.22% |  77.122 us |       1.19% |   -7.871 us |  -9.26% |   FAST   |
|    I32     |     I32      |      I64      |      2^28      |   0.201   |   1.005 ms |       0.12% | 874.635 us |       0.12% | -130.773 us | -13.01% |   FAST   |
|    I32     |     I64      |      I32      |      2^16      |     1     |  18.962 us |       4.73% |  19.808 us |       4.80% |    0.847 us |   4.47% |   SAME   |
|    I32     |     I64      |      I32      |      2^24      |     1     | 122.573 us |       0.57% | 112.100 us |       0.90% |  -10.473 us |  -8.54% |   FAST   |
|    I32     |     I64      |      I32      |      2^28      |     1     |   1.605 ms |       0.09% |   1.436 ms |       0.10% | -168.988 us | -10.53% |   FAST   |
|    I32     |     I64      |      I32      |      2^24      |   0.201   | 106.903 us |       0.82% |  97.969 us |       0.74% |   -8.934 us |  -8.36% |   FAST   |
|    I32     |     I64      |      I32      |      2^28      |   0.201   |   1.372 ms |       0.11% |   1.219 ms |       0.12% | -152.604 us | -11.13% |   FAST   |
|    I32     |     I64      |      I64      |      2^16      |     1     |  18.876 us |       4.83% |  19.641 us |       5.06% |    0.764 us |   4.05% |   SAME   |
|    I32     |     I64      |      I64      |      2^24      |     1     | 122.613 us |       0.46% | 112.415 us |       0.52% |  -10.197 us |  -8.32% |   FAST   |
|    I32     |     I64      |      I64      |      2^28      |     1     |   1.604 ms |       0.09% |   1.443 ms |       0.09% | -161.256 us | -10.05% |   FAST   |
|    I32     |     I64      |      I64      |      2^24      |   0.201   | 106.296 us |       0.55% |  98.090 us |       0.62% |   -8.206 us |  -7.72% |   FAST   |
|    I32     |     I64      |      I64      |      2^28      |   0.201   |   1.363 ms |       0.12% |   1.221 ms |       0.13% | -142.110 us | -10.42% |   FAST   |
|    I64     |      I8      |      I32      |      2^20      |     1     |  31.664 us |       3.19% |  29.826 us |       3.37% |   -1.838 us |  -5.81% |   FAST   |
|    I64     |      I8      |      I32      |      2^24      |     1     | 131.597 us |       0.69% | 105.242 us |       0.95% |  -26.355 us | -20.03% |   FAST   |
|    I64     |      I8      |      I32      |      2^28      |     1     |   1.729 ms |       0.07% |   1.321 ms |       0.07% | -407.953 us | -23.60% |   FAST   |
|    I64     |      I8      |      I32      |      2^20      |   0.201   |  28.916 us |       2.62% |  26.649 us |       1.41% |   -2.267 us |  -7.84% |   FAST   |
|    I64     |      I8      |      I32      |      2^24      |   0.201   | 114.515 us |       0.35% |  90.018 us |       0.40% |  -24.497 us | -21.39% |   FAST   |
|    I64     |      I8      |      I32      |      2^28      |   0.201   |   1.459 ms |       0.12% |   1.055 ms |       0.14% | -403.731 us | -27.68% |   FAST   |
|    I64     |      I8      |      I64      |      2^16      |     1     |  19.776 us |       5.13% |  20.516 us |       3.98% |    0.740 us |   3.74% |   SAME   |
|    I64     |      I8      |      I64      |      2^24      |     1     | 120.285 us |       0.77% | 106.497 us |       0.30% |  -13.788 us | -11.46% |   FAST   |
|    I64     |      I8      |      I64      |      2^28      |     1     |   1.549 ms |       0.08% |   1.346 ms |       0.10% | -202.347 us | -13.07% |   FAST   |
|    I64     |      I8      |      I64      |      2^24      |   0.201   | 101.968 us |       0.66% |  90.029 us |       0.48% |  -11.939 us | -11.71% |   FAST   |
|    I64     |      I8      |      I64      |      2^28      |   0.201   |   1.236 ms |       0.10% |   1.069 ms |       0.12% | -167.531 us | -13.55% |   FAST   |
|    I64     |     I16      |      I32      |      2^16      |     1     |  19.996 us |       4.36% |  20.467 us |       0.22% |    0.471 us |   2.36% |   SLOW   |
|    I64     |     I16      |      I32      |      2^20      |     1     |  31.783 us |       3.18% |  29.154 us |       3.00% |   -2.628 us |  -8.27% |   FAST   |
|    I64     |     I16      |      I32      |      2^24      |     1     | 134.125 us |       0.76% | 106.525 us |       0.29% |  -27.600 us | -20.58% |   FAST   |
|    I64     |     I16      |      I32      |      2^28      |     1     |   1.763 ms |       0.08% |   1.340 ms |       0.08% | -422.710 us | -23.97% |   FAST   |
|    I64     |     I16      |      I32      |      2^20      |   0.201   |  28.714 us |       1.73% |  26.877 us |       2.63% |   -1.837 us |  -6.40% |   FAST   |
|    I64     |     I16      |      I32      |      2^24      |   0.201   | 116.979 us |       0.80% |  90.539 us |       1.06% |  -26.440 us | -22.60% |   FAST   |
|    I64     |     I16      |      I32      |      2^28      |   0.201   |   1.518 ms |       0.18% |   1.074 ms |       0.13% | -444.063 us | -29.24% |   FAST   |
|    I64     |     I16      |      I64      |      2^20      |     1     |  32.366 us |       2.14% |  30.452 us |       2.06% |   -1.914 us |  -5.91% |   FAST   |
|    I64     |     I16      |      I64      |      2^24      |     1     | 134.852 us |       0.61% | 108.477 us |       0.27% |  -26.375 us | -19.56% |   FAST   |
|    I64     |     I16      |      I64      |      2^28      |     1     |   1.768 ms |       0.08% |   1.366 ms |       0.10% | -402.583 us | -22.77% |   FAST   |
|    I64     |     I16      |      I64      |      2^20      |   0.201   |  29.209 us |       3.35% |  27.044 us |       3.36% |   -2.165 us |  -7.41% |   FAST   |
|    I64     |     I16      |      I64      |      2^24      |   0.201   | 116.491 us |       0.27% |  91.720 us |       0.81% |  -24.770 us | -21.26% |   FAST   |
|    I64     |     I16      |      I64      |      2^28      |   0.201   |   1.505 ms |       0.12% |   1.086 ms |       0.13% | -418.995 us | -27.84% |   FAST   |
|    I64     |     I32      |      I32      |      2^24      |     1     | 129.058 us |       0.44% | 116.881 us |       0.47% |  -12.177 us |  -9.44% |   FAST   |
|    I64     |     I32      |      I32      |      2^28      |     1     |   1.660 ms |       0.10% |   1.515 ms |       0.09% | -145.268 us |  -8.75% |   FAST   |
|    I64     |     I32      |      I32      |      2^16      |   0.201   |  18.609 us |       3.74% |  19.154 us |       5.98% |    0.545 us |   2.93% |   SAME   |
|    I64     |     I32      |      I32      |      2^24      |   0.201   | 108.134 us |       0.72% |  98.311 us |       0.36% |   -9.823 us |  -9.08% |   FAST   |
|    I64     |     I32      |      I32      |      2^28      |   0.201   |   1.363 ms |       0.15% |   1.209 ms |       0.10% | -153.879 us | -11.29% |   FAST   |
|    I64     |     I32      |      I64      |      2^24      |     1     | 129.525 us |       0.78% | 118.045 us |       1.15% |  -11.480 us |  -8.86% |   FAST   |
|    I64     |     I32      |      I64      |      2^28      |     1     |   1.678 ms |       0.11% |   1.524 ms |       0.07% | -154.136 us |  -9.18% |   FAST   |
|    I64     |     I32      |      I64      |      2^16      |   0.201   |  18.606 us |       4.22% |  19.009 us |       5.58% |    0.403 us |   2.17% |   SAME   |
|    I64     |     I32      |      I64      |      2^24      |   0.201   | 108.199 us |       0.49% |  98.947 us |       1.03% |   -9.252 us |  -8.55% |   FAST   |
|    I64     |     I32      |      I64      |      2^28      |   0.201   |   1.365 ms |       0.14% |   1.220 ms |       0.11% | -144.678 us | -10.60% |   FAST   |
|    I64     |     I64      |      I32      |      2^24      |     1     | 155.699 us |       0.27% | 147.647 us |       0.42% |   -8.052 us |  -5.17% |   FAST   |
|    I64     |     I64      |      I32      |      2^28      |     1     |   2.176 ms |       0.06% |   2.046 ms |       0.07% | -130.582 us |  -6.00% |   FAST   |
|    I64     |     I64      |      I64      |      2^24      |     1     | 157.026 us |       0.98% | 148.084 us |       0.72% |   -8.941 us |  -5.69% |   FAST   |
|    I64     |     I64      |      I64      |      2^28      |     1     |   2.186 ms |       0.06% |   2.049 ms |       0.07% | -136.471 us |  -6.24% |   FAST   |
|    I128    |      I8      |      I32      |      2^20      |     1     |  37.189 us |       2.07% |  35.751 us |       2.87% |   -1.438 us |  -3.87% |   FAST   |
|    I128    |      I8      |      I32      |      2^24      |     1     | 223.422 us |       0.28% | 190.270 us |       0.32% |  -33.151 us | -14.84% |   FAST   |
|    I128    |      I8      |      I32      |      2^28      |     1     |   3.212 ms |       0.04% |   2.676 ms |       0.05% | -536.036 us | -16.69% |   FAST   |
|    I128    |      I8      |      I32      |      2^24      |   0.201   | 159.707 us |       0.14% | 140.878 us |       0.58% |  -18.828 us | -11.79% |   FAST   |
|    I128    |      I8      |      I32      |      2^28      |   0.201   |   2.209 ms |       0.06% |   1.896 ms |       0.08% | -312.155 us | -14.13% |   FAST   |
|    I128    |      I8      |      I64      |      2^20      |     1     |  37.042 us |       1.83% |  36.167 us |       2.63% |   -0.875 us |  -2.36% |   FAST   |
|    I128    |      I8      |      I64      |      2^24      |     1     | 222.963 us |       0.39% | 192.605 us |       0.29% |  -30.358 us | -13.62% |   FAST   |
|    I128    |      I8      |      I64      |      2^28      |     1     |   3.194 ms |       0.04% |   2.718 ms |       0.05% | -476.407 us | -14.92% |   FAST   |
|    I128    |      I8      |      I64      |      2^24      |   0.201   | 157.702 us |       0.39% | 141.144 us |       0.50% |  -16.559 us | -10.50% |   FAST   |
|    I128    |      I8      |      I64      |      2^28      |   0.201   |   2.182 ms |       0.06% |   1.901 ms |       0.09% | -281.300 us | -12.89% |   FAST   |
|    I128    |     I16      |      I32      |      2^24      |     1     | 219.193 us |       0.19% | 192.799 us |       0.38% |  -26.394 us | -12.04% |   FAST   |
|    I128    |     I16      |      I32      |      2^28      |     1     |   3.157 ms |       0.04% |   2.721 ms |       0.05% | -436.103 us | -13.81% |   FAST   |
|    I128    |     I16      |      I32      |      2^24      |   0.201   | 153.741 us |       0.39% | 143.955 us |       0.67% |   -9.786 us |  -6.37% |   FAST   |
|    I128    |     I16      |      I32      |      2^28      |   0.201   |   2.127 ms |       0.07% |   1.953 ms |       0.08% | -174.162 us |  -8.19% |   FAST   |
|    I128    |     I16      |      I64      |      2^20      |     1     |  37.547 us |       2.68% |  36.543 us |       1.93% |   -1.004 us |  -2.67% |   FAST   |
|    I128    |     I16      |      I64      |      2^24      |     1     | 225.461 us |       0.30% | 196.294 us |       0.44% |  -29.167 us | -12.94% |   FAST   |
|    I128    |     I16      |      I64      |      2^28      |     1     |   3.233 ms |       0.05% |   2.772 ms |       0.05% | -460.942 us | -14.26% |   FAST   |
|    I128    |     I16      |      I64      |      2^24      |   0.201   | 164.141 us |       0.55% | 144.812 us |       0.67% |  -19.328 us | -11.78% |   FAST   |
|    I128    |     I16      |      I64      |      2^28      |   0.201   |   2.279 ms |       0.06% |   1.962 ms |       0.08% | -317.626 us | -13.94% |   FAST   |
|    I128    |     I32      |      I32      |      2^24      |     1     | 220.541 us |       0.43% | 193.572 us |       0.55% |  -26.969 us | -12.23% |   FAST   |
|    I128    |     I32      |      I32      |      2^28      |     1     |   3.173 ms |       0.04% |   2.727 ms |       0.05% | -445.994 us | -14.06% |   FAST   |
|    I128    |     I32      |      I32      |      2^24      |   0.201   | 155.699 us |       0.31% | 147.480 us |       0.26% |   -8.218 us |  -5.28% |   FAST   |
|    I128    |     I32      |      I32      |      2^28      |   0.201   |   2.156 ms |       0.07% |   2.011 ms |       0.08% | -145.576 us |  -6.75% |   FAST   |
|    I128    |     I32      |      I64      |      2^20      |     1     |  38.384 us |       2.32% |  36.828 us |       0.63% |   -1.556 us |  -4.05% |   FAST   |
|    I128    |     I32      |      I64      |      2^24      |     1     | 227.016 us |       0.39% | 196.694 us |       0.27% |  -30.322 us | -13.36% |   FAST   |
|    I128    |     I32      |      I64      |      2^28      |     1     |   3.257 ms |       0.04% |   2.782 ms |       0.04% | -474.905 us | -14.58% |   FAST   |
|    I128    |     I32      |      I64      |      2^24      |   0.201   | 167.248 us |       0.66% | 147.755 us |       0.60% |  -19.494 us | -11.66% |   FAST   |
|    I128    |     I32      |      I64      |      2^28      |   0.201   |   2.324 ms |       0.07% |   2.018 ms |       0.07% | -305.884 us | -13.16% |   FAST   |
|    I128    |     I64      |      I32      |      2^16      |     1     |  21.380 us |       5.38% |  21.930 us |       4.42% |    0.550 us |   2.57% |   SAME   |
|    I128    |     I64      |      I32      |      2^24      |     1     | 250.733 us |       0.41% | 223.047 us |       0.26% |  -27.686 us | -11.04% |   FAST   |
|    I128    |     I64      |      I32      |      2^28      |     1     |   3.686 ms |       0.03% |   3.224 ms |       0.04% | -462.241 us | -12.54% |   FAST   |
|    I128    |     I64      |      I32      |      2^24      |   0.201   | 199.965 us |       0.48% | 184.498 us |       0.34% |  -15.467 us |  -7.73% |   FAST   |
|    I128    |     I64      |      I32      |      2^28      |   0.201   |   2.874 ms |       0.06% |   2.625 ms |       0.07% | -248.361 us |  -8.64% |   FAST   |
|    I128    |     I64      |      I64      |      2^16      |     1     |  21.414 us |       4.84% |  22.076 us |       3.81% |    0.662 us |   3.09% |   SAME   |
|    I128    |     I64      |      I64      |      2^24      |     1     | 251.269 us |       0.37% | 225.189 us |       0.16% |  -26.081 us | -10.38% |   FAST   |
|    I128    |     I64      |      I64      |      2^28      |     1     |   3.692 ms |       0.04% |   3.269 ms |       0.04% | -423.148 us | -11.46% |   FAST   |
|    I128    |     I64      |      I64      |      2^24      |   0.201   | 199.078 us |       0.48% | 186.413 us |       0.33% |  -12.665 us |  -6.36% |   FAST   |
|    I128    |     I64      |      I64      |      2^28      |   0.201   |   2.867 ms |       0.06% |   2.658 ms |       0.07% | -209.102 us |  -7.29% |   FAST   |
|    F32     |      I8      |      I32      |      2^20      |     1     |  27.895 us |       3.56% |  26.591 us |       0.97% |   -1.303 us |  -4.67% |   FAST   |
|    F32     |      I8      |      I32      |      2^24      |     1     |  94.743 us |       1.08% |  87.179 us |       1.47% |   -7.563 us |  -7.98% |   FAST   |
|    F32     |      I8      |      I32      |      2^28      |     1     |   1.123 ms |       0.08% |   1.016 ms |       0.08% | -107.065 us |  -9.54% |   FAST   |
|    F32     |      I8      |      I32      |      2^16      |   0.201   |  19.477 us |       5.10% |  18.981 us |       7.46% |   -0.496 us |  -2.55% |   SAME   |
|    F32     |      I8      |      I32      |      2^20      |   0.201   |  26.488 us |       1.44% |  25.930 us |       3.86% |   -0.558 us |  -2.11% |   FAST   |
|    F32     |      I8      |      I32      |      2^24      |   0.201   |  86.484 us |       1.04% |  78.114 us |       0.97% |   -8.370 us |  -9.68% |   FAST   |
|    F32     |      I8      |      I32      |      2^28      |   0.201   |   1.002 ms |       0.08% | 878.173 us |       0.13% | -123.461 us | -12.33% |   FAST   |
|    F32     |      I8      |      I64      |      2^20      |     1     |  27.627 us |       3.73% |  26.586 us |       0.74% |   -1.041 us |  -3.77% |   FAST   |
|    F32     |      I8      |      I64      |      2^24      |     1     |  94.961 us |       1.04% |  90.289 us |       1.33% |   -4.672 us |  -4.92% |   FAST   |
|    F32     |      I8      |      I64      |      2^28      |     1     |   1.121 ms |       0.10% |   1.056 ms |       0.13% |  -64.365 us |  -5.74% |   FAST   |
|    F32     |      I8      |      I64      |      2^16      |   0.201   |  20.024 us |       3.30% |  19.493 us |       6.76% |   -0.531 us |  -2.65% |   SAME   |
|    F32     |      I8      |      I64      |      2^24      |   0.201   |  86.058 us |       0.71% |  81.256 us |       1.12% |   -4.802 us |  -5.58% |   FAST   |
|    F32     |      I8      |      I64      |      2^28      |   0.201   |   1.004 ms |       0.10% | 911.899 us |       0.16% |  -92.161 us |  -9.18% |   FAST   |
|    F32     |     I16      |      I32      |      2^24      |     1     | 100.588 us |       0.68% |  80.031 us |       0.71% |  -20.558 us | -20.44% |   FAST   |
|    F32     |     I16      |      I32      |      2^28      |     1     |   1.231 ms |       0.08% | 904.324 us |       0.15% | -327.022 us | -26.56% |   FAST   |
|    F32     |     I16      |      I32      |      2^16      |   0.201   |  19.879 us |       4.37% |  19.366 us |       7.04% |   -0.513 us |  -2.58% |   SAME   |
|    F32     |     I16      |      I32      |      2^20      |   0.201   |  26.433 us |       1.63% |  24.550 us |       3.11% |   -1.883 us |  -7.12% |   FAST   |
|    F32     |     I16      |      I32      |      2^24      |   0.201   |  91.892 us |       0.65% |  71.714 us |       0.61% |  -20.179 us | -21.96% |   FAST   |
|    F32     |     I16      |      I32      |      2^28      |   0.201   |   1.089 ms |       0.11% | 776.567 us |       0.17% | -312.859 us | -28.72% |   FAST   |
|    F32     |     I16      |      I64      |      2^24      |     1     | 100.724 us |       1.45% |  80.635 us |       1.27% |  -20.089 us | -19.94% |   FAST   |
|    F32     |     I16      |      I64      |      2^28      |     1     |   1.224 ms |       0.10% | 910.008 us |       0.12% | -314.436 us | -25.68% |   FAST   |
|    F32     |     I16      |      I64      |      2^20      |   0.201   |  25.773 us |       3.70% |  24.589 us |       2.41% |   -1.185 us |  -4.60% |   FAST   |
|    F32     |     I16      |      I64      |      2^24      |   0.201   |  90.007 us |       0.47% |  71.817 us |       1.02% |  -18.190 us | -20.21% |   FAST   |
|    F32     |     I16      |      I64      |      2^28      |   0.201   |   1.085 ms |       0.10% | 780.736 us |       0.15% | -304.235 us | -28.04% |   FAST   |
|    F32     |     I32      |      I32      |      2^20      |     1     |  25.994 us |       3.60% |  26.611 us |       0.18% |    0.617 us |   2.37% |   SLOW   |
|    F32     |     I32      |      I32      |      2^24      |     1     |  94.847 us |       1.08% |  92.689 us |       0.97% |   -2.158 us |  -2.27% |   FAST   |
|    F32     |     I32      |      I32      |      2^28      |     1     |   1.167 ms |       0.09% |   1.129 ms |       0.12% |  -37.576 us |  -3.22% |   FAST   |
|    F32     |     I32      |      I32      |      2^28      |   0.201   | 992.757 us |       0.16% | 970.429 us |       0.16% |  -22.329 us |  -2.25% |   FAST   |
|    F32     |     I32      |      I64      |      2^24      |     1     |  95.898 us |       1.33% |  87.205 us |       1.49% |   -8.693 us |  -9.07% |   FAST   |
|    F32     |     I32      |      I64      |      2^28      |     1     |   1.169 ms |       0.08% |   1.044 ms |       0.06% | -124.965 us | -10.69% |   FAST   |
|    F32     |     I32      |      I64      |      2^24      |   0.201   |  85.134 us |       1.16% |  77.359 us |       1.02% |   -7.776 us |  -9.13% |   FAST   |
|    F32     |     I32      |      I64      |      2^28      |   0.201   |   1.005 ms |       0.14% | 874.911 us |       0.15% | -130.504 us | -12.98% |   FAST   |
|    F32     |     I64      |      I32      |      2^16      |     1     |  19.028 us |       4.92% |  19.643 us |       5.10% |    0.615 us |   3.23% |   SAME   |
|    F32     |     I64      |      I32      |      2^20      |     1     |  27.335 us |       3.58% |  26.715 us |       1.67% |   -0.621 us |  -2.27% |   FAST   |
|    F32     |     I64      |      I32      |      2^24      |     1     | 122.585 us |       0.58% | 112.038 us |       0.88% |  -10.547 us |  -8.60% |   FAST   |
|    F32     |     I64      |      I32      |      2^28      |     1     |   1.606 ms |       0.08% |   1.438 ms |       0.07% | -168.483 us | -10.49% |   FAST   |
|    F32     |     I64      |      I32      |      2^24      |   0.201   | 106.737 us |       0.70% |  97.792 us |       0.89% |   -8.945 us |  -8.38% |   FAST   |
|    F32     |     I64      |      I32      |      2^28      |   0.201   |   1.372 ms |       0.10% |   1.220 ms |       0.10% | -152.622 us | -11.12% |   FAST   |
|    F32     |     I64      |      I64      |      2^16      |     1     |  18.877 us |       4.76% |  19.697 us |       5.03% |    0.820 us |   4.34% |   SAME   |
|    F32     |     I64      |      I64      |      2^24      |     1     | 122.786 us |       0.10% | 112.455 us |       0.43% |  -10.331 us |  -8.41% |   FAST   |
|    F32     |     I64      |      I64      |      2^28      |     1     |   1.605 ms |       0.08% |   1.444 ms |       0.08% | -160.859 us | -10.02% |   FAST   |
|    F32     |     I64      |      I64      |      2^24      |   0.201   | 106.364 us |       0.39% |  98.129 us |       0.46% |   -8.235 us |  -7.74% |   FAST   |
|    F32     |     I64      |      I64      |      2^28      |   0.201   |   1.364 ms |       0.10% |   1.222 ms |       0.09% | -141.943 us | -10.41% |   FAST   |
|    F64     |      I8      |      I32      |      2^20      |     1     |  32.031 us |       3.02% |  29.530 us |       3.41% |   -2.501 us |  -7.81% |   FAST   |
|    F64     |      I8      |      I32      |      2^24      |     1     | 131.024 us |       0.16% | 104.483 us |       0.33% |  -26.541 us | -20.26% |   FAST   |
|    F64     |      I8      |      I32      |      2^28      |     1     |   1.717 ms |       0.06% |   1.310 ms |       0.09% | -406.341 us | -23.67% |   FAST   |
|    F64     |      I8      |      I32      |      2^20      |   0.201   |  28.602 us |       2.79% |  26.660 us |       1.45% |   -1.943 us |  -6.79% |   FAST   |
|    F64     |      I8      |      I32      |      2^24      |   0.201   | 113.030 us |       0.81% |  88.841 us |       1.20% |  -24.189 us | -21.40% |   FAST   |
|    F64     |      I8      |      I32      |      2^28      |   0.201   |   1.435 ms |       0.10% |   1.045 ms |       0.12% | -389.697 us | -27.15% |   FAST   |
|    F64     |      I8      |      I64      |      2^24      |     1     | 120.116 us |       0.82% | 106.470 us |       0.19% |  -13.646 us | -11.36% |   FAST   |
|    F64     |      I8      |      I64      |      2^28      |     1     |   1.549 ms |       0.08% |   1.337 ms |       0.07% | -211.855 us | -13.67% |   FAST   |
|    F64     |      I8      |      I64      |      2^24      |   0.201   | 101.026 us |       1.04% |  89.890 us |       0.37% |  -11.136 us | -11.02% |   FAST   |
|    F64     |      I8      |      I64      |      2^28      |   0.201   |   1.230 ms |       0.08% |   1.057 ms |       0.13% | -173.165 us | -14.08% |   FAST   |
|    F64     |     I16      |      I32      |      2^16      |     1     |  20.017 us |       5.01% |  20.463 us |       0.54% |    0.446 us |   2.23% |   SLOW   |
|    F64     |     I16      |      I32      |      2^20      |     1     |  31.693 us |       3.17% |  29.149 us |       2.97% |   -2.544 us |  -8.03% |   FAST   |
|    F64     |     I16      |      I32      |      2^24      |     1     | 134.570 us |       0.68% | 106.493 us |       0.16% |  -28.077 us | -20.86% |   FAST   |
|    F64     |     I16      |      I32      |      2^28      |     1     |   1.764 ms |       0.07% |   1.333 ms |       0.09% | -431.658 us | -24.46% |   FAST   |
|    F64     |     I16      |      I32      |      2^20      |   0.201   |  28.867 us |       2.38% |  26.698 us |       1.86% |   -2.169 us |  -7.51% |   FAST   |
|    F64     |     I16      |      I32      |      2^24      |   0.201   | 115.593 us |       0.94% |  90.044 us |       0.78% |  -25.549 us | -22.10% |   FAST   |
|    F64     |     I16      |      I32      |      2^28      |   0.201   |   1.486 ms |       0.17% |   1.067 ms |       0.12% | -419.490 us | -28.23% |   FAST   |
|    F64     |     I16      |      I64      |      2^20      |     1     |  32.011 us |       2.88% |  30.342 us |       2.40% |   -1.669 us |  -5.21% |   FAST   |
|    F64     |     I16      |      I64      |      2^24      |     1     | 134.165 us |       0.78% | 108.362 us |       0.55% |  -25.802 us | -19.23% |   FAST   |
|    F64     |     I16      |      I64      |      2^28      |     1     |   1.762 ms |       0.08% |   1.357 ms |       0.07% | -404.484 us | -22.96% |   FAST   |
|    F64     |     I16      |      I64      |      2^20      |   0.201   |  29.049 us |       3.01% |  27.079 us |       3.38% |   -1.970 us |  -6.78% |   FAST   |
|    F64     |     I16      |      I64      |      2^24      |   0.201   | 115.108 us |       0.81% |  90.948 us |       1.13% |  -24.161 us | -20.99% |   FAST   |
|    F64     |     I16      |      I64      |      2^28      |   0.201   |   1.475 ms |       0.12% |   1.075 ms |       0.13% | -399.550 us | -27.10% |   FAST   |
|    F64     |     I32      |      I32      |      2^16      |     1     |  19.828 us |       4.92% |  20.511 us |       1.56% |    0.683 us |   3.45% |   SLOW   |
|    F64     |     I32      |      I32      |      2^24      |     1     | 129.188 us |       0.70% | 116.921 us |       0.52% |  -12.268 us |  -9.50% |   FAST   |
|    F64     |     I32      |      I32      |      2^28      |     1     |   1.660 ms |       0.09% |   1.505 ms |       0.05% | -155.143 us |  -9.34% |   FAST   |
|    F64     |     I32      |      I32      |      2^16      |   0.201   |  18.490 us |       3.00% |  18.899 us |       4.75% |    0.408 us |   2.21% |   SAME   |
|    F64     |     I32      |      I32      |      2^24      |   0.201   | 107.153 us |       0.94% |  98.222 us |       0.30% |   -8.932 us |  -8.34% |   FAST   |
|    F64     |     I32      |      I32      |      2^28      |   0.201   |   1.352 ms |       0.17% |   1.198 ms |       0.09% | -154.330 us | -11.41% |   FAST   |
|    F64     |     I32      |      I64      |      2^16      |     1     |  20.039 us |       6.72% |  20.622 us |       3.54% |    0.583 us |   2.91% |   SAME   |
|    F64     |     I32      |      I64      |      2^24      |     1     | 129.722 us |       0.81% | 117.428 us |       0.86% |  -12.294 us |  -9.48% |   FAST   |
|    F64     |     I32      |      I64      |      2^28      |     1     |   1.674 ms |       0.10% |   1.510 ms |       0.06% | -164.075 us |  -9.80% |   FAST   |
|    F64     |     I32      |      I64      |      2^16      |   0.201   |  18.273 us |       1.68% |  18.777 us |       5.10% |    0.505 us |   2.76% |   SLOW   |
|    F64     |     I32      |      I64      |      2^24      |   0.201   | 107.516 us |       0.95% |  98.173 us |       0.26% |   -9.343 us |  -8.69% |   FAST   |
|    F64     |     I32      |      I64      |      2^28      |   0.201   |   1.355 ms |       0.13% |   1.209 ms |       0.10% | -145.981 us | -10.77% |   FAST   |
|    F64     |     I64      |      I32      |      2^24      |     1     | 155.815 us |       0.37% | 147.470 us |       0.61% |   -8.345 us |  -5.36% |   FAST   |
|    F64     |     I64      |      I32      |      2^28      |     1     |   2.173 ms |       0.06% |   2.027 ms |       0.06% | -145.707 us |  -6.71% |   FAST   |
|    F64     |     I64      |      I64      |      2^24      |     1     | 157.045 us |       0.92% | 147.932 us |       0.60% |   -9.113 us |  -5.80% |   FAST   |
|    F64     |     I64      |      I64      |      2^28      |     1     |   2.187 ms |       0.06% |   2.037 ms |       0.07% | -150.588 us |  -6.89% |   FAST   |

I think we can unconditionally enable BlockLoadToShared for Blackwell. Great!

bernhardmgruber · 2025-11-05T13:47:25Z

Added some manual tunings for Ampere.

keys on A100 80GB PCIe, main vs this PR, 2% threshold:

|  KeyT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|------------|---------------|----------------|-----------|------------|-------------|------------|-------------|-------------|---------|----------|
|     I8     |      I32      |      2^24      |     1     |  80.364 us |       0.68% |  71.093 us |       0.80% |   -9.271 us | -11.54% |   FAST   |
|     I8     |      I32      |      2^28      |     1     | 902.653 us |       0.37% | 737.943 us |       0.42% | -164.710 us | -18.25% |   FAST   |
|     I8     |      I32      |      2^20      |   0.201   |  23.254 us |       2.12% |  22.663 us |       1.77% |   -0.592 us |  -2.54% |   FAST   |
|     I8     |      I32      |      2^24      |   0.201   |  78.144 us |       0.64% |  67.893 us |       0.79% |  -10.251 us | -13.12% |   FAST   |
|     I8     |      I32      |      2^28      |   0.201   | 889.346 us |       0.23% | 719.774 us |       0.32% | -169.572 us | -19.07% |   FAST   |
|     I8     |      I64      |      2^24      |     1     |  81.176 us |       0.71% |  71.528 us |       0.83% |   -9.648 us | -11.89% |   FAST   |
|     I8     |      I64      |      2^28      |     1     | 907.867 us |       0.33% | 745.176 us |       0.40% | -162.691 us | -17.92% |   FAST   |
|     I8     |      I64      |      2^20      |   0.201   |  23.677 us |       1.75% |  23.135 us |       2.23% |   -0.542 us |  -2.29% |   FAST   |
|     I8     |      I64      |      2^24      |   0.201   |  78.679 us |       0.64% |  68.512 us |       0.75% |  -10.167 us | -12.92% |   FAST   |
|     I8     |      I64      |      2^28      |   0.201   | 895.626 us |       0.27% | 723.920 us |       0.31% | -171.705 us | -19.17% |   FAST   |
|    I16     |      I32      |      2^16      |     1     |  18.644 us |       2.94% |  18.210 us |       3.20% |   -0.434 us |  -2.33% |   SAME   |
|    I16     |      I32      |      2^24      |     1     |  96.635 us |       0.56% |  90.967 us |       0.64% |   -5.667 us |  -5.86% |   FAST   |
|    I16     |      I32      |      2^28      |     1     |   1.169 ms |       0.32% |   1.029 ms |       0.38% | -140.569 us | -12.02% |   FAST   |
|    I16     |      I32      |      2^24      |   0.201   |  86.548 us |       0.65% |  83.345 us |       0.87% |   -3.203 us |  -3.70% |   FAST   |
|    I16     |      I32      |      2^28      |   0.201   |   1.015 ms |       0.43% | 884.982 us |       0.49% | -129.619 us | -12.78% |   FAST   |
|    I16     |      I64      |      2^28      |     1     |   1.174 ms |       0.53% |   1.071 ms |       0.25% | -102.706 us |  -8.75% |   FAST   |
|    I16     |      I64      |      2^28      |   0.201   |   1.010 ms |       0.45% | 923.604 us |       0.40% |  -86.007 us |  -8.52% |   FAST   |
|    I64     |      I32      |      2^16      |   0.201   |  20.735 us |       2.75% |  21.231 us |       4.00% |    0.497 us |   2.40% |   SAME   |
|    F32     |      I32      |      2^16      |   0.201   |  19.279 us |       5.59% |  19.741 us |       4.17% |    0.461 us |   2.39% |   SAME   |

Pure improvement.

pairs on A100 80GB PCIe, main vs this PR, 2% threshold:

|  KeyT{ct}  |  ValueT{ct}  |  OffsetT{ct}  |  Elements{io}  |  Entropy  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|------------|--------------|---------------|----------------|-----------|------------|-------------|------------|-------------|-------------|---------|----------|
|     I8     |      I8      |      I32      |      2^16      |     1     |  20.282 us |       2.18% |  19.580 us |       2.09% |   -0.702 us |  -3.46% |   FAST   |
|     I8     |      I8      |      I32      |      2^20      |     1     |  32.104 us |       1.70% |  28.537 us |       1.51% |   -3.567 us | -11.11% |   FAST   |
|     I8     |      I8      |      I32      |      2^24      |     1     | 161.035 us |       0.36% | 118.483 us |       0.48% |  -42.552 us | -26.42% |   FAST   |
|     I8     |      I8      |      I32      |      2^28      |     1     |   2.163 ms |       0.23% |   1.459 ms |       0.25% | -703.660 us | -32.53% |   FAST   |
|     I8     |      I8      |      I32      |      2^20      |   0.201   |  31.052 us |       1.83% |  26.865 us |       1.75% |   -4.187 us | -13.48% |   FAST   |
|     I8     |      I8      |      I32      |      2^24      |   0.201   | 159.106 us |       0.35% | 113.472 us |       0.50% |  -45.633 us | -28.68% |   FAST   |
|     I8     |      I8      |      I32      |      2^28      |   0.201   |   2.150 ms |       0.21% |   1.420 ms |       0.16% | -730.505 us | -33.97% |   FAST   |
|     I8     |      I8      |      I64      |      2^24      |     1     | 161.846 us |       0.32% | 151.091 us |       0.38% |  -10.755 us |  -6.65% |   FAST   |
|     I8     |      I8      |      I64      |      2^28      |     1     |   2.179 ms |       0.23% |   1.970 ms |       0.19% | -209.667 us |  -9.62% |   FAST   |
|     I8     |      I8      |      I64      |      2^20      |   0.201   |  31.347 us |       1.68% |  29.540 us |       1.55% |   -1.808 us |  -5.77% |   FAST   |
|     I8     |      I8      |      I64      |      2^24      |   0.201   | 159.498 us |       0.33% | 145.160 us |       0.40% |  -14.339 us |  -8.99% |   FAST   |
|     I8     |      I8      |      I64      |      2^28      |   0.201   |   2.166 ms |       0.21% |   1.960 ms |       0.23% | -205.808 us |  -9.50% |   FAST   |
|     I8     |     I16      |      I32      |      2^16      |     1     |  20.995 us |       2.54% |  20.141 us |       3.02% |   -0.854 us |  -4.07% |   FAST   |
|     I8     |     I16      |      I32      |      2^20      |     1     |  33.768 us |       1.80% |  29.061 us |       1.80% |   -4.707 us | -13.94% |   FAST   |
|     I8     |     I16      |      I32      |      2^24      |     1     | 164.668 us |       0.33% | 121.838 us |       0.47% |  -42.830 us | -26.01% |   FAST   |
|     I8     |     I16      |      I32      |      2^28      |     1     |   2.197 ms |       0.21% |   1.511 ms |       0.47% | -686.198 us | -31.23% |   FAST   |
|     I8     |     I16      |      I32      |      2^20      |   0.201   |  32.071 us |       1.77% |  27.459 us |       1.77% |   -4.612 us | -14.38% |   FAST   |
|     I8     |     I16      |      I32      |      2^24      |   0.201   | 162.614 us |       0.36% | 118.417 us |       0.57% |  -44.197 us | -27.18% |   FAST   |
|     I8     |     I16      |      I32      |      2^28      |   0.201   |   2.183 ms |       0.20% |   1.497 ms |       0.69% | -686.262 us | -31.44% |   FAST   |
|     I8     |     I16      |      I64      |      2^20      |     1     |  33.744 us |       1.46% |  29.142 us |       1.77% |   -4.602 us | -13.64% |   FAST   |
|     I8     |     I16      |      I64      |      2^24      |     1     | 167.752 us |       0.37% | 122.880 us |       0.00% |  -44.872 us | -26.75% |   ????   |
|     I8     |     I16      |      I64      |      2^28      |     1     |   2.242 ms |       0.22% |   1.524 ms |       0.46% | -718.262 us | -32.04% |   FAST   |
|     I8     |     I16      |      I64      |      2^20      |   0.201   |  32.317 us |       1.72% |  27.648 us |       0.00% |   -4.669 us | -14.45% |   FAST   |
|     I8     |     I16      |      I64      |      2^24      |   0.201   | 165.386 us |       0.35% | 118.987 us |       0.50% |  -46.399 us | -28.06% |   FAST   |
|     I8     |     I16      |      I64      |      2^28      |   0.201   |   2.229 ms |       0.19% |   1.522 ms |       0.75% | -706.684 us | -31.70% |   FAST   |
|     I8     |     I32      |      I32      |      2^16      |     1     |  19.891 us |       2.62% |  20.622 us |       2.34% |    0.731 us |   3.68% |   SLOW   |
|     I8     |     I32      |      I32      |      2^20      |     1     |  33.168 us |       1.54% |  29.696 us |       0.00% |   -3.472 us | -10.47% |   FAST   |
|     I8     |     I32      |      I32      |      2^24      |     1     | 183.805 us |       0.53% | 149.423 us |       0.72% |  -34.382 us | -18.71% |   FAST   |
|     I8     |     I32      |      I32      |      2^28      |     1     |   2.634 ms |       0.82% |   1.995 ms |       0.65% | -639.454 us | -24.28% |   FAST   |
|     I8     |     I32      |      I32      |      2^16      |   0.201   |  19.256 us |       2.63% |  20.072 us |       2.60% |    0.816 us |   4.24% |   SLOW   |
|     I8     |     I32      |      I32      |      2^20      |   0.201   |  31.058 us |       1.64% |  28.441 us |       2.55% |   -2.617 us |  -8.42% |   FAST   |
|     I8     |     I32      |      I32      |      2^24      |   0.201   | 181.953 us |       0.86% | 148.979 us |       0.70% |  -32.974 us | -18.12% |   FAST   |
|     I8     |     I32      |      I32      |      2^28      |   0.201   |   2.573 ms |       0.78% |   2.004 ms |       0.63% | -569.467 us | -22.13% |   FAST   |
|     I8     |     I32      |      I64      |      2^16      |     1     |  19.897 us |       2.62% |  20.575 us |       2.21% |    0.678 us |   3.41% |   SLOW   |
|     I8     |     I32      |      I64      |      2^20      |     1     |  33.698 us |       2.10% |  30.491 us |       2.17% |   -3.207 us |  -9.52% |   FAST   |
|     I8     |     I32      |      I64      |      2^24      |     1     | 186.378 us |       0.51% | 149.311 us |       0.77% |  -37.066 us | -19.89% |   FAST   |
|     I8     |     I32      |      I64      |      2^28      |     1     |   2.668 ms |       0.85% |   2.001 ms |       0.66% | -666.238 us | -24.97% |   FAST   |
|     I8     |     I32      |      I64      |      2^16      |   0.201   |  19.537 us |       2.42% |  20.251 us |       2.58% |    0.715 us |   3.66% |   SLOW   |
|     I8     |     I32      |      I64      |      2^20      |   0.201   |  31.567 us |       2.24% |  28.507 us |       2.57% |   -3.060 us |  -9.69% |   FAST   |
|     I8     |     I32      |      I64      |      2^24      |   0.201   | 185.032 us |       1.01% | 149.702 us |       0.70% |  -35.330 us | -19.09% |   FAST   |
|     I8     |     I32      |      I64      |      2^28      |   0.201   |   2.627 ms |       0.67% |   2.006 ms |       0.65% | -620.599 us | -23.62% |   FAST   |
|     I8     |     I64      |      I32      |      2^16      |     1     |  18.769 us |       2.63% |  19.808 us |       2.52% |    1.039 us |   5.53% |   SLOW   |
|     I8     |     I64      |      I32      |      2^24      |     1     | 218.797 us |       0.45% | 210.295 us |       0.41% |   -8.502 us |  -3.89% |   FAST   |
|     I8     |     I64      |      I32      |      2^28      |     1     |   3.190 ms |       0.64% |   3.045 ms |       0.84% | -145.115 us |  -4.55% |   FAST   |
|     I8     |     I64      |      I32      |      2^16      |   0.201   |  18.638 us |       2.55% |  19.789 us |       2.61% |    1.152 us |   6.18% |   SLOW   |
|     I8     |     I64      |      I32      |      2^20      |   0.201   |  35.175 us |       1.74% |  34.375 us |       1.70% |   -0.801 us |  -2.28% |   FAST   |
|     I8     |     I64      |      I32      |      2^24      |   0.201   | 217.714 us |       0.80% | 211.461 us |       1.07% |   -6.253 us |  -2.87% |   FAST   |
|     I8     |     I64      |      I32      |      2^28      |   0.201   |   3.141 ms |       0.68% |   3.022 ms |       0.83% | -118.767 us |  -3.78% |   FAST   |
|     I8     |     I64      |      I64      |      2^16      |     1     |  19.354 us |       2.72% |  20.215 us |       2.56% |    0.861 us |   4.45% |   SLOW   |
|     I8     |     I64      |      I64      |      2^24      |     1     | 223.698 us |       0.76% | 211.403 us |       0.90% |  -12.295 us |  -5.50% |   FAST   |
|     I8     |     I64      |      I64      |      2^28      |     1     |   3.272 ms |       0.66% |   3.047 ms |       0.83% | -224.798 us |  -6.87% |   FAST   |
|     I8     |     I64      |      I64      |      2^16      |   0.201   |  19.059 us |       2.92% |  20.212 us |       2.50% |    1.153 us |   6.05% |   SLOW   |
|     I8     |     I64      |      I64      |      2^20      |   0.201   |  35.666 us |       1.71% |  34.856 us |       1.67% |   -0.810 us |  -2.27% |   FAST   |
|     I8     |     I64      |      I64      |      2^24      |   0.201   | 222.842 us |       0.84% | 211.424 us |       0.95% |  -11.418 us |  -5.12% |   FAST   |
|     I8     |     I64      |      I64      |      2^28      |   0.201   |   3.232 ms |       0.64% |   3.024 ms |       0.82% | -208.009 us |  -6.44% |   FAST   |
|    I16     |      I8      |      I32      |      2^20      |     1     |  36.308 us |       1.57% |  32.901 us |       1.27% |   -3.406 us |  -9.38% |   FAST   |
|    I16     |      I8      |      I32      |      2^24      |     1     | 184.709 us |       0.47% | 149.719 us |       0.64% |  -34.990 us | -18.94% |   FAST   |
|    I16     |      I8      |      I32      |      2^28      |     1     |   2.516 ms |       0.23% |   1.942 ms |       0.57% | -574.287 us | -22.82% |   FAST   |
|    I16     |      I8      |      I32      |      2^16      |   0.201   |  21.072 us |       2.65% |  21.501 us |       2.95% |    0.429 us |   2.03% |   SAME   |
|    I16     |      I8      |      I32      |      2^20      |   0.201   |  32.948 us |       1.68% |  29.696 us |       0.00% |   -3.252 us |  -9.87% |   FAST   |
|    I16     |      I8      |      I32      |      2^24      |   0.201   | 171.032 us |       0.35% | 135.508 us |       0.56% |  -35.524 us | -20.77% |   FAST   |
|    I16     |      I8      |      I32      |      2^28      |   0.201   |   2.301 ms |       0.20% |   1.778 ms |       0.71% | -523.596 us | -22.75% |   FAST   |
|    I16     |      I8      |      I64      |      2^16      |     1     |  22.398 us |       2.17% |  21.797 us |       2.34% |   -0.602 us |  -2.69% |   FAST   |
|    I16     |      I8      |      I64      |      2^20      |     1     |  35.626 us |       1.43% |  31.553 us |       1.44% |   -4.073 us | -11.43% |   FAST   |
|    I16     |      I8      |      I64      |      2^24      |     1     | 188.856 us |       0.35% | 150.053 us |       0.52% |  -38.802 us | -20.55% |   FAST   |
|    I16     |      I8      |      I64      |      2^28      |     1     |   2.552 ms |       0.23% |   1.971 ms |       0.65% | -581.175 us | -22.77% |   FAST   |
|    I16     |      I8      |      I64      |      2^20      |   0.201   |  33.097 us |       1.65% |  29.351 us |       1.74% |   -3.746 us | -11.32% |   FAST   |
|    I16     |      I8      |      I64      |      2^24      |   0.201   | 173.633 us |       0.36% | 135.022 us |       0.48% |  -38.611 us | -22.24% |   FAST   |
|    I16     |      I8      |      I64      |      2^28      |   0.201   |   2.337 ms |       0.20% |   1.785 ms |       0.56% | -552.367 us | -23.63% |   FAST   |
|    I16     |     I16      |      I32      |      2^20      |     1     |  36.550 us |       1.46% |  32.460 us |       1.47% |   -4.090 us | -11.19% |   FAST   |
|    I16     |     I16      |      I32      |      2^24      |     1     | 192.861 us |       0.39% | 157.267 us |       0.72% |  -35.595 us | -18.46% |   FAST   |
|    I16     |     I16      |      I32      |      2^28      |     1     |   2.594 ms |       0.27% |   2.057 ms |       1.11% | -537.354 us | -20.71% |   FAST   |
|    I16     |     I16      |      I32      |      2^20      |   0.201   |  33.761 us |       1.91% |  30.172 us |       1.88% |   -3.589 us | -10.63% |   FAST   |
|    I16     |     I16      |      I32      |      2^24      |   0.201   | 177.679 us |       0.51% | 147.763 us |       0.80% |  -29.916 us | -16.84% |   FAST   |
|    I16     |     I16      |      I32      |      2^28      |   0.201   |   2.378 ms |       0.24% |   2.000 ms |       0.56% | -377.930 us | -15.89% |   FAST   |
|    I16     |     I16      |      I64      |      2^20      |     1     |  36.150 us |       1.48% |  32.830 us |       1.79% |   -3.320 us |  -9.18% |   FAST   |
|    I16     |     I16      |      I64      |      2^24      |     1     | 195.031 us |       0.43% | 158.294 us |       0.54% |  -36.737 us | -18.84% |   FAST   |
|    I16     |     I16      |      I64      |      2^28      |     1     |   2.629 ms |       0.24% |   2.085 ms |       0.69% | -544.147 us | -20.70% |   FAST   |
|    I16     |     I16      |      I64      |      2^20      |   0.201   |  33.994 us |       1.93% |  30.730 us |       1.87% |   -3.264 us |  -9.60% |   FAST   |
|    I16     |     I16      |      I64      |      2^24      |   0.201   | 180.368 us |       0.43% | 146.843 us |       0.86% |  -33.525 us | -18.59% |   FAST   |
|    I16     |     I16      |      I64      |      2^28      |   0.201   |   2.418 ms |       0.23% |   1.991 ms |       0.57% | -427.296 us | -17.67% |   FAST   |
|    I32     |      I8      |      I32      |      2^16      |     1     |  21.357 us |       2.49% |  22.041 us |       2.64% |    0.683 us |   3.20% |   SLOW   |
|    I32     |      I8      |      I32      |      2^20      |     1     |  35.099 us |       1.41% |  33.290 us |       1.71% |   -1.809 us |  -5.15% |   FAST   |
|    I32     |      I8      |      I32      |      2^24      |     1     | 199.674 us |       0.75% | 178.790 us |       1.14% |  -20.883 us | -10.46% |   FAST   |
|    I32     |      I8      |      I32      |      2^28      |     1     |   2.768 ms |       0.52% |   2.513 ms |       0.65% | -255.752 us |  -9.24% |   FAST   |
|    I32     |      I8      |      I32      |      2^16      |   0.201   |  20.507 us |       2.27% |  21.337 us |       2.48% |    0.830 us |   4.05% |   SLOW   |
|    I32     |      I8      |      I32      |      2^20      |   0.201   |  32.585 us |       2.25% |  30.658 us |       2.59% |   -1.926 us |  -5.91% |   FAST   |
|    I32     |      I8      |      I32      |      2^24      |   0.201   | 180.803 us |       0.53% | 155.812 us |       0.73% |  -24.991 us | -13.82% |   FAST   |
|    I32     |      I8      |      I32      |      2^28      |   0.201   |   2.490 ms |       0.75% |   2.175 ms |       0.72% | -314.478 us | -12.63% |   FAST   |
|    I32     |      I8      |      I64      |      2^16      |     1     |  21.703 us |       2.30% |  22.307 us |       2.23% |    0.604 us |   2.78% |   SLOW   |
|    I32     |      I8      |      I64      |      2^20      |     1     |  34.666 us |       2.21% |  32.393 us |       2.30% |   -2.272 us |  -6.56% |   FAST   |
|    I32     |      I8      |      I64      |      2^24      |     1     | 201.213 us |       0.76% | 178.604 us |       0.57% |  -22.608 us | -11.24% |   FAST   |
|    I32     |      I8      |      I64      |      2^28      |     1     |   2.822 ms |       0.77% |   2.570 ms |       0.64% | -252.418 us |  -8.94% |   FAST   |
|    I32     |      I8      |      I64      |      2^16      |   0.201   |  21.039 us |       3.38% |  21.941 us |       2.71% |    0.902 us |   4.29% |   SLOW   |
|    I32     |      I8      |      I64      |      2^20      |   0.201   |  32.997 us |       2.25% |  30.089 us |       2.10% |   -2.908 us |  -8.81% |   FAST   |
|    I32     |      I8      |      I64      |      2^24      |   0.201   | 184.133 us |       0.59% | 158.973 us |       0.62% |  -25.160 us | -13.66% |   FAST   |
|    I32     |      I8      |      I64      |      2^28      |   0.201   |   2.539 ms |       0.36% |   2.252 ms |       0.73% | -286.717 us | -11.29% |   FAST   |
|    I32     |     I16      |      I32      |      2^16      |     1     |  20.828 us |       2.59% |  21.562 us |       2.38% |    0.734 us |   3.52% |   SLOW   |
|    I32     |     I16      |      I32      |      2^20      |     1     |  30.818 us |       2.27% |  31.578 us |       2.52% |    0.760 us |   2.47% |   SLOW   |
|    I32     |     I16      |      I64      |      2^16      |     1     |  22.132 us |       2.71% |  21.434 us |       2.58% |   -0.698 us |  -3.15% |   FAST   |
|    I32     |     I64      |      I32      |      2^16      |     1     |  21.879 us |       2.61% |  21.235 us |       2.89% |   -0.644 us |  -2.94% |   FAST   |
|    I32     |     I64      |      I64      |      2^16      |     1     |  21.081 us |       2.56% |  20.645 us |       2.53% |   -0.436 us |  -2.07% |   SAME   |
|    I64     |     I32      |      I32      |      2^16      |   0.201   |  21.919 us |       4.06% |  22.409 us |       3.91% |    0.490 us |   2.23% |   SAME   |
|    I64     |     I32      |      I64      |      2^20      |   0.201   |  39.912 us |       1.48% |  40.762 us |       1.61% |    0.850 us |   2.13% |   SLOW   |
|    I64     |     I64      |      I32      |      2^16      |   0.201   |  21.502 us |       3.39% |  21.986 us |       3.82% |    0.484 us |   2.25% |   SAME   |
|    I128    |      I8      |      I32      |      2^16      |     1     |  24.205 us |       2.45% |  24.749 us |       2.58% |    0.544 us |   2.25% |   SAME   |
|    I128    |      I8      |      I32      |      2^16      |   0.201   |  23.895 us |       3.82% |  23.369 us |       3.77% |   -0.526 us |  -2.20% |   SAME   |
|    I128    |     I64      |      I64      |      2^16      |     1     |  24.959 us |       3.05% |  24.312 us |       2.86% |   -0.647 us |  -2.59% |   SAME   |
|    I128    |     I64      |      I64      |      2^16      |   0.201   |  23.942 us |       3.89% |  23.115 us |       3.62% |   -0.827 us |  -3.46% |   SAME   |
|    F32     |      I8      |      I32      |      2^20      |     1     |  35.464 us |       1.59% |  32.896 us |       1.50% |   -2.568 us |  -7.24% |   FAST   |
|    F32     |      I8      |      I32      |      2^24      |     1     | 208.268 us |       2.36% | 185.919 us |       1.66% |  -22.349 us | -10.73% |   FAST   |
|    F32     |      I8      |      I32      |      2^28      |     1     |   3.031 ms |       1.31% |   2.624 ms |       0.79% | -406.770 us | -13.42% |   FAST   |
|    F32     |      I8      |      I32      |      2^16      |   0.201   |  21.271 us |       3.19% |  22.294 us |       3.74% |    1.023 us |   4.81% |   SLOW   |
|    F32     |      I8      |      I32      |      2^20      |   0.201   |  32.611 us |       2.08% |  30.150 us |       2.43% |   -2.461 us |  -7.55% |   FAST   |
|    F32     |      I8      |      I32      |      2^24      |   0.201   | 185.041 us |       1.21% | 160.256 us |       1.76% |  -24.785 us | -13.39% |   FAST   |
|    F32     |      I8      |      I32      |      2^28      |   0.201   |   2.610 ms |       0.60% |   2.265 ms |       0.69% | -344.751 us | -13.21% |   FAST   |
|    F32     |      I8      |      I64      |      2^16      |     1     |  21.998 us |       2.66% |  23.316 us |       3.47% |    1.318 us |   5.99% |   SLOW   |
|    F32     |      I8      |      I64      |      2^20      |     1     |  34.833 us |       2.11% |  32.444 us |       2.34% |   -2.388 us |  -6.86% |   FAST   |
|    F32     |      I8      |      I64      |      2^24      |     1     | 204.562 us |       1.17% | 183.990 us |       1.42% |  -20.573 us | -10.06% |   FAST   |
|    F32     |      I8      |      I64      |      2^28      |     1     |   3.018 ms |       0.53% |   2.660 ms |       0.65% | -358.108 us | -11.86% |   FAST   |
|    F32     |      I8      |      I64      |      2^16      |   0.201   |  21.277 us |       2.91% |  21.883 us |       3.17% |    0.606 us |   2.85% |   SAME   |
|    F32     |      I8      |      I64      |      2^20      |   0.201   |  32.689 us |       2.29% |  30.367 us |       2.37% |   -2.322 us |  -7.10% |   FAST   |
|    F32     |      I8      |      I64      |      2^24      |   0.201   | 186.896 us |       1.11% | 165.151 us |       1.63% |  -21.745 us | -11.63% |   FAST   |
|    F32     |      I8      |      I64      |      2^28      |   0.201   |   2.693 ms |       2.02% |   2.354 ms |       0.58% | -338.457 us | -12.57% |   FAST   |
|    F32     |     I16      |      I32      |      2^16      |     1     |  21.222 us |       3.15% |  21.993 us |       2.91% |    0.771 us |   3.63% |   SLOW   |
|    F32     |     I16      |      I64      |      2^16      |     1     |  22.188 us |       3.18% |  22.753 us |       2.64% |    0.565 us |   2.55% |   SAME   |
|    F32     |     I16      |      I64      |      2^20      |     1     |  31.649 us |       2.40% |  32.601 us |       1.35% |    0.953 us |   3.01% |   SLOW   |
|    F32     |     I32      |      I64      |      2^20      |     1     |  36.192 us |       1.42% |  35.436 us |       1.88% |   -0.755 us |  -2.09% |   FAST   |
|    F32     |     I32      |      I64      |      2^16      |   0.201   |  21.243 us |       3.78% |  20.755 us |       3.05% |   -0.489 us |  -2.30% |   SAME   |
|    F32     |     I64      |      I32      |      2^16      |     1     |  21.874 us |       3.19% |  21.318 us |       2.83% |   -0.557 us |  -2.54% |   SAME   |
|    F32     |     I64      |      I64      |      2^16      |   0.201   |  22.528 us |       0.00% |  21.304 us |       3.60% |   -1.224 us |  -5.43% |   FAST   |
|    F64     |      I8      |      I64      |      2^16      |     1     |  23.654 us |       1.64% |  22.940 us |       4.29% |   -0.715 us |  -3.02% |   FAST   |
|    F64     |      I8      |      I64      |      2^20      |     1     |  42.057 us |       1.86% |  41.116 us |       1.46% |   -0.941 us |  -2.24% |   FAST   |
|    F64     |     I32      |      I64      |      2^16      |     1     |  23.477 us |       4.31% |  22.899 us |       2.87% |   -0.578 us |  -2.46% |   SAME   |
|    F64     |     I32      |      I64      |      2^16      |   0.201   |  22.804 us |       5.02% |  22.016 us |       4.50% |   -0.788 us |  -3.46% |   SAME   |

Good enough

This reverts commit a9fd762.

bernhardmgruber · 2025-11-05T14:00:19Z

/ok to test 23ebab4

github-actions · 2025-11-06T09:15:32Z

🥳 CI Workflow Results

🟩 Finished in 16h 15m: Pass: 100%/81 | Total: 4d 12h | Max: 4h 25m | Hits: 67%/72769

See results here.

bernhardmgruber · 2025-11-06T09:18:10Z

Since this PR is approved and the CI passed, I merged all changes of this PR into #6077 now. (GH does not let me merge into #6077 directly)

github-project-automation bot added this to CCCL Nov 3, 2025

github-project-automation bot moved this to Todo in CCCL Nov 3, 2025

cccl-authenticator-app bot moved this from Todo to In Progress in CCCL Nov 3, 2025

copy-pr-bot bot force-pushed the pull-request/6077 branch from cf6dfa4 to 25d09e9 Compare November 3, 2025 20:27

bernhardmgruber force-pushed the device_merge_tma_bg branch from f88e45e to a9fd762 Compare November 3, 2025 21:17

bernhardmgruber force-pushed the device_merge_tma_bg branch 5 times, most recently from 91f84ee to 500a706 Compare November 5, 2025 01:14

bernhardmgruber force-pushed the device_merge_tma_bg branch from 500a706 to 07c70de Compare November 5, 2025 11:02

copy-pr-bot bot force-pushed the pull-request/6077 branch from 25d09e9 to c2c52f0 Compare November 5, 2025 13:57

bernhardmgruber mentioned this pull request Nov 5, 2025

[CUB] Use BlockLoadToShared in DeviceMerge #6077

Merged

3 tasks

bernhardmgruber added 9 commits November 5, 2025 14:59

Simplify temp storage

9dcd217

unwrap when we know when we can

e69fa5c

Simplify temp storage size calculation

3d0fd08

Skip span creation

a783ea0

use buffer_t only for TMA path

784624c

Fixes

a128cf2

Translate indices when iterating

9767db6

Shorten some names

c0cd666

Revert "Translate indices when iterating"

3f63ca6

This reverts commit a9fd762.

bernhardmgruber added 3 commits November 5, 2025 14:59

Move mbarrier to the back of SMEM

3723065

Allow tuning to enable bl2sh

22cacd4

Manual tuning

23ebab4

bernhardmgruber force-pushed the device_merge_tma_bg branch from 4025063 to 23ebab4 Compare November 5, 2025 14:00

bernhardmgruber marked this pull request as ready for review November 5, 2025 14:01

bernhardmgruber requested review from a team as code owners November 5, 2025 14:01

bernhardmgruber requested review from davebayer and shwina and removed request for a team November 5, 2025 14:01

cccl-authenticator-app bot moved this from In Progress to In Review in CCCL Nov 5, 2025

elstehle approved these changes Nov 5, 2025

View reviewed changes

fix MSVC

e82ea6f

This comment has been minimized.

Sign in to view

bernhardmgruber closed this Nov 6, 2025

github-project-automation bot moved this from In Review to Done in CCCL Nov 6, 2025

bernhardmgruber deleted the device_merge_tma_bg branch November 6, 2025 09:18

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Improvements to: Use BlockLoadToShared in DeviceMerge #6077 #6460

Improvements to: Use BlockLoadToShared in DeviceMerge #6077 #6460

Uh oh!

bernhardmgruber commented Nov 3, 2025 •

edited

Loading

Uh oh!

copy-pr-bot bot commented Nov 3, 2025

Uh oh!

bernhardmgruber commented Nov 4, 2025 •

edited

Loading

Uh oh!

copy-pr-bot bot commented Nov 5, 2025

Uh oh!

bernhardmgruber commented Nov 5, 2025 •

edited

Loading

Uh oh!

bernhardmgruber commented Nov 5, 2025 •

edited

Loading

Uh oh!

bernhardmgruber commented Nov 5, 2025 •

edited

Loading

Uh oh!

bernhardmgruber commented Nov 5, 2025

Uh oh!

bernhardmgruber commented Nov 5, 2025 •

edited

Loading

Uh oh!

bernhardmgruber commented Nov 5, 2025

Uh oh!

This comment has been minimized.

github-actions bot commented Nov 6, 2025

Uh oh!

bernhardmgruber commented Nov 6, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Improvements to: Use BlockLoadToShared in DeviceMerge #6077 #6460

Improvements to: Use BlockLoadToShared in DeviceMerge #6077 #6460

Uh oh!

Conversation

bernhardmgruber commented Nov 3, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

copy-pr-bot bot commented Nov 3, 2025

Uh oh!

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

Uh oh!

Uh oh!

copy-pr-bot bot commented Nov 5, 2025

Uh oh!

bernhardmgruber commented Nov 5, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bernhardmgruber commented Nov 5, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bernhardmgruber commented Nov 5, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bernhardmgruber commented Nov 5, 2025

Uh oh!

bernhardmgruber commented Nov 5, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bernhardmgruber commented Nov 5, 2025

Uh oh!

This comment has been minimized.

github-actions bot commented Nov 6, 2025

🥳 CI Workflow Results

🟩 Finished in 16h 15m: Pass: 100%/81 | Total: 4d 12h | Max: 4h 25m | Hits: 67%/72769

Uh oh!

bernhardmgruber commented Nov 6, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

bernhardmgruber commented Nov 3, 2025 •

edited

Loading

bernhardmgruber commented Nov 4, 2025 •

edited

Loading

bernhardmgruber commented Nov 5, 2025 •

edited

Loading

bernhardmgruber commented Nov 5, 2025 •

edited

Loading

bernhardmgruber commented Nov 5, 2025 •

edited

Loading

bernhardmgruber commented Nov 5, 2025 •

edited

Loading