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

[Bojan Aleksovski]

Update to new upstream version 7.1.0

[Bojan Aleksovski (b0b0a)]

uploaded package to this ppa: https://launchpad.net/~b0b0a/+archive/ubuntu/rocthrust-2139676/

(-proposed and amd64, amd64v3, arm64 archs enabled)

Gonna trigger autopkgtest when it is built and published

[Bojan Aleksovski (b0b0a)]

`reverse-depends --arch ppc64el src:rocthrust -x` returns libstdgpu-hip-dev (src:stdgpu) reverse dependency for for librocthrust-dev.

[Bojan Aleksovski (b0b0a)]

Sebastien Bacher at Debcrafters channel: stdgpu has been removed from Debian testing (https://bugs.debian.org/cgi-bin/bugreport.cgi?bug=1119033) and will be removed from Ubuntu as well until FTBFS gets fixed.

[Andreas Hasenack (ahasenack)]

+ * d/p/0002-Allow-system-sqlite.patch: remove patch

Via grep I see that tests use sqlite3 (like ./test/bitwise_repro/bwr_db.hpp). Inspecting the librocthrust-tests package that is in the archive now, I don't see a dependency on sqlite3, even though it's a build-depends of src:rocthrust. Perhaps it's statically linked? My question here is just to be sure we are not dropping an external sqlite dependency in favor of something vendored. Or, asking in another way, why was this patch dropped?

review: Needs Information

[Andreas Hasenack (ahasenack)]

Ah, I think I see it. USE_SYSTEM_SQLITE was a debian change in the previous package, now upstream has a proper SQLITE_USE_SYSTEM_PACKAGE, is that it?

[Bojan Aleksovski (b0b0a)]

Yes https://github.com/ROCm/rocThrust/releases/tag/rocm-7.1.0 - upstream 7.1.0 release notes:
- Added a new CMake option -DSQLITE_USE_SYSTEM_PACKAGE to allow SQLite to be provided by the system.

[Andreas Hasenack (ahasenack)]

Could you please create a valid d/watch file? I know it's not something this PR broke, and was like this already in ubuntu, but since we are updating to a new usptream version, it's important to have a working d/watch file.

review: Needs Fixing

f4d6881... by Bojan Aleksovski on 2026-02-05

d/changelog: update changelog with watch file entry

292a41d... by Bojan Aleksovski on 2026-02-05

d/watch: create a valid watch file

[Andreas Hasenack (ahasenack)]

awesome, +1

review: Approve

[Andreas Hasenack (ahasenack)]

Sponsored:

Uploading rocthrust_7.1.0-0ubuntu2.dsc
Uploading rocthrust_7.1.0.orig.tar.gz
Uploading rocthrust_7.1.0-0ubuntu2.debian.tar.xz
Uploading rocthrust_7.1.0-0ubuntu2_source.buildinfo
Uploading rocthrust_7.1.0-0ubuntu2_source.changes

[Bojan Aleksovski (b0b0a)]

Thank you Andreas!

Unmerged commits

f4d6881... by Bojan Aleksovski on 2026-02-05

d/changelog: update changelog with watch file entry

292a41d... by Bojan Aleksovski on 2026-02-05

d/watch: create a valid watch file

diff --git a/.clang-format b/.clang-format
index d3b0fef..ec5070f 100644
--- a/.clang-format
+++ b/.clang-format
@@ -1,164 +1,229 @@
-# Style file for MLSE Libraries based on the modified rocBLAS style
-
-# Common settings
-BasedOnStyle:  LLVM
-TabWidth:        2
-IndentWidth:     2
-UseTab:          Never
-ColumnLimit: 120
-
-# Other languages JavaScript, Proto
-
----
-Language: Cpp
-
-# http://releases.llvm.org/6.0.1/tools/clang/docs/ClangFormatStyleOptions.html#disabling-formatting-on-a-piece-of-code
-# int formatted_code;
-# // clang-format off
-#     void    unformatted_code  ;
-# // clang-format on
-# void formatted_code_again;
-
-DisableFormat:   false
-Standard: c++14
+# Note that we don't specify the language in this file because some files are
+# detected as Cpp, but others are detected as ObjC and we want this formatting
+# to apply to all types of files.
+BasedOnStyle: LLVM
 AccessModifierOffset: -2
-AlignAfterOpenBracket: true
-AlignConsecutiveAssignments: true
-AlignConsecutiveDeclarations: true
-AlignEscapedNewlinesLeft: true
-AlignOperands:   true
+AlignAfterOpenBracket: Align
+AlignConsecutiveAssignments: Consecutive
+AlignConsecutiveBitFields: Consecutive
+AlignConsecutiveMacros: Consecutive
+AlignEscapedNewlines: Left
+AlignOperands: AlignAfterOperator
 AllowAllArgumentsOnNextLine: true
-AlignTrailingComments: false
+AlignTrailingComments:
+  Kind: Never
 AllowAllParametersOfDeclarationOnNextLine: true
+AllowAllConstructorInitializersOnNextLine: true
 AllowShortBlocksOnASingleLine: false
 AllowShortCaseLabelsOnASingleLine: false
 AllowShortFunctionsOnASingleLine: Empty
-AllowShortIfStatementsOnASingleLine: false
+AllowShortIfStatementsOnASingleLine: Never
+AllowShortLambdasOnASingleLine: Empty
 AllowShortLoopsOnASingleLine: false
-AlwaysBreakAfterDefinitionReturnType: false
 AlwaysBreakAfterReturnType: None
-AlwaysBreakBeforeMultilineStrings: false
-AlwaysBreakTemplateDeclarations: true
+AlwaysBreakTemplateDeclarations: Yes
 AttributeMacros: [
+                  # rocThrust
                   'THRUST_DEVICE',
                   'THRUST_FORCEINLINE',
                   'THRUST_HOST_DEVICE',
                   'THRUST_HOST',
+                  'THRUST_NODISCARD',
+                  'THRUST_DEPRECATED',
+                  'THRUST_INLINE_CONSTANT',
+                  # thrust
+                  '_CCCL_ALIGNAS_TYPE',
+                  '_CCCL_ALIGNAS',
+                  '_CCCL_CONSTEXPR_CXX14',
+                  '_CCCL_CONSTEXPR_CXX17',
+                  '_CCCL_CONSTEXPR_CXX20',
+                  '_CCCL_CONSTEXPR_CXX23',
                   '_CCCL_DEVICE',
+                  '_CCCL_FALLTHROUGH',
                   '_CCCL_FORCEINLINE',
                   '_CCCL_HOST_DEVICE',
                   '_CCCL_HOST',
+                  '_CCCL_NO_UNIQUE_ADDRESS',
+                  '_CCCL_NODISCARD_FRIEND',
+                  '_CCCL_NODISCARD',
+                  '_CCCL_NORETURN',
+                  '_CCCL_VISIBILITY_HIDDEN',
+                  'CUB_RUNTIME_FUNCTION',
+                  'CUB_DETAIL_KERNEL_ATTRIBUTES',
                   'THRUST_RUNTIME_FUNCTION',
                   'THRUST_DETAIL_KERNEL_ATTRIBUTES',
+                  '_LIBCUDACXX_ALIGNOF',
+                  '_LIBCUDACXX_ALWAYS_INLINE',
+                  '_LIBCUDACXX_AVAILABILITY_THROW_BAD_VARIANT_ACCESS',
+                  '_LIBCUDACXX_CONSTINIT',
+                  '_LIBCUDACXX_DEPRECATED_IN_CXX11',
+                  '_LIBCUDACXX_DEPRECATED_IN_CXX14',
+                  '_LIBCUDACXX_DEPRECATED_IN_CXX17',
+                  '_LIBCUDACXX_DEPRECATED_IN_CXX20',
+                  '_LIBCUDACXX_DEPRECATED',
+                  '_LIBCUDACXX_DISABLE_EXTENTSION_WARNING',
+                  '_LIBCUDACXX_EXCLUDE_FROM_EXPLICIT_INSTANTIATION',
+                  '_LIBCUDACXX_EXPORTED_FROM_ABI',
+                  '_LIBCUDACXX_EXTERN_TEMPLATE_TYPE_VIS',
+                  '_LIBCUDACXX_HIDDEN',
+                  '_LIBCUDACXX_HIDE_FROM_ABI_AFTER_V1',
+                  '_LIBCUDACXX_HIDE_FROM_ABI',
+                  '_LIBCUDACXX_INLINE_VISIBILITY',
+                  '_LIBCUDACXX_INTERNAL_LINKAGE',
+                  '_LIBCUDACXX_METHOD_TEMPLATE_IMPLICIT_INSTANTIATION_VIS',
+                  '_LIBCUDACXX_NO_DESTROY',
+                  '_LIBCUDACXX_NO_SANITIZE',
+                  '_LIBCUDACXX_NOALIAS',
+                  '_LIBCUDACXX_OVERRIDABLE_FUNC_VIS',
+                  '_LIBCUDACXX_STANDALONE_DEBUG',
+                  '_LIBCUDACXX_TEMPLATE_DATA_VIS',
+                  '_LIBCUDACXX_TEMPLATE_VIS',
+                  '_LIBCUDACXX_THREAD_SAFETY_ANNOTATION',
+                  '_LIBCUDACXX_USING_IF_EXISTS',
+                  '_LIBCUDACXX_WEAK',
                  ]
 BinPackArguments: false
 BinPackParameters: false
-
-# Configure each individual brace in BraceWrapping
 BreakBeforeBraces: Custom
-# Control of individual brace wrapping cases
-BraceWrapping: {
-    AfterCaseLabel: 'false'
-    AfterClass: 'true'
-    AfterControlStatement: 'true'
-    AfterEnum : 'true'
-    AfterFunction : 'true'
-    AfterNamespace : 'true'
-    AfterStruct : 'true'
-    AfterUnion : 'true'
-    BeforeCatch : 'true'
-    BeforeElse : 'true'
-    IndentBraces : 'false'
-    SplitEmptyFunction: 'false'
-    SplitEmptyRecord: 'false'
-}
-
+BraceWrapping:
+  AfterCaseLabel: false
+  AfterClass: true
+  AfterControlStatement: true
+  AfterEnum: true
+  AfterFunction: true
+  AfterNamespace: true
+  AfterStruct: true
+  AfterUnion: true
+  BeforeCatch: true
+  BeforeElse: true
+  IndentBraces: false
+  SplitEmptyFunction: false
+  SplitEmptyRecord: false
 BreakBeforeConceptDeclarations: true
 BreakBeforeBinaryOperators: NonAssignment
 BreakBeforeTernaryOperators: true
 BreakConstructorInitializers: BeforeComma
 BreakInheritanceList: BeforeComma
+ColumnLimit: 120
+CompactNamespaces: false
+ContinuationIndentWidth: 2
 EmptyLineAfterAccessModifier: Never
 EmptyLineBeforeAccessModifier: Always
-
+FixNamespaceComments: true
+IfMacros: [
+  # rocThrust
+  'THRUST_IF_CONSTEXPR',
+  # thrust
+  '_CCCL_IF_CONSTEXPR',
+  '_CCCL_ELSE_IF_CONSTEXPR',
+]
+IndentWrappedFunctionNames: false
+IncludeBlocks:   Regroup
+IncludeCategories:
+  # rocThrust
+  - Regex:           '^<(rocprim/rocprim.hpp|thrust/detail/config.h|thrust/system/hip/config.h)'
+    Priority:            0
+    SortPriority:        0
+  - Regex:           '^<rocprim/'
+    Priority:            2
+    SortPriority:        1
+  # Use the priority set by thrust (below)
+  # - Regex:           '^<thrust/'
+  #   Priority:            3
+  #   SortPriority:        2
+  - Regex:           '^<hip/'
+    Priority:            4
+    SortPriority:        3
+  - Regex:           '^<[a-z_]*>$'
+    Priority:            5
+    SortPriority:        4
+  - Regex:           '^<hip'
+    Priority:            0
+    SortPriority:        0
+  # thrust
+  - Regex:           '^<(cuda/std/detail/__config|cub/config.cuh|thrust/detail/config.h|thrust/system/cuda/config.h)'
+    Priority:            0
+    SortPriority:        0
+  - Regex:           '^<cub/'
+    Priority:            2
+    SortPriority:        1
+  - Regex:           '^<thrust/'
+    Priority:            3
+    SortPriority:        2
+  - Regex:           '^<cuda/'
+    Priority:            4
+    SortPriority:        3
+  - Regex:           '^<[a-z_]*>$'
+    Priority:            5
+    SortPriority:        4
+  - Regex:           '^<cuda'
+    Priority:            0
+    SortPriority:        0
 InsertBraces: true
+IndentCaseLabels: true
 InsertNewlineAtEOF: true
 InsertTrailingCommas: Wrapped
 IndentRequires: true
 IndentPPDirectives: AfterHash
-PackConstructorInitializers: Never
-PenaltyBreakAssignment: 30
-PenaltyBreakTemplateDeclaration: 0
-PenaltyIndentedWhitespace: 2
-RemoveSemicolon: false
-SpaceAfterLogicalNot: false
-SpaceAfterTemplateKeyword: true
-SpaceBeforeCtorInitializerColon: true
-SpaceBeforeInheritanceColon: true
-SpaceBeforeRangeBasedForLoopColon: true
-
-
-CommentPragmas:  '^ IWYU pragma:'
-CompactNamespaces: false
-ConstructorInitializerAllOnOneLineOrOnePerLine: false
-ConstructorInitializerIndentWidth: 4
-ContinuationIndentWidth: 2
-Cpp11BracedListStyle: true
-SpaceBeforeCpp11BracedList: false
-ExperimentalAutoDetectBinPacking: false
-ForEachMacros:   [ foreach, Q_FOREACH, BOOST_FOREACH ]
-IndentCaseLabels: true
-FixNamespaceComments: true
-IndentWrappedFunctionNames: false
+IndentWidth: 2
 KeepEmptyLinesAtTheStartOfBlocks: false
-MacroBlockBegin: ''
-MacroBlockEnd:   ''
-#JavaScriptQuotes: Double
 MaxEmptyLinesToKeep: 1
+Macros:
+- _LIBCUDACXX_TEMPLATE(...)=template<...>
+- _LIBCUDACXX_REQUIRES(...)=requires (...)
 NamespaceIndentation: None
-ObjCBlockIndentWidth: 4
-#ObjCSpaceAfterProperty: true
-#ObjCSpaceBeforeProtocolList: true
+PackConstructorInitializers: Never
+PenaltyBreakAssignment: 30
 PenaltyBreakBeforeFirstCallParameter: 50
 PenaltyBreakComment: 0
 PenaltyBreakFirstLessLess: 0
 PenaltyBreakString: 70
+PenaltyBreakTemplateDeclaration: 0
 PenaltyExcessCharacter: 100
 PenaltyReturnTypeOnItsOwnLine: 90
+PenaltyIndentedWhitespace: 2
 PointerAlignment: Left
+ReflowComments: true
+RemoveSemicolon: false
+SortIncludes: CaseInsensitive
 SpaceAfterCStyleCast: true
+SpaceAfterLogicalNot: false
+SpaceAfterTemplateKeyword: true
 SpaceBeforeAssignmentOperators: true
+SpaceBeforeCpp11BracedList: false
+SpaceBeforeCtorInitializerColon: true
+SpaceBeforeInheritanceColon: true
 SpaceBeforeParens: ControlStatements
+SpaceBeforeRangeBasedForLoopColon: true
 SpaceInEmptyParentheses: false
 SpacesBeforeTrailingComments: 1
 SpacesInAngles: Never
-SpacesInContainerLiterals: true
 SpacesInCStyleCastParentheses: false
 SpacesInParentheses: false
 SpacesInSquareBrackets: false
-#SpaceAfterTemplateKeyword: true
-#SpaceBeforeInheritanceColon: true
-
-#SortUsingDeclarations: true
-SortIncludes: CaseInsensitive
-
-ReflowComments: true
-
-#IncludeBlocks: Preserve
-#IndentPPDirectives: AfterHash
-
+# Standard: c++20 not supported by rocThrust
+Standard: c++17
 StatementMacros: [
+  # rocThrust
   'THRUST_EXEC_CHECK_DISABLE',
-  'THRUST_NAMESPACE_BEGIN',
-  'THRUST_NAMESPACE_END',
-  'THRUST_EXEC_CHECK_DISABLE',
+  # thrust
+  '_CCCL_EXEC_CHECK_DISABLE',
   'CUB_NAMESPACE_BEGIN',
   'CUB_NAMESPACE_END',
   'THRUST_NAMESPACE_BEGIN',
   'THRUST_NAMESPACE_END',
   '_LIBCUDACXX_BEGIN_NAMESPACE_STD',
   '_LIBCUDACXX_END_NAMESPACE_STD',
+  '_LIBCUDACXX_BEGIN_NAMESPACE_STD_NOVERSION',
+  '_LIBCUDACXX_END_NAMESPACE_STD_NOVERSION',
+  '_LIBCUDACXX_BEGIN_NAMESPACE_RANGES',
+  '_LIBCUDACXX_END_NAMESPACE_RANGES',
+  '_LIBCUDACXX_BEGIN_NAMESPACE_RANGES_ABI',
+  '_LIBCUDACXX_END_NAMESPACE_RANGES_ABI',
+  '_LIBCUDACXX_BEGIN_NAMESPACE_CPO',
+  '_LIBCUDACXX_END_NAMESPACE_CPO',
+  '_LIBCUDACXX_BEGIN_NAMESPACE_VIEWS',
+  '_LIBCUDACXX_END_NAMESPACE_VIEWS',
 ]
 TabWidth: 2
 UseTab: Never
----
diff --git a/.githooks/install b/.githooks/install
index a98d897..31506e6 100755
--- a/.githooks/install
+++ b/.githooks/install
@@ -3,6 +3,6 @@
 cd "$(git rev-parse --git-dir)"
 cd hooks
 
-echo "Installing hooks..." 
+echo "Installing hooks..."
 ln -s ../../.githooks/pre-commit pre-commit
 echo "Done!"
diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index 809934d..430fdb8 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 27245ad..95422e6 100644
--- a/.gitlab-ci.yml
+++ b/.gitlab-ci.yml
@@ -1,5 +1,5 @@
 # ########################################################################
-# Copyright 2019-2024 Advanced Micro Devices, Inc.
+# Copyright 2019-2025 Advanced Micro Devices, Inc.
 # ########################################################################
 
 include:
@@ -9,6 +9,7 @@ include:
       - /defaults.yaml
       - /deps-cmake.yaml
       - /deps-docs.yaml
+      - /deps-format.yaml
       - /deps-rocm.yaml
       - /deps-windows.yaml
       - /deps-nvcc.yaml
@@ -23,12 +24,20 @@ stages:
   - test      # Tests if unit tests are passing (CTest)
   - benchmark # Runs the non-internal benchmarks (Google Benchmark)
 
+workflow:
+  rules:
+    - if: $CI_MERGE_REQUEST_LABELS !~ /CI Skip/
+
 variables:
   # Helper variables
   PACKAGE_DIR: $BUILD_DIR/package
   ROCPRIM_GIT_BRANCH: develop_stream
   ROCPRIM_DIR: ${CI_PROJECT_DIR}/rocPRIM
 
+clang-format:
+  extends:
+    - .lint:clang-format
+
 copyright-date:
   extends:
     - .deps:rocm
@@ -74,7 +83,6 @@ copyright-date:
       -D CMAKE_CXX_COMPILER=hipcc
       -D CMAKE_BUILD_TYPE=Release
       -D BUILD_TEST=OFF
-      -D BUILD_HIPSTDPAR_TEST=OFF
       -D BUILD_EXAMPLE=OFF
       -D ROCM_DEP_ROCMCORE=OFF
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
@@ -116,7 +124,6 @@ copyright-date:
       -D CMAKE_BUILD_TYPE=$BUILD_TYPE
       -D BUILD_$BUILD_TARGET=ON
       -D GPU_TARGETS=$GPU_TARGETS
-      -D AMDGPU_TEST_TARGETS=$GPU_TARGETS
       -D RNG_SEED_COUNT=$rng_seed_count
       -D PRNG_SEEDS=$prng_seeds
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
@@ -128,7 +135,7 @@ copyright-date:
     - cmake --build $CI_PROJECT_DIR/build
   artifacts:
     paths:
-    - $CI_PROJECT_DIR/build/benchmarks/*
+    - $CI_PROJECT_DIR/build/benchmark/*
     - $CI_PROJECT_DIR/build/test/*
     - $CI_PROJECT_DIR/build/testing/*
     - $CI_PROJECT_DIR/build/deps/*
@@ -137,7 +144,7 @@ copyright-date:
     - $CI_PROJECT_DIR/build/.ninja_log
     exclude:
     - $CI_PROJECT_DIR/build/**/*.o
-    expire_in: 2 weeks
+    expire_in: 1 day
 
 build:cmake-latest:
   stage: build
@@ -147,8 +154,8 @@ build:cmake-latest:
   parallel:
     matrix:
       - BUILD_TYPE: Release
-        BUILD_TARGET: [BENCHMARKS, TEST, EXAMPLES]
-        BUILD_VERSION: [14, 17]
+        BUILD_TARGET: [BENCHMARK, TEST, EXAMPLE]
+        BUILD_VERSION: 17
 
 build:cmake-minimum:
   stage: build
@@ -158,8 +165,8 @@ build:cmake-minimum:
   parallel:
     matrix:
       - BUILD_TYPE: Release
-        BUILD_TARGET: [BENCHMARKS, TEST, EXAMPLES]
-        BUILD_VERSION: 14
+        BUILD_TARGET: [BENCHMARK, TEST, EXAMPLE]
+        BUILD_VERSION: 17
 
 build:package:
   stage: build
@@ -183,7 +190,7 @@ build:package:
     paths:
       - $PACKAGE_DIR/rocthrust*.deb
       - $PACKAGE_DIR/rocthrust*.zip
-    expire_in: 2 weeks
+    expire_in: 1 day
 
 build:windows:
   stage: build
@@ -219,19 +226,20 @@ build:windows:
       -D CMAKE_BUILD_TYPE=Release
       -D GPU_TARGETS=$GPU_TARGET
       -D BUILD_TEST=ON
-      -D BUILD_EXAMPLES=OFF
-      -D BUILD_BENCHMARKS=OFF
+      -D BUILD_EXAMPLE=OFF
+      -D BUILD_BENCHMARK=OFF
+      -D CMAKE_TOOLCHAIN_FILE=$CI_PROJECT_DIR/toolchain-windows.cmake
       -D CMAKE_CXX_FLAGS=-Wno-deprecated-declarations
       -D CMAKE_CXX_COMPILER:FILEPATH="${env:HIP_PATH}/bin/clang++.exe"
       -D CMAKE_INSTALL_PREFIX:PATH="$CI_PROJECT_DIR/build/install"
-      -D CMAKE_CXX_STANDARD=14
+      -D CMAKE_CXX_STANDARD=17
       -D CMAKE_PREFIX_PATH:PATH="$ROCPRIM_DIR/build/install;${env:HIP_PATH}" *>&1
     - \& cmake --build "$CI_PROJECT_DIR/build" *>&1
   artifacts:
     paths:
       - $CI_PROJECT_DIR/build/
       - $ROCPRIM_DIR/build/install
-    expire_in: 2 weeks
+    expire_in: 1 day
 
 test:package:
   stage: test
@@ -268,20 +276,25 @@ test:doc:
   extends:
     - .build:docs
     - .rules:test
+  artifacts:
+    paths:
+      - $DOCS_DIR/_build/html/
+    expire_in: 2 weeks
 
-test:
+.test:rocm:
   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 $CI_PROJECT_DIR/build
     - cmake
@@ -294,12 +307,80 @@ test:
     - HSA_ENABLE_SDMA=0 ctest
       --output-on-failure
       --repeat-until-fail 2
-      --tests-regex $GPU_TARGET
       --resource-spec-file ./resources.json
       --parallel $PARALLEL_JOBS
 
+test:rocm-any-gpu:
+  variables:
+    GPU: ""
+    PARALLEL_JOBS: 1
+  extends:
+    - .test:rocm
+  rules:
+    - if: $CI_MERGE_REQUEST_TITLE =~ /Draft:/ && $CI_MERGE_REQUEST_LABELS !~ /Arch::/
+
+test:rocm-label-arch:
+  extends:
+    - .gpus:rocm
+    - .test:rocm
+    - .rules:arch-labels
+
+test:rocm-all-gpus:
+  variables:
+    SHOULD_BE_UNDRAFTED: "true"
+  extends:
+    - .gpus:rocm
+    - .test:rocm
+    - .rules:test
+
+test-bitwise-repro-generate:
+  stage: test
+  extends:
+    - .cmake-minimum
+    - .rules:manual
+    - .gpus:rocm
+  needs:
+    - job: build:cmake-minimum
+      parallel:
+        matrix:
+          - BUILD_TYPE: Release
+            BUILD_TARGET: TEST
+            BUILD_VERSION: 17
+  cache:
+    key: database
+    paths:
+      - $CI_PROJECT_DIR/bitwise.db
+    policy: push
+  script:
+    - cd $CI_PROJECT_DIR/build/test/
+    - ROCTHRUST_BWR_PATH=$CI_PROJECT_DIR/bitwise.db ROCTHRUST_BWR_GENERATE=1 ./reproducibility.hip
+
+test-bitwise-repro:
+  stage: test
+  extends:
+    - .cmake-minimum
+    - .rules:test
+    - .gpus:rocm
+  needs:
+    - job: build:cmake-minimum
+      parallel:
+        matrix:
+          - BUILD_TYPE: Release
+            BUILD_TARGET: TEST
+            BUILD_VERSION: 17
+  cache:
+    key: database
+    paths:
+      - $CI_PROJECT_DIR/bitwise.db
+    policy: pull
+  script:
+    - cd $CI_PROJECT_DIR/build/test/
+    - ROCTHRUST_BWR_PATH=$CI_PROJECT_DIR/bitwise.db ./reproducibility.hip
+
 .rocm-windows:test:
   extends:
+    - .deps:rocm-windows
+    - .deps:visual-studio-devshell
     - .gpus:rocm-windows
     - .rules:test
   stage: test
@@ -358,19 +439,21 @@ build:cuda-and-omp:
   tags:
     - build
   variables:
-    CCCL_GIT_BRANCH: v2.5.0
+    CCCL_GIT_BRANCH: v2.8.5
     CCCL_DIR: ${CI_PROJECT_DIR}/cccl
   needs: []
   script:
     - git clone -b $CCCL_GIT_BRANCH https://github.com/NVIDIA/cccl.git $CCCL_DIR
     # Replace CCCL Thrust headers with rocThrust headers
-    - rm -R $CCCL_DIR/thrust/thrust
+    - rm -r $CCCL_DIR/thrust/thrust
     - cp -r $CI_PROJECT_DIR/thrust $CCCL_DIR/thrust
+    - sed -i '/include(cmake\/ThrustAddSubdir.cmake)/i \  include(../cmake/CCCLAddSubdirHelper.cmake)' $CCCL_DIR/thrust/CMakeLists.txt
     # Build tests and examples from CCCL Thrust
     - cmake
       -G Ninja
       -D CMAKE_BUILD_TYPE=Release
       -D CMAKE_CUDA_ARCHITECTURES="$GPU_TARGETS"
+      -D CCCL_SOURCE_DIR="$CCCL_DIR"
       -D THRUST_ENABLE_TESTING=ON
       -D THRUST_ENABLE_EXAMPLES=ON
       -D THRUST_ENABLE_BENCHMARKS=OFF
@@ -379,6 +462,7 @@ build:cuda-and-omp:
       -D THRUST_MULTICONFIG_ENABLE_SYSTEM_CUDA=ON
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
+      -D THRUST_MULTICONFIG_ENABLE_DIALECT_CPP14=FALSE
       -D CMAKE_CUDA_COMPILER_LAUNCHER=phc_sccache_cuda
       -B $CI_PROJECT_DIR/build
       -S $CCCL_DIR/thrust
@@ -402,16 +486,17 @@ build:cuda-and-omp:
       - $CCCL_DIR/thrust/cmake/ThrustRunTest.cmake
       - $CCCL_DIR/thrust/cmake/ThrustRunExample.cmake
       - $CI_PROJECT_DIR/build/.ninja_log
-    expire_in: 1 week
+    expire_in: 1 day
 
-test:cuda-and-omp:
+.test:cuda-and-omp:
   stage: test
+  tags:
+    - nvcc
+    - $GPU
   needs:
     - build:cuda-and-omp
   extends:
     - .nvcc
-    - .gpus:nvcc
-    - .rules:test
   before_script:
     # This is only needed because of the legacy before_script in .gpus:nvcc would otherwise overwrite before_script
     - !reference [.nvcc, before_script]
@@ -420,6 +505,23 @@ test:cuda-and-omp:
     # These tests are executed on the build stage because they require sources
     - ctest --output-on-failure --exclude-regex "thrust.example.cmake.add_subdir|thrust.test.cmake.check_source_files"
 
+test:cuda-and-omp-any-gpu:
+  variables:
+    GPU: ""
+    PARALLEL_JOBS: 1
+  extends:
+    - .test:cuda-and-omp
+  rules:
+    - if: $CI_MERGE_REQUEST_TITLE =~ /Draft:/
+
+test:cuda-and-omp-all-gpus:
+  variables:
+    SHOULD_BE_UNDRAFTED: "true"
+  extends:
+    - .gpus:nvcc
+    - .test:cuda-and-omp
+    - .rules:test
+
 .benchmark-base:
   stage: benchmark
   extends:
@@ -430,7 +532,12 @@ test:cuda-and-omp:
 
 benchmark:
   needs:
-    - build:cmake-minimum
+    - job: build:cmake-minimum
+      parallel:
+        matrix:
+          - BUILD_TYPE: Release
+            BUILD_TARGET: BENCHMARK
+            BUILD_VERSION: 17
   extends:
     - .cmake-minimum
     - .gpus:rocm
@@ -445,7 +552,7 @@ benchmark:
     - mkdir -p "${BENCHMARK_RESULT_DIR}"
     - python3
       .gitlab/run_benchmarks.py
-      --benchmark_dir "${CI_PROJECT_DIR}/build/benchmarks"
+      --benchmark_dir "${CI_PROJECT_DIR}/build/benchmark"
       --benchmark_gpu_architecture "${GPU_TARGET}"
       --benchmark_output_dir "${BENCHMARK_RESULT_DIR}"
       --benchmark_filename_regex "${BENCHMARK_FILENAME_REGEX}"
@@ -453,4 +560,4 @@ benchmark:
   artifacts:
     paths:
       - ${BENCHMARK_RESULT_DIR}
-    expire_in: 1 week
+    expire_in: 1 day
diff --git a/.gitlab/run_benchmarks.py b/.gitlab/run_benchmarks.py
index a3fe982..6a94c25 100755
--- a/.gitlab/run_benchmarks.py
+++ b/.gitlab/run_benchmarks.py
@@ -66,7 +66,7 @@ def run_benchmarks(benchmark_context):
             subprocess.check_call(args)
             end_time = time.time()
             duration = end_time - start_time
-            
+
             print(f'Benchmark {benchmark_name} took {duration:.3f} seconds to run', file=sys.stderr, flush=True)
         except subprocess.CalledProcessError as error:
             print(f'Could not run benchmark at {benchmark_path}. Error: "{error}"', file=sys.stderr, flush=True)
diff --git a/.jenkins/common.groovy b/.jenkins/common.groovy
deleted file mode 100644
index 27b0a24..0000000
--- a/.jenkins/common.groovy
+++ /dev/null
@@ -1,78 +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, boolean sameOrg=true)
-{
-    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 getDependenciesCommand = ""
-    if (project.installLibraryDependenciesFromCI)
-    {
-        project.libraryDependencies.each
-        {
-            libraryName ->
-            getDependenciesCommand += auxiliary.getLibrary(libraryName, platform.jenkinsLabel, 'develop', sameOrg)
-        }
-    }
-
-    def command = """#!/usr/bin/env bash
-                set -x
-                ${getDependenciesCommand}
-                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 hmmTestCommand = ''
-    // Note: temporarily disable scan tests below while waiting for a compiler fix
-    def excludeRegex = /(reduce_by_key.hip|scan)/
-    testCommandExclude = "--exclude-regex \"${excludeRegex}\""
-
-    if (platform.jenkinsLabel.contains('gfx90a'))
-    {
-        hmmTestCommand = ""
-                        // temporarily disable hmm testing
-                        //  """
-                        //     export HSA_XNACK=1
-                        //     export ROCTHRUST_USE_HMM=1
-                        //     ${testCommand} ${testCommandExclude}
-                        //  """
-    }
-
-    def command = """
-                    #!/usr/bin/env bash
-                    set -x
-                    cd ${project.paths.project_build_prefix}
-                    cd ${project.testDirectory}
-                    ${testCommand} ${testCommandExclude}
-                    ${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 df19050..0000000
--- a/.jenkins/precheckin.groovy
+++ /dev/null
@@ -1,93 +0,0 @@
-#!/usr/bin/env groovy
-// This shared library is available at https://github.com/ROCmSoftwarePlatform/rocJENKINS/
-@Library('rocJenkins@pong') _
-
-// This file is for internal AMD use.
-// If you are interested in running your own Jenkins, please raise a github issue for assistance.
-
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path;
-
-def runCI =
-{
-    nodeDetails, jobName->
-
-    def prj = new rocProject('rocThrust', 'precheckin')
-
-    prj.defaults.ccache = true
-    prj.timeout.compile = 420
-    prj.libraryDependencies = ["rocPRIM"]
-
-    // Define test architectures, optional rocm version argument is available
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    boolean formatCheck = false
-
-    def commonGroovy
-
-    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)
-
-    Set standardJobNameSet = ["compute-rocm-dkms-no-npi", "compute-rocm-dkms-no-npi-hipclang", "rocm-docker"]
-
-    def jobNameList = ["compute-rocm-dkms-no-npi":([ubuntu16:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx908']]),
-                       "compute-rocm-dkms-no-npi-hipclang":([ubuntu16:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx908']]),
-                       "rocm-docker":([ubuntu16:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx908']])]
-    jobNameList = auxiliary.appendJobNameList(jobNameList)
-
-    auxiliary.registerDependencyBranchParameter(["rocPRIM"])
-
-    propertyList.each
-    {
-        jobName, property->
-        if (urlJobName == jobName)
-            properties(auxiliary.addCommonProperties(property))
-    }
-
-    Set seenJobNames = []
-    jobNameList.each
-    {
-        jobName, nodeDetails->
-        seenJobNames.add(jobName)
-        if (urlJobName == jobName)
-            runCI(nodeDetails, jobName)
-    }
-
-    // For url job names that are outside of the standardJobNameSet i.e. compute-rocm-dkms-no-npi-1901
-    if(!seenJobNames.contains(urlJobName))
-    {
-        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * *')])]))
-        runCI([ubuntu16:['gfx906']], urlJobName)
-    }
-}
diff --git a/.jenkins/staticanalysis.groovy b/.jenkins/staticanalysis.groovy
deleted file mode 100644
index 7d2cf06..0000000
--- a/.jenkins/staticanalysis.groovy
+++ /dev/null
@@ -1,46 +0,0 @@
-#!/usr/bin/env groovy
-// This shared library is available at https://github.com/ROCmSoftwarePlatform/rocJENKINS/
-@Library('rocJenkins@pong') _
-
-// This is file for internal AMD use.
-// If you are interested in running your own Jenkins, please raise a github issue for assistance.
-
-import com.amd.project.*
-import com.amd.docker.*
-import java.nio.file.Path
-
-def runCompileCommand(platform, project, jobName, boolean debug=false)
-{
-    project.paths.construct_build_prefix()
-}
-
-def runCI = 
-{
-    nodeDetails, jobName->
-
-    def prj  = new rocProject('rocThrust', 'StaticAnalysis')
-
-    // Define test architectures, optional rocm version argument is available
-    def nodes = new dockerNodes(nodeDetails, jobName, prj)
-
-    boolean formatCheck = false
-    boolean staticAnalysis = true
-
-    def compileCommand =
-    {
-        platform, project->
-
-        runCompileCommand(platform, project, jobName, false)
-    }
-
-    buildProject(prj , formatCheck, nodes.dockerArray, compileCommand, null, null, staticAnalysis)
-}
-
-ci: { 
-    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
-
-    properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 1 * * 6')])]))
-    stage(urlJobName) {
-        runCI([ubuntu20:['any']], urlJobName)
-    }
-}
diff --git a/.jenkins/staticlibrary.groovy b/.jenkins/staticlibrary.groovy
deleted file mode 100644
index 08f9e3c..0000000
--- a/.jenkins/staticlibrary.groovy
+++ /dev/null
@@ -1,88 +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('rocThrust', 'Static Library PreCheckin')
-
-    prj.defaults.ccache = true
-    prj.timeout.compile = 420
-    prj.libraryDependencies = ["rocPRIM"]
-
-    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, false, true)
-    }
-
-
-    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":([ubuntu16:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx908']]),
-                       "compute-rocm-dkms-no-npi-hipclang":([ubuntu16:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx908']]),
-                       "rocm-docker":([ubuntu16:['gfx900'],centos7:['gfx906'],sles15sp1:['gfx908']])]
-    jobNameList = auxiliary.appendJobNameList(jobNameList)
-
-    auxiliary.registerDependencyBranchParameter(["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 35d3225..a5e4ced 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -3,6 +3,79 @@
 Documentation for rocThrust available at
 [https://rocm.docs.amd.com/projects/rocThrust/en/latest/](https://rocm.docs.amd.com/projects/rocThrust/en/latest/).
 
+## rocThrust 4.1.0 for ROCm 7.1
+
+### Added
+
+* Added a new CMake option `-DSQLITE_USE_SYSTEM_PACKAGE` to allow SQLite to be provided by the system.
+* Introduced `libhipcxx` as a soft depedency. When `liphipcxx` can be included, rocthrust, may use structs and methods defined in `libhipcxx`. This allows for a more complete behaviour parity with CCCL and mirrors CCCL's thrust own depedency on `libcudacxx`.
+* Added a new CMake option `-DUSE_SYSTEM_LIB` to allow tests to be built from `ROCm` libraries provided by the system.
+
+### Known Issues
+* `event` test is failing on CI and local runs on MI300, MI250 and MI210.
+
+* rocThrust, as well as its dependencies rocPRIM and rocRAND have been moved into the new rocm-libraries "monorepo" repository (https://github.com/ROCm/rocm-libraries). This repository contains a number of ROCm libraries that are frequently used together.
+  * The repository migration requires a few changes to the way that rocThrust's ROCm library dependencies are fetched.
+  * There are new cmake options for obtaining rocPRIM and (optionally, if BUILD_BENCHMARKS is enabled) rocRAND.
+  * cmake build options `ROCPRIM_FETCH_METHOD` and `ROCRAND_FETCH_METHOD` may be set to one of the following:
+    * `PACKAGE` - (default) searches for a preinstalled packaged version of the dependency. If it is not found, the build will fall back using option `DOWNLOAD`, described below.
+    * `DOWNLOAD` - downloads the dependency from the rocm-libraries repository. If git >= 2.25 is present, this option uses a sparse checkout that avoids downloading more than it needs to. If not, the whole monorepo is downloaded (this may take some time).
+    * `MONOREPO` - this options is intended to be used if you are building rocThrust from within a copy of the rocm-libraries repository that you have cloned (and therefore already contains the dependencies rocPRIM and rocRAND). When selected, the build will try find the dependency in the local repository tree. If it cannot be found, the build will attempt to add it to the local tree using a sparse-checkout. If that also fails, it will fall back to using the `DOWNLOAD` option.
+
+### Changed
+
+* The previously hidden cmake build option `FORCE_DEPENDENCIES_DOWNLOAD` has been unhidden and renamed `EXTERNAL_DEPS_FORCE_DOWNLOAD` to differentiate it from the new rocPRIM and rocRAND dependency options described above. It's behaviour remains the same - it forces non-ROCm dependencies (Google Benchmark, Google Test, and SQLite) to be downloaded instead of searching for existing installed packages. This option defaults to `OFF`.
+
+### Removed
+  * The previous dependency-related build options `DOWNLOAD_ROCPRIM` and `DOWNLOAD_ROCRAND` have been removed. Please use `ROCPRIM_FETCH_METHOD=DOWNLOAD` and `ROCRAND_FETCH_METHOD=DOWNLOAD` instead.
+
+## rocThrust 4.0.0 for ROCm 7.0
+
+### Changed
+
+* Updated the required version of Google Benchmark from 1.8.0 to 1.9.0.
+* Drop `c++14` support for rocthrust.
+* Renamed `cpp14_required.h` to `cpp_version_check.h`
+* Refactored `test_header.hpp` into separte modules `test_param_fixtures.hpp`, `test_real_assertions.hpp`, `test_imag_assertions.hpp`, and `test_utils.hpp`.
+  * This is done to prevent unit tests from having access to modules that they're not testing. This will improve the accuracy of code coverage reports.
+
+### Added
+* Additional unit tests for:
+  * binary_search
+  * complex
+  * c99math
+  * catrig
+  * ccosh
+  * cexp
+  * clog
+  * csin
+  * csqrt
+  * ctan
+* Added `test_param_fixtures.hpp` to store all the parameters for typed test suites.
+* Added `test_real_assertions.hpp` to handle unit test assertions for real numbers.
+* Added `test_imag_assertions.hpp` to handle unit test assertions for imaginary numbers.
+* `clang++` is now used to compile google benchmarks on Windows.
+* Added gfx950 support.
+* Merged changes from upstream CCCL/thrust 2.6.0
+
+### Removed
+
+* `device_malloc_allocator.h` has been removed. This header file was unused and should not impact users.
+* Removed C++14 support, only C++17 is supported.
+* `test_header.hpp` has been removed. The `HIP_CHECK` function, as well as the `test` and `inter_run_bwr` namespaces, have been moved to `test_utils.hpp`.
+* `test_assertions.hpp` has been split into `test_real_assertions.hpp` and `test_imag_assertions.hpp`.
+
+### Upcoming changes
+
+* `thrust::device_malloc_allocator` is deprecated as of this version. It will be removed in an upcoming version.
+
+### Resolved issues
+
+* Fixed an issue with internal calls to unqualified `distance()` which would be ambigious due to also visibile implementation through ADL.
+
+### Known Issues
+* The order of the values being compared by thrust::exclusive_scan_by_key and thrust::inclusive_scan_by_key can change between runs when integers are being compared. This can cause incorrect output when a non-commutative operator such as division is being used.
+
 ## rocThrust 3.3.0 for ROCm 6.4
 
 ### Added
@@ -26,11 +99,10 @@ Documentation for rocThrust available at
 * `--test|-t` is no longer a required flag for `rtest.py`. Instead, the user can use either `--emulation|-e` or `--test|-t`, but not both.
 * Split the contents of HIPSTDPAR's forwarding header into several implementation headers.
 * Fixed `copy_if` to work with large data types (512 bytes)
-* Updated the required version of Google Benchmark from 1.8.0 to 1.9.0.
 
 ### Known Issues
-*  `thrust::inclusive_scan_by_key` might produce incorrect results when it's used with -O2 or -O3 optimization.  
-  - The error is caused by a recent compiler change. There is a fix available that will be released at a later date. 
+*  `thrust::inclusive_scan_by_key` might produce incorrect results when it's used with -O2 or -O3 optimization.
+  - The error is caused by a recent compiler change. There is a fix available that will be released at a later date.
 
 ## rocThrust 3.2.0 for ROCm 6.3
 
@@ -109,7 +181,6 @@ Documentation for rocThrust available at
 * Updated `docs` directory structure to match the standard of [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core).
 * Removed references to and workarounds for deprecated hcc
 
-
 ## rocThrust 2.17.0 for ROCm 5.5
 
 ### Additions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 887c1e5..5413549 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,9 +1,22 @@
 # ########################################################################
-# Copyright 2019-2024 Advanced Micro Devices, Inc.
+# Copyright 2019-2025 Advanced Micro Devices, Inc.
 # ########################################################################
 
 cmake_minimum_required(VERSION 3.10.2 FATAL_ERROR)
 
+# --------------------------------------
+# Update these variables at release time
+#
+# Set the library version
+set(VERSION_STRING "4.1.0")
+# Set the minimum required rocPRIM version
+set(MIN_ROCPRIM_PACKAGE_VERSION "4.0.0" CACHE STRING "Minimum version of rocPRIM to search for when ROCPRIM_FETCH_METHOD is set to PACKAGE.")
+# Set the minimum required rocRAND version
+set(MIN_ROCRAND_PACKAGE_VERSION "4.0.0" CACHE STRING "Minimum version of rocRAND to search for when ROCRAND_FETCH_METHOD is set to PACKAGE.")
+# Set download branch for dependencies rocPRIM and rocRAND
+set(ROCM_DEP_RELEASE_BRANCH "release/rocm-rel-7.1" CACHE STRING "Download branch for ROCm dependencies")
+# --------------------------------------
+
 # Install prefix
 if(WIN32)
   set(CMAKE_INSTALL_PREFIX ${PROJECT_BINARY_DIR}/package CACHE PATH "Install path prefix, prepended onto install directories")
@@ -22,6 +35,12 @@ endif()
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
 set(CMAKE_CXX_EXTENSIONS OFF)
 
+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")
+endif()
+
 # Set HIP flags
 set(CMAKE_HIP_STANDARD 17)
 set(CMAKE_HIP_STANDARD_REQUIRED ON)
@@ -34,20 +53,23 @@ include(CMakeDependentOption)
 # Disable -Werror
 option(DISABLE_WERROR "Disable building with Werror" ON)
 option(BUILD_TEST "Build tests" OFF)
+option(CODE_COVERAGE "Enable code coverage" OFF)
 option(BUILD_HIPSTDPAR_TEST "Build hipstdpar tests" OFF)
 option(BUILD_HIPSTDPAR_TEST_WITH_TBB "Build hipstdpar tests with TBB" OFF)
-option(BUILD_EXAMPLES "Build examples" OFF)
-option(BUILD_BENCHMARKS "Build benchmarks" OFF)
-option(DOWNLOAD_ROCPRIM "Download rocPRIM and do not search for rocPRIM package" OFF)
-option(DOWNLOAD_ROCRAND "Download rocRAND and do not search for rocRAND package" OFF)
+option(BUILD_EXAMPLE "Build example" OFF)
+option(BUILD_BENCHMARK "Build benchmark" OFF)
 option(BUILD_ADDRESS_SANITIZER "Build with address sanitizer enabled" OFF)
 cmake_dependent_option(ENABLE_UPSTREAM_TESTS "Enable upstream (thrust) tests" ON BUILD_TEST OFF)
-#Set the header wrapper OFF by default.
-option(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY "Build with file/folder reorg with backward compatibility enabled" OFF)
+cmake_dependent_option(USE_SYSTEM_LIB "Use existing system ROCm library installation when building tests" OFF BUILD_TEST OFF)
+option(EXTERNAL_DEPS_FORCE_DOWNLOAD "Force download of non-ROCm dependencies (eg. Google Test, Google Benchmark)" OFF)
 
 check_language(HIP)
 cmake_dependent_option(USE_HIPCXX "Use CMake HIP language support" OFF CMAKE_HIP_COMPILER OFF)
 
+# Allow the user to optionally select offset type dispatch to fixed 32 or 64 bit types
+set(THRUST_DISPATCH_TYPE "Dynamic" CACHE STRING "Select Thrust offset type dispatch." FORCE)
+set_property(CACHE THRUST_DISPATCH_TYPE PROPERTY STRINGS "Dynamic" "Force32bit" "Force64bit")
+
 #Adding CMAKE_PREFIX_PATH
 list( APPEND CMAKE_PREFIX_PATH /opt/rocm/llvm /opt/rocm ${ROCM_PATH} )
 
@@ -59,7 +81,7 @@ list(APPEND CMAKE_MODULE_PATH
 )
 
 # Set a default build type if none was specified
-if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
+if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES AND NOT CODE_COVERAGE)
   message(STATUS "Setting build type to 'Release' as none was specified.")
   set(CMAKE_BUILD_TYPE "Release" CACHE STRING "Choose the type of build." FORCE)
   set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "" "Debug" "Release" "MinSizeRel" "RelWithDebInfo")
@@ -94,17 +116,37 @@ 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()
 
+# Set up options for obtaining dependencies rocPRIM and rocRAND.
+# PACKAGE: Search for an install package that contains the dependency.
+# MONOREPO: Assume this is a monorepo checkout and search for the dependency in the directory at ../../projects/.
+# DOWNLOAD: Download the dependency from the monorepo.
+set(FETCH_METHOD_OPTIONS "PACKAGE" "MONOREPO" "DOWNLOAD")
+
+set(ROCPRIM_FETCH_METHOD "PACKAGE" CACHE STRING "How to obtain the rocPRIM dependency")
+set(ROCRAND_FETCH_METHOD "PACKAGE" CACHE STRING "How to obtain the rocRAND dependency")
+
+# This function checks to see if the fetch method variable it's passed is defined, and contains a valid value.
+# If it does not contain a valid value, it issues a fatal failure with an error message.
+function(check_fetch_method method)
+  if (DEFINED ${method} AND NOT ${${method}} IN_LIST FETCH_METHOD_OPTIONS)
+    message(FATAL_ERROR "Unrecognized ${method}: \"${${method}}\". Valid options are: ${FETCH_METHOD_OPTIONS}.")
+  endif()
+endfunction()
+
+check_fetch_method(ROCPRIM_FETCH_METHOD)
+check_fetch_method(ROCRAND_FETCH_METHOD)
+
 # Get dependencies
 include(cmake/Dependencies.cmake)
 
@@ -130,16 +172,13 @@ if (NOT THRUST_HOST_SYSTEM IN_LIST THRUST_HOST_SYSTEM_OPTIONS)
   )
 endif ()
 
-if(DISABLE_WERROR)
-  add_compile_options(-Wall -Wextra)
-else()
-  add_compile_options(-Wall -Wextra -Werror)
+set(COMPILE_OPTIONS -Wall -Wextra)
+if(NOT DISABLE_WERROR)
+  list(APPEND COMPILE_OPTIONS -Werror)
 endif()
 
-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 (WIN32)
@@ -149,31 +188,43 @@ endif()
 
 # Address Sanitizer
 if(BUILD_ADDRESS_SANITIZER)
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsanitize=address -shared-libasan")
-  set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -fsanitize=address -shared-libasan")
+  list(APPEND COMPILE_OPTIONS -fsanitize=address -shared-libasan)
   add_link_options(-fuse-ld=lld)
 endif()
 
 # Setup VERSION
-rocm_setup_version(VERSION "3.3.0")
-
-# Print configuration summary
-include(cmake/Summary.cmake)
-print_configuration_summary()
+rocm_setup_version(VERSION ${VERSION_STRING})
+math(EXPR rocthrust_VERSION_NUMBER "${rocthrust_VERSION_MAJOR} * 100000 + ${rocthrust_VERSION_MINOR} * 100 + ${rocthrust_VERSION_PATCH}")
 
 # Thrust (with HIP backend)
 add_subdirectory(thrust)
 
-if(BUILD_TEST OR BUILD_BENCHMARKS OR BUILD_HIPSTDPAR_TEST)
+if(BUILD_TEST OR BUILD_BENCHMARK OR BUILD_HIPSTDPAR_TEST)
   rocm_package_setup_component(clients)
 endif()
 
+if(CODE_COVERAGE)
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -O0 -fprofile-instr-generate -fcoverage-mapping")
+endif()
+
 # Tests
+if(BUILD_TEST AND USE_SYSTEM_LIB)
+  find_package(rocprim REQUIRED CONFIG PATHS "/opt/rocm/rocprim")
+  if (${rocprim_VERSION} VERSION_LESS ${MIN_ROCPRIM_PACKAGE_VERSION})
+    message(WARNING "The installed rocprim version, ${rocprim_VERSION}, is less than the minimum required version ${MIN_ROCPRIM_PACKAGE_VERSION}. Building tests with USE_SYSTEM_LIB=ON may not work properly.")
+  endif()
+  find_package(rocthrust REQUIRED CONFIG PATHS "/opt/rocm/rocthrust")
+  if (NOT ${rocthrust_VERSION} VERSION_EQUAL ${VERSION_STRING})
+    message(WARNING "The installed rocthrust version, ${rocthrust_VERSION}, does not match project version ${VERSION_STRING}. Building tests with USE_SYSTEM_LIB=ON may not work properly.")
+  endif()
+endif()
+
 if(BUILD_TEST OR BUILD_HIPSTDPAR_TEST)
   rocm_package_setup_client_component(tests)
   if (ENABLE_UPSTREAM_TESTS)
     enable_testing()
   endif()
+
   # We still want the testing to be compiled to catch some errors
   #TODO: Get testing folder working with HIP on Windows
   if (NOT WIN32 AND BUILD_TEST)
@@ -184,28 +235,15 @@ if(BUILD_TEST OR BUILD_HIPSTDPAR_TEST)
 endif()
 
 # Examples
-if(BUILD_EXAMPLES)
+if(BUILD_EXAMPLE)
   add_subdirectory(examples)
 endif()
 
-# Benchmarks
-if(BUILD_BENCHMARKS)
-  add_subdirectory(benchmarks)
-  add_subdirectory(internal/benchmark)
+# Benchmark
+if(BUILD_BENCHMARK)
+  add_subdirectory(benchmark)
 endif()
 
-#Create header wrapper for backward compatibility
-if(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY AND NOT WIN32)
-  rocm_wrap_header_dir(
-      ${CMAKE_SOURCE_DIR}/thrust
-      PATTERNS "*.h" "*.inl" "*.cuh" "*.hpp"
-      HEADER_LOCATION include/thrust
-      GUARDS SYMLINK WRAPPER
-      WRAPPER_LOCATIONS rocthrust/${CMAKE_INSTALL_INCLUDEDIR}/thrust
-      OUTPUT_LOCATIONS rocthrust/wrapper/include/thrust
-  )
-endif( )
-
 set(THRUST_OPTIONS_DEBUG ${THRUST_OPTIONS_WARNINGS})
 set(THRUST_OPTIONS_RELEASE ${THRUST_OPTIONS_WARNINGS})
 
@@ -230,7 +268,11 @@ set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "\${CPACK_PACKAGING_INSTALL_PR
 
 rocm_create_package(
   NAME rocthrust
-  DESCRIPTION "Radeon Open Compute Thrust library"
+  DESCRIPTION "rocThrust is a ROCm port of the Thrust library, written in HIP"
   MAINTAINER "[e-mail]"
   HEADER_ONLY
 )
+
+# Print configuration summary
+include(cmake/Summary.cmake)
+print_configuration_summary()
diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md
index dfe3344..e985798 100644
--- a/CONTRIBUTING.md
+++ b/CONTRIBUTING.md
@@ -6,7 +6,7 @@
 
 # Contributing to rocThrust #
 
-We welcome contributions to rocThrust.  Please follow these details to help ensure your contributions will be successfully accepted.
+We welcome contributions to rocThrust. Please follow these details to help ensure your contributions will be successfully accepted.
 
 ## Issue Discussion ##
 
@@ -97,4 +97,4 @@ During code reviews, another developer will take a look through your proposed ch
 needed), they may leave a comment. You can follow up and respond to the comment, and/or create comments of your own if you have questions or ideas.
 When a modification request has been completed, the conversation thread about it will be marked as resolved.
 
-To update the code in your PR (eg. in response to a code review discussion), you can simply push another commit to the branch used in your pull request.
\ No newline at end of file
+To update the code in your PR (eg. in response to a code review discussion), you can simply push another commit to the branch used in your pull request.
diff --git a/NOTICES.txt b/NOTICES.txt
index dc4f99e..b91ff6f 100644
--- a/NOTICES.txt
+++ b/NOTICES.txt
@@ -2,7 +2,7 @@ Notices and licenses file
 _________________________
 
 
-AMD copyrighted code (Apache 2.0) 
+AMD copyrighted code (Apache 2.0)
 Copyright © 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
 
 Licensed under the Apache License, Version 2.0 (the "License");
@@ -18,33 +18,44 @@ See the License for the specific language governing permissions and
 limitations under the License.
 
 
-crascit-downloadproject v-u (MIT) 
+crascit-downloadproject v-u (MIT)
 # Distributed under the OSI-approved MIT License.  See accompanying
 # file LICENSE or https://github.com/Crascit/DownloadProject for details.
 
-Dependencies on scipy-scipy v-u (MIT) 
-Copyright (C) 2003-2013 SciPy Developers. 
-Modifications Copyright (C) 2019 Advanced Micro Devices, Inc. All rights reserved. 
-
-Redistribution and use in source and binary forms, with or without modification, are permitted provided that the 
-following conditions are met: 
- 
-Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. 
-Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer 
-in the documentation and/or other materials provided with the distribution. 
-Neither the name of Enthought nor the names of the SciPy Developers may be used to endorse or promote products derived from 
-this software without specific prior written permission. 
- 
-THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS “AS IS� AND ANY EXPRESS OR IMPLIED WARRANTIES, 
-INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. 
-IN NO EVENT SHALL THE REGENTS OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, 
-OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, 
-OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, 
-OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE 
-POSSIBILITY OF SUCH DAMAGE. 
-
-
-Dependencies on thrust-thrust v1.9.2 (Apache 2.0) 
+Dependencies on scipy-scipy v-u (BSD-3-Clause)
+Copyright (C) 2003-2013 SciPy Developers.
+Modifications Copyright (C) 2019 Advanced Micro Devices, Inc. All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions
+are met:
+
+1. Redistributions of source code must retain the above copyright
+   notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above
+   copyright notice, this list of conditions and the following
+   disclaimer in the documentation and/or other materials provided
+   with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its
+   contributors may be used to endorse or promote products derived
+   from this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+"AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
+LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
+A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
+LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+
+Dependencies on thrust-thrust v1.9.2 (Apache 2.0)
 Copyright 2008-2013 NVIDIA Corporation
 Modifications Copyright (C) 2019 Advanced Micro Devices, Inc. All rights reserved.
 
@@ -61,7 +72,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 
 
-rocmsoftwareplatform-rocfft v-u (MIT) 
+rocmsoftwareplatform-rocfft v-u (MIT)
 Copyright © 2016 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 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:
@@ -71,7 +82,7 @@ The above copyright notice and this permission notice shall be included in all c
 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
 
 
-thrust-thrust v1.9.2 (Apache 2.0) 
+thrust-thrust v1.9.2 (Apache 2.0)
 Copyright 2008-2013 NVIDIA Corporation
 
 Licensed under the Apache License, Version 2.0 (the "License");
@@ -87,7 +98,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 
 
-thrust-thrust v1.9.2 (BSD3) 
+thrust-thrust v1.9.2 (BSD3)
 Copyright (c) 2011, Duane Merrill.  All rights reserved.
 Copyright (c) 2011-2018, NVIDIA CORPORATION.  All rights reserved.
 
diff --git a/README.md b/README.md
index 8332574..8a74370 100644
--- a/README.md
+++ b/README.md
@@ -1,12 +1,12 @@
 # rocThrust
 
 > [!NOTE]
-> The published documentation is available at [rocThrust](https://rocm.docs.amd.com/projects/rocThrust/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 rocThrust documentation is available [here](https://rocm.docs.amd.com/projects/rocThrust/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).
 
 Thrust is a parallel algorithm library. It has been ported to
 [HIP](https://github.com/ROCm/HIP) and [ROCm](https://www.github.com/ROCm/ROCm), which use
 the [rocPRIM](https://github.com/ROCm/rocPRIM) library. The HIP-ported library
-works on HIP and ROCm software. Currently there is no CUDA backend in place.
+works on HIP and ROCm software.
 
 ## Requirements
 
@@ -35,46 +35,49 @@ For ROCm hardware requirements, refer to:
 * [Linux support](https://rocm.docs.amd.com/projects/install-on-linux/en/latest/reference/system-requirements.html)
 * [Windows support](https://rocm.docs.amd.com/projects/install-on-windows/en/latest/reference/system-requirements.html)
 
-## Documentation
-
-Documentation for rocThrust available at
-[https://rocm.docs.amd.com/projects/rocThrust/en/latest/](https://rocm.docs.amd.com/projects/rocThrust/en/latest/).
-
-You can build our documentation locally using the following commands:
-
-```shell
-# Go to rocThrust docs directory
-cd rocThrust; cd docs
+## Build and install
 
-# Install Python dependencies
-python3 -m pip install -r sphinx/requirements.txt
+### Obtaining the source code
 
-# Build the documentation
-python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html
+rocThrust can be cloned in two ways:
+1.  Clone rocThrust along with other ROCm libraries that are frequently used together (note that this may take some time to complete):
+```sh
+git clone https://github.com/ROCm/rocm-libraries.git
+cd rocm-libraries
+```
 
-# For e.g. serve the HTML docs locally
-cd _build/html
-python3 -m http.server
+2. To clone rocThrust individually (faster, but requires git version 2.25+):
+```sh
+git clone --no-checkout --depth=1 --filter=tree:0 https://github.com/ROCm/rocm-libraries.git
+cd rocm-libraries
+git sparse-checkout init --cone
+git sparse-checkout set projects/rocthrust
+git checkout develop
 ```
 
-## Build and install
+### Building the library
 
 ```sh
-git clone https://github.com/ROCm/rocThrust
+# Go to the rocthrust directory
+cd projects/rocthrust
 
-# Go to rocThrust directory, create and go to the build directory.
-cd rocThrust; mkdir build; cd build
+# Create a directory for the build and go to it.
+mkdir build; cd build
 
-# Configure rocThrust, setup options for your system.
 # Build options:
-#   DISABLE_WERROR        - ON  by default, This flag disable the -Werror compiler flag
-#   BUILD_TEST            - OFF by default,
-#   BUILD_HIPSTDPAR_TEST  - OFF by default,
-#   BUILD_EXAMPLES        - OFF by default,
-#   BUILD_BENCHMARKS      - OFF by default,
-#   DOWNLOAD_ROCPRIM      - OFF by default, when ON rocPRIM will be downloaded to the build folder,
-#   RNG_SEED_COUNT        - 0 by default, controls non-repeatable random dataset count
-#   PRNG_SEEDS            - 1 by default, reproducible seeds to generate random data
+#   DISABLE_WERROR               - ON  by default, this flag disables the -Werror compiler flag
+#   BUILD_TEST                   - OFF by default,
+#   BUILD_HIPSTDPAR_TEST         - OFF by default,
+#   BUILD_EXAMPLE                - OFF by default,
+#   BUILD_BENCHMARK              - OFF by default,
+#   ROCPRIM_FETCH_METHOD         - PACKAGE is the default, see below for a description of available options
+#   ROCRAND_FETCH_METHOD         - PACKAGE is the default, see below for a description of available options
+#   EXTERNAL_DEPS_FORCE_DOWNLOAD - OFF by default, forces download for non-ROCm dependencies (eg. Google Test / Benchmark)
+#   BUILD_ADDRESS_SANITIZER      - OFF by default, builds with clang address sanitizer enabled.
+#   RNG_SEED_COUNT               - 0 by default, controls non-repeatable random dataset count
+#   PRNG_SEEDS                   - 1 by default, reproducible seeds to generate random data
+#   USE_HIPCXX                   - OFF by default, builds with CMake HIP language support. This eliminates the need to set CXX.
+#   USE_SYSTEM_LIB               - OFF by default, builds tests using the installed ROCm libs provided by the system. This only takes effect when BUILD_TEST is ON.
 #
 # ! IMPORTANT !
 # On ROCm platform set C++ compiler to HipCC. You can do it by adding 'CXX=<path-to-hipcc>'
@@ -95,19 +98,42 @@ make package
 [sudo] make install
 ```
 
+### Build Options for Fetching Dependencies
+
+rocThrust can (optionally) automatically fetch a number of dependencies for you at cmake configuration time.
+Alternatively, it can seach for an existing installation on your system.
+The following cmake build options control how dependencies are located.
+
+- `FORCE_DEPENDENCIES_DOWNLOAD` (default: `OFF`) - when set to `ON`, non-ROCm dependencies (Google Test, Google Benchmark) will always be downloaded, even if they are already installed ony your system. When set to `OFF`, rocThrust first searches for existing installations of the dependencies on your system, and only downloads them if they cannot be found.
+
+- `ROCPRIM_FETCH_METHOD` (default: `PACKAGE`) - controls the way that the rocPRIM dependency is fetched. This option must be set to one of:
+  - `PACKAGE` - searches for an existing installation of the dependency. If it is not found, rocThrust will fall back to using the `DOWNLOAD` setting (below).
+  - `DOWNLOAD` - downloads the dependency from the rocm-libraries repository. If git version 2.25+ is present, uses a [sparse checkout](https://git-scm.com/docs/git-sparse-checkout) to pull only rocThrust files. If not, the whole [rocm-libraries](https://github.com/ROCm/rocm-libraries/) repository will be downloaded (this may take some time).
+  - `MONOREPO` - this option is useful if you are building rocThrust from within a checkout of the rocm-libraries repository (which already includes rocPRIM and rocRAND). When enabled, rocThrust will try to find the dependency in the local repository tree. If it cannot be found, rocThrust will fall back to usign the `DOWNLOAD` option (above).
+- `ROCRAND_FETCH_METHOD` (default: `PACKAGE`) - this option is only considered when the `BUILD_BENCHMARKS` option is set to `ON`. It controls the way that the rocRAND dependency is fetched. See `ROCPRIM_FETCH_METHOD` (above) for available options.
+
+To specify an option, add it to your camke command, prefixed with the `-D` switch, (eg. `[CXX=hipcc] cmake -DFORCE_DEPENDENCIES_DOWNLOAD=ON -DROCPRIM_FETCH_METHOD="MONOREPO" ../.`).
+
 ### HIP on Windows
 
-We've added initial support for HIP on Windows. To install, use the provided `rmake.py` Python script:
+We've added initial support for HIP on Windows.
+To install, first clone rocThrust using the steps described in [obtaining the source code](#obtaining-the-source-code).
+Then, use the provided `rmake.py` Python script:
 
 ```shell
-git clone https://github.com/ROCm/rocThrust.git
-cd rocThrust
+cd projects/rocThrust
 
 # the -i option will install rocPRIM to C:\hipSDK by default
 python rmake.py -i
 
 # the -c option will build all clients including unit tests
 python rmake.py -c
+
+# to build for a specific architecture only, use the -a option
+python rmake.py -ci -a gfx1100
+
+# for a full list of available options, please refer to the help documentation
+python rmake.py -h
 ```
 
 ### Macro options
@@ -139,9 +165,9 @@ target_link_libraries(<your_target> roc::rocthrust)
 
 ```sh
 # Go to rocThrust build directory
-cd rocThrust; cd build
+cd projects/rocthrust; cd build
 
-# Configure with examples flag on
+# Configure with test flag on
 CXX=hipcc cmake -DBUILD_TEST=ON ..
 
 # Build tests
@@ -178,7 +204,7 @@ There is a utility script in the repo that may be called independently:
 
 ```shell
 # Go to rocThrust build directory
-cd rocThrust; cd build
+cd projects/rocthrust; cd build
 
 # Invoke directly or use CMake script mode via cmake -P
 ../cmake/GenerateResourceSpec.cmake
@@ -236,10 +262,10 @@ There are two CMake configuration-time options that control random data fed to u
 
 ```sh
 # Go to rocThrust build directory
-cd rocThrust; cd build
+cd projects/rocthrust; cd build
 
-# Configure with examples flag on
-CXX=hipcc cmake -DBUILD_EXAMPLES=ON ..
+# Configure with example flag on
+CXX=hipcc cmake -DBUILD_EXAMPLE=ON ..
 
 # Build examples
 make -j4
@@ -257,16 +283,16 @@ make -j4
 
 ```sh
 # Go to rocThrust build directory
-cd rocThrust; cd build
+cd projects/rocthrust; cd build
 
-# Configure with benchmarks flag on
-CXX=hipcc cmake -DBUILD_BENCHMARKS=ON ..
+# Configure with benchmark flag on
+CXX=hipcc cmake -DBUILD_BENCHMARK=ON ..
 
 # Build benchmarks
 make -j4
 
 # Run the benchmarks
-./benchmarks/benchmark_thrust_bench
+./benchmark/benchmark_thrust_bench
 ```
 
 ## HIPSTDPAR
@@ -284,11 +310,11 @@ HIPSTDPAR is currently packaged along rocThrust. The `hipstdpar` package is set 
 ### Tests
 rocThrust also includes tests to check the correct building of HIPSTDPAR implementations. They are located in the [tests/hipstdpar](/test/hipstdpar/) folder. When configuring the project with the `BUILD_TEST` option, these tests will not be enabled by default. To enable them, set `BUILD_HIPSTDPAR_TEST=ON`. Additionally, you can configure only HIPSTDPAR's tests by disabling `BUILD_TEST` and enabling `BUILD_HIPSTDPAR_TEST`. In general, the following steps can be followed for building and running the tests:
 
+First, clone rocThrust using the steps described in [obtaining the source code](#obtaining-the-source-code).
+Then, build the tests as follows:
 ```sh
-git clone https://github.com/ROCm/rocThrust
-
 # Go to rocThrust directory, create and go to the build directory.
-cd rocThrust; mkdir build; cd build
+cd projects/rocthrust; mkdir build; cd build
 
 # Configure rocThrust.
 [CXX=hipcc] cmake ../. -D BUILD_TEST=ON # Configure rocThrust's tests.
@@ -309,10 +335,67 @@ ctest --output-on-failure
   * Notice that oneTBB (oneAPI TBB) may fail to compile when libstdc++-9 or -10 is used, due to them using legacy TBB interfaces that are incompatible with the oneTBB ones (see the [release notes](https://www.intel.com/content/www/us/en/developer/articles/release-notes/intel-oneapi-threading-building-blocks-release-notes.html)).
 * CMake (3.10.2 or later)
 
+## 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_rocthrust
+
+# Activate the virtual environment
+pyenv activate venv_rocthrust
+```
+
+### Building
+
+After cloning this repository (see [obtaining the source code](#obtaining-the-source-code)):
+
+```shell
+cd rocm-libraries/projects/rocthrust
+
+# 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.
+
 ## Support
 
 You can report bugs and feature requests through the GitHub
-[issue tracker](https://github.com/ROCm/rocThrust/issues).
+[issue tracker](https://github.com/ROCm/rocm-libraries/issues).
+To help ensure that your issue is seen by the right team more quicly, when creating your issue, please apply the label `project: rocthrust`.
+Similarly, to filter the exising issue list down to only those affecting rocThrust, you can add the filter `label:"project: rocthrust"`,
+or follow [this link](https://github.com/ROCm/rocm-libraries/issues?q=is%3Aissue%20state%3Aopen%20label%3A%22project%3A%20rocthrust%22).
 
 ## License
 
diff --git a/benchmarks/CMakeLists.txt b/benchmark/CMakeLists.txt
similarity index 100%
rename from benchmarks/CMakeLists.txt
rename to benchmark/CMakeLists.txt
diff --git a/benchmarks/bench/adjacent_difference/basic.cu b/benchmark/bench/adjacent_difference/basic.cu
similarity index 56%
rename from benchmarks/bench/adjacent_difference/basic.cu
rename to benchmark/bench/adjacent_difference/basic.cu
index 6733c56..8cf5120 100644
--- a/benchmarks/bench/adjacent_difference/basic.cu
+++ b/benchmark/bench/adjacent_difference/basic.cu
@@ -1,24 +1,30 @@
-// MIT License
-//
-// Copyright (c) 2024 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.
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
 
 // Benchmark utils
 #include "../../bench_utils/bench_utils.hpp"
@@ -32,121 +38,121 @@
 #include <benchmark/benchmark.h>
 
 // STL
-#include <cstdlib>
+#include <cstddef>
 #include <string>
 #include <vector>
 
 struct basic
 {
-    template <typename T, typename Policy>
-    float64_t run(thrust::device_vector<T>& input, thrust::device_vector<T>& output, Policy policy)
-    {
-        bench_utils::gpu_timer d_timer;
+  template <typename T, typename Policy>
+  float64_t run(thrust::device_vector<T>& input, thrust::device_vector<T>& output, Policy policy)
+  {
+    thrust::adjacent_difference(policy, input.cbegin(), input.cend(), output.begin());
 
-        d_timer.start(0);
-        thrust::adjacent_difference(policy, input.cbegin(), input.cend(), output.begin());
-        d_timer.stop(0);
+    bench_utils::gpu_timer d_timer;
 
-        return d_timer.get_duration();
-    }
+    d_timer.start(0);
+    thrust::adjacent_difference(policy, input.cbegin(), input.cend(), output.begin());
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
 };
 
 template <class Benchmark, class T>
 void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string seed_type)
 {
-    // Benchmark object
-    Benchmark benchmark {};
+  // Benchmark object
+  Benchmark benchmark{};
 
-    // GPU times
-    std::vector<double> gpu_times;
+  // GPU times
+  std::vector<double> gpu_times;
 
-    // Generate input
-    thrust::device_vector<T> input = bench_utils::generate(elements, seed_type);
+  // Generate input
+  thrust::device_vector<T> input = bench_utils::generate(elements, seed_type);
 
-    // Output
-    thrust::device_vector<T> output(elements);
+  // Output
+  thrust::device_vector<T> output(elements);
 
-    bench_utils::caching_allocator_t alloc {};
-    thrust::detail::device_t         policy {};
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
 
-    for(auto _ : state)
-    {
-        float64_t duration = benchmark.template run<T>(input, output, policy(alloc));
-        state.SetIterationTime(duration);
-        gpu_times.push_back(duration);
-    }
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(input, output, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
 
-    // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
-    // it will actually be the global memory bandwidth gotten.
-    state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * elements);
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
 
-    const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
-    state.counters["gpu_noise"] = gpu_cv;
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
 }
 
-#define CREATE_BENCHMARK(T, Elements)                                                        \
-    benchmark::RegisterBenchmark(                                                            \
-        bench_utils::bench_naming::format_name("{algo:adjacent_difference,subalgo:" + name   \
-                                               + ",input_type:" #T + ",elements:" #Elements) \
-            .c_str(),                                                                        \
-        run_benchmark<Benchmark, T>,                                                         \
-        Elements,                                                                            \
-        seed_type)
+#define CREATE_BENCHMARK(T, Elements)                                                           \
+  benchmark::RegisterBenchmark(                                                                 \
+    bench_utils::bench_naming::format_name(                                                     \
+      "{algo:adjacent_difference,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements) \
+      .c_str(),                                                                                 \
+    run_benchmark<Benchmark, T>,                                                                \
+    Elements,                                                                                   \
+    seed_type)
 
-#define BENCHMARK_TYPE(type)                                          \
-    CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), \
-        CREATE_BENCHMARK(type, 1 << 24), CREATE_BENCHMARK(type, 1 << 28)
+#define BENCHMARK_TYPE(type)                                                                         \
+  CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), CREATE_BENCHMARK(type, 1 << 24), \
+    CREATE_BENCHMARK(type, 1 << 28)
 
 template <class Benchmark>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    const std::string                             seed_type)
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
 {
-    std::vector<benchmark::internal::Benchmark*> bs = {BENCHMARK_TYPE(int8_t),
-                                                       BENCHMARK_TYPE(int16_t),
-                                                       BENCHMARK_TYPE(int32_t),
-                                                       BENCHMARK_TYPE(int64_t),
-                                                       BENCHMARK_TYPE(float32_t),
-                                                       BENCHMARK_TYPE(float64_t)};
-
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_TYPE(int8_t),
+    BENCHMARK_TYPE(int16_t),
+    BENCHMARK_TYPE(int32_t),
+    BENCHMARK_TYPE(int64_t),
+    BENCHMARK_TYPE(float32_t),
+    BENCHMARK_TYPE(float64_t)};
+
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
 }
 
 int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<std::string>(
-        "name_format", "name_format", "human", "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    bench_utils::bench_naming::set_format(
-        parser.get<std::string>("name_format")); /* either: json,human,txt */
-    const std::string seed_type = parser.get<std::string>("seed");
-
-    // Benchmark info
-    bench_utils::add_common_benchmark_info();
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmark
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<basic>("basic", benchmarks, seed_type);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMicrosecond);
-        b->MinTime(0.4); // in seconds
-    }
-
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
-
-    // Finish
-    benchmark::Shutdown();
-    return 0;
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<basic>("basic", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
 }
diff --git a/benchmarks/bench/adjacent_difference/custom.cu b/benchmark/bench/adjacent_difference/custom.cu
similarity index 53%
rename from benchmarks/bench/adjacent_difference/custom.cu
rename to benchmark/bench/adjacent_difference/custom.cu
index 24d6c03..dd47604 100644
--- a/benchmarks/bench/adjacent_difference/custom.cu
+++ b/benchmark/bench/adjacent_difference/custom.cu
@@ -1,24 +1,30 @@
-// MIT License
-//
-// Copyright (c) 2024 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.
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
 
 // Benchmark utils
 #include "../../bench_utils/bench_utils.hpp"
@@ -32,141 +38,139 @@
 #include <benchmark/benchmark.h>
 
 // STL
-#include <cstdlib>
+#include <cstddef>
 #include <string>
 #include <vector>
 
 template <typename T>
 struct custom_op
 {
-    T val;
+  T val;
 
-    custom_op() = delete;
+  custom_op() = delete;
 
-    explicit custom_op(T val)
-        : val(val)
-    {
-    }
+  explicit custom_op(T val)
+      : val(val)
+  {}
 
-    __device__ T operator()(const T& lhs, const T& rhs)
-    {
-        return lhs * rhs + val;
-    }
+  __device__ T operator()(const T& lhs, const T& rhs)
+  {
+    return lhs * rhs + val; // Hope to gen mad
+  }
 };
 
 template <int Val>
 struct custom
 {
-    template <typename T, typename Policy>
-    float64_t run(thrust::device_vector<T>& input, thrust::device_vector<T>& output, Policy policy)
-    {
-        bench_utils::gpu_timer d_timer;
-
-        d_timer.start(0);
-        thrust::adjacent_difference(
-            policy, input.cbegin(), input.cend(), output.begin(), custom_op<T> {Val});
-        d_timer.stop(0);
-
-        return d_timer.get_duration();
-    }
+  template <typename T, typename Policy>
+  float64_t run(thrust::device_vector<T>& input, thrust::device_vector<T>& output, Policy policy)
+  {
+    thrust::adjacent_difference(policy, input.cbegin(), input.cend(), output.begin(), custom_op<T>{Val});
+
+    bench_utils::gpu_timer d_timer;
+
+    d_timer.start(0);
+    thrust::adjacent_difference(policy, input.cbegin(), input.cend(), output.begin(), custom_op<T>{Val});
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
 };
 
 template <class Benchmark, class T>
 void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string seed_type)
 {
-    // Benchmark object
-    Benchmark benchmark {};
+  // Benchmark object
+  Benchmark benchmark{};
 
-    // GPU times
-    std::vector<double> gpu_times;
+  // GPU times
+  std::vector<double> gpu_times;
 
-    // Generate input
-    thrust::device_vector<T> input = bench_utils::generate(elements, seed_type);
+  // Generate input
+  thrust::device_vector<T> input = bench_utils::generate(elements, seed_type);
 
-    // Output
-    thrust::device_vector<T> output(elements);
+  // Output
+  thrust::device_vector<T> output(elements);
 
-    bench_utils::caching_allocator_t alloc {};
-    thrust::detail::device_t         policy {};
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
 
-    for(auto _ : state)
-    {
-        float64_t duration = benchmark.template run<T>(input, output, policy(alloc));
-        state.SetIterationTime(duration);
-        gpu_times.push_back(duration);
-    }
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(input, output, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
 
-    // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
-    // it will actually be the global memory bandwidth gotten.
-    state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * elements);
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
 
-    const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
-    state.counters["gpu_noise"] = gpu_cv;
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
 }
 
-#define CREATE_BENCHMARK(T, Elements)                                                        \
-    benchmark::RegisterBenchmark(                                                            \
-        bench_utils::bench_naming::format_name("{algo:adjacent_difference,subalgo:" + name   \
-                                               + ",input_type:" #T + ",elements:" #Elements) \
-            .c_str(),                                                                        \
-        run_benchmark<Benchmark<Val>, T>,                                                    \
-        Elements,                                                                            \
-        seed_type)
+#define CREATE_BENCHMARK(T, Elements)                                                           \
+  benchmark::RegisterBenchmark(                                                                 \
+    bench_utils::bench_naming::format_name(                                                     \
+      "{algo:adjacent_difference,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements) \
+      .c_str(),                                                                                 \
+    run_benchmark<Benchmark<Val>, T>,                                                           \
+    Elements,                                                                                   \
+    seed_type)
 
-#define BENCHMARK_TYPE(type)                                          \
-    CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), \
-        CREATE_BENCHMARK(type, 1 << 24), CREATE_BENCHMARK(type, 1 << 28)
+#define BENCHMARK_TYPE(type)                                                                         \
+  CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), CREATE_BENCHMARK(type, 1 << 24), \
+    CREATE_BENCHMARK(type, 1 << 28)
 
 template <template <int> class Benchmark, int Val = 42 /*magic number in Thrust's benchmark*/>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    const std::string                             seed_type)
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
 {
-    std::vector<benchmark::internal::Benchmark*> bs = {BENCHMARK_TYPE(int8_t),
-                                                       BENCHMARK_TYPE(int16_t),
-                                                       BENCHMARK_TYPE(int32_t),
-                                                       BENCHMARK_TYPE(int64_t),
-                                                       BENCHMARK_TYPE(float32_t),
-                                                       BENCHMARK_TYPE(float64_t)};
-
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_TYPE(int8_t),
+    BENCHMARK_TYPE(int16_t),
+    BENCHMARK_TYPE(int32_t),
+    BENCHMARK_TYPE(int64_t),
+    BENCHMARK_TYPE(float32_t),
+    BENCHMARK_TYPE(float64_t)};
+
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
 }
 
 int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<std::string>(
-        "name_format", "name_format", "human", "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    bench_utils::bench_naming::set_format(
-        parser.get<std::string>("name_format")); /* either: json,human,txt */
-    const std::string seed_type = parser.get<std::string>("seed");
-
-    // Benchmark info
-    bench_utils::add_common_benchmark_info();
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmark
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<custom>("custom", benchmarks, seed_type);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMicrosecond);
-        b->MinTime(0.4); // in seconds
-    }
-
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
-
-    // Finish
-    benchmark::Shutdown();
-    return 0;
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<custom>("custom", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
 }
diff --git a/benchmarks/bench/copy/basic.cu b/benchmark/bench/adjacent_difference/in_place.cu
similarity index 53%
rename from benchmarks/bench/copy/basic.cu
rename to benchmark/bench/adjacent_difference/in_place.cu
index 52f22c5..4647152 100644
--- a/benchmarks/bench/copy/basic.cu
+++ b/benchmark/bench/adjacent_difference/in_place.cu
@@ -1,30 +1,36 @@
-// MIT License
-//
-// Copyright (c) 2024 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.
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
 
 // Benchmark utils
 #include "../../bench_utils/bench_utils.hpp"
 
 // rocThrust
-#include <thrust/copy.h>
+#include <thrust/adjacent_difference.h>
 #include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
 
@@ -32,123 +38,118 @@
 #include <benchmark/benchmark.h>
 
 // STL
-#include <cstdlib>
+#include <cstddef>
 #include <string>
 #include <vector>
 
-struct basic
+struct in_place
 {
-    template <typename T, typename Policy>
-    float64_t run(thrust::device_vector<T>& input, thrust::device_vector<T>& output, Policy policy)
-    {
-        bench_utils::gpu_timer d_timer;
+  template <typename T, typename Policy>
+  float64_t run(thrust::device_vector<T>& vec, Policy policy)
+  {
+    thrust::adjacent_difference(policy, vec.begin(), vec.end(), vec.begin());
 
-        d_timer.start(0);
-        thrust::copy(policy, input.cbegin(), input.cend(), output.begin());
-        d_timer.stop(0);
+    bench_utils::gpu_timer d_timer;
 
-        return d_timer.get_duration();
-    }
+    d_timer.start(0);
+    thrust::adjacent_difference(policy, vec.begin(), vec.end(), vec.begin());
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
 };
 
 template <class Benchmark, class T>
-void run_benchmark(benchmark::State& state,
-                   const std::size_t elements,
-                   const std::string /*seed_type*/)
+void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string /*seed_type*/)
 {
-    // Benchmark object
-    Benchmark benchmark {};
-
-    // GPU times
-    std::vector<double> gpu_times;
+  // Benchmark object
+  Benchmark benchmark{};
 
-    // Generate input
-    thrust::device_vector<T> input(elements, 1);
+  // GPU times
+  std::vector<double> gpu_times;
 
-    // Output
-    thrust::device_vector<T> output(elements);
+  // Generate vec
+  thrust::device_vector<T> vec(elements, 0);
 
-    bench_utils::caching_allocator_t alloc;
-    thrust::detail::device_t         policy {};
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
 
-    for(auto _ : state)
-    {
-        float64_t duration = benchmark.template run<T>(input, output, policy(alloc));
-        state.SetIterationTime(duration);
-        gpu_times.push_back(duration);
-    }
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(vec, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
 
-    // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
-    // it will actually be the global memory bandwidth gotten.
-    state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * elements);
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
 
-    const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
-    state.counters["gpu_noise"] = gpu_cv;
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
 }
 
 #define CREATE_BENCHMARK(T, Elements)                                                           \
-    benchmark::RegisterBenchmark(                                                               \
-        bench_utils::bench_naming::format_name("{algo:copy,subalgo:" + name + ",input_type:" #T \
-                                               + ",elements:" #Elements)                        \
-            .c_str(),                                                                           \
-        run_benchmark<Benchmark, T>,                                                            \
-        Elements,                                                                               \
-        seed_type)
-
-#define BENCHMARK_TYPE(type)                                          \
-    CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), \
-        CREATE_BENCHMARK(type, 1 << 24), CREATE_BENCHMARK(type, 1 << 28)
+  benchmark::RegisterBenchmark(                                                                 \
+    bench_utils::bench_naming::format_name(                                                     \
+      "{algo:adjacent_difference,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements) \
+      .c_str(),                                                                                 \
+    run_benchmark<Benchmark, T>,                                                                \
+    Elements,                                                                                   \
+    seed_type)
+
+#define BENCHMARK_TYPE(type)                                                                         \
+  CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), CREATE_BENCHMARK(type, 1 << 24), \
+    CREATE_BENCHMARK(type, 1 << 28)
 
 template <class Benchmark>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    const std::string                             seed_type)
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
 {
-    std::vector<benchmark::internal::Benchmark*> bs = {BENCHMARK_TYPE(int8_t),
-                                                       BENCHMARK_TYPE(int16_t),
-                                                       BENCHMARK_TYPE(int32_t),
-                                                       BENCHMARK_TYPE(int64_t),
-                                                       BENCHMARK_TYPE(float32_t),
-                                                       BENCHMARK_TYPE(float64_t)};
-
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_TYPE(int8_t),
+    BENCHMARK_TYPE(int16_t),
+    BENCHMARK_TYPE(int32_t),
+    BENCHMARK_TYPE(int64_t),
+    BENCHMARK_TYPE(float32_t),
+    BENCHMARK_TYPE(float64_t)};
+
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
 }
 
 int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<std::string>(
-        "name_format", "name_format", "human", "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    bench_utils::bench_naming::set_format(
-        parser.get<std::string>("name_format")); /* either: json,human,txt */
-    const std::string seed_type = parser.get<std::string>("seed");
-
-    // Benchmark info
-    bench_utils::add_common_benchmark_info();
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmark
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<basic>("basic", benchmarks, seed_type);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMicrosecond);
-        b->MinTime(0.4); // in seconds
-    }
-
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
-
-    // Finish
-    benchmark::Shutdown();
-    return 0;
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<in_place>("in_place", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
 }
diff --git a/benchmark/bench/copy/basic.cu b/benchmark/bench/copy/basic.cu
new file mode 100644
index 0000000..42f1a1d
--- /dev/null
+++ b/benchmark/bench/copy/basic.cu
@@ -0,0 +1,191 @@
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+
+// Benchmark utils
+#include "../../bench_utils/bench_utils.hpp"
+
+// rocThrust
+#include <thrust/copy.h>
+#include <thrust/count.h>
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+
+// Google Benchmark
+#include <benchmark/benchmark.h>
+
+// STL
+#include <cstddef>
+#include <string>
+#include <vector>
+#if !_THRUST_HAS_DEVICE_SYSTEM_STD
+#  include <type_traits>
+#endif
+
+struct basic
+{
+  template <typename T, typename Policy>
+  float64_t run(thrust::device_vector<T>& input, thrust::device_vector<T>& output, Policy policy)
+  {
+    thrust::copy(policy, input.cbegin(), input.cend(), output.begin());
+
+    bench_utils::gpu_timer d_timer;
+
+    d_timer.start(0);
+    thrust::copy(policy, input.cbegin(), input.cend(), output.begin());
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
+};
+
+template <class Benchmark, class T>
+void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string /*seed_type*/)
+{
+  // Benchmark object
+  Benchmark benchmark{};
+
+  // GPU times
+  std::vector<double> gpu_times;
+
+  // Generate input
+  thrust::device_vector<T> input(elements, T{1});
+
+  // Output
+  thrust::device_vector<T> output(elements);
+
+  bench_utils::caching_allocator_t alloc;
+  thrust::detail::device_t policy{};
+
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(input, output, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
+
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
+
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
+}
+
+#define CREATE_BENCHMARK(T, Elements)                                                                                 \
+  benchmark::RegisterBenchmark(                                                                                       \
+    bench_utils::bench_naming::format_name("{algo:copy,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements) \
+      .c_str(),                                                                                                       \
+    run_benchmark<Benchmark, T>,                                                                                      \
+    Elements,                                                                                                         \
+    seed_type)
+
+#define BENCHMARK_TYPE(type)                                                                         \
+  CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), CREATE_BENCHMARK(type, 1 << 24), \
+    CREATE_BENCHMARK(type, 1 << 28)
+
+// Non-trivially-copyable/relocatable type which is not allowed to be copied using std::memcpy or cudaMemcpy
+struct non_trivial
+{
+  int a;
+  int b;
+
+  non_trivial() = default;
+
+  THRUST_HOST_DEVICE explicit non_trivial(int i)
+      : a(i)
+      , b(i)
+  {}
+
+  // the user-defined copy constructor prevents the type from being trivially copyable
+  THRUST_HOST_DEVICE non_trivial(const non_trivial& nt)
+      : a(nt.a)
+      , b(nt.b)
+  {}
+
+  non_trivial& operator=(const non_trivial&) = default;
+};
+
+static_assert(!_THRUST_STD::is_trivially_copyable<non_trivial>::value, ""); // as required by the C++ standard
+static_assert(!thrust::is_trivially_relocatable<non_trivial>::value, ""); // thrust uses this check internally
+
+template <class Benchmark>
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
+{
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_TYPE(int8_t),
+    BENCHMARK_TYPE(uint8_t),
+    BENCHMARK_TYPE(int16_t),
+    BENCHMARK_TYPE(uint16_t),
+    BENCHMARK_TYPE(int32_t),
+    BENCHMARK_TYPE(uint32_t),
+    BENCHMARK_TYPE(int64_t),
+    BENCHMARK_TYPE(uint64_t),
+    BENCHMARK_TYPE(float32_t),
+    BENCHMARK_TYPE(float64_t),
+    BENCHMARK_TYPE(non_trivial)};
+
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+}
+
+int main(int argc, char* argv[])
+{
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<basic>("basic", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
+}
diff --git a/benchmark/bench/copy/if.cu b/benchmark/bench/copy/if.cu
new file mode 100644
index 0000000..2c93d05
--- /dev/null
+++ b/benchmark/bench/copy/if.cu
@@ -0,0 +1,203 @@
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+
+// Benchmark utils
+#include "../../bench_utils/bench_utils.hpp"
+
+// rocThrust
+#include <thrust/copy.h>
+#include <thrust/count.h>
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+
+// Google Benchmark
+#include <benchmark/benchmark.h>
+
+// STL
+#include <cstddef>
+#include <string>
+#include <vector>
+
+template <class T>
+struct less_then_t
+{
+  T m_val;
+
+  __host__ __device__ bool operator()(const T& val) const
+  {
+    return val < m_val;
+  }
+};
+
+struct _if
+{
+  template <typename T, typename Policy>
+  float64_t
+  run(thrust::device_vector<T>& input, thrust::device_vector<T>& output, less_then_t<T> select_op, Policy policy)
+  {
+    thrust::copy_if(policy, input.cbegin(), input.cend(), output.begin(), select_op);
+
+    bench_utils::gpu_timer d_timer;
+
+    d_timer.start(0);
+    thrust::copy_if(policy, input.cbegin(), input.cend(), output.begin(), select_op);
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
+};
+
+template <class Benchmark, class T>
+void run_benchmark(
+  benchmark::State& state, const std::size_t elements, const std::string seed_type, const int entropy_reduction)
+{
+  using select_op_t = less_then_t<T>;
+
+  // Benchmark object
+  Benchmark benchmark{};
+
+  // GPU times
+  std::vector<double> gpu_times;
+
+  T val = bench_utils::value_from_entropy<T>(bench_utils::get_entropy_percentage(entropy_reduction));
+  select_op_t select_op{val};
+
+  // Generate input and output
+  thrust::device_vector<T> input;
+  try
+  {
+    input = bench_utils::generate(elements, seed_type);
+  }
+  catch (const ::thrust::system::detail::bad_alloc& e)
+  {
+    (void) hipGetLastError();
+    state.SkipWithError(("thrust::system::detail::bad_alloc: " + std::string(e.what())).c_str());
+    return;
+  }
+  const auto selected_elements = thrust::count_if(input.cbegin(), input.cend(), select_op);
+  thrust::device_vector<T> output;
+  try
+  {
+    output = thrust::device_vector<T>(selected_elements);
+  }
+  catch (const ::thrust::system::detail::bad_alloc& e)
+  {
+    (void) hipGetLastError();
+    state.SkipWithError(("thrust::system::detail::bad_alloc: " + std::string(e.what())).c_str());
+    return;
+  }
+
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
+
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(input, output, select_op, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
+
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * (elements + selected_elements) * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
+
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
+}
+
+#define CREATE_BENCHMARK(T, Elements, EntropyReduction)                                      \
+  benchmark::RegisterBenchmark(                                                              \
+    bench_utils::bench_naming::format_name(                                                  \
+      "{algo:copy,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements              \
+      + ",entropy:" + std::to_string(bench_utils::get_entropy_percentage(EntropyReduction))) \
+      .c_str(),                                                                              \
+    run_benchmark<Benchmark, T>,                                                             \
+    Elements,                                                                                \
+    seed_type,                                                                               \
+    EntropyReduction)
+
+#define BENCHMARK_TYPE_ENTROPY(type, entropy)                                         \
+  CREATE_BENCHMARK(type, 1 << 16, entropy), CREATE_BENCHMARK(type, 1 << 20, entropy), \
+    CREATE_BENCHMARK(type, 1 << 24, entropy), CREATE_BENCHMARK(type, 1 << 28, entropy)
+
+template <class Benchmark>
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
+{
+  constexpr int entropy_reductions[] = {0, 2, 4200}; // 1.000, 0.544, 0.000;
+
+  for (int entropy_reduction : entropy_reductions)
+  {
+    std::vector<benchmark::internal::Benchmark*> bs = {
+      BENCHMARK_TYPE_ENTROPY(int8_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(int16_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(int32_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(int64_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(float, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(double, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(bench_utils::large_data, entropy_reduction)}; // rocThrust issue #565
+    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+  }
+}
+
+int main(int argc, char* argv[])
+{
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<_if>("if", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
+}
diff --git a/benchmark/bench/equal/basic.cu b/benchmark/bench/equal/basic.cu
new file mode 100644
index 0000000..bda269f
--- /dev/null
+++ b/benchmark/bench/equal/basic.cu
@@ -0,0 +1,145 @@
+// SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION. All rights reserved.
+// SPDX-FileCopyrightText: Modifications Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+// Benchmark utils
+#include "../../bench_utils/bench_utils.hpp"
+
+// rocThrust
+#include <thrust/device_vector.h>
+#include <thrust/equal.h>
+#include <thrust/execution_policy.h>
+
+// Google Benchmark
+#include <benchmark/benchmark.h>
+
+// STL
+#include <algorithm>
+#include <cstddef>
+#include <string>
+#include <vector>
+
+struct basic
+{
+  template <typename T, typename Policy>
+  float64_t run(thrust::device_vector<T>& a, thrust::device_vector<T>& b, Policy policy)
+  {
+    bench_utils::do_not_optimize(thrust::equal(policy, a.begin(), a.end(), b.begin()));
+
+    bench_utils::gpu_timer d_timer;
+
+    d_timer.start(0);
+    bench_utils::do_not_optimize(thrust::equal(policy, a.begin(), a.end(), b.begin()));
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
+};
+
+template <class Benchmark, class T>
+void run_benchmark(benchmark::State& state,
+                   const std::size_t elements,
+                   const std::string, /*seed_type*/
+                   const float64_t common_prefix_ratio)
+{
+  // Benchmark object
+  Benchmark benchmark{};
+
+  // GPU times
+  std::vector<double> gpu_times;
+
+  thrust::device_vector<T> a(elements, T{1});
+  thrust::device_vector<T> b(elements, T{1});
+
+  const auto same_elements = std::min(static_cast<std::size_t>(elements * common_prefix_ratio), elements);
+
+  bench_utils::caching_allocator_t alloc;
+  thrust::detail::device_t policy{};
+
+  thrust::fill(policy(alloc), b.begin() + same_elements, b.end(), T{2});
+
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(a, b, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
+
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  // using `same_elements` instead of `elements` corresponds to the
+  // actual elements read in an early exit
+  state.SetBytesProcessed(state.iterations() * 2 * std::max(same_elements, std::size_t(1)) * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * std::max(same_elements, std::size_t(1)));
+
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
+}
+
+#define CREATE_BENCHMARK(T, Elements, CommonPrefixRatio)                                                              \
+  benchmark::RegisterBenchmark(                                                                                       \
+    bench_utils::bench_naming::format_name("{algo:equal,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements \
+                                           + ", common_prefix_ratio:" #CommonPrefixRatio)                             \
+      .c_str(),                                                                                                       \
+    run_benchmark<Benchmark, T>,                                                                                      \
+    Elements,                                                                                                         \
+    seed_type,                                                                                                        \
+    CommonPrefixRatio)
+
+#define BENCHMARK_ELEMENTS(type, elements) \
+  CREATE_BENCHMARK(type, elements, 1.0), CREATE_BENCHMARK(type, elements, 0.5), CREATE_BENCHMARK(type, elements, 0.0)
+
+#define BENCHMARK_TYPE(type)                                                                               \
+  BENCHMARK_ELEMENTS(type, 1 << 16), BENCHMARK_ELEMENTS(type, 1 << 20), BENCHMARK_ELEMENTS(type, 1 << 24), \
+    BENCHMARK_ELEMENTS(type, 1 << 28)
+
+template <class Benchmark>
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
+{
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_TYPE(int8_t),
+    BENCHMARK_TYPE(int16_t),
+    BENCHMARK_TYPE(int32_t),
+    BENCHMARK_TYPE(uint32_t),
+    BENCHMARK_TYPE(int64_t),
+    BENCHMARK_TYPE(uint64_t)};
+
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+}
+
+int main(int argc, char* argv[])
+{
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<basic>("basic", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
+}
diff --git a/benchmarks/bench/fill/basic.cu b/benchmark/bench/fill/basic.cu
similarity index 54%
rename from benchmarks/bench/fill/basic.cu
rename to benchmark/bench/fill/basic.cu
index f88b0fc..62403bd 100644
--- a/benchmarks/bench/fill/basic.cu
+++ b/benchmark/bench/fill/basic.cu
@@ -1,24 +1,30 @@
-// MIT License
-//
-// Copyright (c) 2024 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.
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
 
 // Benchmark utils
 #include "../../bench_utils/bench_utils.hpp"
@@ -32,120 +38,121 @@
 #include <benchmark/benchmark.h>
 
 // STL
-#include <cstdlib>
+#include <cstddef>
 #include <string>
 #include <vector>
 
 template <int Val>
 struct basic
 {
-    template <typename T, typename Policy>
-    float64_t run(thrust::device_vector<T>& output, Policy policy)
-    {
-        bench_utils::gpu_timer d_timer;
+  template <typename T, typename Policy>
+  float64_t run(thrust::device_vector<T>& output, Policy policy)
+  {
+    thrust::fill(policy, output.begin(), output.end(), T{Val});
 
-        d_timer.start(0);
-        thrust::fill(policy, output.begin(), output.end(), T {Val});
-        d_timer.stop(0);
+    bench_utils::gpu_timer d_timer;
 
-        return d_timer.get_duration();
-    }
+    d_timer.start(0);
+    thrust::fill(policy, output.begin(), output.end(), T{Val});
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
 };
 
 template <class Benchmark, class T>
-void run_benchmark(benchmark::State& state,
-                   const std::size_t elements,
-                   const std::string /*seed_type*/)
+void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string /*seed_type*/)
 {
-    // Benchmark object
-    Benchmark benchmark {};
+  // Benchmark object
+  Benchmark benchmark{};
 
-    // GPU times
-    std::vector<double> gpu_times;
+  // GPU times
+  std::vector<double> gpu_times;
 
-    // Output
-    thrust::device_vector<T> output(elements);
+  // Output
+  thrust::device_vector<T> output(elements);
 
-    bench_utils::caching_allocator_t alloc {};
-    thrust::detail::device_t         policy {};
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
 
-    for(auto _ : state)
-    {
-        float64_t duration = benchmark.template run<T>(output, policy(alloc));
-        state.SetIterationTime(duration);
-        gpu_times.push_back(duration);
-    }
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(output, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
 
-    // BytesProcessed is only the number the writes, because thrust::fill does not receive an input
-    state.SetBytesProcessed(state.iterations() * elements * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * elements);
+  // BytesProcessed is only the number the writes, because thrust::fill does not receive an input
+  state.SetBytesProcessed(state.iterations() * elements * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
 
-    const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
-    state.counters["gpu_noise"] = gpu_cv;
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
 }
 
-#define CREATE_BENCHMARK(T, Elements)                                                           \
-    benchmark::RegisterBenchmark(                                                               \
-        bench_utils::bench_naming::format_name("{algo:fill,subalgo:" + name + ",input_type:" #T \
-                                               + ",elements:" #Elements)                        \
-            .c_str(),                                                                           \
-        run_benchmark<Benchmark<Val>, T>,                                                       \
-        Elements,                                                                               \
-        seed_type)
+#define CREATE_BENCHMARK(T, Elements)                                                                                 \
+  benchmark::RegisterBenchmark(                                                                                       \
+    bench_utils::bench_naming::format_name("{algo:fill,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements) \
+      .c_str(),                                                                                                       \
+    run_benchmark<Benchmark<Val>, T>,                                                                                 \
+    Elements,                                                                                                         \
+    seed_type)
 
-#define BENCHMARK_TYPE(type)                                          \
-    CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), \
-        CREATE_BENCHMARK(type, 1 << 24), CREATE_BENCHMARK(type, 1 << 28)
+#define BENCHMARK_TYPE(type)                                                                         \
+  CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), CREATE_BENCHMARK(type, 1 << 24), \
+    CREATE_BENCHMARK(type, 1 << 28)
 
 template <template <int> class Benchmark, int Val = 42 /*magic number in Thrust's benchmark*/>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    const std::string                             seed_type)
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
 {
-    std::vector<benchmark::internal::Benchmark*> bs = {BENCHMARK_TYPE(int8_t),
-                                                       BENCHMARK_TYPE(int16_t),
-                                                       BENCHMARK_TYPE(int32_t),
-                                                       BENCHMARK_TYPE(int64_t),
-                                                       BENCHMARK_TYPE(float32_t),
-                                                       BENCHMARK_TYPE(float64_t)};
-
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_TYPE(int8_t),
+    BENCHMARK_TYPE(int16_t),
+    BENCHMARK_TYPE(int32_t),
+    BENCHMARK_TYPE(int64_t),
+#if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT
+    BENCHMARK_TYPE(int128_t),
+#endif
+    BENCHMARK_TYPE(float32_t),
+    BENCHMARK_TYPE(float64_t)
+  };
+
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
 }
 
 int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<std::string>(
-        "name_format", "name_format", "human", "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    bench_utils::bench_naming::set_format(
-        parser.get<std::string>("name_format")); /* either: json,human,txt */
-    const std::string seed_type = parser.get<std::string>("seed");
-
-    // Benchmark info
-    bench_utils::add_common_benchmark_info();
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmark
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<basic>("basic", benchmarks, seed_type);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMicrosecond);
-        b->MinTime(0.4); // in seconds
-    }
-
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
-
-    // Finish
-    benchmark::Shutdown();
-    return 0;
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<basic>("basic", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
 }
diff --git a/benchmark/bench/for_each/basic.cu b/benchmark/bench/for_each/basic.cu
new file mode 100644
index 0000000..6a5bd58
--- /dev/null
+++ b/benchmark/bench/for_each/basic.cu
@@ -0,0 +1,169 @@
+/******************************************************************************
+ * Copyright (c) 2024, NVIDIA CORPORATION. All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+
+// Benchmark utils
+#include "../../bench_utils/bench_utils.hpp"
+
+// rocThrust
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+#include <thrust/for_each.h>
+
+// Google Benchmark
+#include <benchmark/benchmark.h>
+
+// STL
+#include <cstddef>
+#include <string>
+#include <vector>
+
+template <class T>
+struct square_t
+{
+  __device__ void operator()(T& x) const
+  {
+    x = x * x;
+  }
+};
+
+struct basic
+{
+  template <typename T, typename OpT, typename Policy>
+  float64_t run(thrust::device_vector<T>& in, OpT op, Policy policy)
+  {
+    thrust::for_each(policy, in.begin(), in.end(), op);
+
+    bench_utils::gpu_timer d_timer;
+
+    d_timer.start(0);
+    thrust::for_each(policy, in.begin(), in.end(), op);
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
+};
+
+template <class Benchmark, class T>
+void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string /*seed_type*/)
+{
+  // Benchmark object
+  Benchmark benchmark{};
+
+  // GPU times
+  std::vector<double> gpu_times;
+
+  // Generate input
+  thrust::device_vector<T> in(elements, T{1});
+
+  square_t<T> op{};
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
+
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(in, op, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
+
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
+
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
+}
+
+#define CREATE_BENCHMARK(T, Elements)                                                \
+  benchmark::RegisterBenchmark(                                                      \
+    bench_utils::bench_naming::format_name(                                          \
+      "{algo:for_each,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements) \
+      .c_str(),                                                                      \
+    run_benchmark<Benchmark, T>,                                                     \
+    Elements,                                                                        \
+    seed_type)
+
+#define BENCHMARK_TYPE(type)                                                                         \
+  CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), CREATE_BENCHMARK(type, 1 << 24), \
+    CREATE_BENCHMARK(type, 1 << 28)
+
+template <class Benchmark>
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
+{
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_TYPE(int8_t),
+    BENCHMARK_TYPE(int16_t),
+    BENCHMARK_TYPE(int32_t),
+    BENCHMARK_TYPE(int64_t),
+#if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT
+    BENCHMARK_TYPE(int128_t),
+#endif
+    BENCHMARK_TYPE(float32_t),
+    BENCHMARK_TYPE(float64_t)
+  };
+
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+}
+
+int main(int argc, char* argv[])
+{
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<basic>("basic", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
+}
diff --git a/benchmarks/bench/inner_product/basic.cu b/benchmark/bench/inner_product/basic.cu
similarity index 54%
rename from benchmarks/bench/inner_product/basic.cu
rename to benchmark/bench/inner_product/basic.cu
index 67a9b62..842de89 100644
--- a/benchmarks/bench/inner_product/basic.cu
+++ b/benchmark/bench/inner_product/basic.cu
@@ -1,24 +1,30 @@
-// MIT License
-//
-// Copyright (c) 2024 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.
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
 
 // Benchmark utils
 #include "../../bench_utils/bench_utils.hpp"
@@ -32,126 +38,124 @@
 #include <benchmark/benchmark.h>
 
 // STL
-#include <cstdlib>
+#include <cstddef>
 #include <string>
 #include <vector>
 
 struct basic
 {
-    template <typename T, typename Policy>
-    float64_t run(thrust::device_vector<T>& lhs, thrust::device_vector<T>& rhs, Policy policy)
-    {
-        bench_utils::gpu_timer d_timer;
+  template <typename T, typename Policy>
+  float64_t run(thrust::device_vector<T>& lhs, thrust::device_vector<T>& rhs, Policy policy)
+  {
+    thrust::inner_product(policy, lhs.begin(), lhs.end(), rhs.begin(), T{0});
 
-        d_timer.start(0);
-        thrust::inner_product(policy, lhs.cbegin(), lhs.cend(), rhs.begin(), T {0});
-        d_timer.stop(0);
+    bench_utils::gpu_timer d_timer;
 
-        return d_timer.get_duration();
-    }
+    d_timer.start(0);
+    thrust::inner_product(policy, lhs.begin(), lhs.end(), rhs.begin(), T{0});
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
 };
 
 template <class Benchmark, class T>
 void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string seed_type)
 {
-    // Benchmark object
-    Benchmark benchmark {};
-
-    // GPU times
-    std::vector<double> gpu_times;
-
-    // Generate input
-    thrust::device_vector<T> generator = bench_utils::generate(elements, seed_type);
-    thrust::device_vector<T> lhs       = generator;
-    thrust::device_vector<T> rhs       = generator;
-
-    bench_utils::caching_allocator_t alloc {};
-    thrust::detail::device_t         policy {};
-
-    for(auto _ : state)
-    {
-        float64_t duration = benchmark.template run<T>(lhs, rhs, policy(alloc));
-        state.SetIterationTime(duration);
-        gpu_times.push_back(duration);
-    }
-
-    // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
-    // it will actually be the global memory bandwidth gotten.
-    state.SetBytesProcessed(state.iterations() * (2 * elements + 1) * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * elements);
-
-    const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
-    state.counters["gpu_noise"] = gpu_cv;
+  // Benchmark object
+  Benchmark benchmark{};
+
+  // GPU times
+  std::vector<double> gpu_times;
+
+  // Generate input
+  thrust::device_vector<T> generator = bench_utils::generate(elements, seed_type);
+  thrust::device_vector<T> lhs       = generator;
+  thrust::device_vector<T> rhs       = generator;
+
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
+
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(lhs, rhs, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
+
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * (2 * elements + 1) * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
+
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
 }
 
-#define CREATE_BENCHMARK(T, Elements)                                                        \
-    benchmark::RegisterBenchmark(                                                            \
-        bench_utils::bench_naming::format_name("{algo:inner_product,subalgo:" + name         \
-                                               + ",input_type:" #T + ",elements:" #Elements) \
-            .c_str(),                                                                        \
-        run_benchmark<Benchmark, T>,                                                         \
-        Elements,                                                                            \
-        seed_type)
+#define CREATE_BENCHMARK(T, Elements)                                                     \
+  benchmark::RegisterBenchmark(                                                           \
+    bench_utils::bench_naming::format_name(                                               \
+      "{algo:inner_product,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements) \
+      .c_str(),                                                                           \
+    run_benchmark<Benchmark, T>,                                                          \
+    Elements,                                                                             \
+    seed_type)
 
-#define BENCHMARK_TYPE(type)                                          \
-    CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), \
-        CREATE_BENCHMARK(type, 1 << 24), CREATE_BENCHMARK(type, 1 << 28)
+#define BENCHMARK_TYPE(type)                                                                         \
+  CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), CREATE_BENCHMARK(type, 1 << 24), \
+    CREATE_BENCHMARK(type, 1 << 28)
 
 template <class Benchmark>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    const std::string                             seed_type)
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
 {
-    std::vector<benchmark::internal::Benchmark*> bs
-        = { BENCHMARK_TYPE(int8_t),
-            BENCHMARK_TYPE(int16_t),
-            BENCHMARK_TYPE(int32_t),
-            BENCHMARK_TYPE(int64_t)
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_TYPE(int8_t),
+    BENCHMARK_TYPE(int16_t),
+    BENCHMARK_TYPE(int32_t),
+    BENCHMARK_TYPE(int64_t),
 #if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT
-                ,
-            BENCHMARK_TYPE(int128_t)
+    BENCHMARK_TYPE(int128_t),
 #endif
-                ,
-            BENCHMARK_TYPE(float32_t),
-            BENCHMARK_TYPE(float64_t) };
+    BENCHMARK_TYPE(float32_t),
+    BENCHMARK_TYPE(float64_t)
+  };
 
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
 }
 
 int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<std::string>(
-        "name_format", "name_format", "human", "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    bench_utils::bench_naming::set_format(
-        parser.get<std::string>("name_format")); /* either: json,human,txt */
-    const std::string seed_type = parser.get<std::string>("seed");
-
-    // Benchmark info
-    bench_utils::add_common_benchmark_info();
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmark
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<basic>("basic", benchmarks, seed_type);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMicrosecond);
-        b->MinTime(0.4); // in seconds
-    }
-
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
-
-    // Finish
-    benchmark::Shutdown();
-    return 0;
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<basic>("basic", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
 }
diff --git a/benchmark/bench/merge/basic.cu b/benchmark/bench/merge/basic.cu
new file mode 100644
index 0000000..e8bfb4a
--- /dev/null
+++ b/benchmark/bench/merge/basic.cu
@@ -0,0 +1,187 @@
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+
+// Benchmark utils
+#include "../../bench_utils/bench_utils.hpp"
+
+// rocThrust
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+#include <thrust/merge.h>
+#include <thrust/sort.h>
+
+// Google Benchmark
+#include <benchmark/benchmark.h>
+
+// STL
+#include <cstddef>
+#include <string>
+#include <vector>
+
+struct basic
+{
+  template <typename T, typename Policy>
+  float64_t
+  run(thrust::device_vector<T>& in, thrust::device_vector<T>& out, const std::size_t elements_in_lhs, Policy policy)
+  {
+    thrust::merge(
+      policy, in.cbegin(), in.cbegin() + elements_in_lhs, in.cbegin() + elements_in_lhs, in.cend(), out.begin());
+
+    bench_utils::gpu_timer d_timer;
+
+    d_timer.start(0);
+    thrust::merge(
+      policy, in.cbegin(), in.cbegin() + elements_in_lhs, in.cbegin() + elements_in_lhs, in.cend(), out.begin());
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
+};
+
+template <class Benchmark, class T>
+void run_benchmark(benchmark::State& state,
+                   const std::size_t elements,
+                   const std::string seed_type,
+                   const int entropy_reduction,
+                   const std::size_t input_size_ratio)
+{
+  // Benchmark object
+  Benchmark benchmark{};
+
+  // GPU times
+  std::vector<double> gpu_times;
+
+  // Generate input
+  const auto entropy         = bench_utils::get_entropy_percentage(entropy_reduction) / 100.0f;
+  const auto elements_in_lhs = static_cast<std::size_t>(static_cast<double>(input_size_ratio * elements) / 100.0);
+
+  thrust::device_vector<T> in = bench_utils::generate(elements, seed_type, entropy);
+
+  thrust::sort(in.begin(), in.begin() + elements_in_lhs);
+  thrust::sort(in.begin() + elements_in_lhs, in.end());
+
+  // Output
+  thrust::device_vector<T> out(elements);
+
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
+
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(in, out, elements_in_lhs, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
+
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
+
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
+}
+
+#define CREATE_BENCHMARK(T, Elements, EntropyReduction, InputSizeRatio)                                               \
+  benchmark::RegisterBenchmark(                                                                                       \
+    bench_utils::bench_naming::format_name(                                                                           \
+      "{algo:merge,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements + ",entropy:"                        \
+      + std::to_string(bench_utils::get_entropy_percentage(EntropyReduction)) + ",input_size_ratio:" #InputSizeRatio) \
+      .c_str(),                                                                                                       \
+    run_benchmark<Benchmark, T>,                                                                                      \
+    Elements,                                                                                                         \
+    seed_type,                                                                                                        \
+    EntropyReduction,                                                                                                 \
+    InputSizeRatio)
+
+#define BENCHMARK_ELEMENTS(type, elements, entropy)                                             \
+  CREATE_BENCHMARK(type, elements, entropy, 25), CREATE_BENCHMARK(type, elements, entropy, 50), \
+    CREATE_BENCHMARK(type, elements, entropy, 75)
+
+#define BENCHMARK_TYPE_ENTROPY(type, entropy)                                             \
+  BENCHMARK_ELEMENTS(type, 1 << 16, entropy), BENCHMARK_ELEMENTS(type, 1 << 20, entropy), \
+    BENCHMARK_ELEMENTS(type, 1 << 24, entropy), BENCHMARK_ELEMENTS(type, 1 << 28, entropy)
+
+template <class Benchmark>
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
+{
+  constexpr int entropy_reductions[] = {0, 4}; // 1.000, 0.201;
+
+  for (int entropy_reduction : entropy_reductions)
+  {
+    std::vector<benchmark::internal::Benchmark*> bs = {
+      BENCHMARK_TYPE_ENTROPY(int8_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(int16_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(int32_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(int64_t, entropy_reduction),
+#if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT
+      BENCHMARK_TYPE_ENTROPY(int128_t, entropy_reduction),
+#endif
+      BENCHMARK_TYPE_ENTROPY(float, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(double, entropy_reduction)
+    };
+    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+  }
+}
+
+int main(int argc, char* argv[])
+{
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<basic>("basic", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
+}
diff --git a/benchmarks/bench/partition/basic.cu b/benchmark/bench/partition/basic.cu
similarity index 52%
rename from benchmarks/bench/partition/basic.cu
rename to benchmark/bench/partition/basic.cu
index dfa5de1..9a9aa62 100644
--- a/benchmarks/bench/partition/basic.cu
+++ b/benchmark/bench/partition/basic.cu
@@ -1,189 +1,201 @@
-// MIT License
-//
-// Copyright (c) 2024 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.
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
 
 // Benchmark utils
 #include "../../bench_utils/bench_utils.hpp"
 
 // rocThrust
-#include <thrust/copy.h>
-#include <thrust/count.h>
 #include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
-#include <thrust/iterator/reverse_iterator.h>
+#include <thrust/partition.h>
 
 // Google Benchmark
 #include <benchmark/benchmark.h>
 
 // STL
-#include <cstdlib>
+#include <cstddef>
 #include <string>
 #include <vector>
 
 template <class T>
 struct less_then_t
 {
-    T m_val;
+  T m_val;
 
-    __host__ __device__ bool operator()(const T& val) const
-    {
-        return val < m_val;
-    }
+  __host__ __device__ bool operator()(const T& val) const
+  {
+    return val < m_val;
+  }
 };
 
 struct basic
 {
-    template <typename T, typename Policy>
-    float64_t run(thrust::device_vector<T>& input,
-                  thrust::device_vector<T>& output,
-                  const std::size_t         elements,
-                  less_then_t<T>            select_op,
-                  Policy                    policy)
-    {
-        bench_utils::gpu_timer d_timer;
-
-        d_timer.start(0);
-        thrust::copy_if(policy,
-                        input.cbegin(),
-                        input.cend(),
-                        output.begin(),
-                        thrust::make_reverse_iterator(output.begin() + elements),
-                        select_op);
-        d_timer.stop(0);
-
-        return d_timer.get_duration();
-    }
+  template <typename T, typename Policy>
+  float64_t run(thrust::device_vector<T>& input,
+                thrust::device_vector<T>& output,
+                const std::size_t elements,
+                less_then_t<T> select_op,
+                Policy policy)
+  {
+    thrust::partition_copy(
+      policy,
+      input.cbegin(),
+      input.cend(),
+      output.begin(),
+      thrust::make_reverse_iterator(output.begin() + elements),
+      select_op);
+
+    bench_utils::gpu_timer d_timer;
+
+    d_timer.start(0);
+    thrust::partition_copy(
+      policy,
+      input.cbegin(),
+      input.cend(),
+      output.begin(),
+      thrust::make_reverse_iterator(output.begin() + elements),
+      select_op);
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
 };
 
 template <class Benchmark, class T>
-void run_benchmark(benchmark::State& state,
-                   const std::size_t elements,
-                   const std::string seed_type,
-                   const int         entropy_reduction)
+void run_benchmark(
+  benchmark::State& state, const std::size_t elements, const std::string seed_type, const int entropy_reduction)
 {
-    // Benchmark object
-    Benchmark benchmark {};
+  using select_op_t = less_then_t<T>;
 
-    // GPU times
-    std::vector<double> gpu_times;
+  // Benchmark object
+  Benchmark benchmark{};
 
-    // Generate input
-    T val = bench_utils::value_from_entropy<T>(
-        bench_utils::get_entropy_percentage(entropy_reduction));
-    less_then_t<T> select_op {val};
+  // GPU times
+  std::vector<double> gpu_times;
 
-    thrust::device_vector<T> input = bench_utils::generate(elements, seed_type);
+  // Generate input
+  T val = bench_utils::value_from_entropy<T>(bench_utils::get_entropy_percentage(entropy_reduction));
+  select_op_t select_op{val};
 
-    // Output
-    thrust::device_vector<T> output(elements);
+  thrust::device_vector<T> input = bench_utils::generate(elements, seed_type);
 
-    bench_utils::caching_allocator_t alloc {};
-    thrust::detail::device_t         policy {};
+  // Output
+  thrust::device_vector<T> output(elements);
 
-    for(auto _ : state)
-    {
-        float64_t duration
-            = benchmark.template run<T>(input, output, elements, select_op, policy(alloc));
-        state.SetIterationTime(duration);
-        gpu_times.push_back(duration);
-    }
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
 
-    // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
-    // it will actually be the global memory bandwidth gotten.
-    state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * elements);
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(input, output, elements, select_op, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
 
-    const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
-    state.counters["gpu_noise"] = gpu_cv;
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
+
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
 }
 
-#define CREATE_BENCHMARK(T, Elements, EntropyReduction)                                            \
-    benchmark::RegisterBenchmark(                                                                  \
-        bench_utils::bench_naming::format_name(                                                    \
-            "{algo:partition,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements         \
-            + ",entropy:" + std::to_string(bench_utils::get_entropy_percentage(EntropyReduction))) \
-            .c_str(),                                                                              \
-        run_benchmark<Benchmark, T>,                                                               \
-        Elements,                                                                                  \
-        seed_type,                                                                                 \
-        EntropyReduction)
-
-#define BENCHMARK_TYPE_ENTROPY(type, entropy)                                           \
-    CREATE_BENCHMARK(type, 1 << 16, entropy), CREATE_BENCHMARK(type, 1 << 20, entropy), \
-        CREATE_BENCHMARK(type, 1 << 24, entropy), CREATE_BENCHMARK(type, 1 << 28, entropy)
+#define CREATE_BENCHMARK(T, Elements, EntropyReduction)                                      \
+  benchmark::RegisterBenchmark(                                                              \
+    bench_utils::bench_naming::format_name(                                                  \
+      "{algo:partition,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements         \
+      + ",entropy:" + std::to_string(bench_utils::get_entropy_percentage(EntropyReduction))) \
+      .c_str(),                                                                              \
+    run_benchmark<Benchmark, T>,                                                             \
+    Elements,                                                                                \
+    seed_type,                                                                               \
+    EntropyReduction)
+
+#define BENCHMARK_TYPE_ENTROPY(type, entropy)                                         \
+  CREATE_BENCHMARK(type, 1 << 16, entropy), CREATE_BENCHMARK(type, 1 << 20, entropy), \
+    CREATE_BENCHMARK(type, 1 << 24, entropy), CREATE_BENCHMARK(type, 1 << 28, entropy)
 
 template <class Benchmark>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    const std::string                             seed_type)
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
 {
-    constexpr int entropy_reductions[] = {0, 2, 6}; // 1.000, 0.544, 0.000;
-
-    for(int entropy_reduction : entropy_reductions)
-    {
-        std::vector<benchmark::internal::Benchmark*> bs
-            = {BENCHMARK_TYPE_ENTROPY(int8_t, entropy_reduction),
-               BENCHMARK_TYPE_ENTROPY(int16_t, entropy_reduction),
-               BENCHMARK_TYPE_ENTROPY(int32_t, entropy_reduction),
-               BENCHMARK_TYPE_ENTROPY(int64_t, entropy_reduction),
-               BENCHMARK_TYPE_ENTROPY(float, entropy_reduction),
-               BENCHMARK_TYPE_ENTROPY(double, entropy_reduction)};
-        benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
-    }
+  constexpr int entropy_reductions[] = {0, 2, 4200}; // 1.000, 0.544, 0.000;
+
+  for (int entropy_reduction : entropy_reductions)
+  {
+    std::vector<benchmark::internal::Benchmark*> bs = {
+      BENCHMARK_TYPE_ENTROPY(int8_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(int16_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(int32_t, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(int64_t, entropy_reduction),
+#if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT
+      BENCHMARK_TYPE_ENTROPY(int128_t, entropy_reduction),
+#endif
+      BENCHMARK_TYPE_ENTROPY(float, entropy_reduction),
+      BENCHMARK_TYPE_ENTROPY(double, entropy_reduction)
+    };
+    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+  }
 }
 
 int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<std::string>(
-        "name_format", "name_format", "human", "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    bench_utils::bench_naming::set_format(
-        parser.get<std::string>("name_format")); /* either: json,human,txt */
-    const std::string seed_type = parser.get<std::string>("seed");
-
-    // Benchmark info
-    bench_utils::add_common_benchmark_info();
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmark
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<basic>("basic", benchmarks, seed_type);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMicrosecond);
-        b->MinTime(0.4); // in seconds
-    }
-
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
-
-    // Finish
-    benchmark::Shutdown();
-    return 0;
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<basic>("basic", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
 }
diff --git a/benchmarks/bench/transform_reduce/sum.cu b/benchmark/bench/reduce/basic.cu
similarity index 54%
rename from benchmarks/bench/transform_reduce/sum.cu
rename to benchmark/bench/reduce/basic.cu
index 6dc60ec..f08025b 100644
--- a/benchmarks/bench/transform_reduce/sum.cu
+++ b/benchmark/bench/reduce/basic.cu
@@ -1,24 +1,30 @@
-// MIT License
-//
-// Copyright (c) 2024 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.
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
 
 // Benchmark utils
 #include "../../bench_utils/bench_utils.hpp"
@@ -32,125 +38,121 @@
 #include <benchmark/benchmark.h>
 
 // STL
-#include <cstdlib>
+#include <cstddef>
 #include <string>
 #include <vector>
 
-struct sum
+struct basic
 {
-    template <typename T, typename Policy>
-    float64_t run(thrust::device_vector<T>& input, Policy policy)
-    {
-        bench_utils::gpu_timer d_timer;
-
-        d_timer.start(0);
-        bench_utils::do_not_optimize(
-            thrust::reduce(policy, input.begin(), input.end(), T {}, thrust::plus<T> {}));
-        d_timer.stop(0);
-
-        return d_timer.get_duration();
-    }
+  template <typename T, typename Policy>
+  float64_t run(thrust::device_vector<T>& in, Policy policy)
+  {
+    bench_utils::do_not_optimize(thrust::reduce(policy, in.begin(), in.end()));
+
+    bench_utils::gpu_timer d_timer;
+
+    d_timer.start(0);
+    bench_utils::do_not_optimize(thrust::reduce(policy, in.begin(), in.end()));
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
 };
 
 template <class Benchmark, class T>
 void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string seed_type)
 {
-    // Benchmark object
-    Benchmark benchmark {};
+  // Benchmark object
+  Benchmark benchmark{};
 
-    // GPU times
-    std::vector<double> gpu_times;
+  // GPU times
+  std::vector<double> gpu_times;
 
-    // Generate input
-    thrust::device_vector<T> input = bench_utils::generate(elements, seed_type);
+  // Generate input
+  thrust::device_vector<T> in = bench_utils::generate(elements, seed_type);
 
-    bench_utils::caching_allocator_t alloc {};
-    thrust::detail::device_t         policy {};
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
 
-    for(auto _ : state)
-    {
-        float64_t duration = benchmark.template run<T>(input, policy(alloc));
-        state.SetIterationTime(duration);
-        gpu_times.push_back(duration);
-    }
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<T>(in, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
 
-    // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
-    // it will actually be the global memory bandwidth gotten.
-    state.SetBytesProcessed(state.iterations() * (elements + 1) * sizeof(T));
-    state.SetItemsProcessed(state.iterations() * elements);
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * (elements + 1) * sizeof(T));
+  state.SetItemsProcessed(state.iterations() * elements);
 
-    const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
-    state.counters["gpu_noise"] = gpu_cv;
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
 }
 
-#define CREATE_BENCHMARK(T, Elements)                                                        \
-    benchmark::RegisterBenchmark(                                                            \
-        bench_utils::bench_naming::format_name("{algo:transform_reduce,subalgo:" + name      \
-                                               + ",input_type:" #T + ",elements:" #Elements) \
-            .c_str(),                                                                        \
-        run_benchmark<Benchmark, T>,                                                         \
-        Elements,                                                                            \
-        seed_type)
+#define CREATE_BENCHMARK(T, Elements)                                                                                   \
+  benchmark::RegisterBenchmark(                                                                                         \
+    bench_utils::bench_naming::format_name("{algo:reduce,subalgo:" + name + ",input_type:" #T + ",elements:" #Elements) \
+      .c_str(),                                                                                                         \
+    run_benchmark<Benchmark, T>,                                                                                        \
+    Elements,                                                                                                           \
+    seed_type)
 
-#define BENCHMARK_TYPE(type)                                          \
-    CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), \
-        CREATE_BENCHMARK(type, 1 << 24), CREATE_BENCHMARK(type, 1 << 28)
+#define BENCHMARK_TYPE(type)                                                                         \
+  CREATE_BENCHMARK(type, 1 << 16), CREATE_BENCHMARK(type, 1 << 20), CREATE_BENCHMARK(type, 1 << 24), \
+    CREATE_BENCHMARK(type, 1 << 28)
 
 template <class Benchmark>
-void add_benchmarks(const std::string&                            name,
-                    std::vector<benchmark::internal::Benchmark*>& benchmarks,
-                    const std::string                             seed_type)
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
 {
-    std::vector<benchmark::internal::Benchmark*> bs
-        = { BENCHMARK_TYPE(int8_t),
-            BENCHMARK_TYPE(int16_t),
-            BENCHMARK_TYPE(int32_t),
-            BENCHMARK_TYPE(int64_t)
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_TYPE(int8_t),
+    BENCHMARK_TYPE(int16_t),
+    BENCHMARK_TYPE(int32_t),
+    BENCHMARK_TYPE(int64_t),
 #if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT
-                ,
-            BENCHMARK_TYPE(int128_t)
+    BENCHMARK_TYPE(int128_t),
 #endif
-                ,
-            BENCHMARK_TYPE(float32_t),
-            BENCHMARK_TYPE(float64_t) };
+    BENCHMARK_TYPE(float32_t),
+    BENCHMARK_TYPE(float64_t)
+  };
 
-    benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
 }
 
 int main(int argc, char* argv[])
 {
-    cli::Parser parser(argc, argv);
-    parser.set_optional<std::string>(
-        "name_format", "name_format", "human", "either: json,human,txt");
-    parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
-    parser.run_and_exit_if_error();
-
-    // Parse argv
-    benchmark::Initialize(&argc, argv);
-    bench_utils::bench_naming::set_format(
-        parser.get<std::string>("name_format")); /* either: json,human,txt */
-    const std::string seed_type = parser.get<std::string>("seed");
-
-    // Benchmark info
-    bench_utils::add_common_benchmark_info();
-    benchmark::AddCustomContext("seed", seed_type);
-
-    // Add benchmark
-    std::vector<benchmark::internal::Benchmark*> benchmarks;
-    add_benchmarks<sum>("sum", benchmarks, seed_type);
-
-    // Use manual timing
-    for(auto& b : benchmarks)
-    {
-        b->UseManualTime();
-        b->Unit(benchmark::kMicrosecond);
-        b->MinTime(0.4); // in seconds
-    }
-
-    // Run benchmarks
-    benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
-
-    // Finish
-    benchmark::Shutdown();
-    return 0;
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<basic>("basic", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
 }
diff --git a/benchmark/bench/reduce/by_key.cu b/benchmark/bench/reduce/by_key.cu
new file mode 100644
index 0000000..ea4dee5
--- /dev/null
+++ b/benchmark/bench/reduce/by_key.cu
@@ -0,0 +1,211 @@
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+
+// Benchmark utils
+#include "../../bench_utils/bench_utils.hpp"
+
+// rocThrust
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+#include <thrust/reduce.h>
+#include <thrust/unique.h>
+
+// Google Benchmark
+#include <benchmark/benchmark.h>
+
+// STL
+#include <cstddef>
+#include <string>
+#include <vector>
+
+struct by_key
+{
+  template <typename KeyT, typename ValueT, typename Policy>
+  float64_t run(thrust::device_vector<KeyT>& in_keys,
+                thrust::device_vector<ValueT>& in_vals,
+                thrust::device_vector<KeyT>& out_keys,
+                thrust::device_vector<ValueT>& out_vals,
+                Policy policy)
+  {
+    thrust::reduce_by_key(policy, in_keys.begin(), in_keys.end(), in_vals.begin(), out_keys.begin(), out_vals.begin());
+
+    bench_utils::gpu_timer d_timer;
+
+    d_timer.start(0);
+    thrust::reduce_by_key(policy, in_keys.begin(), in_keys.end(), in_vals.begin(), out_keys.begin(), out_vals.begin());
+    d_timer.stop(0);
+
+    return d_timer.get_duration();
+  }
+};
+
+template <class Benchmark, class KeyT, class ValueT>
+void run_benchmark(
+  benchmark::State& state, const std::size_t elements, const std::string seed_type, const std::size_t max_segment_size)
+{
+  // Benchmark object
+  Benchmark benchmark{};
+
+  // GPU times
+  std::vector<double> gpu_times;
+
+  constexpr std::size_t min_segment_size = 1;
+
+  // Generate input and output
+  thrust::device_vector<KeyT> in_keys;
+  thrust::device_vector<KeyT> out_keys;
+  thrust::device_vector<ValueT> in_vals;
+  try
+  {
+    in_keys  = bench_utils::generate.uniform.key_segments(elements, seed_type, min_segment_size, max_segment_size);
+    out_keys = in_keys;
+    in_vals  = thrust::device_vector<ValueT>(elements);
+  }
+  catch (const ::thrust::system::detail::bad_alloc& e)
+  {
+    (void) hipGetLastError();
+    state.SkipWithError(("thrust::system::detail::bad_alloc: " + std::string(e.what())).c_str());
+    return;
+  }
+  const std::size_t unique_keys = thrust::distance(out_keys.begin(), thrust::unique(out_keys.begin(), out_keys.end()));
+  thrust::device_vector<ValueT> out_vals;
+  try
+  {
+    out_vals = thrust::device_vector<ValueT>(unique_keys);
+  }
+  catch (const ::thrust::system::detail::bad_alloc& e)
+  {
+    (void) hipGetLastError();
+    state.SkipWithError(("thrust::system::detail::bad_alloc: " + std::string(e.what())).c_str());
+    return;
+  }
+
+  bench_utils::caching_allocator_t alloc{};
+  thrust::detail::device_t policy{};
+
+  for (auto _ : state)
+  {
+    float64_t duration = benchmark.template run<KeyT, ValueT>(in_keys, in_vals, out_keys, out_vals, policy(alloc));
+    state.SetIterationTime(duration);
+    gpu_times.push_back(duration);
+  }
+
+  // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported
+  // it will actually be the global memory bandwidth gotten.
+  state.SetBytesProcessed(state.iterations() * ((elements + unique_keys) * (sizeof(KeyT) + sizeof(ValueT))));
+  state.SetItemsProcessed(state.iterations() * elements);
+
+  const double gpu_cv         = bench_utils::StatisticsCV(gpu_times);
+  state.counters["gpu_noise"] = gpu_cv;
+}
+
+#define CREATE_BENCHMARK(KeyT, ValueT, Elements, MaxSegmentSize)                                            \
+  benchmark::RegisterBenchmark(                                                                             \
+    bench_utils::bench_naming::format_name(                                                                 \
+      "{algo:reduce,subalgo:" + name + ",key_type:" #KeyT + ",value_type:" #ValueT + ",elements:" #Elements \
+      + ",max_segment_size:" #MaxSegmentSize)                                                               \
+      .c_str(),                                                                                             \
+    run_benchmark<Benchmark, KeyT, ValueT>,                                                                 \
+    Elements,                                                                                               \
+    seed_type,                                                                                              \
+    MaxSegmentSize)
+
+#define BENCHMARK_ELEMENTS(key_type, value_type, elements)                                                  \
+  CREATE_BENCHMARK(key_type, value_type, elements, 1), CREATE_BENCHMARK(key_type, value_type, elements, 4), \
+    CREATE_BENCHMARK(key_type, value_type, elements, 8)
+
+#define BENCHMARK_VALUE_TYPE(key_type, value_type)                                                      \
+  BENCHMARK_ELEMENTS(key_type, value_type, 1 << 16), BENCHMARK_ELEMENTS(key_type, value_type, 1 << 20), \
+    BENCHMARK_ELEMENTS(key_type, value_type, 1 << 24), BENCHMARK_ELEMENTS(key_type, value_type, 1 << 28)
+
+#if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT
+#  define BENCHMARK_KEY_TYPE(key_type)                                                  \
+    BENCHMARK_VALUE_TYPE(key_type, int8_t), BENCHMARK_VALUE_TYPE(key_type, int16_t),    \
+      BENCHMARK_VALUE_TYPE(key_type, int32_t), BENCHMARK_VALUE_TYPE(key_type, int64_t), \
+      BENCHMARK_VALUE_TYPE(key_type, int128_t), BENCHMARK_VALUE_TYPE(key_type, float),  \
+      BENCHMARK_VALUE_TYPE(key_type, double)
+#else
+#  define BENCHMARK_KEY_TYPE(key_type)                                                  \
+    BENCHMARK_VALUE_TYPE(key_type, int8_t), BENCHMARK_VALUE_TYPE(key_type, int16_t),    \
+      BENCHMARK_VALUE_TYPE(key_type, int32_t), BENCHMARK_VALUE_TYPE(key_type, int64_t), \
+      BENCHMARK_VALUE_TYPE(key_type, float), BENCHMARK_VALUE_TYPE(key_type, double)
+#endif
+
+template <class Benchmark>
+void add_benchmarks(
+  const std::string& name, std::vector<benchmark::internal::Benchmark*>& benchmarks, const std::string seed_type)
+{
+  std::vector<benchmark::internal::Benchmark*> bs = {
+    BENCHMARK_KEY_TYPE(int8_t),
+    BENCHMARK_KEY_TYPE(int16_t),
+    BENCHMARK_KEY_TYPE(int32_t),
+    BENCHMARK_KEY_TYPE(int64_t),
+#if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT
+    BENCHMARK_KEY_TYPE(int128_t),
+#endif
+    BENCHMARK_KEY_TYPE(float32_t),
+    BENCHMARK_KEY_TYPE(float64_t)
+  };
+  benchmarks.insert(benchmarks.end(), bs.begin(), bs.end());
+}
+
+int main(int argc, char* argv[])
+{
+  cli::Parser parser(argc, argv);
+  parser.set_optional<std::string>("name_format", "name_format", "human", "either: json,human,txt");
+  parser.set_optional<std::string>("seed", "seed", "random", bench_utils::get_seed_message());
+  parser.run_and_exit_if_error();
+
+  // Parse argv
+  benchmark::Initialize(&argc, argv);
+  bench_utils::bench_naming::set_format(parser.get<std::string>("name_format")); /* either: json,human,txt */
+  const std::string seed_type = parser.get<std::string>("seed");
+
+  // Benchmark info
+  bench_utils::add_common_benchmark_info();
+  benchmark::AddCustomContext("seed", seed_type);
+
+  // Add benchmark
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  add_benchmarks<by_key>("by_key", benchmarks, seed_type);
+
+  // Use manual timing
+  for (auto& b : benchmarks)
+  {
+    b->UseManualTime();
+    b->Unit(benchmark::kMicrosecond);
+    b->MinTime(0.4); // in seconds
+  }
+
+  // Run benchmarks
+  benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter());
+
+  // Finish
+  benchmark::Shutdown();
+  return 0;
+}
diff --git a/benchmark/bench/scan/exclusive/by_key.cu b/benchmark/bench/scan/exclusive/by_key.cu
new file mode 100644
index 0000000..02a4f52
--- /dev/null
+++ b/benchmark/bench/scan/exclusive/by_key.cu
@@ -0,0 +1,177 @@
+/******************************************************************************
+ * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
+ * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE