Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Add cub::DeviceTransform #2086

Merged
merged 72 commits into from
Sep 8, 2024
Merged

Add cub::DeviceTransform #2086

merged 72 commits into from
Sep 8, 2024

Conversation

bernhardmgruber
Copy link
Contributor

@bernhardmgruber bernhardmgruber commented Jul 26, 2024 •
edited
Loading

This PR creates cub::DeviceTransform and makes the Thrust CUDA backend use it for thrust::transform.

The new cub::DeviceTransform algorithm comes with a set of kernels and corresponding algorithms:

  • fallback_for: reproduces the old behavior of thrust::transform and just forwards to the agent of cub::DeviceFor. This is needed for tuning and to compare against a baseline. We may want to remove this before merging or in a separate PR, once we can cover all cases with better variants.
  • prefetch: does not bitwise-copy any data, but rather prefetches the addresses before starting to pull in data for transformation. this kernel also guarantees stable parameter address to function objects. this forms the new baseline for all GPUs if we cannot copy the parameters before passing them to the function objects (thus, the ultimate baseline kernel).
  • unrolled_staged (bikeshed name?): has two unrolled stages, once for loading ITEMS_PER_THREAD elements from each input stream into registers, followed by a compute stage, that invokes the function object. Ensures the initial data fetch is large and is separate from writing, so the compiler does not detect any aliasing problems. this forms the new baseline for all GPUs if we can copy the parameters before passing them to the function objects.
  • memcpy_async: uses memcpy_async (requires Ampere) to start a bunch for memory loading transactions to shared memory for before crunching multiple iterations on shared memory. Register friendly algorithm and elements per threads are determined at runtime.
  • ublkcp: same as previous but uses ublkcp (requires Hopper). The setup is a bit more complex due to alignment.

All kernels The ublkcp kernel derives their elements per thread from the target minimum number of bytes-in-flight (BIF). Only unrolled_staged does this at compile-time though. However, the unrolling seems to be beneficial over prefetching, especially on the non-data center grade GPUs. From my current understanding, we need the 4 code paths/kernels.

GH200 tuning top 10 
cub.bench.transform.babelstream[T{ct}=I8]:
             variant     score      mins     means      maxs
10  tpb_128.alg_4 ()  1.014491  0.981366  1.014491  1.034848
15  tpb_160.alg_4 ()  1.004240  0.990373  1.004240  1.029778
30  tpb_256.alg_4 ()  1.003050  0.987500  1.003050  1.018750
20  tpb_192.alg_4 ()  1.000345  0.975309  1.000345  1.023465
0            base ()  1.000000  1.000000  1.000000  1.000000
25  tpb_224.alg_4 ()  0.990126  0.970443  0.990126  1.026720
35  tpb_288.alg_4 ()  0.989998  0.976676  0.989998  1.012121
50  tpb_384.alg_4 ()  0.983717  0.972603  0.983717  0.995025
40  tpb_320.alg_4 ()  0.973963  0.946108  0.973963  0.994301
60  tpb_448.alg_4 ()  0.971891  0.957441  0.971891  0.990082
cub.bench.transform.babelstream[T{ct}=I16]:
             variant     score      mins     means      maxs
30  tpb_256.alg_4 ()  1.004151  0.993590  1.004151  1.013575
10  tpb_128.alg_4 ()  1.001833  0.984479  1.001833  1.013353
0            base ()  1.000000  1.000000  1.000000  1.000000
20  tpb_192.alg_4 ()  0.998468  0.977974  0.998468  1.014058
35  tpb_288.alg_4 ()  0.996015  0.986667  0.996015  1.007009
50  tpb_384.alg_4 ()  0.995470  0.980769  0.995470  1.005638
25  tpb_224.alg_4 ()  0.992210  0.973684  0.992210  1.007345
19  tpb_192.alg_3 ()  0.989146  0.928553  0.989146  1.022831
40  tpb_320.alg_4 ()  0.988466  0.965217  0.988466  1.006561
70  tpb_512.alg_4 ()  0.985452  0.969697  0.985452  1.006579
cub.bench.transform.babelstream[T{ct}=F32]:
             variant     score      mins     means      maxs
43  tpb_352.alg_2 ()  1.032837  0.993925  1.032837  1.136364
38  tpb_320.alg_2 ()  1.028563  0.996953  1.028563  1.112676
13  tpb_160.alg_2 ()  1.027531  0.973958  1.027531  1.111111
8   tpb_128.alg_2 ()  1.026094  0.976501  1.026094  1.104895
33  tpb_288.alg_2 ()  1.021513  0.987179  1.021513  1.090909
18  tpb_192.alg_2 ()  1.021103  0.993535  1.021103  1.075862
28  tpb_256.alg_2 ()  1.018935  0.992701  1.018935  1.094891
48  tpb_384.alg_2 ()  1.017742  0.961384  1.017742  1.119403
23  tpb_224.alg_2 ()  1.017260  0.984210  1.017260  1.082192
53  tpb_416.alg_2 ()  1.011580  0.941230  1.011580  1.119403
cub.bench.transform.babelstream[T{ct}=F64]:
             variant     score      mins     means      maxs
33  tpb_288.alg_2 ()  1.038890  0.994334  1.038890  1.153846
18  tpb_192.alg_2 ()  1.030886  0.993129  1.030886  1.144828
63  tpb_480.alg_2 ()  1.025039  0.971300  1.025039  1.145038
28  tpb_256.alg_2 ()  1.023411  0.994737  1.023411  1.099338
23  tpb_224.alg_2 ()  1.022579  0.993530  1.022579  1.106667
8   tpb_128.alg_2 ()  1.022368  0.994334  1.022368  1.099338
13  tpb_160.alg_2 ()  1.021238  0.992727  1.021238  1.092105
43  tpb_352.alg_2 ()  1.020297  0.979730  1.020297  1.093333
38  tpb_320.alg_2 ()  1.018941  0.983051  1.018941  1.084967
48  tpb_384.alg_2 ()  1.016839  0.979424  1.016839  1.099338
cub.bench.transform.babelstream[T{ct}=I128]:
               variant      score      mins      means         maxs
5    tpb_1024.alg_4 ()  82.744400  0.915663  82.744400  1222.974166
28    tpb_256.alg_2 ()   1.011726  0.995138   1.011726     1.077778
33    tpb_288.alg_2 ()   1.008967  0.988839   1.008967     1.071823
48    tpb_384.alg_2 ()   1.007105  0.986637   1.007105     1.071823
78    tpb_576.alg_2 ()   1.006963  0.983974   1.006963     1.065934
68    tpb_512.alg_2 ()   1.006561  0.986637   1.006561     1.065934
73    tpb_544.alg_2 ()   1.006447  0.985959   1.006447     1.065934
138   tpb_960.alg_2 ()   1.005974  0.978025   1.005974     1.065934
25    tpb_224.alg_4 ()   1.005757  0.998374   1.005757     1.028902
3    tpb_1024.alg_2 ()   1.005301  0.977707   1.005301     1.071823
CUB babelstream fallbackfor vs. ublkcp on H200 
# mul

## [0] NVIDIA H200

|  T{ct}  |  OffsetT{ct}  |  Elements{io}  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|---------|---------------|----------------|------------|-------------|------------|-------------|-------------|---------|----------|
|   I8    |      I32      |      2^16      |   4.279 us |       3.45% |   4.846 us |       3.12% |    0.566 us |  13.23% |   FAIL   |
|   I8    |      I32      |      2^20      |   6.265 us |       3.73% |   5.610 us |       3.48% |   -0.654 us | -10.44% |   FAIL   |
|   I8    |      I32      |      2^24      |  37.137 us |       1.02% |  14.292 us |       2.04% |  -22.845 us | -61.52% |   FAIL   |
|   I8    |      I32      |      2^28      | 521.206 us |       0.06% | 134.831 us |       0.65% | -386.376 us | -74.13% |   FAIL   |
|   I8    |      I64      |      2^16      |   4.728 us |       2.90% |   5.190 us |       3.10% |    0.462 us |   9.77% |   FAIL   |
|   I8    |      I64      |      2^20      |   6.380 us |       2.77% |   5.641 us |       3.07% |   -0.740 us | -11.59% |   FAIL   |
|   I8    |      I64      |      2^24      |  37.352 us |       1.05% |  14.358 us |       2.13% |  -22.994 us | -61.56% |   FAIL   |
|   I8    |      I64      |      2^28      | 520.689 us |       0.06% | 135.447 us |       0.65% | -385.241 us | -73.99% |   FAIL   |
|   I16   |      I32      |      2^16      |   4.834 us |       3.04% |   4.826 us |       3.02% |   -0.007 us |  -0.15% |   PASS   |
|   I16   |      I32      |      2^20      |   6.744 us |       2.88% |   5.858 us |       3.48% |   -0.887 us | -13.15% |   FAIL   |
|   I16   |      I32      |      2^24      |  42.113 us |       0.97% |  21.589 us |       2.57% |  -20.525 us | -48.74% |   FAIL   |
|   I16   |      I32      |      2^28      | 586.554 us |       0.12% | 254.011 us |       0.36% | -332.544 us | -56.69% |   FAIL   |
|   I16   |      I64      |      2^16      |   4.806 us |       2.51% |   4.809 us |       2.89% |    0.003 us |   0.07% |   PASS   |
|   I16   |      I64      |      2^20      |   6.811 us |       2.91% |   5.883 us |       2.37% |   -0.928 us | -13.63% |   FAIL   |
|   I16   |      I64      |      2^24      |  42.028 us |       0.96% |  21.574 us |       2.30% |  -20.454 us | -48.67% |   FAIL   |
|   I16   |      I64      |      2^28      | 583.677 us |       0.07% | 253.719 us |       0.41% | -329.958 us | -56.53% |   FAIL   |
|   F32   |      I32      |      2^16      |   4.930 us |       3.22% |   5.135 us |       3.36% |    0.206 us |   4.17% |   FAIL   |
|   F32   |      I32      |      2^20      |   7.546 us |       2.91% |   6.637 us |       3.24% |   -0.908 us | -12.04% |   FAIL   |
|   F32   |      I32      |      2^24      |  49.249 us |       1.09% |  36.985 us |       1.82% |  -12.264 us | -24.90% |   FAIL   |
|   F32   |      I32      |      2^28      | 700.194 us |       0.05% | 509.047 us |       0.18% | -191.147 us | -27.30% |   FAIL   |
|   F32   |      I64      |      2^16      |   4.945 us |       3.88% |   4.803 us |       3.75% |   -0.142 us |  -2.88% |   PASS   |
|   F32   |      I64      |      2^20      |   7.518 us |       2.89% |   6.737 us |       3.54% |   -0.781 us | -10.38% |   FAIL   |
|   F32   |      I64      |      2^24      |  49.108 us |       1.13% |  37.052 us |       1.86% |  -12.056 us | -24.55% |   FAIL   |
|   F32   |      I64      |      2^28      | 697.567 us |       0.06% | 509.116 us |       0.18% | -188.451 us | -27.02% |   FAIL   |
|   F64   |      I32      |      2^16      |   5.169 us |       3.56% |   5.133 us |       3.88% |   -0.036 us |  -0.69% |   PASS   |
|   F64   |      I32      |      2^20      |   9.470 us |       2.65% |   8.944 us |       3.00% |   -0.526 us |  -5.56% |   FAIL   |
|   F64   |      I32      |      2^24      |  73.374 us |       0.88% |  67.611 us |       1.14% |   -5.763 us |  -7.85% |   FAIL   |
|   F64   |      I32      |      2^28      |   1.084 ms |       0.06% |   1.011 ms |       0.12% |  -72.808 us |  -6.72% |   FAIL   |
|   F64   |      I64      |      2^16      |   5.268 us |       4.40% |   4.922 us |       4.57% |   -0.346 us |  -6.57% |   FAIL   |
|   F64   |      I64      |      2^20      |   9.456 us |       2.74% |   8.987 us |       2.98% |   -0.469 us |  -4.96% |   FAIL   |
|   F64   |      I64      |      2^24      |  73.275 us |       0.88% |  67.622 us |       1.10% |   -5.653 us |  -7.71% |   FAIL   |
|   F64   |      I64      |      2^28      |   1.083 ms |       0.06% |   1.011 ms |       0.12% |  -72.089 us |  -6.65% |   FAIL   |
|  I128   |      I32      |      2^16      |   5.569 us |       4.31% |   5.197 us |       4.52% |   -0.372 us |  -6.68% |   FAIL   |
|  I128   |      I32      |      2^20      |  13.558 us |       2.34% |  13.096 us |       1.82% |   -0.462 us |  -3.41% |   FAIL   |
|  I128   |      I32      |      2^24      | 130.880 us |       0.60% | 133.095 us |       0.64% |    2.215 us |   1.69% |   FAIL   |
|  I128   |      I32      |      2^28      |   2.016 ms |       0.06% |   2.052 ms |       0.08% |   35.219 us |   1.75% |   FAIL   |
|  I128   |      I64      |      2^16      |   5.359 us |       3.94% |   5.327 us |       4.01% |   -0.032 us |  -0.60% |   PASS   |
|  I128   |      I64      |      2^20      |  13.467 us |       2.32% |  13.059 us |       1.94% |   -0.408 us |  -3.03% |   FAIL   |
|  I128   |      I64      |      2^24      | 132.284 us |       0.63% | 134.484 us |       0.67% |    2.201 us |   1.66% |   FAIL   |
|  I128   |      I64      |      2^28      |   2.015 ms |       0.31% |   2.054 ms |       0.35% |   39.266 us |   1.95% |   FAIL   |

# add

## [0] NVIDIA H200

|  T{ct}  |  OffsetT{ct}  |  Elements{io}  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|---------|---------------|----------------|------------|-------------|------------|-------------|-------------|---------|----------|
|   I8    |      I32      |      2^16      |   4.361 us |       3.07% |   4.513 us |       3.10% |    0.152 us |   3.48% |   FAIL   |
|   I8    |      I32      |      2^20      |   6.684 us |       3.21% |   5.937 us |       3.61% |   -0.747 us | -11.17% |   FAIL   |
|   I8    |      I32      |      2^24      |  43.103 us |       0.93% |  19.175 us |       1.84% |  -23.927 us | -55.51% |   FAIL   |
|   I8    |      I32      |      2^28      | 577.411 us |       0.08% | 197.937 us |       0.28% | -379.473 us | -65.72% |   FAIL   |
|   I8    |      I64      |      2^16      |   4.831 us |       2.85% |   4.973 us |       2.81% |    0.142 us |   2.94% |   FAIL   |
|   I8    |      I64      |      2^20      |   6.796 us |       3.10% |   5.961 us |       2.66% |   -0.835 us | -12.29% |   FAIL   |
|   I8    |      I64      |      2^24      |  43.399 us |       0.96% |  19.285 us |       1.94% |  -24.115 us | -55.56% |   FAIL   |
|   I8    |      I64      |      2^28      | 581.409 us |       0.07% | 200.064 us |       0.23% | -381.345 us | -65.59% |   FAIL   |
|   I16   |      I32      |      2^16      |   4.869 us |       2.67% |   4.827 us |       2.91% |   -0.042 us |  -0.86% |   PASS   |
|   I16   |      I32      |      2^20      |   7.396 us |       2.78% |   6.647 us |       3.16% |   -0.749 us | -10.12% |   FAIL   |
|   I16   |      I32      |      2^24      |  48.408 us |       0.89% |  30.264 us |       1.65% |  -18.144 us | -37.48% |   FAIL   |
|   I16   |      I32      |      2^28      | 669.724 us |       0.07% | 371.396 us |       0.21% | -298.328 us | -44.54% |   FAIL   |
|   I16   |      I64      |      2^16      |   4.866 us |       2.55% |   4.852 us |       3.08% |   -0.015 us |  -0.30% |   PASS   |
|   I16   |      I64      |      2^20      |   7.132 us |       2.78% |   6.711 us |       2.78% |   -0.421 us |  -5.91% |   FAIL   |
|   I16   |      I64      |      2^24      |  47.909 us |       0.90% |  30.495 us |       1.73% |  -17.414 us | -36.35% |   FAIL   |
|   I16   |      I64      |      2^28      | 663.332 us |       0.08% | 371.926 us |       0.21% | -291.406 us | -43.93% |   FAIL   |
|   F32   |      I32      |      2^16      |   5.043 us |       2.65% |   4.973 us |       3.42% |   -0.070 us |  -1.39% |   PASS   |
|   F32   |      I32      |      2^20      |   8.702 us |       2.58% |   8.210 us |       3.06% |   -0.492 us |  -5.66% |   FAIL   |
|   F32   |      I32      |      2^24      |  62.519 us |       1.01% |  52.920 us |       1.37% |   -9.599 us | -15.35% |   FAIL   |
|   F32   |      I32      |      2^28      | 878.632 us |       0.07% | 730.812 us |       0.11% | -147.819 us | -16.82% |   FAIL   |
|   F32   |      I64      |      2^16      |   5.108 us |       3.62% |   5.054 us |       4.14% |   -0.054 us |  -1.06% |   PASS   |
|   F32   |      I64      |      2^20      |   8.712 us |       2.85% |   8.250 us |       3.13% |   -0.463 us |  -5.31% |   FAIL   |
|   F32   |      I64      |      2^24      |  62.749 us |       0.80% |  52.993 us |       1.39% |   -9.756 us | -15.55% |   FAIL   |
|   F32   |      I64      |      2^28      | 875.608 us |       0.07% | 731.007 us |       0.11% | -144.602 us | -16.51% |   FAIL   |
|   F64   |      I32      |      2^16      |   5.329 us |       3.31% |   5.055 us |       4.14% |   -0.274 us |  -5.14% |   FAIL   |
|   F64   |      I32      |      2^20      |  11.789 us |       2.73% |  11.502 us |       2.81% |   -0.287 us |  -2.44% |   PASS   |
|   F64   |      I32      |      2^24      |  98.967 us |       0.39% |  97.848 us |       0.54% |   -1.119 us |  -1.13% |   FAIL   |
|   F64   |      I32      |      2^28      |   1.463 ms |       0.05% |   1.466 ms |       0.11% |    2.975 us |   0.20% |   FAIL   |
|   F64   |      I64      |      2^16      |   5.412 us |       4.10% |   5.189 us |       4.52% |   -0.223 us |  -4.13% |   FAIL   |
|   F64   |      I64      |      2^20      |  11.768 us |       2.71% |  11.531 us |       2.59% |   -0.237 us |  -2.01% |   PASS   |
|   F64   |      I64      |      2^24      |  98.903 us |       0.42% |  97.919 us |       0.51% |   -0.984 us |  -1.00% |   FAIL   |
|   F64   |      I64      |      2^28      |   1.463 ms |       0.06% |   1.466 ms |       0.11% |    3.664 us |   0.25% |   FAIL   |
|  I128   |      I32      |      2^16      |   5.918 us |       4.57% |   5.657 us |       4.02% |   -0.261 us |  -4.41% |   FAIL   |
|  I128   |      I32      |      2^20      |  18.098 us |       2.08% |  17.622 us |       2.23% |   -0.476 us |  -2.63% |   FAIL   |
|  I128   |      I32      |      2^24      | 190.608 us |       0.31% | 188.985 us |       0.35% |   -1.622 us |  -0.85% |   FAIL   |
|  I128   |      I32      |      2^28      |   2.952 ms |       0.07% |   2.930 ms |       0.07% |  -22.006 us |  -0.75% |   FAIL   |
|  I128   |      I64      |      2^16      |   5.778 us |       3.77% |   5.577 us |       3.24% |   -0.201 us |  -3.47% |   FAIL   |
|  I128   |      I64      |      2^20      |  18.257 us |       2.24% |  17.821 us |       2.30% |   -0.436 us |  -2.39% |   FAIL   |
|  I128   |      I64      |      2^24      | 192.744 us |       0.26% | 190.833 us |       0.32% |   -1.911 us |  -0.99% |   FAIL   |
|  I128   |      I64      |      2^28      |   2.952 ms |       0.41% |   2.934 ms |       0.38% |  -17.396 us |  -0.59% |   FAIL   |

# triad

## [0] NVIDIA H200

