feat(primbench): support an extra backend

projects/rocprim/.gitlab/run_benchmarks.py

@@ -50,8 +50,6 @@ def run_benchmarks(
     max_gpu_temp,
     max_warming_secs,
     max_cooling_secs,
-    output_hip_device_properties_context,
-    output_amdsmi_context,
     output_batches,
     spaces_per_indent,
     stream_blocking_timeout_secs,
@@ -148,10 +146,6 @@ def should_skip(json_out_path):
             args += ["--max-warming-secs", max_warming_secs]
         if max_cooling_secs:
             args += ["--max-cooling-secs", max_cooling_secs]
-        if output_hip_device_properties_context:
-            args += ["--output-hip-device-properties-context"]
-        if output_amdsmi_context:
-            args += ["--output-amdsmi-context"]
         if output_batches:
             args += ["--output-batches"]
         if spaces_per_indent:
@@ -294,20 +288,6 @@ def main():
         default="",
         required=False,
     )
-    parser.add_argument(
-        "--output-hip-device-properties-context",
-        help="Output a `hip_device_properties` object in the context, containing GPU details",
-        default=False,
-        action="store_true",
-        required=False,
-    )
-    parser.add_argument(
-        "--output-amdsmi-context",
-        help="Output an `amdsmi` object in the context, containing GPU details",
-        default=False,
-        action="store_true",
-        required=False,
-    )
     parser.add_argument(
         "--output-batches",
         help="Output a `batches` array for each specialization, containing per-batch details",
@@ -351,8 +331,6 @@ def main():
         args.max_gpu_temp,
         args.max_warming_secs,
         args.max_cooling_secs,
-        args.output_hip_device_properties_context,
-        args.output_amdsmi_context,
         args.output_batches,
         args.spaces_per_indent,
         args.stream_blocking_timeout_secs,

shared/primbench/README.md

@@ -1,6 +1,6 @@
 # primbench
 
-primbench is a single-header HIP benchmarking library.
+primbench is a single-header HIP and CUDA benchmarking library.
 
 ![primbench demo](./assets/primbench.gif)
 
@@ -16,8 +16,8 @@ primbench is a single-header HIP benchmarking library.
 ## Dependencies
 
 primbench has the following dependencies:
-- HIP
-- AMD SMI (for querying live GPU statistics)
+- HIP or CUDA
+- [AMD SMI](https://rocm.docs.amd.com/projects/amdsmi/en/latest/) or [NVML](https://developer.nvidia.com/management-library-nvml) (for querying live GPU statistics)
 - C++17 or later
 
 ## Example
@@ -81,15 +81,15 @@ struct copy_benchmark : public primbench::benchmark_interface
         primbench::log("Allocating device memory");
         T* d_input;
         T* d_output;
-        PRIMBENCH_HIP_CHECK(hipMalloc(&d_input, items * sizeof(T)));
-        PRIMBENCH_HIP_CHECK(hipMalloc(&d_output, items * sizeof(T)));
+        PRIMBENCH_CHECK(hipMalloc(&d_input, items * sizeof(T)));
+        PRIMBENCH_CHECK(hipMalloc(&d_output, items * sizeof(T)));
 
         primbench::log("Copying to device");
-        PRIMBENCH_HIP_CHECK(hipMemcpyAsync(d_input,
-                                           h_input.data(),
-                                           items * sizeof(T),
-                                           hipMemcpyHostToDevice,
-                                           stream));
+        PRIMBENCH_CHECK(hipMemcpyAsync(d_input,
+                                       h_input.data(),
+                                       items * sizeof(T),
+                                       hipMemcpyHostToDevice,
+                                       stream));
 
         dim3 grid(items / items_per_block);
         dim3 block(BlockSize);
@@ -107,8 +107,8 @@ struct copy_benchmark : public primbench::benchmark_interface
                     <<<grid, block, 0, stream>>>(d_input, d_output);
             });
 
-        PRIMBENCH_HIP_CHECK(hipFree(d_input));
-        PRIMBENCH_HIP_CHECK(hipFree(d_output));
+        PRIMBENCH_CHECK(hipFree(d_input));
+        PRIMBENCH_CHECK(hipFree(d_output));
     }
 };
 
@@ -123,29 +123,49 @@ int main(int argc, char* argv[])
 }
 ```
 
-The benchmark is compiled and run like so:
+The HIP benchmark is compiled and run like so:
 
 ```bash
-hipcc -o copy_benchmark examples/copy_benchmark.cpp -I. -lamd_smi && ./copy_benchmark
+hipcc -o copy_benchmark examples/hip/copy_benchmark.cpp -I. -lamd_smi && ./copy_benchmark
+```
+
+And its equivalent CUDA benchmark is compiled and run like so:
+
+```bash
+nvcc -o copy_benchmark examples/cuda/copy_benchmark.cu -I. -lnvidia-ml && ./copy_benchmark
 ```
 
 It outputs this `results.json`:
 ```json
 {
     "context": {
-        "results_version": "2.0.0",
+        "results_version": "3.0.0",
         "general": {
             "algorithm": "copy",
             "specialization_count": 2,
-            "gpu_name": "AMD Instinct MI210",
-            "gpu_arch": "gfx90a",
-            "gpu_pci_bus_id": "0000:83:00.0",
             "library_build_type": "debug",
-            "temp_type": "edge",
+            "gpu": {
+                "name": "AMD Radeon RX 9070 XT",
+                "arch": "gfx1201",
+                "pci_bus_id": "0000:83:00.0"
+            },
+            "backend": {
+                "name": "hip",
+                "hip_version": "6.4.43482-0f2d60242",
+                "runtime_version": "6.4.43482",
+                "driver_version": "6.4.43482",
+                "compiler": {
+                    "name": "clang",
+                    "version": "19.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.4.0 25133 c7fe45cf4b819c5991fe208aaa96edf142730f1d)"
+                },
+                "monitoring": {
+                    "name": "amdsmi",
+                    "version": "25.3.0"
+                }
+            },
+            "temperature_type": "edge",
             "host_name": "host",
-            "date": "2025-11-24T15:01:50+00:00",
-            "hip_version": "6.4.43482-0f2d60242",
-            "clang_version": "19.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.4.0 25133 c7fe45cf4b819c5991fe208aaa96edf142730f1d)"
+            "date": "2025-11-24T15:01:50+00:00"
         },
         "settings": {
             "size": 134217728,
@@ -164,8 +184,6 @@ It outputs this `results.json`:
             "max_gpu_temp": 60,
             "max_warming_secs": 60,
             "max_cooling_secs": 60,
-            "output_hip_device_properties_context": false,
-            "output_amdsmi_context": false,
             "output_batches": false,
             "spaces_per_indent": 4,
             "stream_blocking_timeout_secs": 10
@@ -265,8 +283,6 @@ You can also pass `--help` to benchmarks to print the available options.
 | `--max-gpu-temp`                         | Maximum GPU temperature in °C. Too low slows benchmarks; too high increases noise. (default: 60)                                                                                   |
 | `--max-warming-secs`                     | Maximum seconds allowed for GPU warming before an error is thrown. (default: 60)                                                                                                   |
 | `--max-cooling-secs`                     | Maximum seconds allowed for GPU cooling before an error is thrown. (default: 60)                                                                                                   |
-| `--output-hip-device-properties-context` | Output a `hip_device_properties` object in the context object, containing details about the GPU.                                                                                   |
-| `--output-amdsmi-context`                | Output an `amdsmi` object in the context object, containing details about the GPU.                                                                                                 |
 | `--output-batches`                       | Output a `batches` array for each specialization, containing per-batch details.                                                                                                    |
 | `--spaces-per-indent`                    | Number of spaces per indentation level in JSON output. Set to 0 for no indentation. (default: 4)                                                                                   |
 | `--stream-blocking-timeout-secs`         | Maximum stream blocking duration in seconds before timing out. Stream is blocked while queueing kernel calls. Use `primbench::flags::sync` if kernel is synchronous. (default: 10) |
@@ -344,13 +360,15 @@ Before benchmarking a specialization, primbench:
 
 primbench currently does not cool down the GPU *while* benchmarking a specialization, but this may be changed in the future.
 
-Temperatures are read via AMD SMI. Warming uses short GPU workloads; cooling waits until the GPU naturally drops back into range. If either process takes more than 60 seconds, primbench aborts (`--max-warming-secs`, `--max-cooling-secs`).
+Temperatures are read using AMD SMI/NVML. Warming uses short GPU workloads; cooling waits until the GPU naturally drops back into range. If either process takes more than 60 seconds, primbench aborts (`--max-warming-secs`, `--max-cooling-secs`).
 
 ### GPU Temperature Sensor Selection
 
-primbench supports multiple GPU temperature sensors exposed by AMD SMI, such as *edge* and *hotspot*. At startup, it probes the available sensors and selects the first one that successfully returns a reading. The chosen sensor type is cached and used consistently for warming, cooling, and reporting. If no supported sensor can be read, primbench terminates with an error.
+primbench supports multiple GPU temperature sensors exposed by AMD SMI, such as `edge` and `hotspot`. At startup, it probes the available sensors and selects the first one that successfully returns a reading. The chosen sensor type is cached and used consistently for warming, cooling, and reporting. If no supported sensor can be read, primbench terminates with an error.
+
+The selected temperature sensor type is recorded in the JSON output as `context.general.temperature_type`.
 
-The selected temperature sensor type is recorded in the JSON output as `context.general.temp_type`.
+primbench always uses the GPU temperature sensor called `gpu` for CUDA benchmarks.
 
 ### GPU Cache Clearing
 
@@ -410,7 +428,7 @@ The command below embeds the current branch and commit hash.
 If the repository is in a detached HEAD state (for example in CI or when checking out a specific commit), no branch is active and `BRANCH_NAME` is set to `DETACHED`.
 
 ```bash
-hipcc -o copy_benchmark examples/copy_benchmark.cpp \
+hipcc -o copy_benchmark examples/hip/copy_benchmark.cpp \
   -I. \
   -lamd_smi \
   -DBRANCH_NAME=\"$(git symbolic-ref -q --short HEAD || echo DETACHED)\" \

shared/primbench/examples/cuda/copy_benchmark.cu

@@ -0,0 +1,98 @@
+#include "primbench.hpp"
+
+#include <cuda_runtime.h>
+
+// All benchmarked types must be declared
+// This allows you to for example format `long long` as `int64_t`, or `i64`
+PRIMBENCH_REGISTER_TYPE(char, "char")
+PRIMBENCH_REGISTER_TYPE(long long, "long long")
+
+// Simple copy kernel
+template<typename T, unsigned int BlockSize, unsigned int ItemsPerThread>
+__global__ __launch_bounds__(BlockSize)
+void copy_kernel(const T* input, T* output)
+{
+    unsigned int idx = threadIdx.x + blockIdx.x * BlockSize * ItemsPerThread;
+#pragma unroll
+    for(unsigned int i = 0; i < ItemsPerThread; ++i)
+        output[idx + i * BlockSize] = input[idx + i * BlockSize];
+}
+
+template<typename T>
+struct copy_benchmark : public primbench::benchmark_interface
+{
+    primbench::json meta() const override
+    {
+        // This meta() method must return an `algo` key
+        // The `algo` name must be identical across all queued specializations
+        // primbench::json objects can be nested
+        return primbench::json{}.add("algo", "copy").add("type", primbench::name<T>());
+    }
+
+    // This contains both the setup and running of the benchmark
+    void run(primbench::state& state) override
+    {
+        const auto& stream = state.stream;
+        const auto& bytes  = state.size;
+
+        // primbench::log() calls report progress in gray
+        // They help with discovering slow setup steps
+        primbench::log("Calculating items");
+        size_t                 N              = bytes / sizeof(T);
+        constexpr unsigned int BlockSize      = 256;
+        constexpr unsigned int ItemsPerThread = 4;
+
+        const size_t items_per_block = BlockSize * ItemsPerThread;
+        const size_t items = items_per_block * ((N + items_per_block - 1) / items_per_block);
+
+        primbench::log("Generating input data");
+        std::vector<T> h_input(items);
+        for(size_t i = 0; i < items; ++i)
+            h_input[i] = T(i);
+
+        primbench::log("Allocating output vector");
+        std::vector<T> h_output(items);
+
+        primbench::log("Allocating device memory");
+        T* d_input;
+        T* d_output;
+        PRIMBENCH_CHECK(cudaMalloc(&d_input, items * sizeof(T)));
+        PRIMBENCH_CHECK(cudaMalloc(&d_output, items * sizeof(T)));
+
+        primbench::log("Copying to device");
+        PRIMBENCH_CHECK(cudaMemcpyAsync(d_input,
+                                        h_input.data(),
+                                        items * sizeof(T),
+                                        cudaMemcpyHostToDevice,
+                                        stream));
+
+        dim3 grid(items / items_per_block);
+        dim3 block(BlockSize);
+
+        // primbench uses this to calculates the items/sec and bytes/sec
+        state.set_items(items);
+        state.add_reads<T>(items);
+        state.add_writes<T>(items);
+
+        // This passes a lambda to primbench, which calls it many times
+        // primbench completely handles synchronization
+        state.run(
+            [&] {
+                copy_kernel<T, BlockSize, ItemsPerThread>
+                    <<<grid, block, 0, stream>>>(d_input, d_output);
+            });
+
+        PRIMBENCH_CHECK(cudaFree(d_input));
+        PRIMBENCH_CHECK(cudaFree(d_output));
+    }
+};
+
+int main(int argc, char* argv[])
+{
+    primbench::executor executor(argc, argv);
+
+    executor.queue<copy_benchmark<char>>();
+    executor.queue<copy_benchmark<long long>>();
+
+    executor.run();
+}

shared/primbench/examples/copy_benchmark.cpp → shared/primbench/examples/hip/copy_benchmark.cpp

@@ -54,15 +54,15 @@ struct copy_benchmark : public primbench::benchmark_interface
         primbench::log("Allocating device memory");
         T* d_input;
         T* d_output;
-        PRIMBENCH_HIP_CHECK(hipMalloc(&d_input, items * sizeof(T)));
-        PRIMBENCH_HIP_CHECK(hipMalloc(&d_output, items * sizeof(T)));
+        PRIMBENCH_CHECK(hipMalloc(&d_input, items * sizeof(T)));
+        PRIMBENCH_CHECK(hipMalloc(&d_output, items * sizeof(T)));
 
         primbench::log("Copying to device");
-        PRIMBENCH_HIP_CHECK(hipMemcpyAsync(d_input,
-                                           h_input.data(),
-                                           items * sizeof(T),
-                                           hipMemcpyHostToDevice,
-                                           stream));
+        PRIMBENCH_CHECK(hipMemcpyAsync(d_input,
+                                       h_input.data(),
+                                       items * sizeof(T),
+                                       hipMemcpyHostToDevice,
+                                       stream));
 
         dim3 grid(items / items_per_block);
         dim3 block(BlockSize);
@@ -80,8 +80,8 @@ struct copy_benchmark : public primbench::benchmark_interface
                     <<<grid, block, 0, stream>>>(d_input, d_output);
             });
 
-        PRIMBENCH_HIP_CHECK(hipFree(d_input));
-        PRIMBENCH_HIP_CHECK(hipFree(d_output));
+        PRIMBENCH_CHECK(hipFree(d_input));
+        PRIMBENCH_CHECK(hipFree(d_output));
     }
 };