Revert "feat(rocthrust): CCCL 3.0.x changes"

projects/rocprim/test/rocprim/test_device_segmented_topk.cpp

@@ -910,26 +910,6 @@ void segmented_topk_large_sizes_test(bool debug_synchronous)
 
 TEST(RocprimDeviceSegmentedTopkTests, SegmentedTopkLargeSizesStable)
 {
-    // Temporarily disable this test on gfx115x on Windows until we can determine the root cause.
-    // TODO: remove this after the root cause has been found and fixed properly.
-#if defined(_WIN32)
-    rocprim::detail::target_arch arch = rocprim::detail::target_arch::unknown;
-    if (rocprim::detail::host_target_arch(hipStreamDefault, arch) != HIP_SUCCESS)
-    {
-        std::cerr << "Warning: unable to fetch target architecture for disablement check." << std::endl;
-    }
-
-    const std::set<rocprim::detail::target_arch> disabled_arches = {
-        rocprim::detail::target_arch::gfx1150,
-        rocprim::detail::target_arch::gfx1151,
-        rocprim::detail::target_arch::gfx1152,
-        rocprim::detail::target_arch::gfx1153
-    };
-
-    if (disabled_arches.find(arch) != disabled_arches.end())
-        GTEST_SKIP() << "Skipping test on gfx1151 Windows systems due to known issues.";
-#endif
-
     constexpr bool descending            = false;
     constexpr bool stable                = true;
     using decomposer_t                   = rocprim::identity_decomposer;
@@ -954,26 +934,6 @@ TEST(RocprimDeviceSegmentedTopkTests, SegmentedTopkLargeSizesStable)
 
 TEST(RocprimDeviceSegmentedTopkTests, TopkLargeSizesUnstable)
 {
-    // Temporarily disable this test on gfx115x on Windows until we can determine the root cause.
-    // TODO: remove this after the root cause has been found and fixed properly.
-#if defined(_WIN32)
-    rocprim::detail::target_arch arch = rocprim::detail::target_arch::unknown;
-    if (rocprim::detail::host_target_arch(hipStreamDefault, arch) != HIP_SUCCESS)
-    {
-        std::cerr << "Warning: unable to fetch target architecture for disablement check." << std::endl;
-    }
-
-    const std::set<rocprim::detail::target_arch> disabled_arches = {
-        rocprim::detail::target_arch::gfx1150,
-        rocprim::detail::target_arch::gfx1151,
-        rocprim::detail::target_arch::gfx1152,
-        rocprim::detail::target_arch::gfx1153
-    };
-
-    if (disabled_arches.find(arch) != disabled_arches.end())
-        GTEST_SKIP() << "Skipping test on gfx1151 Windows systems due to known issues.";
-#endif
-
     constexpr bool descending            = false;
     constexpr bool stable                = false;
     using decomposer_t                   = rocprim::identity_decomposer;

projects/rocprim/test/rocprim/test_device_topk.cpp

@@ -919,26 +919,6 @@ void topk_large_sizes_test(bool debug_synchronous)
 
 TEST(RocprimDeviceTopkTests, TopkLargeSizesStable)
 {
-    // Temporarily disable this test on gfx115x on Windows until we can determine the root cause.
-    // TODO: remove this after the root cause has been found and fixed properly.
-#if defined(_WIN32)
-    rocprim::detail::target_arch arch = rocprim::detail::target_arch::unknown;
-    if (rocprim::detail::host_target_arch(hipStreamDefault, arch) != HIP_SUCCESS)
-    {
-        std::cerr << "Warning: unable to fetch target architecture for disablement check." << std::endl;
-    }
-
-    const std::set<rocprim::detail::target_arch> disabled_arches = {
-        rocprim::detail::target_arch::gfx1150,
-        rocprim::detail::target_arch::gfx1151,
-        rocprim::detail::target_arch::gfx1152,
-        rocprim::detail::target_arch::gfx1153
-    };
-
-    if (disabled_arches.find(arch) != disabled_arches.end())
-        GTEST_SKIP() << "Skipping test on gfx1151 Windows systems due to known issues.";
-#endif
-
     constexpr bool descending            = false;
     constexpr bool stable                = true;
     using decomposer_t                   = rocprim::identity_decomposer;
@@ -963,26 +943,6 @@ TEST(RocprimDeviceTopkTests, TopkLargeSizesStable)
 
 TEST(RocprimDeviceTopkTests, TopkLargeSizesUnstable)
 {
-    // Temporarily disable this test on gfx115x on Windows until we can determine the root cause.
-    // TODO: remove this after the root cause has been found and fixed properly.
-#if defined(_WIN32)
-    rocprim::detail::target_arch arch = rocprim::detail::target_arch::unknown;
-    if (rocprim::detail::host_target_arch(hipStreamDefault, arch) != HIP_SUCCESS)
-    {
-        std::cerr << "Warning: unable to fetch target architecture for disablement check." << std::endl;
-    }
-
-    const std::set<rocprim::detail::target_arch> disabled_arches = {
-        rocprim::detail::target_arch::gfx1150,
-        rocprim::detail::target_arch::gfx1151,
-        rocprim::detail::target_arch::gfx1152,
-        rocprim::detail::target_arch::gfx1153
-    };
-
-    if (disabled_arches.find(arch) != disabled_arches.end())
-        GTEST_SKIP() << "Skipping test on gfx1151 Windows systems due to known issues.";
-#endif
-
     constexpr bool descending            = false;
     constexpr bool stable                = false;
     using decomposer_t                   = rocprim::identity_decomposer;

projects/rocthrust/.clang-format

@@ -30,7 +30,7 @@ AttributeMacros: [
                   'THRUST_HOST',
                   'THRUST_NODISCARD',
                   'THRUST_DEPRECATED',
-                  'THRUST_GLOBAL_CONSTANT',
+                  'THRUST_INLINE_CONSTANT',
                   # thrust
                   '_CCCL_ALIGNAS_TYPE',
                   '_CCCL_ALIGNAS',

projects/rocthrust/CHANGELOG.md

@@ -3,15 +3,6 @@
 Documentation for rocThrust available at
 [https://rocm.docs.amd.com/projects/rocThrust/en/latest/](https://rocm.docs.amd.com/projects/rocThrust/en/latest/).
 
-## rocThrust 5.0.0 for ROCm 10.0
-
-### Added
-
-* Largely in feature parity with CCCL/thrust v3.0.3.
-  - `thrust::tuple`, `thrust::pair` and `thrust::zip_iterator` fall back to rocThrust 4.4.0 implementations when a libhipcxx counterpart corresponding to CCCL/libcudacxx >= v3.0.3 is unavailable, ie
-    * `thrust::tuple` and `thrust::pair`: some features may differ from CCCL/thrust v3.0.3.
-    * `thrust::zip_iterator`: some iterator concepts present in CCCL/thrust v3.0.3 are missing.
-
 ## rocThrust 4.5.0 for ROCm 7.14
 
 ### Added

projects/rocthrust/CMakeLists.txt

@@ -8,11 +8,11 @@ cmake_minimum_required(VERSION 3.18 FATAL_ERROR)
 # Update these variables at release time
 #
 # Set the library version
-set(VERSION_STRING "5.0.0")
+set(VERSION_STRING "4.5.0")
 # Set the minimum required rocPRIM version
-set(MIN_ROCPRIM_PACKAGE_VERSION "4.1.0" CACHE STRING "Minimum version of rocPRIM to search for when ROCPRIM_FETCH_METHOD is set to PACKAGE.")
+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.1.0" CACHE STRING "Minimum version of rocRAND to search for when ROCRAND_FETCH_METHOD is set to PACKAGE.")
+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/therock-7.14" CACHE STRING "Download branch for ROCm dependencies")
 # --------------------------------------
@@ -35,9 +35,10 @@ endif()
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
 set(CMAKE_CXX_EXTENSIONS OFF)
 
-set(ALLOWED_CXX_STANDARDS 17 20 23 26)
-if(NOT CMAKE_CXX_STANDARD IN_LIST ALLOWED_CXX_STANDARDS)
-  message(FATAL_ERROR "Only C++17, C++20, C++23 and C++26 are supported")
+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
@@ -201,6 +202,10 @@ if(NOT DISABLE_WERROR)
   list(APPEND COMPILE_OPTIONS -Werror)
 endif()
 
+if(NOT CMAKE_CXX_STANDARD EQUAL 17)
+  message(FATAL_ERROR "Only C++17 is supported")
+endif()
+
 if (WIN32)
   add_compile_options(-xhip)
   add_compile_definitions(THRUST_IGNORE_DEPRECATED_CPP_DIALECT)

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

@@ -1,6 +1,6 @@
 /******************************************************************************
  * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  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:
@@ -110,9 +110,10 @@ void run_benchmark(benchmark::State& state, const std::size_t elements, const st
   }
 
 template <class Benchmark>
-void add_benchmarks(const std::string& name, std::vector<benchmark::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::Benchmark*> bs;
+  std::vector<benchmark::internal::Benchmark*> bs;
   BENCHMARK_TYPE(int8_t)
   BENCHMARK_TYPE(int16_t)
   BENCHMARK_TYPE(int32_t)
@@ -139,7 +140,7 @@ int main(int argc, char* argv[])
   benchmark::AddCustomContext("seed", seed_type);
 
   // Add benchmark
-  std::vector<benchmark::Benchmark*> benchmarks;
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
   add_benchmarks<basic>("basic", benchmarks, seed_type);
 
   // Use manual timing

projects/rocthrust/benchmark/bench/adjacent_difference/custom.cu

@@ -1,6 +1,6 @@
 /******************************************************************************
  * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  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:
@@ -128,9 +128,10 @@ void run_benchmark(benchmark::State& state, const std::size_t elements, const st
   }
 
 template <template <int> class Benchmark, int Val = 42 /*magic number in Thrust's benchmark*/>
-void add_benchmarks(const std::string& name, std::vector<benchmark::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::Benchmark*> bs;
+  std::vector<benchmark::internal::Benchmark*> bs;
   BENCHMARK_TYPE(int8_t)
   BENCHMARK_TYPE(int16_t)
   BENCHMARK_TYPE(int32_t)
@@ -158,7 +159,7 @@ int main(int argc, char* argv[])
   benchmark::AddCustomContext("seed", seed_type);
 
   // Add benchmark
-  std::vector<benchmark::Benchmark*> benchmarks;
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
   add_benchmarks<custom>("custom", benchmarks, seed_type);
 
   // Use manual timing

projects/rocthrust/benchmark/bench/adjacent_difference/in_place.cu

@@ -1,6 +1,6 @@
 /******************************************************************************
  * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  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:
@@ -107,9 +107,10 @@ void run_benchmark(benchmark::State& state, const std::size_t elements, const st
   }
 
 template <class Benchmark>
-void add_benchmarks(const std::string& name, std::vector<benchmark::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::Benchmark*> bs;
+  std::vector<benchmark::internal::Benchmark*> bs;
   BENCHMARK_TYPE(int8_t)
   BENCHMARK_TYPE(int16_t)
   BENCHMARK_TYPE(int32_t)
@@ -137,7 +138,7 @@ int main(int argc, char* argv[])
   benchmark::AddCustomContext("seed", seed_type);
 
   // Add benchmark
-  std::vector<benchmark::Benchmark*> benchmarks;
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
   add_benchmarks<in_place>("in_place", benchmarks, seed_type);
 
   // Use manual timing

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

@@ -1,6 +1,6 @@
 /******************************************************************************
  * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  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:
@@ -139,9 +139,10 @@ static_assert(!_THRUST_STD::is_trivially_copyable<non_trivial>::value, ""); // a
 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::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::Benchmark*> bs;
+  std::vector<benchmark::internal::Benchmark*> bs;
   BENCHMARK_TYPE(int8_t)
   BENCHMARK_TYPE(uint8_t)
   BENCHMARK_TYPE(int16_t)
@@ -174,7 +175,7 @@ int main(int argc, char* argv[])
   benchmark::AddCustomContext("seed", seed_type);
 
   // Add benchmark
-  std::vector<benchmark::Benchmark*> benchmarks;
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
   add_benchmarks<basic>("basic", benchmarks, seed_type);
 
   // Use manual timing

projects/rocthrust/benchmark/bench/copy/if.cu

@@ -1,6 +1,6 @@
 /******************************************************************************
  * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  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:
@@ -32,7 +32,6 @@
 // rocThrust
 #include <thrust/copy.h>
 #include <thrust/count.h>
-#include <thrust/detail/config/namespace.h>
 #include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
 
@@ -94,7 +93,7 @@ void run_benchmark(
   {
     input = bench_utils::generate(elements, seed_type);
   }
-  catch (const THRUST_NS_QUALIFIER::system::detail::bad_alloc& e)
+  catch (const ::thrust::system::detail::bad_alloc& e)
   {
     (void) hipGetLastError();
     state.SkipWithError(("thrust::system::detail::bad_alloc: " + std::string(e.what())).c_str());
@@ -106,7 +105,7 @@ void run_benchmark(
   {
     output = thrust::device_vector<T>(selected_elements);
   }
-  catch (const THRUST_NS_QUALIFIER::system::detail::bad_alloc& e)
+  catch (const ::thrust::system::detail::bad_alloc& e)
   {
     (void) hipGetLastError();
     state.SkipWithError(("thrust::system::detail::bad_alloc: " + std::string(e.what())).c_str());
@@ -154,10 +153,11 @@ void run_benchmark(
   }
 
 template <class Benchmark>
-void add_benchmarks(const std::string& name, std::vector<benchmark::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, 4200}; // 1.000, 0.544, 0.000;
-  std::vector<benchmark::Benchmark*> bs;
+  std::vector<benchmark::internal::Benchmark*> bs;
 
   BENCHMARK_TYPE_ENTROPY(int8_t, entropy_reductions)
   BENCHMARK_TYPE_ENTROPY(int16_t, entropy_reductions)
@@ -187,7 +187,7 @@ int main(int argc, char* argv[])
   benchmark::AddCustomContext("seed", seed_type);
 
   // Add benchmark
-  std::vector<benchmark::Benchmark*> benchmarks;
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
   add_benchmarks<_if>("if", benchmarks, seed_type);
 
   // Use manual timing

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

@@ -1,5 +1,5 @@
 // SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION. All rights reserved.
-// SPDX-FileCopyrightText: Modifications Copyright (c) 2025-2026 Advanced Micro Devices, Inc. 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
@@ -100,9 +100,10 @@ void run_benchmark(benchmark::State& state,
   }
 
 template <class Benchmark>
-void add_benchmarks(const std::string& name, std::vector<benchmark::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::Benchmark*> bs;
+  std::vector<benchmark::internal::Benchmark*> bs;
   BENCHMARK_TYPE(int8_t)
   BENCHMARK_TYPE(int16_t)
   BENCHMARK_TYPE(int32_t)
@@ -130,7 +131,7 @@ int main(int argc, char* argv[])
   benchmark::AddCustomContext("seed", seed_type);
 
   // Add benchmark
-  std::vector<benchmark::Benchmark*> benchmarks;
+  std::vector<benchmark::internal::Benchmark*> benchmarks;
   add_benchmarks<basic>("basic", benchmarks, seed_type);
 
   // Use manual timing