Merge ~bullwinkle-team/ubuntu/+source/rocprim:bullwinkle/llvm-21/ubuntu/devel into ubuntu/+source/rocprim:ubuntu/devel

[Igor Luppi]

Update to upstream version 7.1.0

[Igor Luppi (igorluppi)]

Uploaded to ppa: https://launchpad.net/~igorluppi/+archive/ubuntu/rocprim-lp2140327/+packages

(-proposed and archs enabled)

[Andreas Hasenack (ahasenack)]

--- a/debian/control
+++ b/debian/control
@@ -32,7 +33,7 @@ Description: parallel primitives for GPU-accelerated code - headers

 Package: librocprim-tests
 Section: libdevel
-Architecture: amd64 arm64 ppc64el
+Architecture: amd64 arm64
 XB-X-ROCm-GPU-Architecture: ${rocm:GPU-Architecture}
 Depends: ${misc:Depends}, ${shlibs:Depends},
 Build-Profiles: <!nocheck>

Please investigate the reverse dependencies of librocprim-tests to see what else would have to be dropped in ppc64el.

review: Needs Information

[Igor Luppi (igorluppi)]

Thanks Andreas

$ reverse-depends --arch ppc64el src:rocprim -x
Reverse-Depends
===============
* libhipcub-dev (for librocprim-dev)
* librocthrust-dev (for librocprim-dev)
  * libstdgpu-hip-dev

Its safe to drop since its only required by rocm stack. Also librocthrust-dev was detailed here: https://code.launchpad.net/~bullwinkle-team/ubuntu/+source/rocthrust/+git/rocthrust/+merge/499735
It should be fine.

[Andreas Hasenack (ahasenack)]

+1, sponsored:

Uploading rocprim_7.1.0-0ubuntu3.dsc
Uploading rocprim_7.1.0.orig.tar.gz
Uploading rocprim_7.1.0-0ubuntu3.debian.tar.xz
Uploading rocprim_7.1.0-0ubuntu3_source.buildinfo
Uploading rocprim_7.1.0-0ubuntu3_source.changes

review: Approve

diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index 809934d..2ac616b 100755
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@@ -1,4 +1,4 @@
-* @stanleytsang-amd @umfranzw @RobsonRLemos @lawruble13
+* @stanleytsang-amd @umfranzw @RobsonRLemos 
 # Documentation files
 docs/ @ROCm/rocm-documentation
 *.md @ROCm/rocm-documentation
diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml
index 0f76d8c..ec65005 100644
--- a/.gitlab-ci.yml
+++ b/.gitlab-ci.yml
@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -42,6 +42,17 @@ stages:
   - test
   - benchmark
 
+workflow:
+  rules:
+    - if: $CI_MERGE_REQUEST_LABELS =~ /CI Skip/
+      when: never
+    - if: $CI_MERGE_REQUEST_TITLE !~ /Draft:/
+      variables:
+        ROCPRIM_TEST_RUNS: 1
+    - if: $CI_MERGE_REQUEST_TITLE =~ /Draft:/
+      variables:
+        ROCPRIM_TEST_RUNS: 1
+
 variables:
   PACKAGE_DIR: $BUILD_DIR/package
   AUTOTUNE_CONFIG_DIR: ${CI_PROJECT_DIR}/autotune_config
@@ -110,11 +121,10 @@ copyright-date:
       -D BUILD_TEST=ON
       -D BUILD_EXAMPLE=ON
       -D BUILD_BENCHMARK=ON
-      -D GPU_TARGETS=$GPU_TARGETS
-      -D AMDGPU_TEST_TARGETS=$GPU_TARGETS
+      -D AMDGPU_TARGETS=$GPU_TARGETS
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
-      -D CMAKE_CXX_STANDARD=14
+      -D CMAKE_CXX_STANDARD=17
       -S $CI_PROJECT_DIR
       -B $BUILD_DIR
     - cmake
@@ -164,13 +174,14 @@ build:cmake-minimum-apt:
     - .rules:build
   variables:
     EXTRA_CMAKE_CXX_FLAGS: ""
+    BUILD_TOOL_ARGS: ""
   script:
     - mkdir -p $BUILD_DIR
     - cd $BUILD_DIR
     - | # Add hardened libc++ assertions for tests only
         if [[ $BUILD_TARGET == "TEST" ]]; then
             echo "Configuring with hardened libc++!"
-            EXTRA_CMAKE_CXX_FLAGS+=" -D_GLIBCXX_ASSERTIONS=ON"
+            EXTRA_CMAKE_CXX_FLAGS+=" -D_GLIBCXX_ASSERTIONS=ON -D ROCPRIM_ENABLE_ASSERTS=ON"
         fi
     - cmake
       -G Ninja
@@ -178,15 +189,15 @@ build:cmake-minimum-apt:
       -D CMAKE_CXX_FLAGS="-Wall -Wextra -Werror $EXTRA_CMAKE_CXX_FLAGS"
       -D CMAKE_BUILD_TYPE="$BUILD_TYPE"
       -D BUILD_$BUILD_TARGET=ON
+      -D WITH_ROCRAND=ON
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
       -D BUILD_EXAMPLE=ON
-      -D GPU_TARGETS=$GPU_TARGETS
-      -D AMDGPU_TEST_TARGETS=$GPU_TARGETS
+      -D AMDGPU_TARGETS=$GPU_TARGETS
       -D CMAKE_CXX_STANDARD="$BUILD_VERSION"
       -S $CI_PROJECT_DIR
       -B $BUILD_DIR
-    - cmake --build $BUILD_DIR
+    - cmake --build $BUILD_DIR -- ${BUILD_TOOL_ARGS}
   artifacts:
     paths:
       - $BUILD_DIR/.ninja_log
@@ -197,9 +208,36 @@ build:cmake-minimum-apt:
       - $BUILD_DIR/gtest/
       - $BUILD_DIR/test/CTestTestfile.cmake
       - $BUILD_DIR/test/rocprim/CTestTestfile.cmake
+      - $BUILD_DIR/deps/rocrand/
       - $BUILD_DIR/test/rocprim/test_*
+      - $BUILD_DIR/test/rocprim/libtest_*
       - $BUILD_DIR/test/test_*
-    expire_in: 2 weeks
+    expire_in: 1 day
+
+build:spirv:
+  stage: build
+  needs: []
+  extends:
+    - .cmake-minimum
+    - .build:common
+  variables:
+    # For unknown reasons spir-v builds ignore 'clang diagnostic' pragmas that
+    # we use to ignore internal deprecations.
+    EXTRA_CMAKE_CXX_FLAGS: "-Wno-deprecated-declarations -mf16c -DROCPRIM_EXPERIMENTAL_SPIRV"
+    # Since not all targets are expected to build, do not stop building other
+    # targets when any target fails.
+    BUILD_TOOL_ARGS: "-k 0"
+    GPU_TARGETS: "amdgcnspirv"
+  image: "registry.streamhpc.internal/unstable-rocm:main"
+  allow_failure: true
+  parallel:
+    # Debug builds disabled due to excessive build times for debug test builds
+    matrix:
+      - BUILD_TYPE: Release
+        BUILD_TARGET: [BENCHMARK, TEST]
+        BUILD_VERSION: 17
+  artifacts:
+    when: always
 
 build:cmake-latest:
   stage: build
@@ -212,7 +250,7 @@ build:cmake-latest:
     matrix:
       - BUILD_TYPE: Release
         BUILD_TARGET: [BENCHMARK, TEST]
-        BUILD_VERSION: [14, 17]
+        BUILD_VERSION: 17
 
 build:cmake-minimum:
   needs: []
@@ -223,7 +261,7 @@ build:cmake-minimum:
     matrix:
       - BUILD_TYPE: [Debug, Release]
         BUILD_TARGET: [BENCHMARK, TEST]
-        BUILD_VERSION: 14
+        BUILD_VERSION: 17
 
 build:package:
   stage: build
@@ -240,7 +278,7 @@ build:package:
       -G Ninja
       -D CMAKE_CXX_COMPILER="$AMDCLANG"
       -D CMAKE_BUILD_TYPE=Release
-      -D CMAKE_CXX_STANDARD=14
+      -D CMAKE_CXX_STANDARD=17
       -B $PACKAGE_DIR
       -S $CI_PROJECT_DIR
     - cd $PACKAGE_DIR
@@ -249,7 +287,7 @@ build:package:
     paths:
       - $PACKAGE_DIR/rocprim*.deb
       - $PACKAGE_DIR/rocprim*.zip
-    expire_in: 2 weeks
+    expire_in: 1 day
 
 build:windows:
   stage: build
@@ -261,19 +299,19 @@ build:windows:
     - .deps:visual-studio-devshell
   parallel:
     matrix:
-      - BUILD_TYPE: [Debug, Release]
-        BUILD_TARGET: [BENCHMARK, TEST]
+      - BUILD_TYPE: Release
   script:
     - mkdir -p $CI_PROJECT_DIR/build
     - cmake -G Ninja
       -S $CI_PROJECT_DIR
       -B $CI_PROJECT_DIR/build
-      -D BUILD_$BUILD_TARGET=ON
-      -D GPU_TARGETS=$GPU_TARGET
+      -D BUILD_TEST=ON
+      -D BUILD_BENCHMARK=ON
+      -D AMDGPU_TARGETS=$GPU_TARGET
       -D CMAKE_CXX_COMPILER:PATH="${env:HIP_PATH}\bin\clang++.exe"
       -D CMAKE_PREFIX_PATH:PATH="${env:HIP_PATH}"
       -D CMAKE_BUILD_TYPE="$BUILD_TYPE"
-      -D CMAKE_CXX_STANDARD=14
+      -D CMAKE_CXX_STANDARD=17
     - cmake --build "$CI_PROJECT_DIR/build"
   artifacts:
     paths:
@@ -285,7 +323,7 @@ build:windows:
       - $CI_PROJECT_DIR/build/CMakeCache.txt
       - $CI_PROJECT_DIR/build/.ninja_log
       - $CI_PROJECT_DIR/build/CTestTestfile.cmake
-    expire_in: 2 weeks
+    expire_in: 1 day
 
 autotune:build:
   stage: autotune
@@ -317,10 +355,10 @@ autotune:build:
       -D BUILD_EXAMPLE=OFF
       -D BUILD_BENCHMARK=ON
       -D BENCHMARK_CONFIG_TUNING=ON
-      -D GPU_TARGETS=$GPU_TARGETS
+      -D AMDGPU_TARGETS=$GPU_TARGETS
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
-      -D CMAKE_CXX_STANDARD=14
+      -D CMAKE_CXX_STANDARD=17
     - cmake --build . --target $BENCHMARK_TARGETS
     - 'rm -rf $BUILD_DIR/benchmark/benchmark*.parallel'
     # The autotune benchmarks get very large, above GitLabs upload limit. Fortunately they compress well.
@@ -334,19 +372,20 @@ autotune:build:
       - $BUILD_DIR/deps/googlebenchmark/
     expire_in: 1 week
 
-test:
+.test:common:
   stage: test
+  tags:
+    - rocm
+    - $GPU
   extends:
     - .cmake-minimum
-    - .rules:test
-    - .gpus:rocm
   needs:
     - job: build:cmake-minimum
       parallel:
         matrix:
           - BUILD_TYPE: Release
             BUILD_TARGET: TEST
-            BUILD_VERSION: 14
+            BUILD_VERSION: 17
   script:
     - cd $BUILD_DIR
     - cmake
@@ -359,10 +398,88 @@ test:
     - HSA_ENABLE_SDMA=0 ctest
       --output-on-failure
       --repeat-until-fail 2
-      --tests-regex "hip|$GPU_TARGET"
       --resource-spec-file ./resources.json
       --parallel $PARALLEL_JOBS
 
+test:any-gpu:
+  variables:
+    GPU: ""
+    PARALLEL_JOBS: 1
+  extends:
+    - .test:common
+  rules:
+    - if: $CI_MERGE_REQUEST_TITLE =~ /Draft:/ && $CI_MERGE_REQUEST_LABELS !~ /Arch::/
+
+test:label-arch:
+  extends:
+    - .gpus:rocm
+    - .test:common
+    - .rules:arch-labels
+
+test:all-gpus:
+  variables:
+    SHOULD_BE_UNDRAFTED: "true"
+  extends:
+    - .gpus:rocm
+    - .test:common
+    - .rules:test
+
+.test:common-spirv:
+  stage: test
+  tags:
+    - rocm
+    - $GPU
+  extends:
+    - .cmake-minimum
+  allow_failure: true
+  timeout: 3h
+  needs:
+    - job: build:spirv
+      parallel:
+        matrix:
+          - BUILD_TYPE: Release
+            BUILD_TARGET: TEST
+            BUILD_VERSION: 17
+  image: "registry.streamhpc.internal/unstable-rocm:main"
+  script:
+    - cd $BUILD_DIR
+    - cmake
+      -D CMAKE_PREFIX_PATH=/opt/rocm
+      -P $CI_PROJECT_DIR/cmake/GenerateResourceSpec.cmake
+    - cat ./resources.json
+    # Parallel execution (with other AMDGPU processes) can oversubscribe the SDMA queue.
+    # This causes the hipMemcpy to fail, which is not reported as an error by HIP.
+    # As a temporary workaround, disable the SDMA for test stability.
+    - HSA_ENABLE_SDMA=0 ctest
+      --output-on-failure
+      --repeat-until-fail 2
+      --resource-spec-file ./resources.json
+      --parallel $PARALLEL_JOBS
+      --exclude-regex rocprim.device_partition
+
+test:any-gpu-spirv:
+  variables:
+    GPU: ""
+    PARALLEL_JOBS: 1
+  extends:
+    - .test:common-spirv
+  rules:
+    - if: $CI_MERGE_REQUEST_TITLE =~ /Draft:/ && $CI_MERGE_REQUEST_LABELS !~ /Arch::/
+
+test:label-arch-spirv:
+  extends:
+    - .gpus:rocm
+    - .test:common-spirv
+    - .rules:arch-labels
+
+test:all-gpus-spirv:
+  variables:
+    SHOULD_BE_UNDRAFTED: "true"
+  extends:
+    - .gpus:rocm
+    - .test:common-spirv
+    - .rules:test
+
 .test-windows-base:
   stage: test
   extends:
@@ -394,7 +511,6 @@ test-windows-release:
       parallel:
         matrix:
           - BUILD_TYPE: Release
-            BUILD_TARGET: TEST
 
 .test-package:
   script:
@@ -402,8 +518,8 @@ test-windows-release:
       -G Ninja
       -D CMAKE_CXX_COMPILER="$AMDCLANG"
       -D CMAKE_BUILD_TYPE=Release
-      -D GPU_TARGETS=$GPU_TARGETS
-      -D CMAKE_CXX_STANDARD=14
+      -D AMDGPU_TARGETS=$GPU_TARGETS
+      -D CMAKE_CXX_STANDARD=17
       -S "$CI_PROJECT_DIR/test/extra"
       -B "$CI_PROJECT_DIR/package_test"
     - cmake --build "$CI_PROJECT_DIR/package_test"
@@ -425,10 +541,11 @@ test:install:
       -G Ninja
       -D CMAKE_CXX_COMPILER="$AMDCLANG"
       -D CMAKE_BUILD_TYPE=Release
-      -D CMAKE_CXX_STANDARD=14
+      -D CMAKE_CXX_STANDARD=17
       -B build
       -S $CI_PROJECT_DIR
-    - $SUDO_CMD cmake --build build --target install
+    # Preserve $PATH when sudoing
+    - $SUDO_CMD env PATH="$PATH" cmake --build build --target install
     - !reference [.test-package, script]
 
 test:deb:
@@ -452,6 +569,10 @@ test:docs:
   extends:
     - .rules:test
     - .build:docs
+  artifacts:
+    paths:
+      - $DOCS_DIR/_build/html/
+    expire_in: 2 weeks
 
 .benchmark-base:
   stage: benchmark
@@ -468,7 +589,7 @@ benchmark:
         matrix:
           - BUILD_TYPE: Release
             BUILD_TARGET: BENCHMARK
-            BUILD_VERSION: 14
+            BUILD_VERSION: 17
   extends:
     - .cmake-minimum
     - .gpus:rocm
@@ -489,6 +610,11 @@ benchmark:
       --benchmark_filename_regex "${BENCHMARK_FILENAME_REGEX}"
       --benchmark_filter_regex "${BENCHMARK_ALGORITHM_REGEX}"
       --seed "${BENCHMARK_SEED}"
+    - python3
+      .gitlab/report_noise.py
+      --benchmark_json_dir "${BENCHMARK_RESULT_DIR}"
+      --noise_threshold_percentage 1.0
+      --accept_high_noise
   artifacts:
     paths:
       - ${BENCHMARK_RESULT_DIR}
@@ -577,6 +703,11 @@ autotune:execute-tuning:
       --size="${AUTOTUNE_SIZE}"
       --trials="${AUTOTUNE_TRIALS}"
       --seed=82589933
+    - python3
+      .gitlab/report_noise.py
+      --benchmark_json_dir "${AUTOTUNE_RESULT_DIR}"
+      --noise_threshold_percentage 1.0
+      --accept_high_noise
 
 autotune:generate-config:
   image: python:3.10.5-buster
diff --git a/.gitlab/report_noise.py b/.gitlab/report_noise.py
new file mode 100644
index 0000000..974c2b1
--- /dev/null
+++ b/.gitlab/report_noise.py
@@ -0,0 +1,251 @@
+#!/usr/bin/env python3
+
+# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+# THE SOFTWARE.
+
+
+import argparse
+import json
+import os
+import re
+import stat
+import statistics
+import sys
+
+
+class colors:
+    OK = "\033[92m"
+    FAIL = "\033[91m"
+    END_COLOR = "\033[0m"
+
+
+def print_results(results):
+    # Store the length of the longest value in a column
+    longest = {
+        "name": max(len(result["name"]) for result in results),
+        "noisy_permutations": max(
+            len(result["noisy_permutations"]) for result in results
+        ),
+        "mean": max(len(result["mean"]) for result in results),
+        "median": max(len(result["median"]) for result in results),
+        "max": max(len(result["max"]) for result in results),
+        "batch": max(len(result["batch"]) for result in results),
+        "warmup": max(len(result["warmup"]) for result in results),
+        "bytes": max(len(result["bytes"]) for result in results),
+    }
+
+    # The name of a column can be longer than its values
+    longest = {key: max(value, len(key)) for key, value in longest.items()}
+
+    printed = "name".ljust(longest["name"] + 1)
+    printed += "noisy permutations".ljust(longest["noisy_permutations"] + 1)
+    printed += "mean".ljust(longest["mean"] + 1)
+    printed += "median".ljust(longest["median"] + 1)
+    printed += "max".ljust(longest["max"] + 1)
+    printed += "batch".ljust(longest["batch"] + 1)
+    printed += "warmup".ljust(longest["warmup"] + 1)
+    printed += "bytes".ljust(longest["bytes"] + 1)
+    printed += "seed"
+    print(printed)
+
+    for result in results:
+        printed = result["name"].ljust(longest["name"])
+
+        printed += " "
+        printed += colors.FAIL if result["noisy"] else colors.OK
+        printed += (
+            f'{result["noisy_permutations"].ljust(longest["noisy_permutations"])}'
+        )
+        printed += colors.END_COLOR
+
+        printed += " "
+        printed += colors.FAIL if result["bad_mean"] else colors.OK
+        printed += result["mean"].ljust(longest["mean"])
+        printed += colors.END_COLOR
+
+        printed += " "
+        printed += colors.FAIL if result["bad_median"] else colors.OK
+        printed += result["median"].ljust(longest["median"])
+        printed += colors.END_COLOR
+
+        printed += " "
+        printed += colors.FAIL if result["bad_max"] else colors.OK
+        printed += result["max"].ljust(longest["max"])
+        printed += colors.END_COLOR
+
+        printed += " "
+        printed += colors.FAIL if result["bad_batch"] else colors.OK
+        printed += result["batch"].ljust(longest["batch"])
+        printed += colors.END_COLOR
+
+        printed += " "
+        printed += colors.FAIL if result["bad_warmup"] else colors.OK
+        printed += result["warmup"].ljust(longest["warmup"])
+        printed += colors.END_COLOR
+
+        printed += " "
+        printed += colors.FAIL if result["bad_bytes"] else colors.OK
+        printed += result["bytes"].ljust(longest["bytes"])
+        printed += colors.END_COLOR
+
+        printed += " "
+        printed += colors.FAIL if result["seed"] == "random" else colors.OK
+        printed += result["seed"]
+        printed += colors.END_COLOR
+
+        print(printed)
+
+
+def get_results(benchmarks, threshold):
+    def get_humanized_bytes(size):
+        for unit in ["B", "KiB", "MiB", "GiB", "TiB", "PiB"]:
+            if size < 1024.0 or unit == "PiB":
+                break
+            size /= 1024.0
+        return f"{size:.1f} {unit}"
+
+    success = True
+
+    results = []
+
+    for benchmark in benchmarks:
+        data = benchmark["data"]
+
+        name = benchmark["name"]
+
+        permutations = data["benchmarks"]
+
+        cvs = [permutation["cv"] for permutation in permutations]
+
+        # The cv (coefficient of variation) is a standard way of quantifying noise
+        noises = sum(cv * 100 > threshold for cv in cvs)
+        noisy = noises > 0
+
+        if noisy:
+            success = False
+
+        context = data["context"]
+
+        noisy_permutations = f"{noises}/{len(permutations)}"
+
+        mean = statistics.mean(cvs)
+        median = statistics.median(cvs)
+        max_ = max(cvs)
+
+        batch = context["batch_iterations"]
+        warmup = context["warmup_iterations"]
+
+        bytes_ = int(context["size"])
+        seed = context["seed"]
+
+        results.append(
+            {
+                "name": name,
+                "noisy": noisy,
+                "noisy_permutations": noisy_permutations,
+                "bad_mean": mean * 100 > threshold,
+                "mean": f"{mean:.1%}",
+                "bad_median": median * 100 > threshold,
+                "median": f"{median:.1%}",
+                "bad_max": max_ * 100 > threshold,
+                "max": f"{max_:.1%}",
+                "bad_batch": int(batch) < 10,
+                "batch": batch,
+                "bad_warmup": int(warmup) < 5,
+                "warmup": warmup,
+                "bad_bytes": 0 < bytes_ < 128 * 1024 * 1024,  # 128 MiB
+                "bytes": get_humanized_bytes(int(context["size"])),
+                "seed": seed,
+            }
+        )
+
+    return results, success
+
+
+def load_benchmarks(benchmark_json_dir):
+    def is_benchmark_json(filename):
+        if not re.match(r".*\.json$", filename):
+            return False
+        path = os.path.join(benchmark_json_dir, filename)
+        st_mode = os.stat(path).st_mode
+
+        # we are not interested in permissions, just whether it is a regular file (S_IFREG)
+        return st_mode & stat.S_IFREG
+
+    benchmark_names = [
+        name for name in os.listdir(benchmark_json_dir) if is_benchmark_json(name)
+    ]
+
+    success = True
+    benchmarks = []
+    for benchmark_name in benchmark_names:
+        with open(os.path.join(benchmark_json_dir, benchmark_name)) as f:
+            try:
+                benchmarks.append({"name": benchmark_name, "data": json.load(f)})
+            except json.JSONDecodeError as e:
+                print(
+                    f"{colors.FAIL}Failed to load {benchmark_name}{colors.END_COLOR}: {e}\n",
+                    file=sys.stderr,
+                )
+                success = False
+
+    return benchmarks, success
+
+
+def main():
+    parser = argparse.ArgumentParser()
+    parser.add_argument(
+        "--noise_threshold_percentage",
+        help="The noise threshold percentage, past which benchmark permutations are considered to be too noisy",
+        required=True,
+        type=float,
+    )
+    parser.add_argument(
+        "--benchmark_json_dir",
+        help="The directory of benchmark JSON files, which to report the noise of",
+        required=True,
+    )
+    parser.add_argument(
+        "--accept_high_noise",
+        help="Don't call exit(1) when there is a noisy benchmark permutation",
+        action=argparse.BooleanOptionalAction,
+    )
+    args = parser.parse_args()
+
+    print(f"The noise threshold is {args.noise_threshold_percentage:.1f}%\n")
+
+    benchmarks, load_success = load_benchmarks(args.benchmark_json_dir)
+    results, results_success = get_results(benchmarks, args.noise_threshold_percentage)
+
+    print_results(results)
+
+    if not load_success:
+        return False
+    if args.accept_high_noise:
+        return True
+    return results_success
+
+
+if __name__ == "__main__":
+    success = main()
+    if success:
+        exit(0)
+    else:
+        exit(1)
diff --git a/.gitlab/run_benchmarks.py b/.gitlab/run_benchmarks.py
index f406ea4..d61b4d9 100755
--- a/.gitlab/run_benchmarks.py
+++ b/.gitlab/run_benchmarks.py
@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 
-# Copyright (c) 2022-2024 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2022-2025 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -22,13 +22,14 @@
 
 import argparse
 from collections import namedtuple
+import json
 import os
 import re
 import stat
 import subprocess
 import sys
 
-BenchmarkContext = namedtuple('BenchmarkContext', ['gpu_architecture', 'benchmark_output_dir', 'benchmark_dir', 'benchmark_filename_regex', 'benchmark_filter_regex', 'size', 'trials', 'seed'])
+BenchmarkContext = namedtuple('BenchmarkContext', ['gpu_architecture', 'benchmark_output_dir', 'benchmark_dir', 'benchmark_filename_regex', 'benchmark_filter_regex', 'size', 'trials', 'seed', 'skip_gathered'])
 
 def run_benchmarks(benchmark_context):
     def is_benchmark_executable(filename):
@@ -41,6 +42,18 @@ def run_benchmarks(benchmark_context):
         # and it is a regular file (S_IFREG)
         return (st_mode & (stat.S_IXUSR | stat.S_IXGRP | stat.S_IXOTH)) and (st_mode & stat.S_IFREG)
 
+    def should_skip(results_json_path):
+        if not benchmark_context.skip_gathered:
+            return False
+
+        try:
+            with open(results_json_path) as f:
+                json.load(f)
+        except (FileNotFoundError, json.JSONDecodeError):
+            return False
+
+        return True
+
     success = True
     benchmark_names = [name for name in os.listdir(benchmark_context.benchmark_dir) if is_benchmark_executable(name)]
     print('The following benchmarks will be ran:\n{}'.format('\n'.join(benchmark_names)), file=sys.stderr, flush=True)
@@ -49,11 +62,11 @@ def run_benchmarks(benchmark_context):
 
         benchmark_path = os.path.join(benchmark_context.benchmark_dir, benchmark_name)
         results_json_path = os.path.join(benchmark_context.benchmark_output_dir, results_json_name)
+        if should_skip(results_json_path):
+            print(f'Skipping {benchmark_name}, because its results have already been gathered at {results_json_path}', file=sys.stderr, flush=True)
+            continue
         args = [
             benchmark_path,
-            '--name_format',
-            'json',
-            '--benchmark_out_format=json',
             f'--benchmark_out={results_json_path}',
             f'--benchmark_filter={benchmark_context.benchmark_filter_regex}'
         ]
@@ -103,6 +116,11 @@ def main():
         help='Controls the seed for random number generation for each benchmark case',
         default='',
         required=False)
+    parser.add_argument('--skip_gathered',
+        help='Skip running benchmarks whose JSON data has already been gathered',
+        default=False,
+        action='store_true',
+        required=False)
 
     args = parser.parse_args()
 
@@ -114,7 +132,8 @@ def main():
         args.benchmark_filter_regex,
         args.size,
         args.trials,
-        args.seed)
+        args.seed,
+        args.skip_gathered)
 
     benchmark_run_successful = run_benchmarks(benchmark_context)
 
diff --git a/.jenkins/common.groovy b/.jenkins/common.groovy
deleted file mode 100644
index d49a782..0000000
--- a/.jenkins/common.groovy
+++ /dev/null
@@ -1,76 +0,0 @@
-// This file is for internal AMD use.
-// If you are interested in running your own Jenkins, please raise a github issue for assistance.
-
-def runCompileCommand(platform, project, jobName, boolean debug=false)
-{
-    project.paths.construct_build_prefix()
-
-    String buildTypeArg = debug ? '-DCMAKE_BUILD_TYPE=Debug' : '-DCMAKE_BUILD_TYPE=Release'
-    String buildTypeDir = debug ? 'debug' : 'release'
-    String cmake = platform.jenkinsLabel.contains('centos') ? 'cmake3' : 'cmake'
-    //Set CI node's gfx arch as target if PR, otherwise use default targets of the library
-    String amdgpuTargets = env.BRANCH_NAME.startsWith('PR-') ? '-DAMDGPU_TARGETS=\$gfx_arch' : ''
-
-    def command = """#!/usr/bin/env bash
-                set -x
-                cd ${project.paths.project_build_prefix}
-                mkdir -p build/${buildTypeDir} && cd build/${buildTypeDir}
-                ${auxiliary.gfxTargetParser()}
-                ${cmake} --toolchain=toolchain-linux.cmake ${buildTypeArg} ${amdgpuTargets} -DBUILD_TEST=ON -DBUILD_BENCHMARK=ON ../..
-                make -j\$(nproc)
-                """
-
-    platform.runCommand(this, command)
-}
-
-
-def runTestCommand (platform, project)
-{
-    String sudo = auxiliary.sudo(platform.jenkinsLabel)
-
-    def testCommand = "ctest --output-on-failure "
-    def testCommandExcludeRegex = ''
-    def testCommandExclude = "--exclude-regex \"${testCommandExcludeRegex}\""
-    def hmmExcludeRegex = ''
-    def hmmTestCommandExclude = "--exclude-regex \"${hmmExcludeRegex}\""
-    def hmmTestCommand = ''
-    if (platform.jenkinsLabel.contains('gfx90a'))
-    {
-        echo("HMM TESTS DISABLED")
-        /*hmmTestCommand = """
-                            export HSA_XNACK=1
-                            export ROCPRIM_USE_HMM=1
-                            ${testCommand} ${hmmTestCommandExclude}
-                         """*/
-    }
-    echo(env.JOB_NAME)
-    if (env.JOB_NAME.contains('bleeding-edge'))
-    {
-        testCommand = ''
-        testCommandExclude = ''
-        hmmTestCommand = ''
-        echo("TESTS DISABLED")
-    }
-    def command = """#!/usr/bin/env bash
-                set -x
-                cd ${project.paths.project_build_prefix}
-                cd ${project.testDirectory}
-                ${testCommand} ${testCommandExclude}
-                if (( \$? != 0 )); then
-                    exit 1
-                fi
-                ${hmmTestCommand}
-            """
-
-    platform.runCommand(this, command)
-}
-
-def runPackageCommand(platform, project)
-{
-    def packageHelper = platform.makePackage(platform.jenkinsLabel,"${project.paths.project_build_prefix}/build/release")
-
-    platform.runCommand(this, packageHelper[0])
-        platform.archiveArtifacts(this, packageHelper[1])
-}
-
-return this
diff --git a/.jenkins/precheckin.groovy b/.jenkins/precheckin.groovy
deleted file mode 100644
index dd32cad..0000000
--- a/.jenkins/precheckin.groovy
+++ /dev/null
@@ -1,81 +0,0 @@
-#!/usr/bin/env groovy
-@Library('rocJenkins@pong') _
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path;
-
-def runCI = 
-{
-    nodeDetails, jobName->
-
-    def prj = new rocProject('rocPRIM', 'PreCheckin')
-    prj.paths.build_command = './install -c'
-    prj.timeout.compile = 600
-
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    def commonGroovy
-
-    boolean formatCheck = false
-     
-    def compileCommand =
-    {
-        platform, project->
-
-        commonGroovy = load "${project.paths.project_src_prefix}/.jenkins/common.groovy"
-        commonGroovy.runCompileCommand(platform, project, jobName)
-    }
-
-    def testCommand =
-    {
-        platform, project->
-
-        commonGroovy.runTestCommand(platform, project)
-    }
-
-    def packageCommand =
-    {
-        platform, project->
-
-        commonGroovy.runPackageCommand(platform, project)
-    }
-
-    buildProject(prj, formatCheck, nodes.dockerArray, compileCommand, testCommand, packageCommand)
-}
-
-ci: { 
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    def propertyList = ["compute-rocm-dkms-no-npi":[pipelineTriggers([cron('0 1 * * 0')])], 
-                        "compute-rocm-dkms-no-npi-hipclang":[pipelineTriggers([cron('0 1 * * 0')])],
-                        "rocm-docker":[]]
-    propertyList = auxiliary.appendPropertyList(propertyList)
-
-    def jobNameList = ["compute-rocm-dkms-no-npi-hipclang":([ubuntu18:['gfx900'],centos7:['gfx906'],centos8:['gfx906'],sles15sp1:['gfx908']])]
-    jobNameList = auxiliary.appendJobNameList(jobNameList, 'rocPRIM')
-
-    propertyList.each 
-    {
-        jobName, property->
-        if (urlJobName == jobName)
-            properties(auxiliary.addCommonProperties(property))
-    }
-    
-    jobNameList.each 
-    {
-        jobName, nodeDetails->
-        if (urlJobName == jobName)
-            stage(jobName) {
-                runCI(nodeDetails, jobName)
-            }
-    }
-
-    // For url job names that are not listed by the jobNameList i.e. compute-rocm-dkms-no-npi-1901
-    if(!jobNameList.keySet().contains(urlJobName))
-    {
-        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * *')])]))
-        stage(urlJobName) {
-            runCI([ubuntu16:['gfx906']], urlJobName)
-        }
-    }
-}
diff --git a/.jenkins/static.groovy b/.jenkins/static.groovy
deleted file mode 100644
index 9466b27..0000000
--- a/.jenkins/static.groovy
+++ /dev/null
@@ -1,82 +0,0 @@
-#!/usr/bin/env groovy
-@Library('rocJenkins@pong') _
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path;
-
-def runCI = 
-{
-    nodeDetails, jobName->
-
-    def prj = new rocProject('rocPRIM', 'static')
-    prj.paths.build_command = './install -c -s'
-    prj.timeout.compile = 600
-    prj.timeout.packaging = 120
-
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    def commonGroovy
-
-    boolean formatCheck = false
-     
-    def compileCommand =
-    {
-        platform, project->
-
-        commonGroovy = load "${project.paths.project_src_prefix}/.jenkins/common.groovy"
-        commonGroovy.runCompileCommand(platform, project, jobName)
-    }
-
-    def testCommand =
-    {
-        platform, project->
-
-        commonGroovy.runTestCommand(platform, project)
-    }
-
-    def packageCommand =
-    {
-        platform, project->
-
-        commonGroovy.runPackageCommand(platform, project)
-    }
-
-    buildProject(prj, formatCheck, nodes.dockerArray, compileCommand, testCommand, packageCommand)
-}
-
-ci: { 
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    def propertyList = ["compute-rocm-dkms-no-npi":[pipelineTriggers([cron('0 1 * * 0')])], 
-                        "compute-rocm-dkms-no-npi-hipclang":[pipelineTriggers([cron('0 1 * * 0')])],
-                        "rocm-docker":[]]
-    propertyList = auxiliary.appendPropertyList(propertyList)
-
-    def jobNameList = ["compute-rocm-dkms-no-npi-hipclang":([ubuntu18:['gfx900'],centos7:['gfx906'],centos8:['gfx906'],sles15sp1:['gfx908']])]
-    jobNameList = auxiliary.appendJobNameList(jobNameList, 'rocPRIM')
-
-    propertyList.each 
-    {
-        jobName, property->
-        if (urlJobName == jobName)
-            properties(auxiliary.addCommonProperties(property))
-    }
-    
-    jobNameList.each 
-    {
-        jobName, nodeDetails->
-        if (urlJobName == jobName)
-            stage(jobName) {
-                runCI(nodeDetails, jobName)
-            }
-    }
-
-    // For url job names that are not listed by the jobNameList i.e. compute-rocm-dkms-no-npi-1901
-    if(!jobNameList.keySet().contains(urlJobName))
-    {
-        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * *')])]))
-        stage(urlJobName) {
-            runCI([ubuntu16:['gfx906']], urlJobName)
-        }
-    }
-}
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 624c2f5..8fcaa79 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -2,19 +2,177 @@
 
 Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projects/rocPRIM/en/latest/](https://rocm.docs.amd.com/projects/rocPRIM/en/latest/).
 
+## rocPRIM 4.1.0 for ROCm 7.1
+
+### Added
+
+* Added `get_sreg_lanemask_lt`, `get_sreg_lanemask_le`, `get_sreg_lanemask_gt` and `get_sreg_lanemask_ge`.
+* Added `rocprim::transform_output_iterator` and `rocprim::make_transform_output_iterator`.
+* Added experimental support for SPIR-V, to use the correct tuned config for part of the appliable algorithms.
+* Added a new cmake option, `BUILD_OFFLOAD_COMPRESS`. When rocPRIM is build with this option enabled, the `--offload-compress` switch is passed to the compiler. This causes the compiler to compress the binary that it generates. Compression can be useful in cases where you are compiling for a large number of targets, since this often results in a large binary. Without compression, in some cases, the generated binary may become so large symbols are placed out of range, resulting in linking errors. The new `BUILD_OFFLOAD_COMPRESS` option is set to `ON` by default.
+* Added a new CMake option `-DUSE_SYSTEM_LIB` to allow tests to be built from `ROCm` libraries provided by the system.
+* Added `rocprim::apply` which applies a function to a `rocprim::tuple`.
+
+### Changed
+
+* Changed tests to support `ptr-to-const` output in `/test/rocprim/test_device_batch_memcpy.cpp`.
+
+### Optimizations
+
+* Improved performance of many algorithms, by updating their tuned configs.
+  * 891 specializations have been improved.
+  * 399 specializations have been added.
+
+### Upcoming changes
+
+* Deprecated the `->` operator for the `zip_iterator`.
+
+### Resolved issues
+
+* Fixed `device_select`, `device_merge`, and `device_merge_sort` not allocating the correct amount of virtual shared memory on the host.
+* Fixed the `->` operator for the `transform_iterator`, the `texture_cache_iterator` and the `arg_index_iterator`, by now returning a proxy pointer.
+  * The `arg_index_iterator` also now only returns the internal iterator for the `->`.
+
+## rocPRIM 4.0.1 for ROCm 7.0.2
+
+### Resolved issues
+
+* Fixed compilation issue when using `rocprim::texture_cache_iterator`.
+* Fixed a HIP version check used to determine whether hipStreamLegacy is supported. This resolves runtime errors that occur when hipStreamLegacy is used in versions of ROCm later than 6.4.
+
+## rocPRIM 4.0.0 for ROCm 7.0
+
+### Added
+
+* Added `rocprim::accumulator_t` to ensure parity with CCCL.
+* Added test for `rocprim::accumulator_t`
+* Added `rocprim::invoke_result_r` to ensure parity with CCCL.
+* Added function `is_build_in` into `rocprim::traits::get`.
+* Added virtual shared memory as a fallback option in `rocprim::device_merge` when it exceeds shared memory capacity, similar to `rocprim::device_select`, `rocprim::device_partition`, and `rocprim::device_merge_sort`, which already include this feature.
+* Added initial value support to device level inclusive scans.
+* Added new optimization to the backend for `device_transform` when the input and output are pointers.
+* Added `LoadType` to `transform_config`, which is used for the `device_transform` when the input and output are pointers.
+* Added `rocprim:device_transform` for n-ary transform operations API with as input `n` number of iterators inside a `rocprim::tuple`.
+* Added gfx950 support.
+* Added `rocprim::key_value_pair::operator==`.
+* Added the `rocprim::unrolled_copy` thread function to copy multiple items inside a thread.
+* Added the `rocprim::unrolled_thread_load` function to load multiple items inside a thread using `rocprim::thread_load`.
+* Added `rocprim::int128_t` and `rocprim::uint128_t` to benchmarks for improved performance evaluation on 128-bit integers.
+* Added `rocprim::int128_t` to the supported autotuning types to improve performance for 128-bit integers.
+* Added the `rocprim::merge_inplace` function for merging in-place.
+* Added initial value support for warp- and block-level inclusive scan.
+* Added support for building tests with device-side random data generation, making them finish faster. This requires rocRAND, and is enabled with the `WITH_ROCRAND=ON` build flag.
+* Added tests and documentation to `lookback_scan_state`. It is still in the `detail` namespace.
+
+### Optimizations
+
+* Improved performance of `rocprim::device_select` and `rocprim::device_partition` when using multiple streams on the MI3XX architecture.
+
+### Changed
+
+* Changed the parameters `long_radix_bits` and `LongRadixBits` from `segmented_radix_sort` to `radix_bits` and `RadixBits` respectively.
+* Marked the initialisation constructor of `rocprim::reverse_iterator<Iter>` `explicit`, use `rocprim::make_reverse_iterator`.
+* Merged `radix_key_codec` into type_traits system.
+* Renamed `type_traits_interface.hpp` to `type_traits.hpp`, rename the original `type_traits.hpp` to `type_traits_functions.hpp`.
+* The default scan accumulator types for device-level scan algorithms have changed. This is a breaking change.
+The previous default accumulator types could lead to situations in which unexpected overflow occured, such as
+when the input or inital type was smaller than the output type. 
+  * This is a complete list of affected functions and how their default accumulator types are changing:
+    * `rocprim::inclusive_scan`
+      * Previous default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
+      * Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, typename std::iterator_traits<InputIterator>::value_type>`
+    * `rocprim::deterministic_inclusive_scan`
+      * Previous default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
+      * Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, typename std::iterator_traits<InputIterator>::value_type>`
+    * `rocprim::exclusive_scan`
+      * Previous default: `class AccType = detail::input_type_t<InitValueType>>`
+      * Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, rocprim::detail::input_type_t<InitValueType>>`
+    * `rocprim::deterministic_exclusive_scan`
+      * Previous default: `class AccType = detail::input_type_t<InitValueType>>`
+      * Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, rocprim::detail::input_type_t<InitValueType>>`
+* Undeprecated internal `detail::raw_storage`.
+* A new version of `rocprim::thread_load` and `rocprim::thread_store` replace the deprecated `rocprim::thread_load` and `rocprim::thread_store` functions. The versions avoid inline assembly where possible, and don't hinder the optimizer as much as a result.
+* Renamed `rocprim::load_cs` to `rocprim::load_nontemporal` and `rocprim::store_cs` to `rocprim::store_nontemporal` to express the intent of these load and store methods better.
+* All kernels now have hidden symbol visibility. All symbols now have inline namespaces that include the library version, for example, `rocprim::ROCPRIM_300400_NS::symbol` instead of `rocPRIM::symbol`, letting the user link multiple libraries built with different versions of rocPRIM.
+    
+### Upcoming changes
+
+* `rocprim::invoke_result_binary_op` and `rocprim::invoke_result_binary_op_t` are deprecated. Use `rocprim::accumulator_t` now.
+
+### Removed
+
+* Removed `rocprim::detail::float_bit_mask` and relative tests, use `rocprim::traits::float_bit_mask` instead.
+* Removed `rocprim::traits::is_fundamental`, please use `rocprim::traits::get<T>::is_fundamental()` directly.
+* Removed the deprecated parameters `short_radix_bits` and `ShortRadixBits` from the `segmented_radix_sort` config. They were unused, it is only an API change.
+* Removed the deprecated `operator<<` from the iterators.
+* Removed the deprecated `TwiddleIn` and `TwiddleOut`. Use `radix_key_codec` instead.
+* Removed the deprecated flags API of `block_adjacent_difference`. Use `subtract_left()` or `block_discontinuity::flag_heads()` instead.
+* Removed the deprecated `to_exclusive` functions in the warp scans.
+* Removed the `rocprim::load_cs` from the `cache_load_modifier` enum. Use `rocprim::load_nontemporal` instead.
+* Removed the `rocprim::store_cs` from the `cache_store_modifier` enum. Use `rocprim::store_nontemporal` instead.
+* Removed the deprecated header file `rocprim/detail/match_result_type.hpp`. Include `rocprim/type_traits.hpp` instead.
+  * This header included `rocprim::detail::invoke_result`. Use `rocprim::invoke_result` instead.
+  * This header included `rocprim::detail::invoke_result_binary_op`. Use `rocprim::invoke_result_binary_op` instead.
+  * This header included `rocprim::detail::match_result_type`. Use `rocprim::invoke_result_binary_op_t` instead.
+* Removed the deprecated `rocprim::detail::radix_key_codec` function. Use `rocprim::radix_key_codec` instead.
+* Removed `rocprim/detail/radix_sort.hpp`, functionality can now be found in `rocprim/thread/radix_key_codec.hpp`.
+* Removed C++14 support, only C++17 is supported.
+* Due to the removal of `__AMDGCN_WAVEFRONT_SIZE` in the compiler, the following deprecated warp size-related symbols have been removed:
+  * `rocprim::device_warp_size()`
+    * For compile-time constants, this is replaced with `rocprim::arch::wavefront::min_size()` and `rocprim::arch::wavefront::max_size()`. Use this when allocating global or shared memory.
+    * For run-time constants, this is replaced with `rocprim::arch::wavefront::size().`
+  * `rocprim::warp_size()`
+    * Use `rocprim::host_warp_size()`, `rocprim::arch::wavefront::min_size()` or `rocprim::arch::wavefront::max_size()` instead.
+  * `ROCPRIM_WAVEFRONT_SIZE`
+    * Use `rocprim::arch::wavefront::min_size()` or `rocprim::arch::wavefront::max_size()` instead.
+  * `__AMDGCN_WAVEFRONT_SIZE`
+    * This was a fallback define for the compiler's removed symbol, having the same name. 
+* This release removes support for custom builds on gfx940 and gfx941.
+
+### Resolved issues
+
+* Fixed an issue where `device_batch_memcpy` reported benchmarking throughput being 2x lower than it was in reality.
+* Fixed an issue where `device_segmented_reduce` reported autotuning throughput being 5x lower than it was in reality.
+* Fixed device radix sort not returning the correct required temporary storage when a double buffer contains `nullptr`.
+* Fixed constness of equality operators (`==` and `!=`) in `rocprim::key_value_pair`.
+* Fixed an issue for the comparison operators in `arg_index_iterator` and `texture_cache_iterator`, where `<` and `>` comparators were swapped.
+* Fixed an issue for the `rocprim::thread_reduce` not working correctly with a prefix value.
+
+### Known issues
+* When using `rocprim::deterministic_inclusive_scan_by_key` and `rocprim::deterministic_exclusive_scan_by_key` the intermediate values can change order on Navi3x
+  * However if a commutative scan operator is used then the final scan value (output array) will still always be consistent between runs
+
 ## rocPRIM 3.4.1 for ROCm 6.4.2
 
 ### Upcoming changes
 * Changes to the template parameters of warp and block algorithms will be made in an upcoming release.
 
-### Deprecations
 * Due to an upcoming compiler change the following warp size-related symbols will be removed in the next major release and are thus marked as deprecated:
   * `rocprim::device_warp_size()`
     * For compile-time constants, this is replaced with `rocprim::arch::wavefront::min_size()` and `rocprim::arch::wavefront::max_size()`. Use this when allocating global or shared memory.
     * For run-time constants, this is replaced with `rocprim::arch::wavefront::size().`
   * `rocprim::warp_size()`
-  * `ROCPRIM_WAVEFRONT_SIZE
+  * `ROCPRIM_WAVEFRONT_SIZE`
   
+* The default scan accumulator types for device-level scan algorithms will be changed in an upcoming release, resulting in a breaking change. Previously, the default accumulator type was set to the input type for the inclusive scans and to the initial value type for the exclusive scans. This could lead to unexpected overflow if the input or initial type was smaller than the output type when the accumulator type was't explicitly set using the `AccType` template parameter. The new default accumulator types will be set to the type that results when the input or initial value type is applied to the scan operator.  
+
+    The following is the complete list of affected functions and how their default accumulator types are changing:
+    
+    * `rocprim::inclusive_scan`
+        * current default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
+        * future default: `class AccType = rocprim::invoke_result_binary_op_t<typename std::iterator_traits<InputIterator>::value_type, BinaryFunction>`
+    * `rocprim::deterministic_inclusive_scan`
+        * current default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
+        * future default: `class AccType = rocprim::invoke_result_binary_op_t<typename std::iterator_traits<InputIterator>::value_type, BinaryFunction>`
+    * `rocprim::exclusive_scan`
+        * current default: `class AccType = detail::input_type_t<InitValueType>>`
+        * future default: `class AccType = rocprim::invoke_result_binary_op_t<rocprim::detail::input_type_t<InitValueType>, BinaryFunction>`
+    * `rocprim::deterministic_exclusive_scan`
+        * current default: `class AccType = detail::input_type_t<InitValueType>>`
+        * future default: `class AccType = rocprim::invoke_result_binary_op_t<rocprim::detail::input_type_t<InitValueType>, BinaryFunction>`
+
+* `rocprim::load_cs` and `rocprim::store_cs` are deprecated and will be removed in an upcoming release. Alternatively, you can use `rocprim::load_nontemporal` and `rocprim::store_nontemporal` to load and store values in specific conditions (like bypassing the cache) for `rocprim::thread_load` and `rocprim::thread_store`.
+
 ## rocPRIM 3.4.0 for ROCm 6.4.0
 
 ### Added
@@ -34,8 +192,8 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
 * Added the parallel `search` and `find_end` device functions similar to `std::search` and `std::find_end`, these functions search for the first and last occurrence of the sequence respectively.
 * Added a parallel device-level function, `rocprim::search_n`, similar to the C++ Standard Library `std::search_n` algorithm.
 * Added new constructors and a `base` function, and added `constexpr` specifier to all functions in `rocprim::reverse_iterator` to improve parity with the C++17 `std::reverse_iterator`.
-* Added hipGraph support to device run-length-encode for non trivial runs (`rocprim::run_length_encode_non_trivial_runs`).
-* Added configuration autotuning to device run-length-encode for non trivial runs (`rocprim::run_length_encode_non_trivial_runs`) for improved performance on selected architectures.
+* Added hipGraph support to device run-length-encode for nontrivial runs (`rocprim::run_length_encode_non_trivial_runs`).
+* Added configuration autotuning to device run-length-encode for nontrivial runs (`rocprim::run_length_encode_non_trivial_runs`) for improved performance on selected architectures.
 * Added configuration autotuning to device run-length-encode for trivial runs (`rocprim::run_length_encode`) for improved performance on selected architectures.
 * Added a new type traits interface to enable users to provide additional type trait information to rocPRIM, facilitating better compatibility with custom types.
 
@@ -523,3 +681,5 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
 
 * Switched to HIP-Clang as the default compiler
 * CMake searches for rocPRIM locally first; if t's not found, CMake downloads it from GitHub
+
+
diff --git a/CMakeLists.txt b/CMakeLists.txt
index f6579e5..20999f1 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -41,11 +41,8 @@ set(CMAKE_HIP_STANDARD 14)
 set(CMAKE_HIP_STANDARD_REQUIRED ON)
 set(CMAKE_HIP_EXTENSIONS OFF)
 
-# Set CXX standard
-if (CMAKE_CXX_STANDARD EQUAL 14)
-  message(WARNING "C++14 will be deprecated in the next major release")
-elseif(NOT CMAKE_CXX_STANDARD EQUAL 17)
-  message(FATAL_ERROR "Only C++14 and C++17 are supported")
+if(NOT CMAKE_CXX_STANDARD EQUAL 17)
+  message(FATAL_ERROR "Only C++17 is supported")
 endif()
 
 if (CMAKE_CURRENT_SOURCE_DIR STREQUAL CMAKE_SOURCE_DIR)
@@ -68,16 +65,31 @@ include(CMakeDependentOption)
 # Disables building tests, benchmarks, examples
 option(ONLY_INSTALL "Only install" OFF)
 cmake_dependent_option(BUILD_TEST "Build tests (requires googletest)" OFF "NOT ONLY_INSTALL" OFF)
+option(WITH_ROCRAND "Build tests with device-side data generation(requires rocRAND)" OFF)
 cmake_dependent_option(BUILD_BENCHMARK "Build benchmarks" OFF "NOT ONLY_INSTALL" OFF)
 cmake_dependent_option(BUILD_EXAMPLE "Build examples" OFF "NOT ONLY_INSTALL" OFF)
 option(BUILD_NAIVE_BENCHMARK "Build naive benchmarks" OFF)
 cmake_dependent_option(BUILD_DOCS "Build documentation (requires sphinx)" OFF "NOT ONLY_INSTALL" OFF)
 option(BUILD_CODE_COVERAGE "Build with code coverage enabled" OFF)
 option(ROCPRIM_INSTALL "Enable installation of rocPRIM (projects embedding rocPRIM may want to turn this OFF)" ON)
+option(ROCPRIM_ENABLE_ASSERTS "Enable asserts in release build)" OFF)
+option(BUILD_OFFLOAD_COMPRESS "Build rocPRIM with offload compression" ON)
+cmake_dependent_option(USE_SYSTEM_LIB "Use installed ROCm libs when building tests" OFF BUILD_TEST OFF)
 
 check_language(HIP)
 cmake_dependent_option(USE_HIPCXX "Use CMake HIP language support" OFF CMAKE_HIP_COMPILER OFF)
 
+include(CheckCXXCompilerFlag)
+
+if(BUILD_OFFLOAD_COMPRESS)
+  check_cxx_compiler_flag("--offload-compress" CXX_COMPILER_SUPPORTS_OFFLOAD_COMPRESS)
+  if(CXX_COMPILER_SUPPORTS_OFFLOAD_COMPRESS)
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --offload-compress")
+  else()
+    message(STATUS "Warning: BUILD_OFFLOAD_COMPRESS=ON but flag not supported by compiler. Ignoring option.")
+  endif()
+endif()
+
 if (CMAKE_CURRENT_SOURCE_DIR STREQUAL CMAKE_SOURCE_DIR)
   set(ROCPRIM_PROJECT_IS_TOP_LEVEL TRUE)
 else()
@@ -104,6 +116,21 @@ if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
   set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "" "Debug" "Release" "MinSizeRel" "RelWithDebInfo")
 endif()
 
+if(ROCPRIM_ENABLE_ASSERTS)
+  if(NOT "${CMAKE_BUILD_TYPE}" STREQUAL "")
+      string(TOUPPER ${CMAKE_BUILD_TYPE} BUILD_TYPE)
+      set(BUILD_TYPE_CXX_FLAGS "CMAKE_CXX_FLAGS_${BUILD_TYPE}")
+      set(BUILD_TYPE_C_FLAGS "CMAKE_C_FLAGS_${BUILD_TYPE}")
+  endif()
+
+  string(REGEX REPLACE "-DNDEBUG( |$)" "" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
+  string(REGEX REPLACE "-DNDEBUG( |$)" "" CMAKE_C_FLAGS "${CMAKE_C_FLAGS}")
+  if(NOT "${CMAKE_BUILD_TYPE}" STREQUAL "")
+      string(REGEX REPLACE "-DNDEBUG( |$)" "" ${BUILD_TYPE_CXX_FLAGS} "${${BUILD_TYPE_CXX_FLAGS}}")
+      string(REGEX REPLACE "-DNDEBUG( |$)" "" ${BUILD_TYPE_C_FLAGS} "${${BUILD_TYPE_C_FLAGS}}")
+  endif()
+endif()
+
 set(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE CACHE BOOL "Add paths to linker search and installed rpath")
 
 if(DEFINED BUILD_SHARED_LIBS)
@@ -115,6 +142,7 @@ set(BUILD_SHARED_LIBS OFF) # don't build client dependencies as shared
 
 # Get dependencies (required here to get rocm-cmake)
 include(cmake/Dependencies.cmake)
+
 # Use target ID syntax if supported for GPU_TARGETS
 if(USE_HIPCXX)
   enable_language(HIP)
@@ -130,18 +158,24 @@ else()
     if(BUILD_ADDRESS_SANITIZER)
       # ASAN builds require xnack
       rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
-        TARGETS "gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+"
+        TARGETS "gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+;gfx950:xnack+"
       )
     else()
       rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
-        TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201"
+        TARGETS "gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201"
       )
     endif()
-
     set(GPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for" FORCE)
   endif()
 endif()
 
+# Compressed offload binaries are currently not working with the SPIR-V target
+if("amdgcnspirv" IN_LIST GPU_TARGETS)
+  if(BUILD_OFFLOAD_COMPRESS)
+    message(FATAL_ERROR "Cannot combine SPIR-V and BUILD_OFFLOAD_COMPRESS")
+  endif()
+endif()
+
 # TODO: Fix VerifyCompiler for HIP on Windows
 if (NOT WIN32)
   include(cmake/VerifyCompiler.cmake)
@@ -149,25 +183,15 @@ endif()
 list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH} ${ROCM_PATH}/hip ${ROCM_PATH}/llvm ${ROCM_ROOT}/llvm ${ROCM_ROOT} ${ROCM_ROOT}/hip)
 find_package(hip REQUIRED CONFIG PATHS ${HIP_DIR} ${ROCM_PATH} /opt/rocm)
 
-# FOR HANDLING ENABLE/DISABLE OPTIONAL BACKWARD COMPATIBILITY for FILE/FOLDER REORG
-option(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY "Build with file/folder reorg with backward compatibility enabled" OFF)
-if(ROCPRIM_INSTALL AND BUILD_FILE_REORG_BACKWARD_COMPATIBILITY AND NOT WIN32)
-  rocm_wrap_header_dir(
-    "${PROJECT_SOURCE_DIR}/rocprim/include/rocprim"
-    WRAPPER_LOCATIONS rocprim/include/rocprim
-    OUTPUT_LOCATIONS rocprim/wrapper/include/rocprim
-    PATTERNS *.hpp
-  )
-endif()
-
 if(BUILD_CODE_COVERAGE)
   add_compile_options(-fprofile-arcs -ftest-coverage)
   add_link_options(--coverage)
 endif()
 
 # Setup VERSION
-set(VERSION_STRING "3.4.1")
+set(VERSION_STRING "4.1.0")
 rocm_setup_version(VERSION ${VERSION_STRING})
+math(EXPR rocprim_VERSION_NUMBER "${rocprim_VERSION_MAJOR} * 100000 + ${rocprim_VERSION_MINOR} * 100 + ${rocprim_VERSION_PATCH}")
 
 # Print configuration summary
 include(cmake/Summary.cmake)
@@ -182,6 +206,13 @@ endif()
 
 # Tests
 if(BUILD_TEST)
+  if(USE_SYSTEM_LIB)
+    find_package(rocprim REQUIRED CONFIG PATHS "/opt/rocm/rocprim")
+    if (NOT ${rocprim_VERSION} VERSION_EQUAL ${VERSION_STRING})
+      message(WARNING "The installed rocprim version, ${rocprim_VERSION}, does not match project version ${VERSION_STRING}. Building tests with USE_SYSTEM_LIB=ON may not work properly.")
+    endif()
+  endif()
+
   if (ROCPRIM_PROJECT_IS_TOP_LEVEL)
     rocm_package_setup_client_component(tests)
   endif()
@@ -223,14 +254,14 @@ if (ROCPRIM_PROJECT_IS_TOP_LEVEL)
   rocm_package_add_deb_dependencies(STATIC_DEPENDS "hip-static-dev >= ${HIP_RUNTIME_MINIMUM}")
   rocm_package_add_rpm_dependencies(STATIC_DEPENDS "hip-static-devel >= ${HIP_RUNTIME_MINIMUM}")
 
-  set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.txt")
+  set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.md")
   set(CPACK_RPM_PACKAGE_LICENSE "MIT")
 
   set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "\${CPACK_PACKAGING_INSTALL_PREFIX}" )
 
   rocm_create_package(
     NAME rocprim
-    DESCRIPTION "Radeon Open Compute Parallel Primitives Library"
+    DESCRIPTION "rocPRIM is a header-only library that provides HIP parallel primitives." 
     MAINTAINER "rocPRIM Maintainer <[e-mail]>"
     HEADER_ONLY
   )
diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md
index 77c2847..8769ab6 100644
--- a/CONTRIBUTING.md
+++ b/CONTRIBUTING.md
@@ -41,6 +41,8 @@ and performs well across a variety of input types and sizes. More specifically:
 - Tests and benchmarks must be instantiated with all supported data types.
   - If the algorithm uses multiple data types (for instance, if it uses different types for input and output), a selected and representative few combinations should be tested instead of the full combination matrix.
 
+Any utility needed by the tests **and** benchmarks must be added to the appropriate header within the `common` folder. Non-common utilities may be hosted in the corresponding headers from the `test` or `benchmark` folders. For a more detailed description of the cases to be considered for adding new utilities, please check [common](/common/README.md).
+
 We also employ automated testing and benchmarking via checks that are run when a pull request is created.
 These checks:
 - test all algorithms for correctness across a variety of input configurations (eg. types, sizes, etc.)
diff --git a/LICENSE.txt b/LICENSE.md
similarity index 94%
rename from LICENSE.txt
rename to LICENSE.md
index ba22abe..4d43ac8 100644
--- a/LICENSE.txt
+++ b/LICENSE.md
@@ -1,6 +1,6 @@
 MIT License
 
-Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (C) Advanced Micro Devices, Inc.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -18,4 +18,4 @@ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
-SOFTWARE.
\ No newline at end of file
+SOFTWARE.
diff --git a/README.md b/README.md
index 329b170..8caa96c 100644
--- a/README.md
+++ b/README.md
@@ -1,7 +1,7 @@
 # rocPRIM
 
 > [!NOTE]
-> The published documentation is available at [rocPRIM](https://rocm.docs.amd.com/projects/rocPRIM/en/latest/) in an organized, easy-to-read format, with search and a table of contents. The documentation source files reside in the `docs` folder of this repository. As with all ROCm projects, the documentation is open source. For more information on contributing to the documentation, see [Contribute to ROCm documentation](https://rocm.docs.amd.com/en/latest/contribute/contributing.html).
+> The published rocPRIM documentation is available [here](https://rocm.docs.amd.com/projects/rocPRIM/en/latest/) in an organized, easy-to-read format, with search and a table of contents. The documentation source files reside in the `docs` folder of this repository. As with all ROCm projects, the documentation is open source. For more information on contributing to the documentation, see [Contribute to ROCm documentation](https://rocm.docs.amd.com/en/latest/contribute/contributing.html).
 
 rocPRIM is a header-only library that provides HIP parallel primitives. You can use this library to
 develop performant GPU-accelerated code on AMD ROCm platforms.
@@ -28,55 +28,6 @@ Optional:
   * Required only for benchmarks. Building benchmarks is off by default.
   * This is automatically downloaded and built by the CMake script.
 
-## Documentation
-
-Documentation for rocPRIM is available at
-[https://rocm.docs.amd.com/projects/rocPRIM/en/latest/](https://rocm.docs.amd.com/projects/rocPRIM/en/latest/).
-
-### Build documentation locally
-
-```shell
-# Change directory to rocPRIM docs
-cd rocPRIM; cd docs
-
-# Install documentation dependencies
-python3 -m pip install -r sphinx/requirements.txt
-
-# Build the documentation
-python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html
-
-# To serve the HTML docs locally
-cd _build/html
-python3 -m http.server
-```
-
-### Build documentation via CMake
-
-Install [rocm-cmake](https://github.com/ROCm/rocm-cmake/)
-
-```shell
-# Change directory to rocPRIM
-cd rocPRIM
-
-# Install documentation dependencies
-python3 -m pip install -r docs/sphinx/requirements.txt
-
-# Set C++ compiler
-# This example uses hipcc and assumes it is at the path /usr/bin
-export CXX=hipcc
-export PATH=/usr/bin:$PATH
-
-# Configure the project
-cmake -S . -B ./build -D BUILD_DOCS=ON
-
-# Build the documentation
-cmake --build ./build --target doc
-
-# To serve the HTML docs locally
-cd ./build/docs/html
-python3 -m http.server
-```
-
 ## Build and install
 
 You can build and install rocPRIM on Linux or Windows.
@@ -104,6 +55,7 @@ You can build and install rocPRIM on Linux or Windows.
   #     If you want to detect failures on a per GFX IP basis, setting it to some set of ips will create
   #     separate tests with the ip name embedded into the test name. Building for all, but selecting
   #     tests only of a specific architecture is possible for eg: ctest -R gfx803|gfx900
+  #   USE_SYSTEM_LIB - OFF by default. Setting this flag to ON will build tests from the installed ROCm libs provided by the system. This only takes effect when BUILD_TEST is ON.
   #
   # ! IMPORTANT !
   # Set C++ compiler to HIP-clang. You can do it by adding 'CXX=<path-to-compiler>'
@@ -314,9 +266,8 @@ different types and operations, by passing compile-time configuration structures
 parameter. The main "knobs" are usually the size of the block and the number of items processed by a
 single thread.
 
-rocPRIM has built-in default configurations for each of its primitives. In order to use the included
-configurations, you need to define the macro `ROCPRIM_TARGET_ARCH` as `803` if you want the
-algorithms optimized for gfx803 GCN version, or to `900` for gfx900.
+rocPRIM has built-in default configurations for each of its primitives, these will be used automatically
+based on the input types and the target architecture from the stream used.
 
 ## hipCUB
 
@@ -325,7 +276,86 @@ algorithms optimized for gfx803 GCN version, or to `900` for gfx900.
 [CUB](https://github.com/NVlabs/cub). You can use it to port projects that use the CUB library to the
 [HIP](https://github.com/ROCm/HIP) layer and run them on AMD hardware. In the
 [ROCm](https://rocm.docs.amd.com/en/latest/) environment, hipCUB uses the rocPRIM library as a
-backend; on CUDA platforms, it uses CUB as a backend.
+backend.
+
+## Building the documentation locally
+
+### Requirements
+
+#### Doxygen
+
+The build system uses Doxygen [version 1.9.4](https://github.com/doxygen/doxygen/releases/tag/Release_1_9_4). You can try using a newer version, but that might cause issues.
+
+After you have downloaded Doxygen version 1.9.4:
+
+```shell
+# Add doxygen to your PATH
+echo 'export PATH=<doxygen 1.9.4 path>/bin:$PATH' >> ~/.bashrc
+
+# Apply the updated .bashrc
+source ~/.bashrc
+
+# Confirm that you are using version 1.9.4
+doxygen --version
+```
+
+#### Python
+
+The build system uses Python version 3.10. You can try using a newer version, but that might cause issues.
+
+You can install Python 3.10 alongside your other Python versions using [pyenv](https://github.com/pyenv/pyenv?tab=readme-ov-file#installation):
+
+```shell
+# Install Python 3.10
+pyenv install 3.10
+
+# Create a Python 3.10 virtual environment
+pyenv virtualenv 3.10 venv_rocprim
+
+# Activate the virtual environment
+pyenv activate venv_rocprim
+```
+
+### Building
+
+After cloning this repository, and `cd`ing into it:
+
+```shell
+# Install Python dependencies
+python3 -m pip install -r docs/sphinx/requirements.txt
+
+# Build the documentation
+python3 -m sphinx -T -E -b html -d docs/_build/doctrees -D language=en docs docs/_build/html
+```
+
+You can then open `docs/_build/html/index.html` in your browser to view the documentation.
+
+### Build documentation via CMake
+
+Install [rocm-cmake](https://github.com/ROCm/rocm-cmake/)
+
+```shell
+# Change directory to rocPRIM
+cd rocPRIM
+
+# Install documentation dependencies
+python3 -m pip install -r docs/sphinx/requirements.txt
+
+# Set C++ compiler
+# This example uses hipcc and assumes it is at the path /usr/bin
+export CXX=hipcc
+export PATH=/usr/bin:$PATH
+
+# Configure the project
+cmake -S . -B ./build -D BUILD_DOCS=ON
+
+# Build the documentation
+cmake --build ./build --target doc
+
+# To serve the HTML docs locally
+cd ./build/docs/html
+python3 -m http.server
+```
 
 ## Support
 
diff --git a/benchmark/CMakeLists.txt b/benchmark/CMakeLists.txt
index fd8ee78..a45a6dd 100644
--- a/benchmark/CMakeLists.txt
+++ b/benchmark/CMakeLists.txt
@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -147,6 +147,7 @@ add_rocprim_benchmark(benchmark_device_find_first_of.cpp)
 add_rocprim_benchmark(benchmark_device_find_end.cpp)
 add_rocprim_benchmark(benchmark_device_histogram.cpp)
 add_rocprim_benchmark(benchmark_device_merge.cpp)
+add_rocprim_benchmark(benchmark_device_merge_inplace.cpp)
 add_rocprim_benchmark(benchmark_device_merge_sort.cpp)
 add_rocprim_benchmark(benchmark_device_merge_sort_block_sort.cpp)
 add_rocprim_benchmark(benchmark_device_merge_sort_block_merge.cpp)
@@ -173,6 +174,7 @@ add_rocprim_benchmark(benchmark_device_segmented_radix_sort_keys.cpp)
 add_rocprim_benchmark(benchmark_device_segmented_radix_sort_pairs.cpp)
 add_rocprim_benchmark(benchmark_device_segmented_reduce.cpp)
 add_rocprim_benchmark(benchmark_device_transform.cpp)
+add_rocprim_benchmark(benchmark_device_transform_pointer.cpp)
 add_rocprim_benchmark(benchmark_predicate_iterator.cpp)
 add_rocprim_benchmark(benchmark_warp_exchange.cpp)
 add_rocprim_benchmark(benchmark_warp_reduce.cpp)
diff --git a/benchmark/ConfigAutotuneSettings.cmake b/benchmark/ConfigAutotuneSettings.cmake
index 37f3a97..acd4f12 100644
--- a/benchmark/ConfigAutotuneSettings.cmake
+++ b/benchmark/ConfigAutotuneSettings.cmake
@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2022-2024 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2022-2025 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -21,18 +21,18 @@
 # SOFTWARE.
 
 # All default fallback types as listed in scripts/autotune/fallback_config.json
-set(TUNING_TYPES "int64_t int short int8_t double float rocprim::half")
+set(TUNING_TYPES "rocprim::int128_t int64_t int short int8_t double float rocprim::half")
 # If config selection happens based on two types, the second type has limited fallbacks. The selection is based
 # on the size and it is ignored whether the type is floating-point or integral. The autotuning script uses the
 # benchmarks for the integral types as fallback, hence tuning for the floating-point types is not needed.
-set(LIMITED_TUNING_TYPES "int64_t int short int8_t")
+set(LIMITED_TUNING_TYPES "rocprim::int128_t int64_t int short int8_t")
 
 function(read_config_autotune_settings file list_across_names list_across output_pattern_suffix)
   if(file STREQUAL "benchmark_device_adjacent_difference")
-    set(list_across_names "DataType;Left;InPlace;BlockSize" PARENT_SCOPE)
+    set(list_across_names "DataType;Left;Aliasing;BlockSize" PARENT_SCOPE)
     set(list_across "${TUNING_TYPES};\
-true;false true;32 64 128 256 512 1024" PARENT_SCOPE)
-    set(output_pattern_suffix "@DataType@_@Left@_@InPlace@_@BlockSize@" PARENT_SCOPE)
+true;no_alias in_place;32 64 128 256 512 1024" PARENT_SCOPE)
+    set(output_pattern_suffix "@DataType@_@Left@_@Aliasing@_@BlockSize@" PARENT_SCOPE)
   elseif(file STREQUAL "benchmark_device_adjacent_find")
     set(list_across_names "InputType;BlockSize" PARENT_SCOPE)
     set(list_across "${TUNING_TYPES};64 128 256 512 1024" PARENT_SCOPE)
@@ -44,13 +44,13 @@ true;false true;32 64 128 256 512 1024" PARENT_SCOPE)
   elseif(file STREQUAL "benchmark_device_merge_sort_block_merge")
     set(list_across_names "KeyType;ValueType;BlockSize;UseMergePath" PARENT_SCOPE)
     set(list_across "\
-${TUNING_TYPES};rocprim::empty_type ${LIMITED_TUNING_TYPES} custom_type<char,double>;\
+${TUNING_TYPES};rocprim::empty_type ${LIMITED_TUNING_TYPES};\
 128 256 512 1024;true" PARENT_SCOPE)
     set(output_pattern_suffix "@KeyType@_@ValueType@_@BlockSize@_@UseMergePath@" PARENT_SCOPE)
   elseif(file STREQUAL "benchmark_device_merge_sort_block_sort")
     set(list_across_names "KeyType;ValueType;BlockSize;BlockSortMethod" PARENT_SCOPE)
     set(list_across "\
-${TUNING_TYPES};rocprim::empty_type ${LIMITED_TUNING_TYPES} custom_type<char,double>;\
+${TUNING_TYPES};rocprim::empty_type ${LIMITED_TUNING_TYPES};\
 256 512 1024;rocprim::block_sort_algorithm::stable_merge_sort" PARENT_SCOPE)
     set(output_pattern_suffix "@KeyType@_@ValueType@_@BlockSize@_@BlockSortMethod@" PARENT_SCOPE)
   elseif(file STREQUAL "benchmark_device_radix_sort_block_sort")
@@ -85,24 +85,40 @@ ${TUNING_TYPES};${LIMITED_TUNING_TYPES};using_warp_scan reduce_then_scan" PARENT
     set(list_across "\
 binary_search upper_bound lower_bound;${TUNING_TYPES};${LIMITED_TUNING_TYPES};64 128 256;1 2 4 8 16" PARENT_SCOPE)
     set(output_pattern_suffix "@SubAlgorithm@_@ValueType@_@OutputType@_@BlockSize@_@ItemsPerThread@" PARENT_SCOPE)
+  elseif(file STREQUAL "benchmark_device_search_n")
+    set(list_across_names "InputType;BlockSize;ItemsPerThread;Threshold" PARENT_SCOPE)
+    set(list_across "\
+${TUNING_TYPES};64 128 256 512 1024;1 2 4 8 16;4 8 12 16" PARENT_SCOPE)
+    set(output_pattern_suffix "@InputType@_@BlockSize@_@ItemsPerThread@_@Threshold@" PARENT_SCOPE)
   elseif(file STREQUAL "benchmark_device_segmented_radix_sort_keys")
     set(list_across_names "\
-KeyType;LongBits;BlockSize;ItemsPerThread;WarpSmallLWS;WarpSmallIPT;WarpSmallBS;WarpPartition;WarpMediumLWS;WarpMediumIPT;WarpMediumBS" PARENT_SCOPE)
+KeyType;RadixBits;BlockSize;ItemsPerThread;WarpSmallLWS;WarpSmallIPT;WarpSmallBS;WarpPartition;WarpMediumLWS;WarpMediumIPT;WarpMediumBS" PARENT_SCOPE)
     set(list_across "${TUNING_TYPES};8;256;4 8 16;8;4;256;64;16;8;256" PARENT_SCOPE)
     set(output_pattern_suffix "\
-@KeyType@_@LongBits@_@BlockSize@_@ItemsPerThread@_@WarpSmallLWS@_@WarpSmallIPT@_@WarpSmallBS@_@WarpPartition@_@WarpMediumLWS@_@WarpMediumIPT@_@WarpMediumBS@" PARENT_SCOPE)
+@KeyType@_@RadixBits@_@BlockSize@_@ItemsPerThread@_@WarpSmallLWS@_@WarpSmallIPT@_@WarpSmallBS@_@WarpPartition@_@WarpMediumLWS@_@WarpMediumIPT@_@WarpMediumBS@" PARENT_SCOPE)
   elseif(file STREQUAL "benchmark_device_segmented_radix_sort_pairs")
     set(list_across_names "\
-KeyType;ValueType;LongBits;BlockSize;ItemsPerThread;WarpSmallLWS;WarpSmallIPT;WarpSmallBS;WarpPartition;WarpMediumLWS;WarpMediumIPT;WarpMediumBS" PARENT_SCOPE)
-    set(list_across "${TUNING_TYPES};int8_t;8;256;4 8 16;8;4;256;64;16;8;256" PARENT_SCOPE)
+KeyType;ValueType;RadixBits;BlockSize;ItemsPerThread;WarpSmallLWS;WarpSmallIPT;WarpSmallBS;WarpPartition;WarpMediumLWS;WarpMediumIPT;WarpMediumBS" PARENT_SCOPE)
+    set(list_across "${TUNING_TYPES};${LIMITED_TUNING_TYPES};8;256;4 8 16;8;4;256;64;16;8;256" PARENT_SCOPE)
     set(output_pattern_suffix "\
-@KeyType@_@ValueType@_@LongBits@_@BlockSize@_@ItemsPerThread@_@WarpSmallLWS@_@WarpSmallIPT@_@WarpSmallBS@_@WarpPartition@_@WarpMediumLWS@_@WarpMediumIPT@_@WarpMediumBS@" PARENT_SCOPE)
-  elseif(file STREQUAL "benchmark_device_transform")
+@KeyType@_@ValueType@_@RadixBits@_@BlockSize@_@ItemsPerThread@_@WarpSmallLWS@_@WarpSmallIPT@_@WarpSmallBS@_@WarpPartition@_@WarpMediumLWS@_@WarpMediumIPT@_@WarpMediumBS@" PARENT_SCOPE)
+elseif(file STREQUAL "benchmark_device_segmented_reduce")
+    set(list_across_names "DataType;BlockSize;ItemsPerThread" PARENT_SCOPE)
+    set(list_across "\
+${TUNING_TYPES};64 128 256;1 2 4 8 16" PARENT_SCOPE)
+    set(output_pattern_suffix "@DataType@_@BlockSize@_@ItemsPerThread@" PARENT_SCOPE)
+elseif(file STREQUAL "benchmark_device_transform")
     set(list_across_names "\
 DataType;BlockSize;" PARENT_SCOPE)
     set(list_across "${TUNING_TYPES};64 128 256 512 1024" PARENT_SCOPE)
     set(output_pattern_suffix "\
 @DataType@_@BlockSize@" PARENT_SCOPE)
+elseif(file STREQUAL "benchmark_device_transform_pointer")
+    set(list_across_names "\
+DataType;BlockSize;LoadType" PARENT_SCOPE)
+    set(list_across "${TUNING_TYPES};64 128 256 512 1024;rocprim::load_default rocprim::load_nontemporal" PARENT_SCOPE)
+    set(output_pattern_suffix "\
+@DataType@_@BlockSize@_@LoadType@" PARENT_SCOPE)
   elseif(file STREQUAL "benchmark_device_partition")
     set(list_across_names "DataType;BlockSize" PARENT_SCOPE)
     set(list_across "${TUNING_TYPES};128 192 256 384 512" PARENT_SCOPE)
diff --git a/benchmark/benchmark_block_adjacent_difference.cpp b/benchmark/benchmark_block_adjacent_difference.cpp
index 8d237f7..71af4c1 100644
--- a/benchmark/benchmark_block_adjacent_difference.cpp
+++ b/benchmark/benchmark_block_adjacent_difference.cpp
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2022-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2022-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -21,11 +21,8 @@
 // SOFTWARE.
 
 #include "benchmark_utils.hpp"
-// CmdParser
-#include "cmdparser.hpp"
 
-// Google Benchmark
-#include <benchmark/benchmark.h>
+#include "../common/utils_device_ptr.hpp"
 
 // HIP API
 #include <hip/hip_runtime.h>
@@ -34,83 +31,81 @@
 #include <rocprim/block/block_adjacent_difference.hpp>
 #include <rocprim/block/block_load_func.hpp>
 #include <rocprim/block/block_store_func.hpp>
+#include <rocprim/config.hpp>
+#include <rocprim/functional.hpp>
+#include <rocprim/intrinsics/thread.hpp>
+#include <rocprim/types.hpp>
 
-#include <algorithm>
-#include <iostream>
-#include <limits>
+#include <cstddef>
+#include <stdint.h>
 #include <string>
+#include <type_traits>
 #include <vector>
 
-#include <cstdio>
-#include <cstdlib>
-
-#ifndef DEFAULT_N
-const size_t DEFAULT_BYTES = 1024 * 1024 * 128 * 4;
-#endif
-
-namespace rp = rocprim;
-
-template <class Benchmark,
-          unsigned int BlockSize,
-          unsigned int ItemsPerThread,
-          bool         WithTile,
-          typename... Args>
-__global__ __launch_bounds__(BlockSize) void kernel(Args ...args)
+template<typename Benchmark,
+         unsigned int BlockSize,
+         unsigned int ItemsPerThread,
+         bool         WithTile,
+         typename... Args>
+__global__ __launch_bounds__(BlockSize)
+void kernel(Args... args)
 {
     Benchmark::template run<BlockSize, ItemsPerThread, WithTile>(args...);
 }
 
 struct subtract_left
 {
-    template <unsigned int BlockSize, unsigned int ItemsPerThread, bool WithTile, typename T>
-    __device__ static void run(const T* d_input, T* d_output, unsigned int trials)
+    template<unsigned int BlockSize, unsigned int ItemsPerThread, bool WithTile, typename T>
+    __device__
+    static void run(const T* d_input, T* d_output, unsigned int trials)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
-        using adjacent_diff_t = rp::block_adjacent_difference<T, BlockSize>;
+        using adjacent_diff_t = rocprim::block_adjacent_difference<T, BlockSize>;
         __shared__ typename adjacent_diff_t::storage_type storage;
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < trials; trial++)
+        for(unsigned int trial = 0; trial < trials; ++trial)
         {
             T output[ItemsPerThread];
             if(WithTile)
             {
-                adjacent_diff_t().subtract_left(input, output, rp::minus<>{}, T(123), storage);
+                adjacent_diff_t().subtract_left(input, output, rocprim::minus<>{}, T(123), storage);
             }
             else
             {
-                adjacent_diff_t().subtract_left(input, output, rp::minus<>{}, storage);
+                adjacent_diff_t().subtract_left(input, output, rocprim::minus<>{}, storage);
             }
 
             for(unsigned int i = 0; i < ItemsPerThread; ++i)
             {
                 input[i] += output[i];
             }
-            rp::syncthreads();
+            rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct subtract_left_partial
 {
-    template <unsigned int BlockSize, unsigned int ItemsPerThread, bool WithTile, typename T>
-    __device__ static void
+    template<unsigned int BlockSize, unsigned int ItemsPerThread, bool WithTile, typename T>
+    __device__
+    static void
         run(const T* d_input, const unsigned int* tile_sizes, T* d_output, unsigned int trials)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
-        using adjacent_diff_t = rp::block_adjacent_difference<T, BlockSize>;
+        using adjacent_diff_t = rocprim::block_adjacent_difference<T, BlockSize>;
         __shared__ typename adjacent_diff_t::storage_type storage;
 
         unsigned int tile_size = tile_sizes[blockIdx.x];
@@ -119,16 +114,25 @@ struct subtract_left_partial
         const auto tile_size_diff = (BlockSize * ItemsPerThread) / trials + 1;
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < trials; trial++)
+        for(unsigned int trial = 0; trial < trials; ++trial)
         {
             T output[ItemsPerThread];
             if(WithTile)
             {
-                adjacent_diff_t().subtract_left_partial(input, output, rp::minus<>{}, T(123), tile_size, storage);
+                adjacent_diff_t().subtract_left_partial(input,
+                                                        output,
+                                                        rocprim::minus<>{},
+                                                        T(123),
+                                                        tile_size,
+                                                        storage);
             }
             else
             {
-                adjacent_diff_t().subtract_left_partial(input, output, rp::minus<>{}, tile_size, storage);
+                adjacent_diff_t().subtract_left_partial(input,
+                                                        output,
+                                                        rocprim::minus<>{},
+                                                        tile_size,
+                                                        storage);
             }
 
             for(unsigned int i = 0; i < ItemsPerThread; ++i)
@@ -138,66 +142,69 @@ struct subtract_left_partial
 
             // Change the tile_size to even out the distribution
             tile_size = (tile_size + tile_size_diff) % (BlockSize * ItemsPerThread);
-            rp::syncthreads();
+            rocprim::syncthreads();
         }
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct subtract_right
 {
-    template <unsigned int BlockSize,
-              unsigned int ItemsPerThread,
-              bool         WithTile,
-              typename T>
-    __device__ static void run(const T* d_input, T* d_output, unsigned int trials)
+    template<unsigned int BlockSize, unsigned int ItemsPerThread, bool WithTile, typename T>
+    __device__
+    static void run(const T* d_input, T* d_output, unsigned int trials)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
-        using adjacent_diff_t = rp::block_adjacent_difference<T, BlockSize>;
+        using adjacent_diff_t = rocprim::block_adjacent_difference<T, BlockSize>;
         __shared__ typename adjacent_diff_t::storage_type storage;
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < trials; trial++)
+        for(unsigned int trial = 0; trial < trials; ++trial)
         {
             T output[ItemsPerThread];
             if(WithTile)
             {
-                adjacent_diff_t().subtract_right(input, output, rp::minus<>{}, T(123), storage);
+                adjacent_diff_t().subtract_right(input,
+                                                 output,
+                                                 rocprim::minus<>{},
+                                                 T(123),
+                                                 storage);
             }
             else
             {
-                adjacent_diff_t().subtract_right(input, output, rp::minus<>{}, storage);
+                adjacent_diff_t().subtract_right(input, output, rocprim::minus<>{}, storage);
             }
 
             for(unsigned int i = 0; i < ItemsPerThread; ++i)
             {
                 input[i] += output[i];
             }
-            rp::syncthreads();
+            rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct subtract_right_partial
 {
-    template <unsigned int BlockSize, unsigned int ItemsPerThread, bool WithTile, typename T>
-    __device__ static void
+    template<unsigned int BlockSize, unsigned int ItemsPerThread, bool WithTile, typename T>
+    __device__
+    static void
         run(const T* d_input, const unsigned int* tile_sizes, T* d_output, unsigned int trials)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
-        using adjacent_diff_t = rp::block_adjacent_difference<T, BlockSize>;
+        using adjacent_diff_t = rocprim::block_adjacent_difference<T, BlockSize>;
         __shared__ typename adjacent_diff_t::storage_type storage;
 
         unsigned int tile_size = tile_sizes[blockIdx.x];
@@ -205,10 +212,14 @@ struct subtract_right_partial
         const auto tile_size_diff = (BlockSize * ItemsPerThread) / trials + 1;
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < trials; trial++)
+        for(unsigned int trial = 0; trial < trials; ++trial)
         {
             T output[ItemsPerThread];
-            adjacent_diff_t().subtract_right_partial(input, output, rp::minus<>{}, tile_size, storage);
+            adjacent_diff_t().subtract_right_partial(input,
+                                                     output,
+                                                     rocprim::minus<>{},
+                                                     tile_size,
+                                                     storage);
 
             for(unsigned int i = 0; i < ItemsPerThread; ++i)
             {
@@ -216,27 +227,31 @@ struct subtract_right_partial
             }
             // Change the tile_size to even out the distribution
             tile_size = (tile_size + tile_size_diff) % (BlockSize * ItemsPerThread);
-            rp::syncthreads();
+            rocprim::syncthreads();
         }
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
-template<class Benchmark,
-         class T,
+template<typename Benchmark,
+         typename T,
          unsigned int BlockSize,
          unsigned int ItemsPerThread,
          bool         WithTile,
          unsigned int Trials = 100>
-auto run_benchmark(benchmark::State& state, size_t bytes, const managed_seed& seed, hipStream_t stream)
+auto run_benchmark(benchmark_utils::state&& state)
     -> std::enable_if_t<!std::is_same<Benchmark, subtract_left_partial>::value
                         && !std::is_same<Benchmark, subtract_right_partial>::value>
 {
+    const auto& bytes  = state.bytes;
+    const auto& seed   = state.seed;
+    const auto& stream = state.stream;
+
     // Calculate the number of elements N
     size_t N = bytes / sizeof(T);
-    
+
     constexpr auto items_per_block = BlockSize * ItemsPerThread;
-    const auto num_blocks = (N + items_per_block - 1) / items_per_block;
+    const auto     num_blocks      = (N + items_per_block - 1) / items_per_block;
     // Round up size to the next multiple of items_per_block
     const auto size = num_blocks * items_per_block;
 
@@ -244,70 +259,41 @@ auto run_benchmark(benchmark::State& state, size_t bytes, const managed_seed& se
     const std::vector<T> input
         = get_random_data<T>(size, random_range.first, random_range.second, seed.get_0());
 
-    T* d_input;
-    T* d_output;
-    HIP_CHECK(hipMalloc(&d_input, input.size() * sizeof(input[0])));
-    HIP_CHECK(hipMalloc(&d_output, input.size() * sizeof(T)));
-    HIP_CHECK(
-        hipMemcpy(
-            d_input, input.data(),
-            input.size() * sizeof(input[0]),
-            hipMemcpyHostToDevice
-        )
-    );
-
-    // HIP events creation
-    hipEvent_t start, stop;
-    HIP_CHECK(hipEventCreate(&start));
-    HIP_CHECK(hipEventCreate(&stop));
-
-    for(auto _ : state)
-    {
-        // Record start event
-        HIP_CHECK(hipEventRecord(start, stream));
-
-        hipLaunchKernelGGL(
-            HIP_KERNEL_NAME(kernel<Benchmark, BlockSize, ItemsPerThread, WithTile>),
-            dim3(num_blocks), dim3(BlockSize), 0, stream,
-            d_input, d_output, Trials
-        );
-        HIP_CHECK(hipGetLastError());
-
-        // Record stop event and wait until it completes
-        HIP_CHECK(hipEventRecord(stop, stream));
-        HIP_CHECK(hipEventSynchronize(stop));
-
-        float elapsed_mseconds;
-        HIP_CHECK(hipEventElapsedTime(&elapsed_mseconds, start, stop));
-        state.SetIterationTime(elapsed_mseconds / 1000);
-    }
-
-    // Destroy HIP events
-    HIP_CHECK(hipEventDestroy(start));
-    HIP_CHECK(hipEventDestroy(stop));
+    common::device_ptr<T> d_input(input);
+    common::device_ptr<T> d_output(input.size());
 
-    state.SetBytesProcessed(state.iterations() * Trials * size * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * Trials * size);
+    state.run(
+        [&]
+        {
+            kernel<Benchmark, BlockSize, ItemsPerThread, WithTile>
+                <<<dim3(num_blocks), dim3(BlockSize), 0, stream>>>(d_input.get(),
+                                                                   d_output.get(),
+                                                                   Trials);
+            HIP_CHECK(hipGetLastError());
+        });
 
-    HIP_CHECK(hipFree(d_input));
-    HIP_CHECK(hipFree(d_output));
+    state.set_throughput(size * Trials, sizeof(T));
 }
 
-template<class Benchmark,
-         class T,
+template<typename Benchmark,
+         typename T,
          unsigned int BlockSize,
          unsigned int ItemsPerThread,
          bool         WithTile,
          unsigned int Trials = 100>
-auto run_benchmark(benchmark::State& state, size_t bytes, const managed_seed& seed, hipStream_t stream)
+auto run_benchmark(benchmark_utils::state&& state)
     -> std::enable_if_t<std::is_same<Benchmark, subtract_left_partial>::value
                         || std::is_same<Benchmark, subtract_right_partial>::value>
 {
+    const auto& bytes  = state.bytes;
+    const auto& seed   = state.seed;
+    const auto& stream = state.stream;
+
     // Calculate the number of elements N
     size_t N = bytes / sizeof(T);
 
     static constexpr auto items_per_block = BlockSize * ItemsPerThread;
-    const auto num_blocks = (N + items_per_block - 1) / items_per_block;
+    const auto            num_blocks      = (N + items_per_block - 1) / items_per_block;
     // Round up size to the next multiple of items_per_block
     const auto size = num_blocks * items_per_block;
 
@@ -323,171 +309,73 @@ auto run_benchmark(benchmark::State& state, size_t bytes, const managed_seed& se
                                         random_range_tile_sizes.second,
                                         seed.get_1());
 
-    T*            d_input;
-    unsigned int* d_tile_sizes;
-    T*            d_output;
-    HIP_CHECK(hipMalloc(&d_input, input.size() * sizeof(input[0])));
-    HIP_CHECK(hipMalloc(&d_tile_sizes, tile_sizes.size() * sizeof(tile_sizes[0])));
-    HIP_CHECK(hipMalloc(&d_output, input.size() * sizeof(input[0])));
-    HIP_CHECK(
-        hipMemcpy(
-            d_input, input.data(),
-            input.size() * sizeof(input[0]),
-            hipMemcpyHostToDevice
-        )
-    );
-    HIP_CHECK(
-        hipMemcpy(
-            d_tile_sizes, tile_sizes.data(),
-            tile_sizes.size() * sizeof(tile_sizes[0]),
-            hipMemcpyHostToDevice
-        )
-    );
-
-    // HIP events creation
-    hipEvent_t start, stop;
-    HIP_CHECK(hipEventCreate(&start));
-    HIP_CHECK(hipEventCreate(&stop));
-
-    for(auto _ : state)
-    {
-        // Record start event
-        HIP_CHECK(hipEventRecord(start, stream));
-
-        hipLaunchKernelGGL(
-            HIP_KERNEL_NAME(kernel<Benchmark, BlockSize, ItemsPerThread, WithTile>),
-            dim3(num_blocks), dim3(BlockSize), 0, stream,
-            d_input, d_tile_sizes, d_output, Trials
-        );
-        HIP_CHECK(hipGetLastError());
-
-        // Record stop event and wait until it completes
-        HIP_CHECK(hipEventRecord(stop, stream));
-        HIP_CHECK(hipEventSynchronize(stop));
-
-        float elapsed_mseconds;
-        HIP_CHECK(hipEventElapsedTime(&elapsed_mseconds, start, stop));
-        state.SetIterationTime(elapsed_mseconds / 1000);
-    }
+    common::device_ptr<T>            d_input(input);
+    common::device_ptr<unsigned int> d_tile_sizes(tile_sizes);
+    common::device_ptr<T>            d_output(input.size());
 
-    // Destroy HIP events
-    HIP_CHECK(hipEventDestroy(start));
-    HIP_CHECK(hipEventDestroy(stop));
-
-    state.SetBytesProcessed(state.iterations() * Trials * size * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * Trials * size);
+    state.run(
+        [&]
+        {
+            kernel<Benchmark, BlockSize, ItemsPerThread, WithTile>
+                <<<dim3(num_blocks), dim3(BlockSize), 0, stream>>>(d_input.get(),
+                                                                   d_tile_sizes.get(),
+                                                                   d_output.get(),
+                                                                   Trials);
+            HIP_CHECK(hipGetLastError());
+        });
 
-    HIP_CHECK(hipFree(d_input));
-    HIP_CHECK(hipFree(d_tile_sizes));
-    HIP_CHECK(hipFree(d_output));
+    state.set_throughput(size * Trials, sizeof(T));
 }
 
 #define CREATE_BENCHMARK(T, BS, IPT, WITH_TILE)                                         \
-    benchmark::RegisterBenchmark(                                                       \
+    executor.queue_fn(                                                                  \
         bench_naming::format_name("{lvl:block,algo:adjacent_difference,subalgo:" + name \
                                   + ",key_type:" #T ",cfg:{bs:" #BS ",ipt:" #IPT        \
                                     ",with_tile:" #WITH_TILE "}}")                      \
             .c_str(),                                                                   \
-        run_benchmark<Benchmark, T, BS, IPT, WITH_TILE>,                                \
-        bytes,                                                                           \
-        seed,                                                                           \
-        stream)
-
-#define BENCHMARK_TYPE(type, block, with_tile)    \
-    CREATE_BENCHMARK(type, block, 1,  with_tile), \
-    CREATE_BENCHMARK(type, block, 3,  with_tile), \
-    CREATE_BENCHMARK(type, block, 4,  with_tile), \
-    CREATE_BENCHMARK(type, block, 8,  with_tile), \
-    CREATE_BENCHMARK(type, block, 16, with_tile), \
+        run_benchmark<Benchmark, T, BS, IPT, WITH_TILE>);
+
+#define BENCHMARK_TYPE(type, block, with_tile)   \
+    CREATE_BENCHMARK(type, block, 1, with_tile)  \
+    CREATE_BENCHMARK(type, block, 3, with_tile)  \
+    CREATE_BENCHMARK(type, block, 4, with_tile)  \
+    CREATE_BENCHMARK(type, block, 8, with_tile)  \
+    CREATE_BENCHMARK(type, block, 16, with_tile) \
     CREATE_BENCHMARK(type, block, 32, with_tile)
 
-template<class Benchmark>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    size_t                                        bytes,
-                    const managed_seed&                           seed,
-                    hipStream_t                                   stream)
+template<typename Benchmark>
+void add_benchmarks(const std::string& name, benchmark_utils::executor& executor)
 {
-    std::vector<benchmark::internal::Benchmark*> bs =
+    BENCHMARK_TYPE(int, 256, false)
+    BENCHMARK_TYPE(float, 256, false)
+    BENCHMARK_TYPE(int8_t, 256, false)
+    BENCHMARK_TYPE(rocprim::half, 256, false)
+    BENCHMARK_TYPE(long long, 256, false)
+    BENCHMARK_TYPE(double, 256, false)
+    BENCHMARK_TYPE(rocprim::int128_t, 256, false)
+    BENCHMARK_TYPE(rocprim::uint128_t, 256, false)
+
+    if(!std::is_same<Benchmark, subtract_right_partial>::value)
     {
-        BENCHMARK_TYPE(int, 256, false),
-        BENCHMARK_TYPE(float, 256, false),
-        BENCHMARK_TYPE(int8_t, 256, false),
-        BENCHMARK_TYPE(rocprim::half, 256, false),
-        BENCHMARK_TYPE(long long, 256, false),
-        BENCHMARK_TYPE(double, 256, false)
-    };
-
-    if(!std::is_same<Benchmark, subtract_right_partial>::value) {
-        bs.insert(bs.end(), {
-            BENCHMARK_TYPE(int, 256, true),
-            BENCHMARK_TYPE(float, 256, true),
-            BENCHMARK_TYPE(int8_t, 256, true),
-            BENCHMARK_TYPE(rocprim::half, 256, true),
-            BENCHMARK_TYPE(long long, 256, true),
-            BENCHMARK_TYPE(double, 256, true)
-        });
+        BENCHMARK_TYPE(int, 256, true)
+        BENCHMARK_TYPE(float, 256, true)
+        BENCHMARK_TYPE(int8_t, 256, true)
+        BENCHMARK_TYPE(rocprim::half, 256, true)
+        BENCHMARK_TYPE(long long, 256, true)
+        BENCHMARK_TYPE(double, 256, true)
+        BENCHMARK_TYPE(rocprim::int128_t, 256, true)
+        BENCHMARK_TYPE(rocprim::uint128_t, 256, true)
     }
-
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
 }
 
-int main(int argc, char *argv[])
+int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<size_t>("size", "size", DEFAULT_BYTES, "number of bytes");
-    parser.set_optional<int>("trials", "trials", -1, "number of iterations");
-    parser.set_optional<std::string>("name_format",
-                                     "name_format",
-                                     "human",
-                                     "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    const size_t bytes = parser.get<size_t>("size");
-    const int trials = parser.get<int>("trials");
-    bench_naming::set_format(parser.get<std::string>("name_format"));
-    const std::string  seed_type = parser.get<std::string>("seed");
-    const managed_seed seed(seed_type);
-
-    // HIP
-    hipStream_t stream = 0; // default
-
-    // Benchmark info
-    add_common_benchmark_info();
-    benchmark::AddCustomContext("bytes", std::to_string(bytes));
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmarks
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<subtract_left>("subtract_left", benchmarks, bytes, seed, stream);
-    add_benchmarks<subtract_right>("subtract_right", benchmarks, bytes, seed, stream);
-    add_benchmarks<subtract_left_partial>("subtract_left_partial", benchmarks, bytes, seed, stream);
-    add_benchmarks<subtract_right_partial>("subtract_right_partial",
-                                           benchmarks,
-                                           bytes,
-                                           seed,
-                                           stream);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMillisecond);
-    }
+    benchmark_utils::executor executor(argc, argv, 512 * benchmark_utils::MiB, 1, 0);
 
-    // Force number of iterations
-    if(trials > 0)
-    {
-        for(auto& b : benchmarks)
-        {
-            b->Iterations(trials);
-        }
-    }
+    add_benchmarks<subtract_left>("subtract_left", executor);
+    add_benchmarks<subtract_right>("subtract_right", executor);
+    add_benchmarks<subtract_left_partial>("subtract_left_partial", executor);
+    add_benchmarks<subtract_right_partial>("subtract_right_partial", executor);
 
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks();
-    return 0;
+    executor.run();
 }
diff --git a/benchmark/benchmark_block_discontinuity.cpp b/benchmark/benchmark_block_discontinuity.cpp
index 1d07cdb..62fd01a 100644
--- a/benchmark/benchmark_block_discontinuity.cpp
+++ b/benchmark/benchmark_block_discontinuity.cpp
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -20,12 +20,7 @@
 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 // SOFTWARE.
 
-// CmdParser
 #include "benchmark_utils.hpp"
-#include "cmdparser.hpp"
-
-// Google Benchmark
-#include <benchmark/benchmark.h>
 
 // HIP API
 #include <hip/hip_runtime.h>
@@ -34,327 +29,243 @@
 #include <rocprim/block/block_discontinuity.hpp>
 #include <rocprim/block/block_load_func.hpp>
 #include <rocprim/block/block_store_func.hpp>
+#include <rocprim/config.hpp>
+#include <rocprim/functional.hpp>
+#include <rocprim/intrinsics/thread.hpp>
+#include <rocprim/types.hpp>
 
-#include <iostream>
-#include <limits>
+#include <cstddef>
+#include <stdint.h>
 #include <string>
 #include <vector>
 
-#include <cstdio>
-#include <cstdlib>
-
-#ifndef DEFAULT_N
-const size_t DEFAULT_BYTES = 1024 * 1024 * 128 * 4;
-#endif
-
-namespace rp = rocprim;
-
-template<
-    class Runner,
-    class T,
-    unsigned int BlockSize,
-    unsigned int ItemsPerThread,
-    bool WithTile,
-    unsigned int Trials
->
-__global__
-__launch_bounds__(BlockSize)
-void kernel(const T * d_input, T * d_output)
+template<typename Runner,
+         typename T,
+         unsigned int BlockSize,
+         unsigned int ItemsPerThread,
+         bool         WithTile,
+         unsigned int Trials>
+__global__ __launch_bounds__(BlockSize)
+void kernel(const T* d_input, T* d_output)
 {
     Runner::template run<T, BlockSize, ItemsPerThread, WithTile, Trials>(d_input, d_output);
 }
 
 struct flag_heads
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        bool WithTile,
-        unsigned int Trials
-    >
+    template<typename T,
+             unsigned int BlockSize,
+             unsigned int ItemsPerThread,
+             bool         WithTile,
+             unsigned int Trials>
     __device__
-    static void run(const T * d_input, T * d_output)
+    static void run(const T* d_input, T* d_output)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
-            rp::block_discontinuity<T, BlockSize> bdiscontinuity;
-            bool head_flags[ItemsPerThread];
+            rocprim::block_discontinuity<T, BlockSize> bdiscontinuity;
+            bool                                       head_flags[ItemsPerThread];
             if(WithTile)
             {
-                bdiscontinuity.flag_heads(head_flags, T(123), input, rp::equal_to<T>());
+                bdiscontinuity.flag_heads(head_flags, T(123), input, rocprim::equal_to<T>());
             }
             else
             {
-                bdiscontinuity.flag_heads(head_flags, input, rp::equal_to<T>());
+                bdiscontinuity.flag_heads(head_flags, input, rocprim::equal_to<T>());
             }
 
-            for(unsigned int i = 0; i < ItemsPerThread; i++)
+            for(unsigned int i = 0; i < ItemsPerThread; ++i)
             {
                 input[i] += head_flags[i];
             }
-            rp::syncthreads();
+            rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct flag_tails
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        bool WithTile,
-        unsigned int Trials
-    >
+    template<typename T,
+             unsigned int BlockSize,
+             unsigned int ItemsPerThread,
+             bool         WithTile,
+             unsigned int Trials>
     __device__
-    static void run(const T * d_input, T * d_output)
+    static void run(const T* d_input, T* d_output)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
-            rp::block_discontinuity<T, BlockSize> bdiscontinuity;
-            bool tail_flags[ItemsPerThread];
+            rocprim::block_discontinuity<T, BlockSize> bdiscontinuity;
+            bool                                       tail_flags[ItemsPerThread];
             if(WithTile)
             {
-                bdiscontinuity.flag_tails(tail_flags, T(123), input, rp::equal_to<T>());
+                bdiscontinuity.flag_tails(tail_flags, T(123), input, rocprim::equal_to<T>());
             }
             else
             {
-                bdiscontinuity.flag_tails(tail_flags, input, rp::equal_to<T>());
+                bdiscontinuity.flag_tails(tail_flags, input, rocprim::equal_to<T>());
             }
 
-            for(unsigned int i = 0; i < ItemsPerThread; i++)
+            for(unsigned int i = 0; i < ItemsPerThread; ++i)
             {
                 input[i] += tail_flags[i];
             }
-            rp::syncthreads();
+            rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct flag_heads_and_tails
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        bool WithTile,
-        unsigned int Trials
-    >
+    template<typename T,
+             unsigned int BlockSize,
+             unsigned int ItemsPerThread,
+             bool         WithTile,
+             unsigned int Trials>
     __device__
-    static void run(const T * d_input, T * d_output)
+    static void run(const T* d_input, T* d_output)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
-            rp::block_discontinuity<T, BlockSize> bdiscontinuity;
-            bool head_flags[ItemsPerThread];
-            bool tail_flags[ItemsPerThread];
+            rocprim::block_discontinuity<T, BlockSize> bdiscontinuity;
+            bool                                       head_flags[ItemsPerThread];
+            bool                                       tail_flags[ItemsPerThread];
             if(WithTile)
             {
-                bdiscontinuity.flag_heads_and_tails(head_flags, T(123), tail_flags, T(234), input, rp::equal_to<T>());
+                bdiscontinuity.flag_heads_and_tails(head_flags,
+                                                    T(123),
+                                                    tail_flags,
+                                                    T(234),
+                                                    input,
+                                                    rocprim::equal_to<T>());
             }
             else
             {
-                bdiscontinuity.flag_heads_and_tails(head_flags, tail_flags, input, rp::equal_to<T>());
+                bdiscontinuity.flag_heads_and_tails(head_flags,
+                                                    tail_flags,
+                                                    input,
+                                                    rocprim::equal_to<T>());
             }
 
-            for(unsigned int i = 0; i < ItemsPerThread; i++)
+            for(unsigned int i = 0; i < ItemsPerThread; ++i)
             {
                 input[i] += head_flags[i];
                 input[i] += tail_flags[i];
             }
-            rp::syncthreads();
+            rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
-template<class Benchmark,
-         class T,
+template<typename Benchmark,
+         typename T,
          unsigned int BlockSize,
          unsigned int ItemsPerThread,
          bool         WithTile,
          unsigned int Trials = 100>
-void run_benchmark(benchmark::State& state, size_t bytes, const managed_seed& seed, hipStream_t stream)
+void run_benchmark(benchmark_utils::state&& state)
 {
+    const auto& bytes  = state.bytes;
+    const auto& seed   = state.seed;
+    const auto& stream = state.stream;
+
     // Calculate the number of elements N
     size_t N = bytes / sizeof(T);
 
     constexpr auto items_per_block = BlockSize * ItemsPerThread;
-    const auto size = items_per_block * ((N + items_per_block - 1)/items_per_block);
+    const auto     size = items_per_block * ((N + items_per_block - 1) / items_per_block);
 
     const auto     random_range = limit_random_range<T>(0, 10);
     std::vector<T> input
         = get_random_data<T>(size, random_range.first, random_range.second, seed.get_0());
-    T * d_input;
-    T * d_output;
+    T* d_input;
+    T* d_output;
     HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_input), size * sizeof(T)));
     HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_output), size * sizeof(T)));
-    HIP_CHECK(
-        hipMemcpy(
-            d_input, input.data(),
-            size * sizeof(T),
-            hipMemcpyHostToDevice
-        )
-    );
+    HIP_CHECK(hipMemcpy(d_input, input.data(), size * sizeof(T), hipMemcpyHostToDevice));
     HIP_CHECK(hipDeviceSynchronize());
 
-    // HIP events creation
-    hipEvent_t start, stop;
-    HIP_CHECK(hipEventCreate(&start));
-    HIP_CHECK(hipEventCreate(&stop));
-
-    for(auto _ : state)
-    {
-        // Record start event
-        HIP_CHECK(hipEventRecord(start, stream));
-
-        hipLaunchKernelGGL(
-            HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, WithTile, Trials>),
-            dim3(size/items_per_block), dim3(BlockSize), 0, stream,
-            d_input, d_output
-        );
-        HIP_CHECK(hipGetLastError());
-
-        // Record stop event and wait until it completes
-        HIP_CHECK(hipEventRecord(stop, stream));
-        HIP_CHECK(hipEventSynchronize(stop));
-
-        float elapsed_mseconds;
-        HIP_CHECK(hipEventElapsedTime(&elapsed_mseconds, start, stop));
-        state.SetIterationTime(elapsed_mseconds / 1000);
-    }
-
-    // Destroy HIP events
-    HIP_CHECK(hipEventDestroy(start));
-    HIP_CHECK(hipEventDestroy(stop));
+    state.run(
+        [&]
+        {
+            kernel<Benchmark, T, BlockSize, ItemsPerThread, WithTile, Trials>
+                <<<dim3(size / items_per_block), dim3(BlockSize), 0, stream>>>(d_input, d_output);
+            HIP_CHECK(hipGetLastError());
+        });
 
-    state.SetBytesProcessed(state.iterations() * Trials * size * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * Trials * size);
+    state.set_throughput(size * Trials, sizeof(T));
 
     HIP_CHECK(hipFree(d_input));
     HIP_CHECK(hipFree(d_output));
 }
 
-#define CREATE_BENCHMARK(T, BS, IPT, WITH_TILE)                                   \
-    benchmark::RegisterBenchmark(                                                 \
-        bench_naming::format_name("{lvl:block,algo:discontinuity,subalgo:" + name \
-                                  + ",key_type:" #T ",cfg:{bs:" #BS ",ipt:" #IPT  \
-                                    ",with_tile:" #WITH_TILE "}}")                \
-            .c_str(),                                                             \
-        run_benchmark<Benchmark, T, BS, IPT, WITH_TILE>,                          \
-        bytes,                                                                     \
-        seed,                                                                     \
-        stream)
-
-#define BENCHMARK_TYPE(type, block, bool) \
-    CREATE_BENCHMARK(type, block, 1, bool), \
-    CREATE_BENCHMARK(type, block, 2, bool), \
-    CREATE_BENCHMARK(type, block, 3, bool), \
-    CREATE_BENCHMARK(type, block, 4, bool), \
+#define CREATE_BENCHMARK(T, BS, IPT, WITH_TILE)                                                 \
+    executor.queue_fn(bench_naming::format_name("{lvl:block,algo:discontinuity,subalgo:" + name \
+                                                + ",key_type:" #T ",cfg:{bs:" #BS ",ipt:" #IPT  \
+                                                  ",with_tile:" #WITH_TILE "}}")                \
+                          .c_str(),                                                             \
+                      run_benchmark<Benchmark, T, BS, IPT, WITH_TILE>);
+
+#define BENCHMARK_TYPE(type, block, bool)  \
+    CREATE_BENCHMARK(type, block, 1, bool) \
+    CREATE_BENCHMARK(type, block, 2, bool) \
+    CREATE_BENCHMARK(type, block, 3, bool) \
+    CREATE_BENCHMARK(type, block, 4, bool) \
     CREATE_BENCHMARK(type, block, 8, bool)
 
-template<class Benchmark>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    size_t                                        bytes,
-                    const managed_seed&                           seed,
-                    hipStream_t                                   stream)
+template<typename Benchmark>
+void add_benchmarks(const std::string& name, benchmark_utils::executor& executor)
 {
-    std::vector<benchmark::internal::Benchmark*> bs =
-    {
-        BENCHMARK_TYPE(int, 256, false),
-        BENCHMARK_TYPE(int, 256, true),
-        BENCHMARK_TYPE(int8_t, 256, false),
-        BENCHMARK_TYPE(int8_t, 256, true),
-        BENCHMARK_TYPE(uint8_t, 256, false),
-        BENCHMARK_TYPE(uint8_t, 256, true),
-        BENCHMARK_TYPE(rocprim::half, 256, false),
-        BENCHMARK_TYPE(rocprim::half, 256, true),
-        BENCHMARK_TYPE(long long, 256, false),
-        BENCHMARK_TYPE(long long, 256, true),
-    };
-
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+    BENCHMARK_TYPE(int, 256, false)
+    BENCHMARK_TYPE(int, 256, true)
+    BENCHMARK_TYPE(int8_t, 256, false)
+    BENCHMARK_TYPE(int8_t, 256, true)
+    BENCHMARK_TYPE(uint8_t, 256, false)
+    BENCHMARK_TYPE(uint8_t, 256, true)
+    BENCHMARK_TYPE(rocprim::half, 256, false)
+    BENCHMARK_TYPE(rocprim::half, 256, true)
+    BENCHMARK_TYPE(long long, 256, false)
+    BENCHMARK_TYPE(long long, 256, true)
+    BENCHMARK_TYPE(rocprim::int128_t, 256, false)
+    BENCHMARK_TYPE(rocprim::int128_t, 256, true)
+    BENCHMARK_TYPE(rocprim::uint128_t, 256, false)
+    BENCHMARK_TYPE(rocprim::uint128_t, 256, true)
 }
 
-int main(int argc, char *argv[])
+int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<size_t>("size", "size", DEFAULT_BYTES, "number of bytes");
-    parser.set_optional<int>("trials", "trials", -1, "number of iterations");
-    parser.set_optional<std::string>("name_format",
-                                     "name_format",
-                                     "human",
-                                     "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    const size_t bytes = parser.get<size_t>("size");
-    const int trials = parser.get<int>("trials");
-    bench_naming::set_format(parser.get<std::string>("name_format"));
-    const std::string  seed_type = parser.get<std::string>("seed");
-    const managed_seed seed(seed_type);
-
-    // HIP
-    hipStream_t stream = 0; // default
-
-    // Benchmark info
-    add_common_benchmark_info();
-    benchmark::AddCustomContext("bytes", std::to_string(bytes));
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmarks
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<flag_heads>("flag_heads", benchmarks, bytes, seed, stream);
-    add_benchmarks<flag_tails>("flag_tails", benchmarks, bytes, seed, stream);
-    add_benchmarks<flag_heads_and_tails>("flag_heads_and_tails", benchmarks, bytes, seed, stream);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMillisecond);
-    }
+    benchmark_utils::executor executor(argc, argv, 512 * benchmark_utils::MiB, 1, 0);
 
-    // Force number of iterations
-    if(trials > 0)
-    {
-        for(auto& b : benchmarks)
-        {
-            b->Iterations(trials);
-        }
-    }
+    add_benchmarks<flag_heads>("flag_heads", executor);
+    add_benchmarks<flag_tails>("flag_tails", executor);
+    add_benchmarks<flag_heads_and_tails>("flag_heads_and_tails", executor);
 
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks();
-    return 0;
+    executor.run();
 }
diff --git a/benchmark/benchmark_block_exchange.cpp b/benchmark/benchmark_block_exchange.cpp
index ec3c95f..889fc78 100644
--- a/benchmark/benchmark_block_exchange.cpp
+++ b/benchmark/benchmark_block_exchange.cpp
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -20,12 +20,10 @@
 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 // SOFTWARE.
 
-// CmdParser
-#include "cmdparser.hpp"
 #include "benchmark_utils.hpp"
 
-// Google Benchmark
-#include <benchmark/benchmark.h>
+#include "../common/utils_custom_type.hpp"
+#include "../common/utils_device_ptr.hpp"
 
 // HIP API
 #include <hip/hip_runtime.h>
@@ -34,407 +32,270 @@
 #include <rocprim/block/block_exchange.hpp>
 #include <rocprim/block/block_load_func.hpp>
 #include <rocprim/block/block_store_func.hpp>
-
-#include <iostream>
-#include <limits>
+#include <rocprim/config.hpp>
+#include <rocprim/intrinsics/thread.hpp>
+#include <rocprim/types.hpp>
+
+#include <algorithm>
+#include <cstddef>
+#include <numeric>
+#include <stdint.h>
 #include <string>
 #include <vector>
 
-#include <cstdio>
-#include <cstdlib>
-
-#ifndef DEFAULT_N
-const size_t DEFAULT_BYTES = 1024 * 1024 * 32 * 4;
-#endif
-
-namespace rp = rocprim;
-
-template<
-    class Runner,
-    class T,
-    unsigned int BlockSize,
-    unsigned int ItemsPerThread,
-    unsigned int Trials
->
-__global__
-__launch_bounds__(BlockSize)
-void kernel(const T * d_input, const unsigned int * d_ranks, T * d_output)
+template<typename Runner,
+         typename T,
+         unsigned int BlockSize,
+         unsigned int ItemsPerThread,
+         unsigned int Trials>
+__global__ __launch_bounds__(BlockSize)
+void kernel(const T* d_input, const unsigned int* d_ranks, T* d_output)
 {
     Runner::template run<T, BlockSize, ItemsPerThread, Trials>(d_input, d_ranks, d_output);
 }
 
 struct blocked_to_striped
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        unsigned int Trials
-    >
+    template<typename T, unsigned int BlockSize, unsigned int ItemsPerThread, unsigned int Trials>
     __device__
-    static void run(const T * d_input, const unsigned int *, T * d_output)
+    static void run(const T* d_input, const unsigned int*, T* d_output)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
-            rp::block_exchange<T, BlockSize, ItemsPerThread> exchange;
+            rocprim::block_exchange<T, BlockSize, ItemsPerThread> exchange;
             exchange.blocked_to_striped(input, input);
             ::rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct striped_to_blocked
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        unsigned int Trials
-    >
+    template<typename T, unsigned int BlockSize, unsigned int ItemsPerThread, unsigned int Trials>
     __device__
-    static void run(const T * d_input, const unsigned int *, T * d_output)
+    static void run(const T* d_input, const unsigned int*, T* d_output)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
-            rp::block_exchange<T, BlockSize, ItemsPerThread> exchange;
+            rocprim::block_exchange<T, BlockSize, ItemsPerThread> exchange;
             exchange.striped_to_blocked(input, input);
             ::rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct blocked_to_warp_striped
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        unsigned int Trials
-    >
+    template<typename T, unsigned int BlockSize, unsigned int ItemsPerThread, unsigned int Trials>
     __device__
-    static void run(const T * d_input, const unsigned int *, T * d_output)
+    static void run(const T* d_input, const unsigned int*, T* d_output)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
-            rp::block_exchange<T, BlockSize, ItemsPerThread> exchange;
+            rocprim::block_exchange<T, BlockSize, ItemsPerThread> exchange;
             exchange.blocked_to_warp_striped(input, input);
             ::rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct warp_striped_to_blocked
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        unsigned int Trials
-    >
+    template<typename T, unsigned int BlockSize, unsigned int ItemsPerThread, unsigned int Trials>
     __device__
-    static void run(const T * d_input, const unsigned int *, T * d_output)
+    static void run(const T* d_input, const unsigned int*, T* d_output)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
         T input[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
-            rp::block_exchange<T, BlockSize, ItemsPerThread> exchange;
+            rocprim::block_exchange<T, BlockSize, ItemsPerThread> exchange;
             exchange.warp_striped_to_blocked(input, input);
             ::rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct scatter_to_blocked
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        unsigned int Trials
-    >
+    template<typename T, unsigned int BlockSize, unsigned int ItemsPerThread, unsigned int Trials>
     __device__
-    static void run(const T * d_input, const unsigned int * d_ranks, T * d_output)
+    static void run(const T* d_input, const unsigned int* d_ranks, T* d_output)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
-        T input[ItemsPerThread];
+        T            input[ItemsPerThread];
         unsigned int ranks[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
-        rp::block_load_direct_striped<BlockSize>(lid, d_ranks + block_offset, ranks);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_ranks + block_offset, ranks);
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
-            rp::block_exchange<T, BlockSize, ItemsPerThread> exchange;
+            rocprim::block_exchange<T, BlockSize, ItemsPerThread> exchange;
             exchange.scatter_to_blocked(input, input, ranks);
             ::rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
 struct scatter_to_striped
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        unsigned int Trials
-    >
+    template<typename T, unsigned int BlockSize, unsigned int ItemsPerThread, unsigned int Trials>
     __device__
-    static void run(const T * d_input, const unsigned int * d_ranks, T * d_output)
+    static void run(const T* d_input, const unsigned int* d_ranks, T* d_output)
     {
-        const unsigned int lid = threadIdx.x;
+        const unsigned int lid          = threadIdx.x;
         const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
-        T input[ItemsPerThread];
+        T            input[ItemsPerThread];
         unsigned int ranks[ItemsPerThread];
-        rp::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
-        rp::block_load_direct_striped<BlockSize>(lid, d_ranks + block_offset, ranks);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_input + block_offset, input);
+        rocprim::block_load_direct_striped<BlockSize>(lid, d_ranks + block_offset, ranks);
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
-            rp::block_exchange<T, BlockSize, ItemsPerThread> exchange;
+            rocprim::block_exchange<T, BlockSize, ItemsPerThread> exchange;
             exchange.scatter_to_striped(input, input, ranks);
             ::rocprim::syncthreads();
         }
 
-        rp::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
+        rocprim::block_store_direct_striped<BlockSize>(lid, d_output + block_offset, input);
     }
 };
 
-template<class Benchmark,
-         class T,
+template<typename Benchmark,
+         typename T,
          unsigned int BlockSize,
          unsigned int ItemsPerThread,
          unsigned int Trials = 100>
-void run_benchmark(benchmark::State& state, size_t bytes, const managed_seed& seed, hipStream_t stream)
+void run_benchmark(benchmark_utils::state&& state)
 {
+    const auto& bytes  = state.bytes;
+    const auto& seed   = state.seed;
+    const auto& stream = state.stream;
+
     // Calculate the number of elements N
     size_t N = bytes / sizeof(T);
-    
+
     constexpr auto items_per_block = BlockSize * ItemsPerThread;
-    const auto size = items_per_block * ((N + items_per_block - 1)/items_per_block);
+    const auto     size = items_per_block * ((N + items_per_block - 1) / items_per_block);
 
     std::vector<T> input(size);
     // Fill input
-    for(size_t i = 0; i < size; i++)
+    for(size_t i = 0; i < size; ++i)
     {
         input[i] = T(i);
     }
     std::vector<unsigned int> ranks(size);
     // Fill ranks (for scatter operations)
     engine_type gen(seed.get_0());
-    for(size_t bi = 0; bi < size / items_per_block; bi++)
+    for(size_t bi = 0; bi < size / items_per_block; ++bi)
     {
         auto block_ranks = ranks.begin() + bi * items_per_block;
         std::iota(block_ranks, block_ranks + items_per_block, 0);
         std::shuffle(block_ranks, block_ranks + items_per_block, gen);
     }
-    T * d_input;
-    unsigned int * d_ranks;
-    T * d_output;
-    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_input), size * sizeof(T)));
-    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_ranks), size * sizeof(unsigned int)));
-    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_output), size * sizeof(T)));
-    HIP_CHECK(
-        hipMemcpy(
-            d_input, input.data(),
-            size * sizeof(T),
-            hipMemcpyHostToDevice
-        )
-    );
-    HIP_CHECK(
-        hipMemcpy(
-            d_ranks, ranks.data(),
-            size * sizeof(unsigned int),
-            hipMemcpyHostToDevice
-        )
-    );
+    common::device_ptr<T>            d_input(input);
+    common::device_ptr<unsigned int> d_ranks(ranks);
+    common::device_ptr<T>            d_output(size);
     HIP_CHECK(hipDeviceSynchronize());
 
-    // HIP events creation
-    hipEvent_t start, stop;
-    HIP_CHECK(hipEventCreate(&start));
-    HIP_CHECK(hipEventCreate(&stop));
-
-    for(auto _ : state)
-    {
-        // Record start event
-        HIP_CHECK(hipEventRecord(start, stream));
-
-        hipLaunchKernelGGL(
-            HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>),
-            dim3(size/items_per_block), dim3(BlockSize), 0, stream,
-            d_input, d_ranks, d_output
-        );
-        HIP_CHECK(hipGetLastError());
-
-        // Record stop event and wait until it completes
-        HIP_CHECK(hipEventRecord(stop, stream));
-        HIP_CHECK(hipEventSynchronize(stop));
-
-        float elapsed_mseconds;
-        HIP_CHECK(hipEventElapsedTime(&elapsed_mseconds, start, stop));
-        state.SetIterationTime(elapsed_mseconds / 1000);
-    }
-
-    // Destroy HIP events
-    HIP_CHECK(hipEventDestroy(start));
-    HIP_CHECK(hipEventDestroy(stop));
-
-    state.SetBytesProcessed(state.iterations() * Trials * size * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * Trials * size);
-
-    HIP_CHECK(hipFree(d_input));
-    HIP_CHECK(hipFree(d_ranks));
-    HIP_CHECK(hipFree(d_output));
+    state.run(
+        [&]
+        {
+            kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>
+                <<<dim3(size / items_per_block), dim3(BlockSize), 0, stream>>>(d_input.get(),
+                                                                               d_ranks.get(),
+                                                                               d_output.get());
+            HIP_CHECK(hipGetLastError());
+        });
+
+    state.set_throughput(size * Trials, sizeof(T));
 }
 
-#define CREATE_BENCHMARK(T, BS, IPT)                                                   \
-    benchmark::RegisterBenchmark(                                                      \
-        bench_naming::format_name("{lvl:block,algo:exchange,subalgo:" + name           \
-                                  + ",key_type:" #T ",cfg:{bs:" #BS ",ipt:" #IPT "}}") \
-            .c_str(),                                                                  \
-        run_benchmark<Benchmark, T, BS, IPT>,                                          \
-        bytes,                                                                          \
-        seed,                                                                          \
-        stream)
-
-#define BENCHMARK_TYPE(type, block) \
-    CREATE_BENCHMARK(type, block, 1), \
-    CREATE_BENCHMARK(type, block, 2), \
-    CREATE_BENCHMARK(type, block, 3), \
-    CREATE_BENCHMARK(type, block, 4), \
-    CREATE_BENCHMARK(type, block, 7), \
+#define CREATE_BENCHMARK(T, BS, IPT)                                                           \
+    executor.queue_fn(bench_naming::format_name("{lvl:block,algo:exchange,subalgo:" + name     \
+                                                + ",key_type:" #T ",cfg:{bs:" #BS ",ipt:" #IPT \
+                                                  "}}")                                        \
+                          .c_str(),                                                            \
+                      run_benchmark<Benchmark, T, BS, IPT>);
+
+#define BENCHMARK_TYPE(type, block)  \
+    CREATE_BENCHMARK(type, block, 1) \
+    CREATE_BENCHMARK(type, block, 2) \
+    CREATE_BENCHMARK(type, block, 3) \
+    CREATE_BENCHMARK(type, block, 4) \
+    CREATE_BENCHMARK(type, block, 7) \
     CREATE_BENCHMARK(type, block, 8)
 
-template<class Benchmark>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    size_t                                        bytes,
-                    const managed_seed&                           seed,
-                    hipStream_t                                   stream)
+template<typename Benchmark>
+void add_benchmarks(const std::string& name, benchmark_utils::executor& executor)
 {
-    using custom_float2 = custom_type<float, float>;
-    using custom_double2 = custom_type<double, double>;
-
-    std::vector<benchmark::internal::Benchmark*> bs =
-    {
-        BENCHMARK_TYPE(int, 256),
-        BENCHMARK_TYPE(int8_t, 256),
-        BENCHMARK_TYPE(rocprim::half, 256),
-        BENCHMARK_TYPE(long long, 256),
-        BENCHMARK_TYPE(custom_float2, 256),
-        BENCHMARK_TYPE(float2, 256),
-        BENCHMARK_TYPE(custom_double2, 256),
-        BENCHMARK_TYPE(double2, 256),
-        BENCHMARK_TYPE(float4, 256),
-    };
-
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+    using custom_float2  = common::custom_type<float, float>;
+    using custom_double2 = common::custom_type<double, double>;
+
+    BENCHMARK_TYPE(int, 256)
+    BENCHMARK_TYPE(int8_t, 256)
+    BENCHMARK_TYPE(rocprim::half, 256)
+    BENCHMARK_TYPE(long long, 256)
+    BENCHMARK_TYPE(custom_float2, 256)
+    BENCHMARK_TYPE(float2, 256)
+    BENCHMARK_TYPE(custom_double2, 256)
+    BENCHMARK_TYPE(double2, 256)
+    BENCHMARK_TYPE(float4, 256)
+    BENCHMARK_TYPE(rocprim::int128_t, 256)
+    BENCHMARK_TYPE(rocprim::uint128_t, 256)
 }
 
-int main(int argc, char *argv[])
+int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<size_t>("bytes", "bytes", DEFAULT_BYTES, "number of values");
-    parser.set_optional<int>("trials", "trials", -1, "number of iterations");
-    parser.set_optional<std::string>("name_format",
-                                     "name_format",
-                                     "human",
-                                     "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    const size_t bytes = parser.get<size_t>("bytes");
-    const int trials = parser.get<int>("trials");
-    bench_naming::set_format(parser.get<std::string>("name_format"));
-    const std::string  seed_type = parser.get<std::string>("seed");
-    const managed_seed seed(seed_type);
-
-    // HIP
-    hipStream_t stream = 0; // default
-
-    // Benchmark info
-    add_common_benchmark_info();
-    benchmark::AddCustomContext("bytes", std::to_string(bytes));
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmarks
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<blocked_to_striped>("blocked_to_striped", benchmarks, bytes, seed, stream);
-    add_benchmarks<striped_to_blocked>("striped_to_blocked", benchmarks, bytes, seed, stream);
-    add_benchmarks<blocked_to_warp_striped>("blocked_to_warp_striped",
-                                            benchmarks,
-                                            bytes,
-                                            seed,
-                                            stream);
-    add_benchmarks<warp_striped_to_blocked>("warp_striped_to_blocked",
-                                            benchmarks,
-                                            bytes,
-                                            seed,
-                                            stream);
-    add_benchmarks<scatter_to_blocked>("scatter_to_blocked", benchmarks, bytes, seed, stream);
-    add_benchmarks<scatter_to_striped>("scatter_to_striped", benchmarks, bytes, seed, stream);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMillisecond);
-    }
+    benchmark_utils::executor executor(argc, argv, 128 * benchmark_utils::MiB, 1, 0);
 
-    // Force number of iterations
-    if(trials > 0)
-    {
-        for(auto& b : benchmarks)
-        {
-            b->Iterations(trials);
-        }
-    }
+    add_benchmarks<blocked_to_striped>("blocked_to_striped", executor);
+    add_benchmarks<striped_to_blocked>("striped_to_blocked", executor);
+    add_benchmarks<blocked_to_warp_striped>("blocked_to_warp_striped", executor);
+    add_benchmarks<warp_striped_to_blocked>("warp_striped_to_blocked", executor);
+    add_benchmarks<scatter_to_blocked>("scatter_to_blocked", executor);
+    add_benchmarks<scatter_to_striped>("scatter_to_striped", executor);
 
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks();
-    return 0;
+    executor.run();
 }
diff --git a/benchmark/benchmark_block_histogram.cpp b/benchmark/benchmark_block_histogram.cpp
index 7dc2827..b108314 100644
--- a/benchmark/benchmark_block_histogram.cpp
+++ b/benchmark/benchmark_block_histogram.cpp
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -21,44 +21,24 @@
 // SOFTWARE.
 
 #include "benchmark_utils.hpp"
-// CmdParser
-#include "cmdparser.hpp"
 
-// Google Benchmark
-#include <benchmark/benchmark.h>
+#include "../common/utils_device_ptr.hpp"
 
 // HIP API
 #include <hip/hip_runtime.h>
 
 // rocPRIM
 #include <rocprim/block/block_histogram.hpp>
-#include <rocprim/block/block_load_func.hpp>
-#include <rocprim/block/block_store_func.hpp>
-
-#include <iostream>
-#include <limits>
-#include <string>
-#include <vector>
-
-#include <cstdio>
-#include <cstdlib>
-
-#ifndef DEFAULT_N
-const size_t DEFAULT_BYTES = 1024 * 1024 * 128 * 4;
-#endif
-
-namespace rp = rocprim;
-
-template<
-    class Runner,
-    class T,
-    unsigned int BlockSize,
-    unsigned int ItemsPerThread,
-    unsigned int BinSize,
-    unsigned int Trials
->
-__global__
-__launch_bounds__(BlockSize)
+#include <rocprim/config.hpp>
+#include <rocprim/types.hpp>
+
+template<typename Runner,
+         typename T,
+         unsigned int BlockSize,
+         unsigned int ItemsPerThread,
+         unsigned int BinSize,
+         unsigned int Trials>
+__global__ __launch_bounds__(BlockSize)
 void kernel(const T* input, T* output)
 {
     Runner::template run<T, BlockSize, ItemsPerThread, BinSize, Trials>(input, output);
@@ -67,42 +47,42 @@ void kernel(const T* input, T* output)
 template<rocprim::block_histogram_algorithm algorithm>
 struct histogram
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        unsigned int BinSize,
-        unsigned int Trials
-    >
+    static constexpr auto algorithm_type = algorithm;
+    template<typename T,
+             unsigned int BlockSize,
+             unsigned int ItemsPerThread,
+             unsigned int BinSize,
+             unsigned int Trials>
     __device__
     static void run(const T* input, T* output)
     {
         // TODO: Move global_offset into final loop
         const unsigned int index = ((blockIdx.x * BlockSize) + threadIdx.x) * ItemsPerThread;
-        unsigned int global_offset = blockIdx.x * BinSize;
+        unsigned int       global_offset = blockIdx.x * BinSize;
 
         T values[ItemsPerThread];
-        for(unsigned int k = 0; k < ItemsPerThread; k++)
+        for(unsigned int k = 0; k < ItemsPerThread; ++k)
         {
             values[k] = input[index + k];
         }
 
-        using bhistogram_t =  rp::block_histogram<T, BlockSize, ItemsPerThread, BinSize, algorithm>;
+        using bhistogram_t
+            = rocprim::block_histogram<T, BlockSize, ItemsPerThread, BinSize, algorithm>;
         __shared__ T histogram[BinSize];
         __shared__ typename bhistogram_t::storage_type storage;
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
             bhistogram_t().histogram(values, histogram, storage);
-            for(unsigned int k = 0; k < ItemsPerThread; k++)
+            for(unsigned int k = 0; k < ItemsPerThread; ++k)
             {
                 values[k] = BinSize - 1 - values[k];
             }
         }
 
         ROCPRIM_UNROLL
-        for (unsigned int offset = 0; offset < BinSize; offset += BlockSize)
+        for(unsigned int offset = 0; offset < BinSize; offset += BlockSize)
         {
             if(offset + threadIdx.x < BinSize)
             {
@@ -113,161 +93,94 @@ struct histogram
     }
 };
 
-template<
-    class Benchmark,
-    class T,
-    unsigned int BlockSize,
-    unsigned int ItemsPerThread,
-    unsigned int BinSize = BlockSize,
-    unsigned int Trials = 100
->
-void run_benchmark(benchmark::State& state, hipStream_t stream, size_t bytes)
+template<typename Benchmark,
+         typename T,
+         unsigned int BlockSize,
+         unsigned int ItemsPerThread,
+         unsigned int BinSize = BlockSize,
+         unsigned int Trials  = 100>
+void run_benchmark(benchmark_utils::state&& state)
 {
+    const auto& stream = state.stream;
+    const auto& bytes  = state.bytes;
+
     // Calculate the number of elements N
     size_t N = bytes / sizeof(T);
     // Make sure size is a multiple of BlockSize
     constexpr auto items_per_block = BlockSize * ItemsPerThread;
-    const auto size = items_per_block * ((N + items_per_block - 1)/items_per_block);
-    const auto bin_size = BinSize * ((N + items_per_block - 1)/items_per_block);
+    const auto     size     = items_per_block * ((N + items_per_block - 1) / items_per_block);
+    const auto     bin_size = BinSize * ((N + items_per_block - 1) / items_per_block);
     // Allocate and fill memory
     std::vector<T> input(size, 0.0f);
-    T * d_input;
-    T * d_output;
-    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_input), size * sizeof(T)));
-    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_output), bin_size * sizeof(T)));
-    HIP_CHECK(
-        hipMemcpy(
-            d_input, input.data(),
-            size * sizeof(T),
-            hipMemcpyHostToDevice
-        )
-    );
+    common::device_ptr<T> d_input(input);
+    common::device_ptr<T> d_output(bin_size);
     HIP_CHECK(hipDeviceSynchronize());
 
-    // HIP events creation
-    hipEvent_t start, stop;
-    HIP_CHECK(hipEventCreate(&start));
-    HIP_CHECK(hipEventCreate(&stop));
-
-    for (auto _ : state)
-    {
-        // Record start event
-        HIP_CHECK(hipEventRecord(start, stream));
-
-        hipLaunchKernelGGL(
-            HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, BinSize, Trials>),
-            dim3(size/items_per_block), dim3(BlockSize), 0, stream,
-            d_input, d_output
-        );
-        HIP_CHECK(hipGetLastError());
-
-        // Record stop event and wait until it completes
-        HIP_CHECK(hipEventRecord(stop, stream));
-        HIP_CHECK(hipEventSynchronize(stop));
-
-        float elapsed_mseconds;
-        HIP_CHECK(hipEventElapsedTime(&elapsed_mseconds, start, stop));
-        state.SetIterationTime(elapsed_mseconds / 1000);
-    }
-
-    // Destroy HIP events
-    HIP_CHECK(hipEventDestroy(start));
-    HIP_CHECK(hipEventDestroy(stop));
-
-    state.SetBytesProcessed(state.iterations() * size * sizeof(T) * Trials);
-    state.SetItemsProcessed(state.iterations() * size * Trials);
-
-    HIP_CHECK(hipFree(d_input));
-    HIP_CHECK(hipFree(d_output));
-}
-
-// IPT - items per thread
-#define CREATE_BENCHMARK(T, BS, IPT)                                                       \
-    benchmark::RegisterBenchmark(                                                          \
-        bench_naming::format_name("{lvl:block,algo:histogram,key_type:" #T ",cfg:{bs:" #BS \
-                                  ",ipt:" #IPT ",method:"                                  \
-                                  + method_name + "}}")                                    \
-            .c_str(),                                                                      \
-        run_benchmark<Benchmark, T, BS, IPT>,                                              \
-        stream,                                                                            \
-        bytes)
-
-#define BENCHMARK_TYPE(type, block) \
-    CREATE_BENCHMARK(type, block, 1), \
-    CREATE_BENCHMARK(type, block, 2), \
-    CREATE_BENCHMARK(type, block, 3), \
-    CREATE_BENCHMARK(type, block, 4), \
-    CREATE_BENCHMARK(type, block, 8), \
-    CREATE_BENCHMARK(type, block, 16)
-
-template<class Benchmark>
-void add_benchmarks(std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    const std::string&                            method_name,
-                    hipStream_t                                   stream,
-                    size_t                                        bytes)
-{
-    std::vector<benchmark::internal::Benchmark*> new_benchmarks =
-    {
-        BENCHMARK_TYPE(int, 256),
-        BENCHMARK_TYPE(int, 320),
-        BENCHMARK_TYPE(int, 512),
+    state.run(
+        [&]
+        {
+            kernel<Benchmark, T, BlockSize, ItemsPerThread, BinSize, Trials>
+                <<<dim3(size / items_per_block), dim3(BlockSize), 0, stream>>>(d_input.get(),
+                                                                               d_output.get());
+            HIP_CHECK(hipGetLastError());
+        });
 
-        BENCHMARK_TYPE(unsigned long long, 256),
-        BENCHMARK_TYPE(unsigned long long, 320)
-    };
-    benchmarks.insert(benchmarks.end(), new_benchmarks.begin(), new_benchmarks.end());
+    state.set_throughput(size * Trials, sizeof(T));
 }
 
-int main(int argc, char *argv[])
+#define CREATE_BENCHMARK(Benchmark, method, T, BS, IPT)                                  \
+    executor.queue_fn(bench_naming::format_name("{lvl:block,algo:histogram,key_type:" #T \
+                                                ",cfg:{bs:" #BS ",ipt:" #IPT ",method:"  \
+                                                + std::string(method) + "}}")            \
+                          .c_str(),                                                      \
+                      run_benchmark<Benchmark, T, BS, IPT>);
+
+#define BENCHMARK_TYPE(Benchmark, method, T, BS)  \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 1) \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 2) \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 3) \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 4) \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 8) \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 16)
+
+#define BENCHMARK_TYPE_128(Benchmark, method, T, BS) \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 1)    \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 2)    \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 3)    \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 4)    \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 8)    \
+    CREATE_BENCHMARK(Benchmark, method, T, BS, 12)
+
+#define BENCHMARK_ATOMIC()                                                      \
+    BENCHMARK_TYPE(histogram_atomic_t, "using_atomic", int, 256)                \
+    BENCHMARK_TYPE(histogram_atomic_t, "using_atomic", int, 320)                \
+    BENCHMARK_TYPE(histogram_atomic_t, "using_atomic", int, 512)                \
+                                                                                \
+    BENCHMARK_TYPE(histogram_atomic_t, "using_atomic", unsigned long long, 256) \
+    BENCHMARK_TYPE(histogram_atomic_t, "using_atomic", unsigned long long, 320)
+
+#define BENCHMARK_SORT()                                                       \
+    BENCHMARK_TYPE(histogram_sort_t, "using_sort", int, 256)                   \
+    BENCHMARK_TYPE(histogram_sort_t, "using_sort", int, 320)                   \
+    BENCHMARK_TYPE(histogram_sort_t, "using_sort", int, 512)                   \
+                                                                               \
+    BENCHMARK_TYPE(histogram_sort_t, "using_sort", unsigned long long, 256)    \
+    BENCHMARK_TYPE(histogram_sort_t, "using_sort", unsigned long long, 320)    \
+                                                                               \
+    BENCHMARK_TYPE_128(histogram_sort_t, "using_sort", rocprim::int128_t, 256) \
+    BENCHMARK_TYPE_128(histogram_sort_t, "using_sort", rocprim::uint128_t, 256)
+
+int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<size_t>("size", "size", DEFAULT_BYTES, "number of bytes");
-    parser.set_optional<int>("trials", "trials", -1, "number of iterations");
-    parser.set_optional<std::string>("name_format",
-                                     "name_format",
-                                     "human",
-                                     "either: json,human,txt");
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    const size_t bytes = parser.get<size_t>("size");
-    const int trials = parser.get<int>("trials");
-    bench_naming::set_format(parser.get<std::string>("name_format"));
-
-    // HIP
-    hipStream_t stream = 0; // default
-
-    // Benchmark info
-    add_common_benchmark_info();
-    benchmark::AddCustomContext("bytes", std::to_string(bytes));
+    benchmark_utils::executor executor(argc, argv, 512 * benchmark_utils::MiB, 1, 0);
 
-    // Add benchmarks
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    // using_atomic
-    using histogram_a_t = histogram<rocprim::block_histogram_algorithm::using_atomic>;
-    add_benchmarks<histogram_a_t>(benchmarks, "using_atomic", stream, bytes);
-    // using_sort
-    using histogram_s_t = histogram<rocprim::block_histogram_algorithm::using_sort>;
-    add_benchmarks<histogram_s_t>(benchmarks, "using_sort", stream, bytes);
+#ifndef BENCHMARK_CONFIG_TUNING
+    using histogram_atomic_t = histogram<rocprim::block_histogram_algorithm::using_atomic>;
+    using histogram_sort_t   = histogram<rocprim::block_histogram_algorithm::using_sort>;
 
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMillisecond);
-    }
-
-    // Force number of iterations
-    if(trials > 0)
-    {
-        for(auto& b : benchmarks)
-        {
-            b->Iterations(trials);
-        }
-    }
+    BENCHMARK_ATOMIC()
+    BENCHMARK_SORT()
+#endif
 
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks();
-    return 0;
+    executor.run();
 }
diff --git a/benchmark/benchmark_block_radix_rank.cpp b/benchmark/benchmark_block_radix_rank.cpp
index faee6d3..a708037 100644
--- a/benchmark/benchmark_block_radix_rank.cpp
+++ b/benchmark/benchmark_block_radix_rank.cpp
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2022-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2022-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -21,11 +21,9 @@
 // SOFTWARE.
 
 #include "benchmark_utils.hpp"
-// CmdParser
-#include "cmdparser.hpp"
 
-// Google Benchmark
-#include <benchmark/benchmark.h>
+#include "../common/utils_data_generation.hpp"
+#include "../common/utils_device_ptr.hpp"
 
 // HIP API
 #include <hip/hip_runtime.h>
@@ -34,34 +32,31 @@
 #include <rocprim/block/block_load_func.hpp>
 #include <rocprim/block/block_radix_rank.hpp>
 #include <rocprim/block/block_store_func.hpp>
+#include <rocprim/config.hpp>
+#include <rocprim/types.hpp>
 
 #include <chrono>
-#include <limits>
+#include <stdint.h>
 #include <string>
-
-#ifndef DEFAULT_N
-const size_t DEFAULT_BYTES = 1024 * 1024 * 128 * 4;
-#endif
-
-namespace rp = rocprim;
+#include <vector>
 
 template<typename T,
-         unsigned int                   BlockSize,
-         unsigned int                   ItemsPerThread,
-         unsigned int                   RadixBits,
-         bool                           Descending,
-         rp::block_radix_rank_algorithm Algorithm,
-         unsigned int                   Trials>
-__global__ __launch_bounds__(BlockSize) void rank_kernel(const T*      keys_input,
-                                                         unsigned int* ranks_output)
+         unsigned int                        BlockSize,
+         unsigned int                        ItemsPerThread,
+         unsigned int                        RadixBits,
+         bool                                Descending,
+         rocprim::block_radix_rank_algorithm Algorithm,
+         unsigned int                        Trials>
+__global__ __launch_bounds__(BlockSize)
+void rank_kernel(const T* keys_input, unsigned int* ranks_output)
 {
-    using rank_type = rp::block_radix_rank<BlockSize, RadixBits, Algorithm>;
+    using rank_type = rocprim::block_radix_rank<BlockSize, RadixBits, Algorithm>;
 
     const unsigned int lid          = threadIdx.x;
     const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
     T keys[ItemsPerThread];
-    rp::block_load_direct_striped<BlockSize>(lid, keys_input + block_offset, keys);
+    rocprim::block_load_direct_striped<BlockSize>(lid, keys_input + block_offset, keys);
 
     unsigned int ranks[ItemsPerThread];
 
@@ -75,7 +70,7 @@ __global__ __launch_bounds__(BlockSize) void rank_kernel(const T*      keys_inpu
         while(begin_bit < end_bit)
         {
             const unsigned pass_bits = min(RadixBits, end_bit - begin_bit);
-            if ROCPRIM_IF_CONSTEXPR(Descending)
+            if constexpr(Descending)
             {
                 rank_type().rank_keys_desc(keys, ranks, storage, begin_bit, pass_bits);
             }
@@ -87,164 +82,93 @@ __global__ __launch_bounds__(BlockSize) void rank_kernel(const T*      keys_inpu
         }
     }
 
-    rp::block_store_direct_striped<BlockSize>(lid, ranks_output + block_offset, ranks);
+    rocprim::block_store_direct_striped<BlockSize>(lid, ranks_output + block_offset, ranks);
 }
 
 template<typename T,
-         unsigned int                   BlockSize,
-         unsigned int                   ItemsPerThread,
-         rp::block_radix_rank_algorithm Algorithm,
-         unsigned int                   RadixBits  = 4,
-         bool                           Descending = false,
-         unsigned int                   Trials     = 10>
-void run_benchmark(benchmark::State& state, size_t bytes, const managed_seed& seed, hipStream_t stream)
+         size_t                              BlockSize,
+         size_t                              ItemsPerThread,
+         rocprim::block_radix_rank_algorithm Algorithm,
+         size_t                              RadixBits  = 4,
+         bool                                Descending = false,
+         size_t                              Trials     = 10>
+void run_benchmark(benchmark_utils::state&& state)
 {
+    const auto& bytes  = state.bytes;
+    const auto& seed   = state.seed;
+    const auto& stream = state.stream;
+
     // Calculate the number of elements N
-    size_t N = bytes / sizeof(T);
-    constexpr unsigned int items_per_block = BlockSize * ItemsPerThread;
-    const unsigned int     grid_size       = ((N + items_per_block - 1) / items_per_block);
-    const unsigned int     size            = items_per_block * grid_size;
+    size_t           N               = bytes / sizeof(T);
+    constexpr size_t items_per_block = BlockSize * ItemsPerThread;
+    const size_t     grid_size       = ((N + items_per_block - 1) / items_per_block);
+    const size_t     size            = items_per_block * grid_size;
 
     std::vector<T> input = get_random_data<T>(size,
-                                              generate_limits<T>::min(),
-                                              generate_limits<T>::max(),
+                                              common::generate_limits<T>::min(),
+                                              common::generate_limits<T>::max(),
                                               seed.get_0());
 
-    T*            d_input;
-    unsigned int* d_output;
-    HIP_CHECK(hipMalloc(&d_input, size * sizeof(T)));
-    HIP_CHECK(hipMalloc(&d_output, size * sizeof(unsigned int)));
-    HIP_CHECK(hipMemcpy(d_input, input.data(), size * sizeof(T), hipMemcpyHostToDevice));
+    common::device_ptr<T>            d_input(input);
+    common::device_ptr<unsigned int> d_output(size);
     HIP_CHECK(hipDeviceSynchronize());
 
-    for(auto _ : state)
-    {
-        auto start = std::chrono::steady_clock::now();
-
-        hipLaunchKernelGGL(HIP_KERNEL_NAME(rank_kernel<T,
-                                                       BlockSize,
-                                                       ItemsPerThread,
-                                                       RadixBits,
-                                                       Descending,
-                                                       Algorithm,
-                                                       Trials>),
-                           dim3(grid_size),
-                           dim3(BlockSize),
-                           0,
-                           stream,
-                           d_input,
-                           d_output);
-        HIP_CHECK(hipPeekAtLastError());
-        HIP_CHECK(hipDeviceSynchronize());
-
-        auto end = std::chrono::steady_clock::now();
-        auto elapsed_seconds
-            = std::chrono::duration_cast<std::chrono::duration<double>>(end - start);
-        state.SetIterationTime(elapsed_seconds.count());
-    }
-    state.SetBytesProcessed(state.iterations() * Trials * size * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * Trials * size);
+    state.run(
+        [&]
+        {
+            rank_kernel<T, BlockSize, ItemsPerThread, RadixBits, Descending, Algorithm, Trials>
+                <<<dim3(grid_size), dim3(BlockSize), 0, stream>>>(d_input.get(), d_output.get());
+            HIP_CHECK(hipGetLastError());
+        });
 
-    HIP_CHECK(hipFree(d_input));
-    HIP_CHECK(hipFree(d_output));
+    state.set_throughput(size * Trials, sizeof(T));
 }
 
-#define CREATE_BENCHMARK(T, BS, IPT, KIND)                                                  \
-    benchmark::RegisterBenchmark(                                                           \
-        bench_naming::format_name("{lvl:block,algo:radix_rank,key_type:" #T ",cfg:{bs:" #BS \
-                                  ",ipt:" #IPT ",method:" #KIND "}}")                       \
-            .c_str(),                                                                       \
-        run_benchmark<T, BS, IPT, KIND>,                                                    \
-        bytes,                                                                               \
-        seed,                                                                               \
-        stream)
+#define CREATE_BENCHMARK(T, BS, IPT, KIND)                                                    \
+    executor.queue_fn(bench_naming::format_name("{lvl:block,algo:radix_rank,key_type:" #T     \
+                                                ",cfg:{bs:" #BS ",ipt:" #IPT ",method:" #KIND \
+                                                "}}")                                         \
+                          .c_str(),                                                           \
+                      run_benchmark<T, BS, IPT, KIND>);
 
 // clang-format off
-#define CREATE_BENCHMARK_KINDS(type, block, ipt)                                       \
-    CREATE_BENCHMARK(type, block, ipt, rp::block_radix_rank_algorithm::basic),         \
-    CREATE_BENCHMARK(type, block, ipt, rp::block_radix_rank_algorithm::basic_memoize), \
-    CREATE_BENCHMARK(type, block, ipt, rp::block_radix_rank_algorithm::match)
-
-#define BENCHMARK_TYPE(type, block)          \
-    CREATE_BENCHMARK_KINDS(type, block, 1),  \
-    CREATE_BENCHMARK_KINDS(type, block, 4),  \
-    CREATE_BENCHMARK_KINDS(type, block, 8),  \
-    CREATE_BENCHMARK_KINDS(type, block, 12), \
-    CREATE_BENCHMARK_KINDS(type, block, 16), \
+#define CREATE_BENCHMARK_KINDS(type, block, ipt) \
+    CREATE_BENCHMARK(type, block, ipt, rocprim::block_radix_rank_algorithm::basic) \
+    CREATE_BENCHMARK(type, block, ipt, rocprim::block_radix_rank_algorithm::basic_memoize) \
+    CREATE_BENCHMARK(type, block, ipt, rocprim::block_radix_rank_algorithm::match)
+
+#define BENCHMARK_TYPE(type, block) \
+    CREATE_BENCHMARK_KINDS(type, block, 1) \
+    CREATE_BENCHMARK_KINDS(type, block, 4) \
+    CREATE_BENCHMARK_KINDS(type, block, 8) \
+    CREATE_BENCHMARK_KINDS(type, block, 12) \
+    CREATE_BENCHMARK_KINDS(type, block, 16) \
     CREATE_BENCHMARK_KINDS(type, block, 20)
 // clang-format on
 
-void add_benchmarks(std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    size_t                                        bytes,
-                    const managed_seed&                           seed,
-                    hipStream_t                                   stream)
+int main(int argc, char* argv[])
 {
-    std::vector<benchmark::internal::Benchmark*> bs = {
-        BENCHMARK_TYPE(int, 128),
-        BENCHMARK_TYPE(int, 256),
-        BENCHMARK_TYPE(int, 512),
+    benchmark_utils::executor executor(argc, argv, 512 * benchmark_utils::MiB, 1, 0);
 
-        BENCHMARK_TYPE(uint8_t, 128),
-        BENCHMARK_TYPE(uint8_t, 256),
-        BENCHMARK_TYPE(uint8_t, 512),
+    BENCHMARK_TYPE(int, 128)
+    BENCHMARK_TYPE(int, 256)
+    BENCHMARK_TYPE(int, 512)
 
-        BENCHMARK_TYPE(long long, 128),
-        BENCHMARK_TYPE(long long, 256),
-        BENCHMARK_TYPE(long long, 512),
-    };
+    BENCHMARK_TYPE(uint8_t, 128)
+    BENCHMARK_TYPE(uint8_t, 256)
+    BENCHMARK_TYPE(uint8_t, 512)
 
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
-}
+    BENCHMARK_TYPE(long long, 128)
+    BENCHMARK_TYPE(long long, 256)
+    BENCHMARK_TYPE(long long, 512)
 
-int main(int argc, char* argv[])
-{
-    cli::Parser parser(argc, argv);
-    parser.set_optional<size_t>("size", "size", DEFAULT_BYTES, "number of bytes");
-    parser.set_optional<int>("trials", "trials", -1, "number of iterations");
-    parser.set_optional<std::string>("name_format",
-                                     "name_format",
-                                     "human",
-                                     "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    const size_t bytes   = parser.get<size_t>("size");
-    const int    trials = parser.get<int>("trials");
-    bench_naming::set_format(parser.get<std::string>("name_format"));
-    const std::string  seed_type = parser.get<std::string>("seed");
-    const managed_seed seed(seed_type);
-
-    // HIP
-    hipStream_t stream = 0; // default
-
-    // Benchmark info
-    add_common_benchmark_info();
-    benchmark::AddCustomContext("bytes", std::to_string(bytes));
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmarks
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks(benchmarks, bytes, seed, stream);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMillisecond);
-    }
+    BENCHMARK_TYPE(rocprim::int128_t, 128)
+    BENCHMARK_TYPE(rocprim::int128_t, 256)
+    BENCHMARK_TYPE(rocprim::int128_t, 512)
 
-    // Force number of iterations
-    if(trials > 0)
-    {
-        for(auto& b : benchmarks)
-        {
-            b->Iterations(trials);
-        }
-    }
+    BENCHMARK_TYPE(rocprim::uint128_t, 128)
+    BENCHMARK_TYPE(rocprim::uint128_t, 256)
+    BENCHMARK_TYPE(rocprim::uint128_t, 512)
 
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks();
-    return 0;
+    executor.run();
 }
diff --git a/benchmark/benchmark_block_radix_sort.cpp b/benchmark/benchmark_block_radix_sort.cpp
index ff14473..331f809 100644
--- a/benchmark/benchmark_block_radix_sort.cpp
+++ b/benchmark/benchmark_block_radix_sort.cpp
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -21,11 +21,10 @@
 // SOFTWARE.
 
 #include "benchmark_utils.hpp"
-// CmdParser
-#include "cmdparser.hpp"
 
-// Google Benchmark
-#include <benchmark/benchmark.h>
+#include "../common/utils_custom_type.hpp"
+#include "../common/utils_data_generation.hpp"
+#include "../common/utils_device_ptr.hpp"
 
 // HIP API
 #include <hip/hip_runtime.h>
@@ -34,325 +33,244 @@
 #include <rocprim/block/block_load_func.hpp>
 #include <rocprim/block/block_radix_sort.hpp>
 #include <rocprim/block/block_store_func.hpp>
+#include <rocprim/config.hpp>
+#include <rocprim/types.hpp>
 
-#include <iostream>
-#include <limits>
+#include <cstddef>
+#include <stdint.h>
 #include <string>
 #include <type_traits>
 #include <vector>
 
-#include <cstdio>
-#include <cstdlib>
-
-#ifndef DEFAULT_N
-const size_t DEFAULT_BYTES = 1024 * 1024 * 128 * 4;
-#endif
-
 enum class benchmark_kinds
 {
     sort_keys,
     sort_pairs
 };
 
-namespace rp = rocprim;
-
-template<class T>
-using select_decomposer_t = std::
-    conditional_t<is_custom_type<T>::value, custom_type_decomposer<T>, rp::identity_decomposer>;
+template<typename T>
+using select_decomposer_t = std::conditional_t<common::is_custom_type<T>::value,
+                                               custom_type_decomposer<T>,
+                                               rocprim::identity_decomposer>;
 
-template<class T,
+template<typename T,
          unsigned int BlockSize,
          unsigned int RadixBitsPerPass,
          unsigned int ItemsPerThread,
          unsigned int Trials>
-__global__ __launch_bounds__(BlockSize) void sort_keys_kernel(const T* input, T* output)
+__global__ __launch_bounds__(BlockSize)
+void sort_keys_kernel(const T* input, T* output)
 {
-    const unsigned int lid = threadIdx.x;
+    const unsigned int lid          = threadIdx.x;
     const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
     T keys[ItemsPerThread];
-    rp::block_load_direct_striped<BlockSize>(lid, input + block_offset, keys);
+    rocprim::block_load_direct_striped<BlockSize>(lid, input + block_offset, keys);
 
     ROCPRIM_NO_UNROLL
-    for(unsigned int trial = 0; trial < Trials; trial++)
+    for(unsigned int trial = 0; trial < Trials; ++trial)
     {
-        rp::block_radix_sort<T,
-                             BlockSize,
-                             ItemsPerThread,
-                             rocprim::empty_type,
-                             1,
-                             1,
-                             RadixBitsPerPass>
+        rocprim::block_radix_sort<T,
+                                  BlockSize,
+                                  ItemsPerThread,
+                                  rocprim::empty_type,
+                                  1,
+                                  1,
+                                  RadixBitsPerPass>
             sort;
         sort.sort(keys, 0, sizeof(T) * 8, select_decomposer_t<T>{});
     }
 
-    rp::block_store_direct_striped<BlockSize>(lid, output + block_offset, keys);
+    rocprim::block_store_direct_striped<BlockSize>(lid, output + block_offset, keys);
 }
 
-template<class T,
+template<typename T,
          unsigned int BlockSize,
          unsigned int RadixBitsPerPass,
          unsigned int ItemsPerThread,
          unsigned int Trials>
-__global__ __launch_bounds__(BlockSize) void sort_pairs_kernel(const T* input, T* output)
+__global__ __launch_bounds__(BlockSize)
+void sort_pairs_kernel(const T* input, T* output)
 {
-    const unsigned int lid = threadIdx.x;
+    const unsigned int lid          = threadIdx.x;
     const unsigned int block_offset = blockIdx.x * ItemsPerThread * BlockSize;
 
     T keys[ItemsPerThread];
     T values[ItemsPerThread];
-    rp::block_load_direct_striped<BlockSize>(lid, input + block_offset, keys);
-    for(unsigned int i = 0; i < ItemsPerThread; i++)
+    rocprim::block_load_direct_striped<BlockSize>(lid, input + block_offset, keys);
+    for(unsigned int i = 0; i < ItemsPerThread; ++i)
     {
         values[i] = keys[i] + T(1);
     }
 
     ROCPRIM_NO_UNROLL
-    for(unsigned int trial = 0; trial < Trials; trial++)
+    for(unsigned int trial = 0; trial < Trials; ++trial)
     {
-        rp::block_radix_sort<T, BlockSize, ItemsPerThread, T, 1, 1, RadixBitsPerPass> sort;
+        rocprim::block_radix_sort<T, BlockSize, ItemsPerThread, T, 1, 1, RadixBitsPerPass> sort;
         sort.sort(keys, values, 0, sizeof(T) * 8, select_decomposer_t<T>{});
     }
 
-    for(unsigned int i = 0; i < ItemsPerThread; i++)
+    for(unsigned int i = 0; i < ItemsPerThread; ++i)
     {
         keys[i] += values[i];
     }
-    rp::block_store_direct_striped<BlockSize>(lid, output + block_offset, keys);
+    rocprim::block_store_direct_striped<BlockSize>(lid, output + block_offset, keys);
 }
 
-template<class T,
-         unsigned int BlockSize,
-         unsigned int RadixBitsPerPass,
-         unsigned int ItemsPerThread,
-         unsigned int Trials = 10>
-void run_benchmark(benchmark::State&   state,
-                   benchmark_kinds     benchmark_kind,
-                   size_t              bytes,
-                   const managed_seed& seed,
-                   hipStream_t         stream)
+template<typename T,
+         benchmark_kinds BenchmarkKind,
+         unsigned int    BlockSize,
+         unsigned int    RadixBitsPerPass,
+         unsigned int    ItemsPerThread,
+         unsigned int    Trials = 10>
+void run_benchmark(benchmark_utils::state&& state)
 {
+    const auto& bytes  = state.bytes;
+    const auto& seed   = state.seed;
+    const auto& stream = state.stream;
+
     // Calculate the number of elements N
     size_t N = bytes / sizeof(T);
-    
+
     constexpr auto items_per_block = BlockSize * ItemsPerThread;
-    const auto size = items_per_block * ((N + items_per_block - 1)/items_per_block);
+    const auto     size = items_per_block * ((N + items_per_block - 1) / items_per_block);
 
     std::vector<T> input = get_random_data<T>(size,
-                                              generate_limits<T>::min(),
-                                              generate_limits<T>::max(),
+                                              common::generate_limits<T>::min(),
+                                              common::generate_limits<T>::max(),
                                               seed.get_0());
 
-    T*  d_input;
-    T * d_output;
-    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_input), size * sizeof(T)));
-    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_output), size * sizeof(T)));
-    HIP_CHECK(
-        hipMemcpy(
-            d_input, input.data(),
-            size * sizeof(T),
-            hipMemcpyHostToDevice
-        )
-    );
+    common::device_ptr<T> d_input(input);
+    common::device_ptr<T> d_output(size);
     HIP_CHECK(hipDeviceSynchronize());
 
-    // HIP events creation
-    hipEvent_t start, stop;
-    HIP_CHECK(hipEventCreate(&start));
-    HIP_CHECK(hipEventCreate(&stop));
-
-    for(auto _ : state)
-    {
-        // Record start event
-        HIP_CHECK(hipEventRecord(start, stream));
-
-        if(benchmark_kind == benchmark_kinds::sort_keys)
-        {
-            hipLaunchKernelGGL(
-                HIP_KERNEL_NAME(
-                    sort_keys_kernel<T, BlockSize, RadixBitsPerPass, ItemsPerThread, Trials>),
-                dim3(size / items_per_block),
-                dim3(BlockSize),
-                0,
-                stream,
-                d_input,
-                d_output);
-        }
-        else if(benchmark_kind == benchmark_kinds::sort_pairs)
+    state.run(
+        [&]
         {
-            hipLaunchKernelGGL(
-                HIP_KERNEL_NAME(
-                    sort_pairs_kernel<T, BlockSize, RadixBitsPerPass, ItemsPerThread, Trials>),
-                dim3(size / items_per_block),
-                dim3(BlockSize),
-                0,
-                stream,
-                d_input,
-                d_output);
-        }
-        HIP_CHECK(hipGetLastError());
-
-        // Record stop event and wait until it completes
-        HIP_CHECK(hipEventRecord(stop, stream));
-        HIP_CHECK(hipEventSynchronize(stop));
-
-        float elapsed_mseconds;
-        HIP_CHECK(hipEventElapsedTime(&elapsed_mseconds, start, stop));
-        state.SetIterationTime(elapsed_mseconds / 1000);
-    }
-
-    // Destroy HIP events
-    HIP_CHECK(hipEventDestroy(start));
-    HIP_CHECK(hipEventDestroy(stop));
-
-    state.SetBytesProcessed(state.iterations() * Trials * size * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * Trials * size);
-
-    HIP_CHECK(hipFree(d_input));
-    HIP_CHECK(hipFree(d_output));
+            if constexpr(BenchmarkKind == benchmark_kinds::sort_keys)
+            {
+                sort_keys_kernel<T, BlockSize, RadixBitsPerPass, ItemsPerThread, Trials>
+                    <<<dim3(size / items_per_block), dim3(BlockSize), 0, stream>>>(d_input.get(),
+                                                                                   d_output.get());
+            }
+            else if constexpr(BenchmarkKind == benchmark_kinds::sort_pairs)
+            {
+                sort_pairs_kernel<T, BlockSize, RadixBitsPerPass, ItemsPerThread, Trials>
+                    <<<dim3(size / items_per_block), dim3(BlockSize), 0, stream>>>(d_input.get(),
+                                                                                   d_output.get());
+            }
+            HIP_CHECK(hipGetLastError());
+        });
+
+    state.set_throughput(size * Trials, sizeof(T));
 }
 
 #define CREATE_BENCHMARK(T, BS, RB, IPT)                                                       \
-    benchmark::RegisterBenchmark(                                                              \
+    executor.queue_fn(                                                                         \
         bench_naming::format_name("{lvl:block,algo:radix_sort,key_type:" #T ",subalgo:" + name \
                                   + ",cfg:{bs:" #BS ",rb:" #RB ",ipt:" #IPT "}}")              \
             .c_str(),                                                                          \
-        run_benchmark<T, BS, RB, IPT>,                                                         \
-        benchmark_kind,                                                                        \
-        bytes,                                                                                  \
-        seed,                                                                                  \
-        stream)
-
-#define BENCHMARK_TYPE(type, block, radix_bits)                                                 \
-    CREATE_BENCHMARK(type, block, radix_bits, 1), CREATE_BENCHMARK(type, block, radix_bits, 2), \
-        CREATE_BENCHMARK(type, block, radix_bits, 3),                                           \
-        CREATE_BENCHMARK(type, block, radix_bits, 4), CREATE_BENCHMARK(type, block, radix_bits, 8)
-
-void add_benchmarks(benchmark_kinds                               benchmark_kind,
-                    const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    size_t                                        bytes,
-                    const managed_seed&                           seed,
-                    hipStream_t                                   stream)
+        run_benchmark<T, BenchmarkKind, BS, RB, IPT>);
+
+#define BENCHMARK_TYPE(type, block, radix_bits)  \
+    CREATE_BENCHMARK(type, block, radix_bits, 1) \
+    CREATE_BENCHMARK(type, block, radix_bits, 2) \
+    CREATE_BENCHMARK(type, block, radix_bits, 3) \
+    CREATE_BENCHMARK(type, block, radix_bits, 4) \
+    CREATE_BENCHMARK(type, block, radix_bits, 8)
+
+template<benchmark_kinds BenchmarkKind>
+void add_benchmarks(const std::string& name, benchmark_utils::executor& executor)
 {
-    using custom_int_type = custom_type<int, int>;
-
-    std::vector<benchmark::internal::Benchmark*> bs = {
-        BENCHMARK_TYPE(int, 64, 3),
-        BENCHMARK_TYPE(int, 512, 3),
-
-        BENCHMARK_TYPE(int, 64, 4),
-        BENCHMARK_TYPE(int, 128, 4),
-        BENCHMARK_TYPE(int, 192, 4),
-        BENCHMARK_TYPE(int, 256, 4),
-        BENCHMARK_TYPE(int, 320, 4),
-        BENCHMARK_TYPE(int, 512, 4),
-
-        BENCHMARK_TYPE(int8_t, 64, 3),
-        BENCHMARK_TYPE(int8_t, 512, 3),
-
-        BENCHMARK_TYPE(int8_t, 64, 4),
-        BENCHMARK_TYPE(int8_t, 128, 4),
-        BENCHMARK_TYPE(int8_t, 192, 4),
-        BENCHMARK_TYPE(int8_t, 256, 4),
-        BENCHMARK_TYPE(int8_t, 320, 4),
-        BENCHMARK_TYPE(int8_t, 512, 4),
-
-        BENCHMARK_TYPE(uint8_t, 64, 3),
-        BENCHMARK_TYPE(uint8_t, 512, 3),
-
-        BENCHMARK_TYPE(uint8_t, 64, 4),
-        BENCHMARK_TYPE(uint8_t, 128, 4),
-        BENCHMARK_TYPE(uint8_t, 192, 4),
-        BENCHMARK_TYPE(uint8_t, 256, 4),
-        BENCHMARK_TYPE(uint8_t, 320, 4),
-        BENCHMARK_TYPE(uint8_t, 512, 4),
-
-        BENCHMARK_TYPE(rocprim::half, 64, 3),
-        BENCHMARK_TYPE(rocprim::half, 512, 3),
-
-        BENCHMARK_TYPE(rocprim::half, 64, 4),
-        BENCHMARK_TYPE(rocprim::half, 128, 4),
-        BENCHMARK_TYPE(rocprim::half, 192, 4),
-        BENCHMARK_TYPE(rocprim::half, 256, 4),
-        BENCHMARK_TYPE(rocprim::half, 320, 4),
-        BENCHMARK_TYPE(rocprim::half, 512, 4),
-
-        BENCHMARK_TYPE(long long, 64, 3),
-        BENCHMARK_TYPE(long long, 512, 3),
-
-        BENCHMARK_TYPE(long long, 64, 4),
-        BENCHMARK_TYPE(long long, 128, 4),
-        BENCHMARK_TYPE(long long, 192, 4),
-        BENCHMARK_TYPE(long long, 256, 4),
-        BENCHMARK_TYPE(long long, 320, 4),
-        BENCHMARK_TYPE(long long, 512, 4),
-
-        BENCHMARK_TYPE(custom_int_type, 64, 3),
-        BENCHMARK_TYPE(custom_int_type, 512, 3),
-
-        BENCHMARK_TYPE(custom_int_type, 64, 4),
-        BENCHMARK_TYPE(custom_int_type, 128, 4),
-        BENCHMARK_TYPE(custom_int_type, 192, 4),
-        BENCHMARK_TYPE(custom_int_type, 256, 4),
-        BENCHMARK_TYPE(custom_int_type, 320, 4),
-        BENCHMARK_TYPE(custom_int_type, 512, 4),
-    };
-
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+    using custom_int_type = common::custom_type<int, int>;
+
+    BENCHMARK_TYPE(int, 64, 3)
+    BENCHMARK_TYPE(int, 512, 3)
+
+    BENCHMARK_TYPE(int, 64, 4)
+    BENCHMARK_TYPE(int, 128, 4)
+    BENCHMARK_TYPE(int, 192, 4)
+    BENCHMARK_TYPE(int, 256, 4)
+    BENCHMARK_TYPE(int, 320, 4)
+    BENCHMARK_TYPE(int, 512, 4)
+
+    BENCHMARK_TYPE(int8_t, 64, 3)
+    BENCHMARK_TYPE(int8_t, 512, 3)
+
+    BENCHMARK_TYPE(int8_t, 64, 4)
+    BENCHMARK_TYPE(int8_t, 128, 4)
+    BENCHMARK_TYPE(int8_t, 192, 4)
+    BENCHMARK_TYPE(int8_t, 256, 4)
+    BENCHMARK_TYPE(int8_t, 320, 4)
+    BENCHMARK_TYPE(int8_t, 512, 4)
+
+    BENCHMARK_TYPE(uint8_t, 64, 3)
+    BENCHMARK_TYPE(uint8_t, 512, 3)
+
+    BENCHMARK_TYPE(uint8_t, 64, 4)
+    BENCHMARK_TYPE(uint8_t, 128, 4)
+    BENCHMARK_TYPE(uint8_t, 192, 4)
+    BENCHMARK_TYPE(uint8_t, 256, 4)
+    BENCHMARK_TYPE(uint8_t, 320, 4)
+    BENCHMARK_TYPE(uint8_t, 512, 4)
+
+    BENCHMARK_TYPE(rocprim::half, 64, 3)
+    BENCHMARK_TYPE(rocprim::half, 512, 3)
+
+    BENCHMARK_TYPE(rocprim::half, 64, 4)
+    BENCHMARK_TYPE(rocprim::half, 128, 4)
+    BENCHMARK_TYPE(rocprim::half, 192, 4)
+    BENCHMARK_TYPE(rocprim::half, 256, 4)
+    BENCHMARK_TYPE(rocprim::half, 320, 4)
+    BENCHMARK_TYPE(rocprim::half, 512, 4)
+
+    BENCHMARK_TYPE(long long, 64, 3)
+    BENCHMARK_TYPE(long long, 512, 3)
+
+    BENCHMARK_TYPE(long long, 64, 4)
+    BENCHMARK_TYPE(long long, 128, 4)
+    BENCHMARK_TYPE(long long, 192, 4)
+    BENCHMARK_TYPE(long long, 256, 4)
+    BENCHMARK_TYPE(long long, 320, 4)
+    BENCHMARK_TYPE(long long, 512, 4)
+
+    BENCHMARK_TYPE(custom_int_type, 64, 3)
+    BENCHMARK_TYPE(custom_int_type, 512, 3)
+
+    BENCHMARK_TYPE(custom_int_type, 64, 4)
+    BENCHMARK_TYPE(custom_int_type, 128, 4)
+    BENCHMARK_TYPE(custom_int_type, 192, 4)
+    BENCHMARK_TYPE(custom_int_type, 256, 4)
+    BENCHMARK_TYPE(custom_int_type, 320, 4)
+    BENCHMARK_TYPE(custom_int_type, 512, 4)
+
+    BENCHMARK_TYPE(rocprim::int128_t, 64, 3)
+    BENCHMARK_TYPE(rocprim::int128_t, 512, 3)
+
+    BENCHMARK_TYPE(rocprim::int128_t, 64, 4)
+    BENCHMARK_TYPE(rocprim::int128_t, 128, 4)
+    BENCHMARK_TYPE(rocprim::int128_t, 192, 4)
+    BENCHMARK_TYPE(rocprim::int128_t, 256, 4)
+    BENCHMARK_TYPE(rocprim::int128_t, 320, 4)
+    BENCHMARK_TYPE(rocprim::int128_t, 512, 4)
+
+    BENCHMARK_TYPE(rocprim::uint128_t, 64, 3)
+    BENCHMARK_TYPE(rocprim::uint128_t, 512, 3)
+
+    BENCHMARK_TYPE(rocprim::uint128_t, 64, 4)
+    BENCHMARK_TYPE(rocprim::uint128_t, 128, 4)
+    BENCHMARK_TYPE(rocprim::uint128_t, 192, 4)
+    BENCHMARK_TYPE(rocprim::uint128_t, 256, 4)
+    BENCHMARK_TYPE(rocprim::uint128_t, 320, 4)
+    BENCHMARK_TYPE(rocprim::uint128_t, 512, 4)
 }
 
-int main(int argc, char *argv[])
+int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<size_t>("size", "size", DEFAULT_BYTES, "number of bytes");
-    parser.set_optional<int>("trials", "trials", -1, "number of iterations");
-    parser.set_optional<std::string>("name_format",
-                                     "name_format",
-                                     "human",
-                                     "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    const size_t bytes = parser.get<size_t>("size");
-    const int trials = parser.get<int>("trials");
-    bench_naming::set_format(parser.get<std::string>("name_format"));
-    const std::string  seed_type = parser.get<std::string>("seed");
-    const managed_seed seed(seed_type);
-
-    // HIP
-    hipStream_t stream = 0; // default
-
-    // Benchmark info
-    add_common_benchmark_info();
-    benchmark::AddCustomContext("bytes", std::to_string(bytes));
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmarks
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks(benchmark_kinds::sort_keys, "keys", benchmarks, bytes, seed, stream);
-    add_benchmarks(benchmark_kinds::sort_pairs, "pairs", benchmarks, bytes, seed, stream);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMillisecond);
-    }
+    benchmark_utils::executor executor(argc, argv, 512 * benchmark_utils::MiB, 1, 0);
 
-    // Force number of iterations
-    if(trials > 0)
-    {
-        for(auto& b : benchmarks)
-        {
-            b->Iterations(trials);
-        }
-    }
+    add_benchmarks<benchmark_kinds::sort_keys>("keys", executor);
+    add_benchmarks<benchmark_kinds::sort_pairs>("pairs", executor);
 
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks();
-    return 0;
+    executor.run();
 }
diff --git a/benchmark/benchmark_block_reduce.cpp b/benchmark/benchmark_block_reduce.cpp
index 2571654..27678d0 100644
--- a/benchmark/benchmark_block_reduce.cpp
+++ b/benchmark/benchmark_block_reduce.cpp
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -21,41 +21,29 @@
 // SOFTWARE.
 
 #include "benchmark_utils.hpp"
-// CmdParser
-#include "cmdparser.hpp"
 
-// Google Benchmark
-#include <benchmark/benchmark.h>
+#include "../common/utils_custom_type.hpp"
+#include "../common/utils_device_ptr.hpp"
 
 // HIP API
 #include <hip/hip_runtime.h>
 
 // rocPRIM
 #include <rocprim/block/block_reduce.hpp>
+#include <rocprim/config.hpp>
+#include <rocprim/types.hpp>
 
-#include <iostream>
-#include <limits>
+#include <cstddef>
+#include <stdint.h>
 #include <string>
 #include <vector>
 
-#include <cstdio>
-#include <cstdlib>
-
-#ifndef DEFAULT_N
-const size_t DEFAULT_BYTES = 1024 * 1024 * 32 * 4;
-#endif
-
-namespace rp = rocprim;
-
-template<
-    class Runner,
-    class T,
-    unsigned int BlockSize,
-    unsigned int ItemsPerThread,
-    unsigned int Trials
->
-__global__
-__launch_bounds__(BlockSize)
+template<typename Runner,
+         typename T,
+         unsigned int BlockSize,
+         unsigned int ItemsPerThread,
+         unsigned int Trials>
+__global__ __launch_bounds__(BlockSize)
 void kernel(const T* input, T* output)
 {
     Runner::template run<T, BlockSize, ItemsPerThread, Trials>(input, output);
@@ -64,12 +52,7 @@ void kernel(const T* input, T* output)
 template<rocprim::block_reduce_algorithm algorithm>
 struct reduce
 {
-    template<
-        class T,
-        unsigned int BlockSize,
-        unsigned int ItemsPerThread,
-        unsigned int Trials
-    >
+    template<typename T, unsigned int BlockSize, unsigned int ItemsPerThread, unsigned int Trials>
     __device__
     static void run(const T* input, T* output)
     {
@@ -77,16 +60,16 @@ struct reduce
 
         T values[ItemsPerThread];
         T reduced_value;
-        for(unsigned int k = 0; k < ItemsPerThread; k++)
+        for(unsigned int k = 0; k < ItemsPerThread; ++k)
         {
             values[k] = input[i * ItemsPerThread + k];
         }
 
-        using breduce_t = rp::block_reduce<T, BlockSize, algorithm>;
+        using breduce_t = rocprim::block_reduce<T, BlockSize, algorithm>;
         __shared__ typename breduce_t::storage_type storage;
 
         ROCPRIM_NO_UNROLL
-        for(unsigned int trial = 0; trial < Trials; trial++)
+        for(unsigned int trial = 0; trial < Trials; ++trial)
         {
             breduce_t().reduce(values, reduced_value, storage);
             values[0] = reduced_value;
@@ -99,193 +82,113 @@ struct reduce
     }
 };
 
-template<
-    class Benchmark,
-    class T,
-    unsigned int BlockSize,
-    unsigned int ItemsPerThread,
-    unsigned int Trials = 100
->
-void run_benchmark(benchmark::State& state, hipStream_t stream, size_t bytes)
+template<typename Benchmark,
+         typename T,
+         unsigned int BlockSize,
+         unsigned int ItemsPerThread,
+         unsigned int Trials = 100>
+void run_benchmark(benchmark_utils::state&& state)
 {
+    const auto& bytes  = state.bytes;
+    const auto& stream = state.stream;
+
     // Calculate the number of elements N
     size_t N = bytes / sizeof(T);
     // Make sure size is a multiple of BlockSize
     constexpr auto items_per_block = BlockSize * ItemsPerThread;
-    const auto size = items_per_block * ((N + items_per_block - 1)/items_per_block);
+    const auto     size = items_per_block * ((N + items_per_block - 1) / items_per_block);
     // Allocate and fill memory
     std::vector<T> input(size, T(1));
-    T * d_input;
-    T * d_output;
-    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_input), size * sizeof(T)));
-    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_output), size * sizeof(T)));
-    HIP_CHECK(
-        hipMemcpy(
-            d_input, input.data(),
-            size * sizeof(T),
-            hipMemcpyHostToDevice
-        )
-    );
+    common::device_ptr<T> d_input(input);
+    common::device_ptr<T> d_output(size);
     HIP_CHECK(hipDeviceSynchronize());
 
-    // HIP events creation
-    hipEvent_t start, stop;
-    HIP_CHECK(hipEventCreate(&start));
-    HIP_CHECK(hipEventCreate(&stop));
-
-    for (auto _ : state)
-    {
-        // Record start event
-        HIP_CHECK(hipEventRecord(start, stream));
-
-        hipLaunchKernelGGL(
-            HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>),
-            dim3(size/items_per_block), dim3(BlockSize), 0, stream,
-            d_input, d_output
-        );
-        HIP_CHECK(hipGetLastError());
-
-        // Record stop event and wait until it completes
-        HIP_CHECK(hipEventRecord(stop, stream));
-        HIP_CHECK(hipEventSynchronize(stop));
-
-        float elapsed_mseconds;
-        HIP_CHECK(hipEventElapsedTime(&elapsed_mseconds, start, stop));
-        state.SetIterationTime(elapsed_mseconds / 1000);
-    }
-
-    // Destroy HIP events
-    HIP_CHECK(hipEventDestroy(start));
-    HIP_CHECK(hipEventDestroy(stop));
-
-    state.SetBytesProcessed(state.iterations() * size * sizeof(T) * Trials);
-    state.SetItemsProcessed(state.iterations() * size * Trials);
+    state.run(
+        [&]
+        {
+            kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>
+                <<<dim3(size / items_per_block), dim3(BlockSize), 0, stream>>>(d_input.get(),
+                                                                               d_output.get());
+            HIP_CHECK(hipGetLastError());
+        });
 
-    HIP_CHECK(hipFree(d_input));
-    HIP_CHECK(hipFree(d_output));
+    state.set_throughput(size * Trials, sizeof(T));
 }
 
-// IPT - items per thread
-#define CREATE_BENCHMARK(T, BS, IPT)                                                               \
-    benchmark::RegisterBenchmark(bench_naming::format_name("{lvl:block,algo:reduce,key_type:" #T   \
-                                                           ",cfg:{bs:" #BS ",ipt:" #IPT ",method:" \
-                                                           + method_name + "}}")                   \
-                                     .c_str(),                                                     \
-                                 run_benchmark<Benchmark, T, BS, IPT>,                             \
-                                 stream,                                                           \
-                                 bytes)
-
-#define BENCHMARK_TYPE(type, block) \
-    CREATE_BENCHMARK(type, block, 1), \
-    CREATE_BENCHMARK(type, block, 2), \
-    CREATE_BENCHMARK(type, block, 3), \
-    CREATE_BENCHMARK(type, block, 4), \
-    CREATE_BENCHMARK(type, block, 8), \
-    CREATE_BENCHMARK(type, block, 11), \
+#define CREATE_BENCHMARK(T, BS, IPT)                                                    \
+    executor.queue_fn(bench_naming::format_name("{lvl:block,algo:reduce,key_type:" #T   \
+                                                ",cfg:{bs:" #BS ",ipt:" #IPT ",method:" \
+                                                + name + "}}")                          \
+                          .c_str(),                                                     \
+                      run_benchmark<Benchmark, T, BS, IPT>);
+
+#define BENCHMARK_TYPE(type, block)   \
+    CREATE_BENCHMARK(type, block, 1)  \
+    CREATE_BENCHMARK(type, block, 2)  \
+    CREATE_BENCHMARK(type, block, 3)  \
+    CREATE_BENCHMARK(type, block, 4)  \
+    CREATE_BENCHMARK(type, block, 8)  \
+    CREATE_BENCHMARK(type, block, 11) \
     CREATE_BENCHMARK(type, block, 16)
 
-template<class Benchmark>
-void add_benchmarks(std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    const std::string&                            method_name,
-                    hipStream_t                                   stream,
-                    size_t                                        bytes)
+template<typename Benchmark>
+void add_benchmarks(const std::string& name, benchmark_utils::executor& executor)
 {
-    using custom_float2 = custom_type<float, float>;
-    using custom_double2 = custom_type<double, double>;
-
-    std::vector<benchmark::internal::Benchmark*> new_benchmarks =
-    {
-        // When block size is less than or equal to warp size
-        BENCHMARK_TYPE(int, 64),
-        BENCHMARK_TYPE(float, 64),
-        BENCHMARK_TYPE(double, 64),
-        BENCHMARK_TYPE(int8_t, 64),
-        BENCHMARK_TYPE(uint8_t, 64),
-        BENCHMARK_TYPE(rocprim::half, 64),
-
-        BENCHMARK_TYPE(int, 256),
-        BENCHMARK_TYPE(float, 256),
-        BENCHMARK_TYPE(double, 256),
-        BENCHMARK_TYPE(int8_t, 256),
-        BENCHMARK_TYPE(uint8_t, 256),
-        BENCHMARK_TYPE(rocprim::half, 256),
-
-        CREATE_BENCHMARK(custom_float2, 256, 1),
-        CREATE_BENCHMARK(custom_float2, 256, 4),
-        CREATE_BENCHMARK(custom_float2, 256, 8),
-
-        CREATE_BENCHMARK(float2, 256, 1),
-        CREATE_BENCHMARK(float2, 256, 4),
-        CREATE_BENCHMARK(float2, 256, 8),
-
-        CREATE_BENCHMARK(custom_double2, 256, 1),
-        CREATE_BENCHMARK(custom_double2, 256, 4),
-        CREATE_BENCHMARK(custom_double2, 256, 8),
-
-        CREATE_BENCHMARK(double2, 256, 1),
-        CREATE_BENCHMARK(double2, 256, 4),
-        CREATE_BENCHMARK(double2, 256, 8),
-
-        CREATE_BENCHMARK(float4, 256, 1),
-        CREATE_BENCHMARK(float4, 256, 4),
-        CREATE_BENCHMARK(float4, 256, 8),
-    };
-    benchmarks.insert(benchmarks.end(), new_benchmarks.begin(), new_benchmarks.end());
+    using custom_float2  = common::custom_type<float, float>;
+    using custom_double2 = common::custom_type<double, double>;
+
+    // When block size is less than or equal to warp size
+    BENCHMARK_TYPE(int, 64)
+    BENCHMARK_TYPE(float, 64)
+    BENCHMARK_TYPE(double, 64)
+    BENCHMARK_TYPE(int8_t, 64)
+    BENCHMARK_TYPE(uint8_t, 64)
+    BENCHMARK_TYPE(rocprim::half, 64)
+    BENCHMARK_TYPE(rocprim::int128_t, 64)
+    BENCHMARK_TYPE(rocprim::uint128_t, 64)
+
+    BENCHMARK_TYPE(int, 256)
+    BENCHMARK_TYPE(float, 256)
+    BENCHMARK_TYPE(double, 256)
+    BENCHMARK_TYPE(int8_t, 256)
+    BENCHMARK_TYPE(uint8_t, 256)
+    BENCHMARK_TYPE(rocprim::half, 256)
+    BENCHMARK_TYPE(rocprim::int128_t, 256)
+    BENCHMARK_TYPE(rocprim::uint128_t, 256)
+
+    CREATE_BENCHMARK(custom_float2, 256, 1)
+    CREATE_BENCHMARK(custom_float2, 256, 4)
+    CREATE_BENCHMARK(custom_float2, 256, 8)
+
+    CREATE_BENCHMARK(float2, 256, 1)
+    CREATE_BENCHMARK(float2, 256, 4)
+    CREATE_BENCHMARK(float2, 256, 8)
+
+    CREATE_BENCHMARK(custom_double2, 256, 1)
+    CREATE_BENCHMARK(custom_double2, 256, 4)
+    CREATE_BENCHMARK(custom_double2, 256, 8)
+
+    CREATE_BENCHMARK(double2, 256, 1)
+    CREATE_BENCHMARK(double2, 256, 4)
+    CREATE_BENCHMARK(double2, 256, 8)
+
+    CREATE_BENCHMARK(float4, 256, 1)
+    CREATE_BENCHMARK(float4, 256, 4)
+    CREATE_BENCHMARK(float4, 256, 8)
 }
 
-int main(int argc, char *argv[])
+int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<size_t>("size", "size", DEFAULT_BYTES, "number of bytes");
-    parser.set_optional<int>("trials", "trials", -1, "number of iterations");
-    parser.set_optional<std::string>("name_format",
-                                     "name_format",
-                                     "human",
-                                     "either: json,human,txt");
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    const size_t bytes = parser.get<size_t>("size");
-    const int trials = parser.get<int>("trials");
-    bench_naming::set_format(parser.get<std::string>("name_format"));
+    benchmark_utils::executor executor(argc, argv, 128 * benchmark_utils::MiB, 1, 0);
 
-    // HIP
-    hipStream_t stream = 0; // default
-
-    // Benchmark info
-    add_common_benchmark_info();
-    benchmark::AddCustomContext("bytes", std::to_string(bytes));
-
-    // Add benchmarks
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    // using_warp_scan
     using reduce_uwr_t = reduce<rocprim::block_reduce_algorithm::using_warp_reduce>;
-    add_benchmarks<reduce_uwr_t>(benchmarks, "using_warp_reduce", stream, bytes);
-    // reduce then scan
-    using reduce_rr_t = reduce<rocprim::block_reduce_algorithm::raking_reduce>;
-    add_benchmarks<reduce_rr_t>(benchmarks, "raking_reduce", stream, bytes);
-    // reduce commutative only
-    using reduce_rrco_t = reduce<rocprim::block_reduce_algorithm::raking_reduce_commutative_only>;
-    add_benchmarks<reduce_rrco_t>(benchmarks, "raking_reduce_commutative_only", stream, bytes);
+    add_benchmarks<reduce_uwr_t>("using_warp_reduce", executor);
 
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMillisecond);
-    }
+    using reduce_rr_t = reduce<rocprim::block_reduce_algorithm::raking_reduce>;
+    add_benchmarks<reduce_rr_t>("raking_reduce", executor);
 
-    // Force number of iterations
-    if(trials > 0)
-    {
-        for(auto& b : benchmarks)
-        {
-            b->Iterations(trials);
-        }
-    }
+    using reduce_rrco_t = reduce<rocprim::block_reduce_algorithm::raking_reduce_commutative_only>;
+    add_benchmarks<reduce_rrco_t>("raking_reduce_commutative_only", executor);
 
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks();
-    return 0;
+    executor.run();
 }
diff --git a/benchmark/benchmark_block_run_length_decode.cpp b/benchmark/benchmark_block_run_length_decode.cpp
index e56bf09..3b20d0f 100644
--- a/benchmark/benchmark_block_run_length_decode.cpp
+++ b/benchmark/benchmark_block_run_length_decode.cpp
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2021-2025 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -21,32 +21,33 @@
 // SOFTWARE.
 
 #include "benchmark_utils.hpp"
-#include "cmdparser.hpp"
 
-#include <benchmark/benchmark.h>
+#include "../common/utils_data_generation.hpp"
+#include "../common/utils_device_ptr.hpp"
 
 #include <rocprim/block/block_load_func.hpp>
 #include <rocprim/block/block_run_length_decode.hpp>
 #include <rocprim/block/block_store_func.hpp>
+#include <rocprim/type_traits.hpp>
+#include <rocprim/types.hpp>
 
-#include <random>
+#include <chrono>
+#include <cstddef>
+#include <string>
+#include <type_traits>
 #include <vector>
 
-#ifndef DEFAULT_N
-const size_t DEFAULT_BYTES = 1024 * 1024 * 32 * 4;
-#endif
-
-template<class ItemT,
-         class OffsetT,
+template<typename ItemT,
+         typename OffsetT,
          unsigned BlockSize,
          unsigned RunsPerThread,
          unsigned DecodedItemsPerThread,
          unsigned Trials>
-__global__
-    __launch_bounds__(BlockSize) void block_run_length_decode_kernel(const ItemT*   d_run_items,
-                                                                     const OffsetT* d_run_offsets,
-                                                                     ItemT*         d_decoded_items,