feat(rocthrust): allow grafting thrust onto benchmarks

projects/rocrand/cmake/SetupNVCC.cmake

@@ -61,7 +61,7 @@ function(hip_cuda_detect_lowest_cc out_variable)
     endif()
 
     if(NOT HIP_CUDA_lowest_cc)
-        set(HIP_CUDA_lowest_cc "35")
+        set(HIP_CUDA_lowest_cc "52")
         set(${out_variable} ${HIP_CUDA_lowest_cc} PARENT_SCOPE)
     else()
         set(${out_variable} ${HIP_CUDA_lowest_cc} PARENT_SCOPE)

projects/rocthrust/CMakeLists.txt

@@ -71,6 +71,8 @@ cmake_dependent_option(USE_SYSTEM_LIB "Use existing system ROCm library installa
 option(EXTERNAL_DEPS_FORCE_DOWNLOAD "Force download of non-ROCm dependencies (eg. Google Test, Google Benchmark)" OFF)
 option(LINK_HIP_DEVICE_LIBS "Links against HIP device libraries. Required for device-side acceleration." ON)
 
+option(GRAFT_THRUST_ONTO_BINARIES "!!! TODO !!!" OFF)
+
 if((BUILD_TEST OR BUILD_BENCHMARK OR BUILD_EXAMPLE) AND NOT (${LINK_HIP_DEVICE_LIBS}))
     message(FATAL_ERROR "rocThrust tests, benchmarks, and examples must be built with LINK_HIP_DEVICE_LIBS=ON since they require device acceleration.")
 endif()
@@ -178,7 +180,7 @@ check_option(FETCH_METHOD_OPTIONS ROCRAND_FETCH_METHOD)
 include(cmake/Dependencies.cmake)
 
 # Verify that supported compilers are used
-if (NOT WIN32)
+if (NOT WIN32 OR NOT GRAFT_THRUST_ONTO_BINARIES)
   include(cmake/VerifyCompiler.cmake)
 endif()
 
@@ -191,11 +193,13 @@ set(THRUST_HOST_SYSTEM CPP CACHE STRING "The host backend to target.")
 check_option(THRUST_HOST_SYSTEM_OPTIONS THRUST_HOST_SYSTEM)
 add_compile_definitions(THRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_${THRUST_HOST_SYSTEM})
 
-# Note CUDA backend is currently not supported, so disallow the option for now.
-set(THRUST_DEVICE_SYSTEM_OPTIONS OMP TBB CPP HIP)
-set(THRUST_DEVICE_SYSTEM HIP CACHE STRING "The host backend to target.")
-check_option(THRUST_DEVICE_SYSTEM_OPTIONS THRUST_DEVICE_SYSTEM)
-add_compile_definitions(THRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_${THRUST_DEVICE_SYSTEM})
+# If we're not grafting existing CUDA onto binaries, we must dissallow CUDA backend.
+if (NOT GRAFT_THRUST_ONTO_BINARIES)
+  set(THRUST_DEVICE_SYSTEM_OPTIONS OMP TBB CPP HIP)
+  set(THRUST_DEVICE_SYSTEM HIP CACHE STRING "The host backend to target.")
+  check_option(THRUST_DEVICE_SYSTEM_OPTIONS THRUST_DEVICE_SYSTEM)
+  add_compile_definitions(THRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_${THRUST_DEVICE_SYSTEM})
+endif()
 
 set(COMPILE_OPTIONS -Wall -Wextra)
 if(NOT DISABLE_WERROR)
@@ -230,8 +234,13 @@ endif()
 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)
+# Thrust (header-only) library
+if(GRAFT_THRUST_ONTO_BINARIES)
+  # If we're grafting CCCL we don't need the actual rocThrust library!
+  find_package(CCCL CONFIG REQUIRED)
+else()
+  add_subdirectory(thrust)
+endif()
 
 if(BUILD_TEST OR BUILD_BENCHMARK OR BUILD_HIPSTDPAR_TEST)
   rocm_package_setup_component(clients)
@@ -242,7 +251,7 @@ if(CODE_COVERAGE)
 endif()
 
 # Tests
-if(BUILD_TEST AND USE_SYSTEM_LIB)
+if(BUILD_TEST AND USE_SYSTEM_LIB AND NOT GRAFT_THRUST_ONTO_BINARIES)
   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.")
@@ -261,7 +270,7 @@ if(BUILD_TEST OR BUILD_HIPSTDPAR_TEST)
 
   # Used by tests to define their RESOURCE_GROUPS
   set(DEFAULT_GPU_RESOURCE "gpus")
-  
+
   # 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)

projects/rocthrust/CMakePresets.json

@@ -0,0 +1,56 @@
+{
+    "version": 3,
+    "cmakeMinimumRequired": {
+        "major": 3,
+        "minor": 28,
+        "patch": 0
+    },
+    "configurePresets": [
+        {
+            "name": "hip",
+            "hidden": true,
+            "generator": "Ninja",
+            "binaryDir": "${sourceDir}/build/${presetName}",
+            "cacheVariables": {
+                "CMAKE_CXX_STANDARD": "17",
+                "CMAKE_BUILD_TYPE": "Release",
+                "CMAKE_EXPORT_COMPILE_COMMANDS": true
+            }
+        },
+        {
+            "inherits": "hip",
+            "name": "hip-amd",
+            "cacheVariables": {
+              "HIP_PLATFORM": "amd"
+            }
+        },
+        {
+            "inherits": "hip-amd",
+            "name": "hip-amd-all",
+            "cacheVariables": {
+              "BUILD_BENCHMARK": true,
+              "BUILD_TEST": true
+            }
+        },
+        {
+            "inherits": "hip-amd-all",
+            "name": "hip-amd-dev",
+            "cacheVariables": {
+              "ROCRAND_FETCH_METHOD": "MONOREPO"
+            }
+        },
+        {
+            "inherits": "hip",
+            "name": "hip-nv-dev",
+            "cacheVariables": {
+                "BUILD_BENCHMARK": true,
+                "CMAKE_HIP_FLAGS": "--extended-lambda -w",
+                "GRAFT_THRUST_ONTO_BINARIES": true,
+                "HIP_PLATFORM": "nvidia",
+                "ROCRAND_FETCH_METHOD": "MONOREPO",
+                "THRUST_DEVICE_SYSTEM": "CUDA",
+                "USE_HIPCXX": true
+            }
+        }
+    ]
+}

projects/rocthrust/benchmark/bench/transform/basic.cu

@@ -32,7 +32,6 @@
 // rocThrust
 #include <thrust/copy.h>
 #include <thrust/count.h>
-#include <thrust/detail/functional/address_stability.h>
 #include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
 #include <thrust/iterator/zip_iterator.h>
@@ -50,6 +49,8 @@
 #  include <utility>
 #endif
 
+#include _THRUST_LIBCXX_INCLUDE(functional)
+
 template <class InT, class OutT>
 struct fib_t
 {
@@ -187,7 +188,7 @@ struct mul
   {
     const T scalar = startScalar;
     return bench_transform(
-      c.begin(), c.end(), b.begin(), ::thrust::detail::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ci) {
+      c.begin(), c.end(), b.begin(), _THRUST_LIBCXX::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ci) {
         return ci * scalar;
       }));
   }
@@ -205,7 +206,7 @@ struct add
       a.end(),
       b.begin(),
       c.begin(),
-      ::thrust::detail::proclaim_copyable_arguments([] THRUST_DEVICE(const T& ai, const T& bi) -> T {
+      _THRUST_LIBCXX::proclaim_copyable_arguments([] THRUST_DEVICE(const T& ai, const T& bi) -> T {
         return ai + bi;
       }));
   }
@@ -225,7 +226,7 @@ struct triad
       b.end(),
       c.begin(),
       a.begin(),
-      ::thrust::detail::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& bi, const T& ci) {
+      _THRUST_LIBCXX::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& bi, const T& ci) {
         return bi + scalar * ci;
       }));
   }
@@ -245,7 +246,7 @@ struct nstream
       thrust::make_zip_iterator(a.end(), b.end(), c.end()),
       a.begin(),
       thrust::make_zip_function(
-        ::thrust::detail::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ai, const T& bi, const T& ci) {
+        _THRUST_LIBCXX::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ai, const T& bi, const T& ci) {
           return ai + bi + scalar * ci;
         })));
   }

projects/rocthrust/benchmark/bench_utils/bench_utils.hpp

@@ -26,10 +26,11 @@
 // Utils
 #include <thrust/execution_policy.h>
 
-#include "cmdparser.hpp"
-#include "common/types.hpp"
-#include "custom_reporter.hpp"
-#include "generation_utils.hpp"
+#include "cmdparser.hpp" // IWYU pragma: export
+#include "common/types.hpp" // IWYU pragma: export
+#include "custom_reporter.hpp" // IWYU pragma: export
+#include "generation_utils.hpp" // IWYU pragma: export
+#include "thrust_compat.hpp" // IWYU pragma: export
 
 // HIP/CUDA
 #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP
@@ -57,7 +58,7 @@
 
 namespace bench_utils
 {
-#if (THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HIP)
+#if (THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HIP) || (THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA)
 
 #  define HIP_CHECK(condition)                                                     \
     {                                                                              \
@@ -280,7 +281,14 @@ inline void add_common_benchmark_info()
   num("hdp_max_shared_memory_per_multi_processor", devProp.maxSharedMemoryPerMultiProcessor);
   num("hdp_is_multi_gpu_board", devProp.isMultiGpuBoard);
   num("hdp_can_map_host_memory", devProp.canMapHostMemory);
+
+#ifdef __NVCC__
+  // NVIDIA GPUs do not have a GCN Architecture Name; devProp.gcnArchName is uninitialized
+  str("hdp_gcn_arch_name", "");
+#else
   str("hdp_gcn_arch_name", devProp.gcnArchName);
+#endif
+
   num("hdp_integrated", devProp.integrated);
   num("hdp_cooperative_launch", devProp.cooperativeLaunch);
   num("hdp_cooperative_multi_device_launch", devProp.cooperativeMultiDeviceLaunch);

projects/rocthrust/benchmark/bench_utils/generation_utils.hpp

@@ -46,9 +46,6 @@
 #include <thrust/sequence.h>
 #include <thrust/tuple.h>
 
-// rocPRIM
-#include <rocprim/rocprim.hpp>
-
 // rocRAND
 #include <rocrand/rocrand.h>
 

projects/rocthrust/benchmark/bench_utils/thrust_compat.hpp

@@ -0,0 +1,38 @@
+// MIT License
+//
+// Copyright (c) 2026 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.
+
+// This header allows running rocThrust benchmarks with CCCL Thrust.
+
+#ifndef ROCTHRUST_BENCHMARKS_BENCH_UTILS_THRUST_COMPAT_HPP_
+#define ROCTHRUST_BENCHMARKS_BENCH_UTILS_THRUST_COMPAT_HPP_
+
+#ifndef _THRUST_HAS_DEVICE_SYSTEM_STD
+
+#define THRUST_HOST_DEVICE __host__ __device__
+#define THRUST_DEVICE __device__
+
+#define _THRUST_LIBCXX_INCLUDE(LIB) <cuda/LIB>
+#define _THRUST_STD ::cuda::std
+#define _THRUST_LIBCXX ::cuda
+
+#endif // _THRUST_HAS_DEVICE_SYSTEM_STD
+#endif // ROCTHRUST_BENCHMARKS_BENCH_UTILS_THRUST_COMPAT_HPP_

projects/rocthrust/cmake/Benchmarks.cmake

@@ -25,11 +25,15 @@ function(add_thrust_benchmark BENCHMARK_NAME BENCHMARK_SOURCE NOT_INTERNAL)
     add_executable(${BENCHMARK_TARGET} ${BENCHMARK_SOURCE})
 
     target_compile_options(${BENCHMARK_TARGET} PRIVATE ${COMPILE_OPTIONS})
-    target_link_libraries(${BENCHMARK_TARGET}
-        PRIVATE
-            rocthrust
-            roc::rocprim_hip
-    )
+    if(GRAFT_THRUST_ONTO_BINARIES)
+        target_link_libraries(${BENCHMARK_TARGET} INTERFACE CCCL::CCCL)
+    else()
+        target_link_libraries(${BENCHMARK_TARGET}
+          PRIVATE
+              rocthrust
+              roc::rocprim_hip
+        )
+    endif()
     # Internal benchmark does not use Google Benchmark nor rocRAND.
     # This can be omited when that benchmark is removed.
     if(NOT_INTERNAL)

projects/rocthrust/cmake/Dependencies.cmake

@@ -244,7 +244,7 @@ function(fetch_dep method repo_name repo_path download_branch)
   endif()
 endfunction()
 
-if(${LINK_HIP_DEVICE_LIBS})
+if(${LINK_HIP_DEVICE_LIBS} AND NOT GRAFT_THRUST_ONTO_BINARIES)
   fetch_dep(ROCPRIM_FETCH_METHOD rocprim ROCPRIM_PATH ROCM_DEP_RELEASE_BRANCH)
 
   if(${ROCPRIM_FETCH_METHOD} STREQUAL "DOWNLOAD" OR ${ROCPRIM_FETCH_METHOD} STREQUAL "MONOREPO")
@@ -415,12 +415,12 @@ if(BUILD_BENCHMARK)
     set(_ROCTHRUST_DISABLE_ROCM_CHECKS FALSE)
 	# Clang on Windows throws the following warnings with Googlebenchmark v1.9.5 (along with Werror):
     # googlebench-src/src/string_util.cc:158:34: error: format string is not a string literal [-Werror,-Wformat-nonliteral]
-	if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang" AND WIN32)  
-	  if(TARGET benchmark)  
-	    target_compile_options(benchmark PRIVATE -Wno-format-nonliteral -Wno-missing-format-attribute -Wno-unused-command-line-argument)  
-	  endif()  
-	  if(TARGET benchmark_main)  
-	    target_compile_options(benchmark_main PRIVATE -Wno-format-nonliteral -Wno-missing-format-attribute -Wno-unused-command-line-argument)	
+	if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang" AND WIN32)
+	  if(TARGET benchmark)
+	    target_compile_options(benchmark PRIVATE -Wno-format-nonliteral -Wno-missing-format-attribute -Wno-unused-command-line-argument)
+	  endif()
+	  if(TARGET benchmark_main)
+	    target_compile_options(benchmark_main PRIVATE -Wno-format-nonliteral -Wno-missing-format-attribute -Wno-unused-command-line-argument)
 	  endif()
       if(NOT TARGET benchmark::benchmark)
         add_library(benchmark::benchmark ALIAS benchmark)
@@ -443,7 +443,7 @@ if(BUILD_BENCHMARK)
     if(CMAKE_CXX_COMPILER_LAUNCHER)
       set(EXTRA_CMAKE_ARGS "${EXTRA_CMAKE_ARGS} -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER}")
     endif()
-    
+
     FetchContent_Declare(
       rocrand
       SOURCE_DIR    ${ROCRAND_PATH}

projects/rocthrust/cmake/Summary.cmake

@@ -83,9 +83,9 @@ endif()
     message(STATUS "  Install prefix        : ${CMAKE_INSTALL_PREFIX}")
 if(HIP_COMPILER STREQUAL "clang")
 if(USE_HIPCXX)
-    message(STATUS "  Device targets        : ${CMAKE_HIP_ARCHITECTURES}")
+    message(STATUS "  HIP_ARCHITECTURES     : ${HIP_ARCHITECTURES}")
 else()
-    message(STATUS "  Device targets        : ${GPU_TARGETS}")
+    message(STATUS "  GPU_TARGETS           : ${GPU_TARGETS}")
 endif()
 endif()
     message(STATUS "")

projects/rocthrust/cmake/VerifyCompiler.cmake

@@ -25,10 +25,6 @@ if(${LINK_HIP_DEVICE_LIBS})
     find_package(hip REQUIRED CONFIG PATHS /opt/rocm)
 endif()
 
-if(HIP_COMPILER STREQUAL "nvcc")
-    message(FATAL_ERROR "rocThrust does not support the CUDA backend.")
-endif()
-
 if(${LINK_HIP_DEVICE_LIBS})
     # When building for HIP, make sure we have a hip-aware clang.
     if(HIP_COMPILER STREQUAL "clang")
@@ -41,7 +37,5 @@ if(${LINK_HIP_DEVICE_LIBS})
                 message(FATAL_ERROR "When LINK_HIP_DEVICE_LIBS is set to 'ON', then 'hipcc' or a HIP-aware Clang must be used as the C++ compiler.")
             endif()
         endif()
-    else()
-        message(FATAL_ERROR "When LINK_HIP_DEVICE_LIBS is set to 'ON', HIP_COMPILER must be `clang` (AMD ROCm platform)")
     endif()
 endif()

projects/rocthrust/test/CMakeLists.txt

@@ -64,8 +64,10 @@ function(add_rocthrust_test TEST)
             ${sqlite_local_SOURCE_DIR}
     )
 
-    if(USE_SYSTEM_LIB)
-        target_link_libraries(${TEST_TARGET} PRIVATE roc::rocthrust)
+    if(GRAFT_THRUST_ONTO_BINARIES)
+        target_link_libraries(${TEST_TARGET} INTERFACE CCCL::CCCL)
+    elseif(USE_SYSTEM_LIB)
+        target_link_libraries(${TEST_TARGET} INTERFACE roc::rocthrust)
     else()
         target_link_libraries(${TEST_TARGET}
             PRIVATE

projects/rocthrust/testing/CMakeLists.txt

@@ -10,7 +10,9 @@ target_include_directories(testing_common
         $<BUILD_INTERFACE:${CMAKE_CURRENT_LIST_DIR}>
 )
 
-if(USE_SYSTEM_LIB)
+if(GRAFT_THRUST_ONTO_BINARIES)
+    target_link_libraries(testing_common INTERFACE CCCL::CCCL)
+elseif(USE_SYSTEM_LIB)
     target_link_libraries(testing_common INTERFACE roc::rocthrust)
 else()
     target_link_libraries(testing_common

projects/rocthrust/testing/allocator.cu

@@ -17,7 +17,6 @@
 
 #include <thrust/detail/config.h>
 
-#include <thrust/detail/nv_target.h>
 #include <thrust/device_malloc_allocator.h>
 #include <thrust/system/cpp/vector.h>