|  T{ct}  |  OffsetT{ct}  |  Elements{io}  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|---------|---------------|----------------|------------|-------------|------------|-------------|-------------|---------|----------|
|   I8    |      I32      |      2^16      |   4.963 us |       3.46% |   5.048 us |       2.79% |    0.085 us |   1.72% |   PASS   |
|   I8    |      I32      |      2^20      |   6.929 us |       3.16% |   5.887 us |       2.96% |   -1.042 us | -15.03% |   FAIL   |
|   I8    |      I32      |      2^24      |  42.744 us |       0.98% |  19.148 us |       2.06% |  -23.597 us | -55.20% |   FAIL   |
|   I8    |      I32      |      2^28      | 572.157 us |       0.33% | 199.072 us |       0.23% | -373.085 us | -65.21% |   FAIL   |
|   I8    |      I64      |      2^16      |   4.908 us |       3.12% |   5.083 us |       3.73% |    0.176 us |   3.58% |   FAIL   |
|   I8    |      I64      |      2^20      |   6.969 us |       3.07% |   5.936 us |       2.77% |   -1.033 us | -14.82% |   FAIL   |
|   I8    |      I64      |      2^24      |  43.314 us |       0.94% |  19.120 us |       2.14% |  -24.193 us | -55.86% |   FAIL   |
|   I8    |      I64      |      2^28      | 588.792 us |       0.06% | 200.097 us |       1.06% | -388.695 us | -66.02% |   FAIL   |
|   I16   |      I32      |      2^16      |   4.999 us |       3.02% |   4.856 us |       2.89% |   -0.143 us |  -2.85% |   PASS   |
|   I16   |      I32      |      2^20      |   7.440 us |       2.45% |   6.682 us |       3.14% |   -0.758 us | -10.18% |   FAIL   |
|   I16   |      I32      |      2^24      |  47.974 us |       0.92% |  30.261 us |       2.16% |  -17.712 us | -36.92% |   FAIL   |
|   I16   |      I32      |      2^28      | 658.497 us |       0.11% | 372.018 us |       0.21% | -286.478 us | -43.50% |   FAIL   |
|   I16   |      I64      |      2^16      |   5.048 us |       2.84% |   4.873 us |       3.25% |   -0.175 us |  -3.46% |   FAIL   |
|   I16   |      I64      |      2^20      |   7.420 us |       3.01% |   6.729 us |       3.07% |   -0.690 us |  -9.31% |   FAIL   |
|   I16   |      I64      |      2^24      |  47.570 us |       0.91% |  30.245 us |       1.85% |  -17.325 us | -36.42% |   FAIL   |
|   I16   |      I64      |      2^28      | 651.567 us |       0.08% | 372.418 us |       0.21% | -279.149 us | -42.84% |   FAIL   |
|   F32   |      I32      |      2^16      |   5.115 us |       3.12% |   4.835 us |       3.04% |   -0.280 us |  -5.47% |   FAIL   |
|   F32   |      I32      |      2^20      |   8.741 us |       2.74% |   8.268 us |       2.69% |   -0.473 us |  -5.41% |   FAIL   |
|   F32   |      I32      |      2^24      |  62.791 us |       0.84% |  53.255 us |       1.34% |   -9.536 us | -15.19% |   FAIL   |
|   F32   |      I32      |      2^28      | 872.096 us |       0.07% | 731.790 us |       0.14% | -140.305 us | -16.09% |   FAIL   |
|   F32   |      I64      |      2^16      |   5.172 us |       4.18% |   4.991 us |       4.65% |   -0.181 us |  -3.51% |   PASS   |
|   F32   |      I64      |      2^20      |   8.724 us |       2.72% |   8.291 us |       2.94% |   -0.433 us |  -4.96% |   FAIL   |
|   F32   |      I64      |      2^24      |  62.719 us |       0.79% |  53.274 us |       1.36% |   -9.445 us | -15.06% |   FAIL   |
|   F32   |      I64      |      2^28      | 869.231 us |       0.08% | 731.674 us |       0.12% | -137.557 us | -15.83% |   FAIL   |
|   F64   |      I32      |      2^16      |   5.409 us |       4.00% |   5.165 us |       4.45% |   -0.244 us |  -4.51% |   FAIL   |
|   F64   |      I32      |      2^20      |  11.867 us |       2.29% |  11.618 us |       2.78% |   -0.249 us |  -2.10% |   PASS   |
|   F64   |      I32      |      2^24      |  98.923 us |       0.39% |  97.578 us |       0.54% |   -1.345 us |  -1.36% |   FAIL   |
|   F64   |      I32      |      2^28      |   1.465 ms |       0.05% |   1.467 ms |       0.10% |    2.133 us |   0.15% |   FAIL   |
|   F64   |      I64      |      2^16      |   5.537 us |       4.74% |   5.300 us |       4.67% |   -0.237 us |  -4.28% |   PASS   |
|   F64   |      I64      |      2^20      |  12.103 us |       3.63% |  11.657 us |       2.56% |   -0.446 us |  -3.69% |   FAIL   |
|   F64   |      I64      |      2^24      |  98.861 us |       0.43% |  97.605 us |       0.54% |   -1.256 us |  -1.27% |   FAIL   |
|   F64   |      I64      |      2^28      |   1.464 ms |       0.05% |   1.467 ms |       0.10% |    2.344 us |   0.16% |   FAIL   |
|  I128   |      I32      |      2^16      |   6.133 us |       4.79% |   5.706 us |       4.94% |   -0.427 us |  -6.97% |   FAIL   |
|  I128   |      I32      |      2^20      |  18.317 us |       2.13% |  17.945 us |       2.34% |   -0.372 us |  -2.03% |   PASS   |
|  I128   |      I32      |      2^24      | 190.521 us |       0.32% | 188.933 us |       0.34% |   -1.587 us |  -0.83% |   FAIL   |
|  I128   |      I32      |      2^28      |   2.953 ms |       0.07% |   2.933 ms |       0.06% |  -19.303 us |  -0.65% |   FAIL   |
|  I128   |      I64      |      2^16      |   5.883 us |       2.92% |   5.576 us |       4.12% |   -0.307 us |  -5.22% |   FAIL   |
|  I128   |      I64      |      2^20      |  18.270 us |       2.12% |  17.970 us |       2.37% |   -0.300 us |  -1.64% |   PASS   |
|  I128   |      I64      |      2^24      | 192.793 us |       0.30% | 191.179 us |       0.32% |   -1.614 us |  -0.84% |   FAIL   |
|  I128   |      I64      |      2^28      |   2.953 ms |       0.36% |   2.933 ms |       0.42% |  -19.444 us |  -0.66% |   FAIL   |

# nstream

## [0] NVIDIA H200

|  T{ct}  |  OffsetT{ct}  |  Elements{io}  |  OverwriteInput  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|---------|---------------|----------------|------------------|------------|-------------|------------|-------------|-------------|---------|----------|
|   I8    |      I32      |      2^16      |        1         |   4.451 us |       3.23% |   4.672 us |       3.06% |    0.222 us |   4.98% |   FAIL   |
|   I8    |      I32      |      2^20      |        1         |   7.042 us |       3.10% |   6.330 us |       2.82% |   -0.712 us | -10.11% |   FAIL   |
|   I8    |      I32      |      2^24      |        1         |  46.278 us |       0.93% |  23.230 us |       2.07% |  -23.048 us | -49.80% |   FAIL   |
|   I8    |      I32      |      2^28      |        1         | 625.269 us |       0.08% | 258.946 us |       0.35% | -366.323 us | -58.59% |   FAIL   |
|   I8    |      I64      |      2^16      |        1         |   4.903 us |       2.61% |   5.176 us |       2.93% |    0.273 us |   5.58% |   FAIL   |
|   I8    |      I64      |      2^20      |        1         |   7.109 us |       2.74% |   6.351 us |       2.49% |   -0.758 us | -10.66% |   FAIL   |
|   I8    |      I64      |      2^24      |        1         |  46.669 us |       0.94% |  23.437 us |       2.15% |  -23.232 us | -49.78% |   FAIL   |
|   I8    |      I64      |      2^28      |        1         | 631.695 us |       0.08% | 265.763 us |       0.35% | -365.933 us | -57.93% |   FAIL   |
|   I16   |      I32      |      2^16      |        1         |   5.022 us |       2.92% |   4.980 us |       3.49% |   -0.041 us |  -0.82% |   PASS   |
|   I16   |      I32      |      2^20      |        1         |   7.983 us |       2.94% |   7.438 us |       2.97% |   -0.545 us |  -6.83% |   FAIL   |
|   I16   |      I32      |      2^24      |        1         |  54.482 us |       0.87% |  38.778 us |       1.72% |  -15.705 us | -28.83% |   FAIL   |
|   I16   |      I32      |      2^28      |        1         | 747.609 us |       0.08% | 502.881 us |       0.09% | -244.728 us | -32.73% |   FAIL   |
|   I16   |      I64      |      2^16      |        1         |   5.054 us |       2.97% |   5.020 us |       3.40% |   -0.033 us |  -0.66% |   PASS   |
|   I16   |      I64      |      2^20      |        1         |   7.930 us |       3.22% |   7.502 us |       3.16% |   -0.429 us |  -5.41% |   FAIL   |
|   I16   |      I64      |      2^24      |        1         |  54.279 us |       0.87% |  39.021 us |       1.70% |  -15.259 us | -28.11% |   FAIL   |
|   I16   |      I64      |      2^28      |        1         | 744.414 us |       0.08% | 510.716 us |       0.14% | -233.698 us | -31.39% |   FAIL   |
|   F32   |      I32      |      2^16      |        1         |   5.284 us |       3.66% |   4.894 us |       3.09% |   -0.390 us |  -7.38% |   FAIL   |
|   F32   |      I32      |      2^20      |        1         |   9.705 us |       2.78% |   9.517 us |       2.58% |   -0.187 us |  -1.93% |   PASS   |
|   F32   |      I32      |      2^24      |        1         |  74.644 us |       0.86% |  68.380 us |       1.03% |   -6.264 us |  -8.39% |   FAIL   |
|   F32   |      I32      |      2^28      |        1         |   1.067 ms |       0.07% | 968.744 us |       0.09% |  -97.846 us |  -9.17% |   FAIL   |
|   F32   |      I64      |      2^16      |        1         |   5.336 us |       4.21% |   4.973 us |       3.76% |   -0.363 us |  -6.80% |   FAIL   |
|   F32   |      I64      |      2^20      |        1         |   9.719 us |       2.88% |   9.570 us |       2.64% |   -0.150 us |  -1.54% |   PASS   |
|   F32   |      I64      |      2^24      |        1         |  74.506 us |       0.88% |  68.469 us |       1.04% |   -6.037 us |  -8.10% |   FAIL   |
|   F32   |      I64      |      2^28      |        1         |   1.064 ms |       0.02% | 969.419 us |       0.09% |  -94.741 us |  -8.90% |   FAIL   |
|   F64   |      I32      |      2^16      |        1         |   5.544 us |       3.95% |   5.283 us |       4.35% |   -0.262 us |  -4.72% |   FAIL   |
|   F64   |      I32      |      2^20      |        1         |  13.956 us |       2.27% |  13.925 us |       2.06% |   -0.031 us |  -0.22% |   PASS   |
|   F64   |      I32      |      2^24      |        1         | 128.653 us |       0.63% | 128.299 us |       0.61% |   -0.354 us |  -0.28% |   PASS   |
|   F64   |      I32      |      2^28      |        1         |   1.935 ms |       0.09% |   1.929 ms |       0.04% |   -5.400 us |  -0.28% |   FAIL   |
|   F64   |      I64      |      2^16      |        1         |   5.681 us |       4.46% |   5.447 us |       4.42% |   -0.235 us |  -4.13% |   PASS   |
|   F64   |      I64      |      2^20      |        1         |  13.998 us |       2.34% |  13.969 us |       2.12% |   -0.029 us |  -0.21% |   PASS   |
|   F64   |      I64      |      2^24      |        1         | 128.722 us |       0.66% | 128.373 us |       0.60% |   -0.349 us |  -0.27% |   PASS   |
|   F64   |      I64      |      2^28      |        1         |   1.935 ms |       0.09% |   1.930 ms |       0.04% |   -4.866 us |  -0.25% |   FAIL   |
|  I128   |      I32      |      2^16      |        1         |   6.320 us |       3.82% |   5.984 us |       4.22% |   -0.336 us |  -5.31% |   FAIL   |
|  I128   |      I32      |      2^20      |        1         |  22.093 us |       2.22% |  21.827 us |       2.19% |   -0.266 us |  -1.20% |   PASS   |
|  I128   |      I32      |      2^24      |        1         | 250.055 us |       0.36% | 248.230 us |       0.35% |   -1.825 us |  -0.73% |   FAIL   |
|  I128   |      I32      |      2^28      |        1         |   3.903 ms |       0.13% |   3.861 ms |       0.04% |  -42.478 us |  -1.09% |   FAIL   |
|  I128   |      I64      |      2^16      |        1         |   6.012 us |       3.34% |   5.840 us |       2.89% |   -0.171 us |  -2.85% |   PASS   |
|  I128   |      I64      |      2^20      |        1         |  21.997 us |       2.13% |  21.835 us |       2.13% |   -0.162 us |  -0.74% |   PASS   |
|  I128   |      I64      |      2^24      |        1         | 249.989 us |       0.35% | 248.256 us |       0.35% |   -1.733 us |  -0.69% |   FAIL   |
|  I128   |      I64      |      2^28      |        1         |   3.903 ms |       0.14% |   3.861 ms |       0.04% |  -42.055 us |  -1.08% |   FAIL   |

Fixes: #2091

@bernhardmgruber bernhardmgruber added thrust For all items related to Thrust. cub For all items related to CUB labels Jul 26, 2024
@bernhardmgruber bernhardmgruber force-pushed the transform branch 2 times, most recently from c48b0e8 to 02c1aff Compare July 26, 2024 16:18
@bernhardmgruber bernhardmgruber linked an issue Jul 26, 2024 that may be closed by this pull request
Create cub::DeviceTransform #2091
Closed
@bernhardmgruber bernhardmgruber force-pushed the transform branch 2 times, most recently from 445a3aa to 84bb0ad Compare July 26, 2024 19:00
gevtushenko
gevtushenko reviewed Jul 28, 2024
View reviewed changes
Copy link
Collaborator

@gevtushenko gevtushenko left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Left a few comments. Regarding the address stability detection, I think it's a good idea to separate it into a different PR. There's no risk in merging only opt-in mechanism for Thrust.

cub/cub/device/dispatch/dispatch_transform.cuh Outdated Show resolved Hide resolved
cub/cub/device/dispatch/dispatch_transform.cuh Outdated Show resolved Hide resolved
cub/cub/device/dispatch/dispatch_transform.cuh Outdated Show resolved Hide resolved
cub/cub/device/dispatch/dispatch_transform.cuh Outdated Show resolved Hide resolved
cub/cub/device/dispatch/dispatch_transform.cuh Outdated Show resolved Hide resolved
cub/cub/device/dispatch/dispatch_transform.cuh Outdated Show resolved Hide resolved
cub/cub/device/device_transform.cuh Outdated Show resolved Hide resolved
cub/benchmarks/bench/transform/babelstream.cu Outdated Show resolved Hide resolved
cub/benchmarks/bench/transform/babelstream.cu Outdated Show resolved Hide resolved
@bernhardmgruber bernhardmgruber force-pushed the transform branch 3 times, most recently from 5ad3a62 to 6efb3d7 Compare July 29, 2024 15:18
@gevtushenko gevtushenko requested a review from elstehle July 29, 2024 15:29
@bernhardmgruber bernhardmgruber force-pushed the transform branch 2 times, most recently from c2ed3ce to e76e241 Compare July 30, 2024 09:39
@bernhardmgruber bernhardmgruber mentioned this pull request Jul 30, 2024
Merged
@bernhardmgruber bernhardmgruber force-pushed the transform branch 2 times, most recently from e39bae3 to a414e36 Compare July 30, 2024 12:24
@bernhardmgruber
Copy link
Contributor Author

I removed all the address stability detection from this PR. I will follow up on this separately.

@bernhardmgruber bernhardmgruber force-pushed the transform branch 3 times, most recently from 5b99eb5 to 5a92ac0 Compare August 2, 2024 22:01
@bernhardmgruber bernhardmgruber marked this pull request as ready for review August 2, 2024 22:15
@bernhardmgruber bernhardmgruber requested review from a team as code owners August 2, 2024 22:15
@bernhardmgruber
Copy link
Contributor Author

Calling the for_each agent from the transform_kernel leads to equivalent SASS. Only the kernel symbol name changed and some offsets:

< 		Function : _ZN3cub17CUB_200700_860_NS6detail8for_each13static_kernelINS2_12policy_hub_t12policy_350_tElNS1_9transform10dispatch_tILb0ElN4cuda3std3__45tupleIJN6thrust20THRUST_200700_860_NS6detail15normal_iteratorINSD_10device_ptrInEEEEEEESI_ZN53_INTERNAL_c645bd8d_15_babelstream1_cu_885ae364_1399973mulInlEEvRN7nvbench5stateENSM_9type_listIJT_T0_EEEEUlRKnE_12policy_hub_tIJSI_EEE28non_contiguous_fallback_op_tIJLm0EEEEEEvSR_T1_
---
> 		Function : _ZN3cub17CUB_200700_860_NS6detail9transform16transform_kernelIN12policy_hub_tIJN6thrust20THRUST_200700_860_NS6detail15normal_iteratorINS6_10device_ptrInEEEEEE10max_policyElZN53_INTERNAL_c645bd8d_15_babelstream1_cu_885ae364_1393953mulInlEEvRN7nvbench5stateENSG_9type_listIJT_T0_EEEEUlRKnE_SB_JPnEEEvSL_iT1_T2_DpNS2_10kernel_argIT3_EE
86c86
<         /*00d0*/                   IADD3 R4, P1, R2.reuse, c[0x0][0x170], RZ ;          /* 0x00005c0002047a10 */
---
>         /*00d0*/                   IADD3 R4, P1, R2.reuse, c[0x0][0x188], RZ ;          /* 0x0000620002047a10 */
88c88
<         /*00e0*/                   IADD3 R2, P2, R2, c[0x0][0x178], RZ ;                /* 0x00005e0002027a10 */
---
>         /*00e0*/                   IADD3 R2, P2, R2, c[0x0][0x180], RZ ;                /* 0x0000600002027a10 */
90c90
<         /*00f0*/                   IADD3.X R5, R0.reuse, c[0x0][0x174], RZ, P1, !PT ;   /* 0x00005d0000057a10 */
---
>         /*00f0*/                   IADD3.X R5, R0.reuse, c[0x0][0x18c], RZ, P1, !PT ;   /* 0x0000630000057a10 */
92c92
<         /*0100*/                   IADD3.X R3, R0, c[0x0][0x17c], RZ, P2, !PT ;         /* 0x00005f0000037a10 */
---
>         /*0100*/                   IADD3.X R3, R0, c[0x0][0x184], RZ, P2, !PT ;         /* 0x0000610000037a10 */
207c207

@bernhardmgruber bernhardmgruber mentioned this pull request Sep 6, 2024
Merged
7 tasks
gevtushenko
gevtushenko approved these changes Sep 6, 2024
View reviewed changes
Copy link
Collaborator

@gevtushenko gevtushenko left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Incredible speedups on H200. Great job!

@github-actions GitHub Actions
Copy link
Contributor

github-actions bot commented Sep 7, 2024

🟩 CI finished in 10h 16m: Pass: 100%/251 | Total: 1d 09h | Avg: 7m 57s | Max: 34m 05s | Hits: 99%/24441

  • 🟩 cub: Pass: 100%/132 | Total: 20h 12m | Avg: 9m 11s | Max: 34m 05s | Hits: 99%/4362 

    🟩 cpu
  🟩 amd64              Pass: 100%/124 | Total: 19h 37m | Avg:  9m 29s | Max: 34m 05s | Hits:  99%/4362  
  🟩 arm64              Pass: 100%/8   | Total: 35m 07s | Avg:  4m 23s | Max:  4m 56s
🟩 ctk
  🟩 11.1               Pass: 100%/15  | Total:  1h 11m | Avg:  4m 46s | Max: 12m 03s | Hits:  99%/727   
  🟩 11.8               Pass: 100%/3   | Total: 16m 44s | Avg:  5m 34s | Max:  6m 02s
  🟩 12.5               Pass: 100%/114 | Total: 18h 44m | Avg:  9m 51s | Max: 34m 05s | Hits:  99%/3635  
🟩 cudacxx
  🟩 ClangCUDA17        Pass: 100%/2   | Total:  8m 59s | Avg:  4m 29s | Max:  4m 37s
  🟩 nvcc11.1           Pass: 100%/15  | Total:  1h 11m | Avg:  4m 46s | Max: 12m 03s | Hits:  99%/727   
  🟩 nvcc11.8           Pass: 100%/3   | Total: 16m 44s | Avg:  5m 34s | Max:  6m 02s
  🟩 nvcc12.5           Pass: 100%/112 | Total: 18h 35m | Avg:  9m 57s | Max: 34m 05s | Hits:  99%/3635  
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total:  8m 59s | Avg:  4m 29s | Max:  4m 37s
  🟩 nvcc               Pass: 100%/130 | Total: 20h 03m | Avg:  9m 15s | Max: 34m 05s | Hits:  99%/4362  
