Add pytorch inspired DeviceTransform benchmark

[bernhardmgruber]

This benchmark is donated by @MatthiasKohl.

Q: How much should be simplify it by adopting primitives from CCCL?

Running locally:

## chained_eltwise_many_in_many_inst

### [0] NVIDIA GeForce RTX 5090

| T{ct} | OffsetT{ct} |   Elements{io}   | Samples |  CPU Time  |  Noise  |  GPU Time  | Noise | Elem/s  | GlobalMem BW | BWUtil  |
|-------|-------------|------------------|---------|------------|---------|------------|-------|---------|--------------|---------|
|   F32 |         I32 |     2^16 = 65536 |    304x | 107.405 us | 786.62% |  22.544 us | 1.45% |  2.907G | 348.838 GB/s |  19.47% |
|   F32 |         I32 |   2^20 = 1048576 |    688x |  95.059 us |   2.57% |  58.624 us | 1.38% | 17.886G |   2.146 TB/s | 119.77% |
|   F32 |         I32 |  2^24 = 16777216 |    430x |   1.134 ms |   0.33% |   1.096 ms | 0.27% | 15.309G |   1.837 TB/s | 102.51% |
|   F32 |         I32 | 2^28 = 268435456 |    440x |  20.510 ms |   0.18% |  20.465 ms | 0.13% | 13.117G |   1.574 TB/s |  87.83% |
|  bf16 |         I32 |     2^16 = 65536 |    410x |  63.858 us |   5.35% |  26.951 us | 2.70% |  2.432G | 145.899 GB/s |   8.14% |
|  bf16 |         I32 |   2^20 = 1048576 |    568x |  85.948 us |   7.41% |  47.075 us | 1.32% | 22.275G |   1.336 TB/s |  74.57% |
|  bf16 |         I32 |  2^24 = 16777216 |    384x | 430.876 us |   0.92% | 392.816 us | 0.28% | 42.710G |   2.563 TB/s | 142.99% |
|  bf16 |         I32 | 2^28 = 268435456 |    350x |  10.303 ms |   0.08% |  10.262 ms | 0.05% | 26.158G |   1.569 TB/s |  87.58% |

## chained_eltwise_many_in_few_inst

### [0] NVIDIA GeForce RTX 5090

| T{ct} | OffsetT{ct} |   Elements{io}   | Samples |  CPU Time  | Noise |  GPU Time  | Noise | Elem/s  | GlobalMem BW | BWUtil  |
|-------|-------------|------------------|---------|------------|-------|------------|-------|---------|--------------|---------|
|   F32 |         I32 |     2^16 = 65536 |    450x |  56.519 us | 4.47% |  18.502 us | 2.95% |  3.542G | 453.399 GB/s |  25.30% |
|   F32 |         I32 |   2^20 = 1048576 |    524x |  94.659 us | 2.23% |  57.380 us | 1.16% | 18.274G |   2.339 TB/s | 130.52% |
|   F32 |         I32 |  2^24 = 16777216 |    334x |   1.241 ms | 0.38% |   1.201 ms | 0.31% | 13.965G |   1.788 TB/s |  99.74% |
|   F32 |         I32 | 2^28 = 268435456 |    472x |  21.788 ms | 0.14% |  21.742 ms | 0.13% | 12.347G |   1.580 TB/s |  88.18% |
|  bf16 |         I32 |     2^16 = 65536 |    348x |  56.106 us | 5.11% |  19.428 us | 3.70% |  3.373G | 215.888 GB/s |  12.05% |
|  bf16 |         I32 |   2^20 = 1048576 |    312x |  77.319 us | 3.15% |  40.861 us | 2.39% | 25.662G |   1.642 TB/s |  91.64% |
|  bf16 |         I32 |  2^24 = 16777216 |    302x | 484.484 us | 0.57% | 446.249 us | 0.26% | 37.596G |   2.406 TB/s | 134.26% |
|  bf16 |         I32 | 2^28 = 268435456 |    444x |  10.922 ms | 0.09% |  10.879 ms | 0.05% | 24.674G |   1.579 TB/s |  88.11% |

## chained_eltwise_few_in_many_inst

### [0] NVIDIA GeForce RTX 5090

| T{ct} | OffsetT{ct} |   Elements{io}   | Samples |  CPU Time  | Noise |  GPU Time  | Noise | Elem/s  | GlobalMem BW | BWUtil  |
|-------|-------------|------------------|---------|------------|-------|------------|-------|---------|--------------|---------|
|   F32 |         I32 |     2^16 = 65536 |    250x |  52.764 us | 4.22% |  16.325 us | 1.74% |  4.015G | 321.164 GB/s |  17.92% |
|   F32 |         I32 |   2^20 = 1048576 |    342x |  70.282 us | 4.77% |  33.342 us | 2.03% | 31.449G |   2.516 TB/s | 140.39% |
|   F32 |         I32 |  2^24 = 16777216 |    436x | 539.024 us | 0.98% | 499.806 us | 0.70% | 33.567G |   2.685 TB/s | 149.84% |
|   F32 |         I32 | 2^28 = 268435456 |    474x |  14.087 ms | 0.07% |  14.046 ms | 0.06% | 19.111G |   1.529 TB/s |  85.31% |
|  bf16 |         I32 |     2^16 = 65536 |    410x |  55.907 us | 6.44% |  19.348 us | 4.27% |  3.387G | 135.487 GB/s |   7.56% |
|  bf16 |         I32 |   2^20 = 1048576 |    312x |  68.972 us | 4.63% |  31.913 us | 1.68% | 32.857G |   1.314 TB/s |  73.34% |
|  bf16 |         I32 |  2^24 = 16777216 |    344x | 266.021 us | 1.66% | 228.682 us | 0.68% | 73.365G |   2.935 TB/s | 163.75% |
|  bf16 |         I32 | 2^28 = 268435456 |    332x |   7.100 ms | 0.31% |   7.059 ms | 0.31% | 38.025G |   1.521 TB/s |  84.87% |

## chained_eltwise_few_in_few_inst

### [0] NVIDIA GeForce RTX 5090

| T{ct} | OffsetT{ct} |   Elements{io}   | Samples |  CPU Time  | Noise |  GPU Time  | Noise |  Elem/s  | GlobalMem BW | BWUtil  |
|-------|-------------|------------------|---------|------------|-------|------------|-------|----------|--------------|---------|
|   F32 |         I32 |     2^16 = 65536 |     12x |  50.535 us | 2.50% |  13.312 us | 0.00% |   4.923G | 393.846 GB/s |  21.98% |
|   F32 |         I32 |   2^20 = 1048576 |    376x |  62.790 us | 3.35% |  26.610 us | 1.80% |  39.406G |   3.152 TB/s | 175.91% |
|   F32 |         I32 |  2^24 = 16777216 |    436x | 539.553 us | 0.87% | 501.387 us | 0.66% |  33.462G |   2.677 TB/s | 149.37% |
|   F32 |         I32 | 2^28 = 268435456 |    356x |  14.087 ms | 0.08% |  14.048 ms | 0.07% |  19.108G |   1.529 TB/s |  85.30% |
|  bf16 |         I32 |     2^16 = 65536 |    510x |  49.647 us | 7.52% |  13.414 us | 2.33% |   4.886G | 195.428 GB/s |  10.90% |
|  bf16 |         I32 |   2^20 = 1048576 |    302x |  58.203 us | 3.05% |  22.280 us | 2.26% |  47.065G |   1.883 TB/s | 105.05% |
|  bf16 |         I32 |  2^24 = 16777216 |    360x | 200.228 us | 1.19% | 163.100 us | 0.56% | 102.864G |   4.115 TB/s | 229.59% |
|  bf16 |         I32 | 2^28 = 268435456 |    336x |   7.065 ms | 0.10% |   7.027 ms | 0.10% |  38.201G |   1.528 TB/s |  85.26% |

Fixes: https://github.com/NVIDIA-dev/cccl_private/issues/639

[coderabbitai]

Note

Reviews paused

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

Use the following commands to manage reviews:

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

Use the checkboxes below for quick actions:

• ▶️ Resume reviews
• 🔍 Trigger review

📝 Walkthrough

Walkthrough

Adds a new NVBench CUDA benchmark with a replicated BFloat16 type, device math helpers, deterministic input generation, and four chained elementwise workloads. Base benchmark CUDA flags also enable --expt-relaxed-constexpr.

Changes

Chained elementwise transform benchmarks

Layer / File(s)	Summary
Benchmark build flags cub/benchmarks/CMakeLists.txt	Adds --expt-relaxed-constexpr alongside --extended-lambda for base benchmark CUDA targets.
BFloat16 and transform infrastructure cub/benchmarks/bench/transform/applications/P1/bfloat16.h, cub/benchmarks/bench/transform/applications/P1/pytorch.cu	Adds BFloat16 conversions, arithmetic, comparisons, device math functors, deterministic normal input generation, tuned transform execution, and NVBench type definitions.
Many-input workloads cub/benchmarks/bench/transform/applications/P1/pytorch.cu	Adds many-input/many-instruction and many-input/few-instruction chained elementwise workloads.
Few-input workloads and registration cub/benchmarks/bench/transform/applications/P1/pytorch.cu	Adds few-input/many-instruction and few-input/few-instruction chains, then registers all four workloads with float/BFloat16 types and the configured element range.

Suggested reviewers: fbusato, shwina, jrhemstad

---

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

[coderabbitai]

Actionable comments posted: 2

🧹 Nitpick comments (1)

cub/benchmarks/bench/transform/applications/P1/pytorch.cu (1)

258-265: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick win

suggestion: these comparators take non-const references, so they can't bind const operands or temporaries — the very thing std::min/std::max (const-ref params) require, contradicting the comment. In this file std::min/std::max are all called on opmath_t (float), so as written these are effectively unused; if ever hit on a const BFloat16 they only compile by falling back to the implicit float conversion. Use const&.

-inline __host__ __device__ bool operator>(BFloat16& lhs, BFloat16& rhs)
+inline __host__ __device__ bool operator>(const BFloat16& lhs, const BFloat16& rhs)
 {
   return float(lhs) > float(rhs);
 }
-inline __host__ __device__ bool operator<(BFloat16& lhs, BFloat16& rhs)
+inline __host__ __device__ bool operator<(const BFloat16& lhs, const BFloat16& rhs)
 {
   return float(lhs) < float(rhs);
 }

As per coding guidelines: "All variables that are not modified must be declared const".

Source: Coding guidelines

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: a9c3b764-29b6-4949-87f1-04bc6015f851

📥 Commits

Reviewing files that changed from the base of the PR and between f51b10a and 4a9f9fe.

📒 Files selected for processing (2)

• cub/benchmarks/CMakeLists.txt
• cub/benchmarks/bench/transform/applications/P1/pytorch.cu

[coderabbitai]

Actionable comments posted: 1

🧹 Nitpick comments (1)

cub/benchmarks/bench/transform/applications/P1/bfloat16.h (1)

6-9: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: This header relies on transitive includes for several symbols: uint16_t/UINT16_C/UINT32_C (<cstdint>), __nv_bfloat16/__float2bfloat16/__bfloat16_as_ushort/__bfloat162float (<cuda_bf16.h>), and NV_IF_TARGET/NV_IS_DEVICE (<nv/target>). Include them directly.

 `#include` <cmath>
+#include <cstdint>
 `#include` <cstring>
 
+#include <cuda_bf16.h>
+#include <nv/target>
+
 `#include` <nvbench/type_strings.cuh>

As per coding guidelines: "Include all headers needed by the symbols being used; do not rely on transitive includes."

Source: Coding guidelines

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: a045ad98-301f-4eef-8321-e761f9ec9a95

📥 Commits

Reviewing files that changed from the base of the PR and between 7159511 and 63f97a2.

📒 Files selected for processing (3)

• cub/benchmarks/CMakeLists.txt
• cub/benchmarks/bench/transform/applications/P1/bfloat16.h
• cub/benchmarks/bench/transform/applications/P1/pytorch.cu

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

• cub/benchmarks/CMakeLists.txt

[github-actions]

😬 CI Workflow Results

🟥 Finished in 1h 06m: Pass: 58%/242 | Total: 1d 08h | Max: 1h 06m | Hits: 99%/87559

See results here.