🟩 cxx
  🟩 Clang9             Pass: 100%/6   | Total: 30m 00s | Avg:  5m 00s | Max:  6m 04s
  🟩 Clang10            Pass: 100%/3   | Total: 18m 59s | Avg:  6m 19s | Max:  6m 51s
  🟩 Clang11            Pass: 100%/4   | Total: 22m 34s | Avg:  5m 38s | Max:  5m 56s
  🟩 Clang12            Pass: 100%/4   | Total: 23m 06s | Avg:  5m 46s | Max:  6m 04s
  🟩 Clang13            Pass: 100%/4   | Total: 21m 09s | Avg:  5m 17s | Max:  5m 30s
  🟩 Clang14            Pass: 100%/4   | Total: 21m 41s | Avg:  5m 25s | Max:  5m 29s
  🟩 Clang15            Pass: 100%/4   | Total: 21m 45s | Avg:  5m 26s | Max:  5m 39s
  🟩 Clang16            Pass: 100%/4   | Total: 22m 31s | Avg:  5m 37s | Max:  5m 46s
  🟩 Clang17            Pass: 100%/26  | Total:  5h 58m | Avg: 13m 46s | Max: 26m 30s
  🟩 GCC6               Pass: 100%/2   | Total:  8m 29s | Avg:  4m 14s | Max:  4m 30s
  🟩 GCC7               Pass: 100%/6   | Total: 28m 16s | Avg:  4m 42s | Max:  5m 35s
  🟩 GCC8               Pass: 100%/6   | Total: 27m 59s | Avg:  4m 39s | Max:  5m 10s
  🟩 GCC9               Pass: 100%/6   | Total: 29m 53s | Avg:  4m 58s | Max:  5m 34s
  🟩 GCC10              Pass: 100%/4   | Total: 21m 09s | Avg:  5m 17s | Max:  5m 30s
  🟩 GCC11              Pass: 100%/7   | Total: 38m 11s | Avg:  5m 27s | Max:  6m 02s
  🟩 GCC12              Pass: 100%/4   | Total: 23m 40s | Avg:  5m 55s | Max:  6m 22s
  🟩 GCC13              Pass: 100%/29  | Total:  6h 42m | Avg: 13m 53s | Max: 34m 05s
  🟩 Intel2023.2.0      Pass: 100%/3   | Total: 18m 12s | Avg:  6m 04s | Max:  6m 09s
  🟩 MSVC14.16          Pass: 100%/1   | Total: 12m 03s | Avg: 12m 03s | Max: 12m 03s | Hits:  99%/727   
  🟩 MSVC14.29          Pass: 100%/2   | Total: 22m 07s | Avg: 11m 03s | Max: 11m 05s | Hits:  99%/1454  
  🟩 MSVC14.39          Pass: 100%/3   | Total: 39m 44s | Avg: 13m 14s | Max: 13m 49s | Hits:  99%/2181  
🟩 cxx_family
  🟩 Clang              Pass: 100%/59  | Total:  8h 59m | Avg:  9m 09s | Max: 26m 30s
  🟩 GCC                Pass: 100%/64  | Total:  9h 40m | Avg:  9m 04s | Max: 34m 05s
  🟩 Intel              Pass: 100%/3   | Total: 18m 12s | Avg:  6m 04s | Max:  6m 09s
  🟩 MSVC               Pass: 100%/6   | Total:  1h 13m | Avg: 12m 19s | Max: 13m 49s | Hits:  99%/4362  
🟩 gpu
  🟩 v100               Pass: 100%/132 | Total: 20h 12m | Avg:  9m 11s | Max: 34m 05s | Hits:  99%/4362  
🟩 jobs
  🟩 Build              Pass: 100%/99  | Total:  9h 18m | Avg:  5m 38s | Max: 13m 49s | Hits:  99%/4362  
  🟩 DeviceLaunch       Pass: 100%/8   | Total:  2h 32m | Avg: 19m 01s | Max: 20m 16s
  🟩 GraphCapture       Pass: 100%/8   | Total:  2h 04m | Avg: 15m 31s | Max: 17m 09s
  🟩 HostLaunch         Pass: 100%/8   | Total:  2h 20m | Avg: 17m 37s | Max: 18m 57s
  🟩 SmallGMem          Pass: 100%/1   | Total: 31m 01s | Avg: 31m 01s | Max: 31m 01s
  🟩 TestGPU            Pass: 100%/8   | Total:  3h 25m | Avg: 25m 44s | Max: 34m 05s
🟩 sm
  🟩 60;70;80;90        Pass: 100%/3   | Total: 16m 44s | Avg:  5m 34s | Max:  6m 02s
  🟩 90a                Pass: 100%/4   | Total: 17m 07s | Avg:  4m 16s | Max:  4m 26s
🟩 std
  🟩 11                 Pass: 100%/34  | Total:  4h 44m | Avg:  8m 22s | Max: 26m 30s
  🟩 14                 Pass: 100%/37  | Total:  5h 21m | Avg:  8m 40s | Max: 24m 54s | Hits:  99%/2181  
  🟩 17                 Pass: 100%/37  | Total:  5h 51m | Avg:  9m 29s | Max: 31m 01s | Hits:  99%/1454  
  🟩 20                 Pass: 100%/24  | Total:  4h 15m | Avg: 10m 39s | Max: 34m 05s | Hits:  99%/727

  • 🟩 thrust: Pass: 100%/118 | Total: 12h 47m | Avg: 6m 30s | Max: 23m 58s | Hits: 99%/20079 

    🟩 cpu
  🟩 amd64              Pass: 100%/110 | Total: 12h 12m | Avg:  6m 39s | Max: 23m 58s | Hits:  99%/20079 
  🟩 arm64              Pass: 100%/8   | Total: 34m 38s | Avg:  4m 19s | Max:  4m 36s
🟩 ctk
  🟩 11.1               Pass: 100%/15  | Total:  1h 11m | Avg:  4m 44s | Max: 16m 55s | Hits:  99%/2231  
  🟩 11.8               Pass: 100%/3   | Total: 14m 13s | Avg:  4m 44s | Max:  4m 59s
  🟩 12.5               Pass: 100%/100 | Total: 11h 22m | Avg:  6m 49s | Max: 23m 58s | Hits:  99%/17848 
🟩 cudacxx
  🟩 ClangCUDA17        Pass: 100%/2   | Total:  9m 14s | Avg:  4m 37s | Max:  4m 37s
  🟩 nvcc11.1           Pass: 100%/15  | Total:  1h 11m | Avg:  4m 44s | Max: 16m 55s | Hits:  99%/2231  
  🟩 nvcc11.8           Pass: 100%/3   | Total: 14m 13s | Avg:  4m 44s | Max:  4m 59s
  🟩 nvcc12.5           Pass: 100%/98  | Total: 11h 12m | Avg:  6m 52s | Max: 23m 58s | Hits:  99%/17848 
🟩 cudacxx_family
  🟩 ClangCUDA          Pass: 100%/2   | Total:  9m 14s | Avg:  4m 37s | Max:  4m 37s
  🟩 nvcc               Pass: 100%/116 | Total: 12h 38m | Avg:  6m 32s | Max: 23m 58s | Hits:  99%/20079 
🟩 cxx
  🟩 Clang9             Pass: 100%/6   | Total: 29m 27s | Avg:  4m 54s | Max:  5m 50s
  🟩 Clang10            Pass: 100%/3   | Total: 17m 48s | Avg:  5m 56s | Max:  6m 25s
  🟩 Clang11            Pass: 100%/4   | Total: 18m 38s | Avg:  4m 39s | Max:  4m 57s
  🟩 Clang12            Pass: 100%/4   | Total: 18m 44s | Avg:  4m 41s | Max:  5m 11s
  🟩 Clang13            Pass: 100%/4   | Total: 18m 12s | Avg:  4m 33s | Max:  4m 58s
  🟩 Clang14            Pass: 100%/4   | Total: 18m 59s | Avg:  4m 44s | Max:  4m 53s
  🟩 Clang15            Pass: 100%/4   | Total: 20m 18s | Avg:  5m 04s | Max:  5m 30s
  🟩 Clang16            Pass: 100%/4   | Total: 20m 01s | Avg:  5m 00s | Max:  5m 18s
  🟩 Clang17            Pass: 100%/18  | Total:  2h 06m | Avg:  7m 02s | Max: 13m 59s
  🟩 GCC6               Pass: 100%/2   | Total:  7m 33s | Avg:  3m 46s | Max:  4m 03s
  🟩 GCC7               Pass: 100%/6   | Total: 24m 20s | Avg:  4m 03s | Max:  4m 58s
  🟩 GCC8               Pass: 100%/6   | Total: 24m 26s | Avg:  4m 04s | Max:  4m 41s
  🟩 GCC9               Pass: 100%/6   | Total: 24m 57s | Avg:  4m 09s | Max:  4m 57s
  🟩 GCC10              Pass: 100%/4   | Total: 19m 49s | Avg:  4m 57s | Max:  5m 55s
  🟩 GCC11              Pass: 100%/7   | Total: 33m 56s | Avg:  4m 50s | Max:  5m 13s
  🟩 GCC12              Pass: 100%/4   | Total: 20m 05s | Avg:  5m 01s | Max:  5m 24s
  🟩 GCC13              Pass: 100%/20  | Total:  2h 20m | Avg:  7m 00s | Max: 14m 46s
  🟩 Intel2023.2.0      Pass: 100%/3   | Total: 17m 47s | Avg:  5m 55s | Max:  6m 30s
  🟩 MSVC14.16          Pass: 100%/1   | Total: 16m 55s | Avg: 16m 55s | Max: 16m 55s | Hits:  99%/2231  
  🟩 MSVC14.29          Pass: 100%/2   | Total: 31m 03s | Avg: 15m 31s | Max: 16m 29s | Hits:  99%/4462  
  🟩 MSVC14.39          Pass: 100%/6   | Total:  1h 57m | Avg: 19m 36s | Max: 23m 58s | Hits:  99%/13386 
🟩 cxx_family
  🟩 Clang              Pass: 100%/51  | Total:  4h 48m | Avg:  5m 39s | Max: 13m 59s
  🟩 GCC                Pass: 100%/55  | Total:  4h 55m | Avg:  5m 22s | Max: 14m 46s
  🟩 Intel              Pass: 100%/3   | Total: 17m 47s | Avg:  5m 55s | Max:  6m 30s
  🟩 MSVC               Pass: 100%/9   | Total:  2h 45m | Avg: 18m 24s | Max: 23m 58s | Hits:  99%/20079 
🟩 gpu
  🟩 v100               Pass: 100%/118 | Total: 12h 47m | Avg:  6m 30s | Max: 23m 58s | Hits:  99%/20079 
🟩 jobs
  🟩 Build              Pass: 100%/99  | Total:  8h 55m | Avg:  5m 24s | Max: 17m 54s | Hits:  99%/13386 
  🟩 TestCPU            Pass: 100%/11  | Total:  2h 04m | Avg: 11m 18s | Max: 23m 58s | Hits:  99%/6693  
  🟩 TestGPU            Pass: 100%/8   | Total:  1h 47m | Avg: 13m 27s | Max: 14m 46s
🟩 sm
  🟩 60;70;80;90        Pass: 100%/3   | Total: 14m 13s | Avg:  4m 44s | Max:  4m 59s
  🟩 90a                Pass: 100%/4   | Total: 16m 10s | Avg:  4m 02s | Max:  4m 26s
🟩 std
  🟩 11                 Pass: 100%/30  | Total:  2h 32m | Avg:  5m 05s | Max: 13m 53s
  🟩 14                 Pass: 100%/34  | Total:  3h 54m | Avg:  6m 53s | Max: 20m 21s | Hits:  99%/8924  
  🟩 17                 Pass: 100%/33  | Total:  3h 46m | Avg:  6m 52s | Max: 20m 57s | Hits:  99%/6693  
  🟩 20                 Pass: 100%/21  | Total:  2h 33m | Avg:  7m 19s | Max: 23m 58s | Hits:  99%/4462

  • 🟩 pycuda: Pass: 100%/1 | Total: 15m 34s | Avg: 15m 34s | Max: 15m 34s

     
    🟩 cpu
  🟩 amd64              Pass: 100%/1   | Total: 15m 34s | Avg: 15m 34s | Max: 15m 34s
🟩 ctk
  🟩 12.5               Pass: 100%/1   | Total: 15m 34s | Avg: 15m 34s | Max: 15m 34s
🟩 cudacxx
  🟩 nvcc12.5           Pass: 100%/1   | Total: 15m 34s | Avg: 15m 34s | Max: 15m 34s
🟩 cudacxx_family
  🟩 nvcc               Pass: 100%/1   | Total: 15m 34s | Avg: 15m 34s | Max: 15m 34s
🟩 cxx
  🟩 GCC13              Pass: 100%/1   | Total: 15m 34s | Avg: 15m 34s | Max: 15m 34s
🟩 cxx_family
  🟩 GCC                Pass: 100%/1   | Total: 15m 34s | Avg: 15m 34s | Max: 15m 34s
🟩 gpu
  🟩 v100               Pass: 100%/1   | Total: 15m 34s | Avg: 15m 34s | Max: 15m 34s
🟩 jobs
  🟩 Test               Pass: 100%/1   | Total: 15m 34s | Avg: 15m 34s | Max: 15m 34s

👃 Inspect Changes

Modifications in project?

Project
CCCL Infrastructure
libcu++
+/- CUB
Thrust
CUDA Experimental
pycuda
CUDA C Core Library

Modifications in project or dependencies?

Project
CCCL Infrastructure
libcu++
+/- CUB
+/- Thrust
CUDA Experimental
+/- pycuda
+/- CUDA C Core Library

🏃‍ Runner counts (total jobs: 251)
# Runner
178 linux-amd64-cpu16
42 linux-amd64-gpu-v100-latest-1
16 linux-arm64-cpu16
15 windows-amd64-cpu16

@bernhardmgruber bernhardmgruber merged commit 71b9f98 into NVIDIA:main Sep 8, 2024
266 checks passed
@bernhardmgruber bernhardmgruber deleted the transform branch September 8, 2024 21:40
@dkolsen-pgi dkolsen-pgi mentioned this pull request Sep 10, 2024
Closed
1 task
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@fbusato fbusato fbusato left review comments

@ahendriksen ahendriksen ahendriksen left review comments

@elstehle elstehle elstehle left review comments

@miscco miscco miscco left review comments

@gevtushenko gevtushenko gevtushenko approved these changes

@wmaxey wmaxey Awaiting requested review from wmaxey wmaxey is a code owner automatically assigned from NVIDIA/cccl-cmake-codeowners

@griwes griwes Awaiting requested review from griwes griwes is a code owner automatically assigned from NVIDIA/cccl-codeowners

Assignees
No one assigned
Labels
cub For all items related to CUB thrust For all items related to Thrust.
Projects
CCCL
Archived in project
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.

Create cub::DeviceTransform
6 participants
@bernhardmgruber @fbusato @miscco @elstehle @ahendriksen @gevtushenko