umfranzw · Jul 26, 2026 · Oct 2, 2025 · Oct 2, 2025 · Oct 2, 2025 · Jul 17, 2026
@@ -910,6 +910,26 @@ 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;
@@ -934,6 +954,26 @@ 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;

@@ -919,6 +919,26 @@ 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;
@@ -943,6 +963,26 @@ 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;

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

@@ -3,6 +3,15 @@
 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

@@ -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 "4.5.0")
+set(VERSION_STRING "5.0.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(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 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(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 download branch for dependencies rocPRIM and rocRAND
 set(ROCM_DEP_RELEASE_BRANCH "release/therock-7.14" CACHE STRING "Download branch for ROCm dependencies")
 # --------------------------------------
@@ -35,10 +35,9 @@ 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")
+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")
 endif()
 
 # Set HIP flags
@@ -202,10 +201,6 @@ 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)

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

@@ -1,6 +1,6 @@
 /******************************************************************************
  * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ * Modifications Copyright (c) 2024-2026, 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,10 +128,9 @@ 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::internal::Benchmark*>& benchmarks, const std::string seed_type)
+void add_benchmarks(const std::string& name, std::vector<benchmark::Benchmark*>& benchmarks, const std::string seed_type)
 {
-  std::vector<benchmark::internal::Benchmark*> bs;
+  std::vector<benchmark::Benchmark*> bs;
   BENCHMARK_TYPE(int8_t)
   BENCHMARK_TYPE(int16_t)
   BENCHMARK_TYPE(int32_t)
@@ -159,7 +158,7 @@ int main(int argc, char* argv[])
   benchmark::AddCustomContext("seed", seed_type);
 
   // Add benchmark
-  std::vector<benchmark::internal::Benchmark*> benchmarks;
+  std::vector<benchmark::Benchmark*> benchmarks;
   add_benchmarks<custom>("custom", benchmarks, seed_type);
 
   // Use manual timing

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

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

@@ -1,6 +1,6 @@
 /******************************************************************************
  * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc.  All rights reserved.
+ * Modifications Copyright (c) 2024-2026, 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,6 +32,7 @@
 // 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>
 
@@ -93,7 +94,7 @@ void run_benchmark(
   {
     input = bench_utils::generate(elements, seed_type);
   }
-  catch (const ::thrust::system::detail::bad_alloc& e)
+  catch (const THRUST_NS_QUALIFIER::system::detail::bad_alloc& e)
   {
     (void) hipGetLastError();
     state.SkipWithError(("thrust::system::detail::bad_alloc: " + std::string(e.what())).c_str());
@@ -105,7 +106,7 @@ void run_benchmark(
   {
     output = thrust::device_vector<T>(selected_elements);
   }
-  catch (const ::thrust::system::detail::bad_alloc& e)
+  catch (const THRUST_NS_QUALIFIER::system::detail::bad_alloc& e)
   {
     (void) hipGetLastError();
     state.SkipWithError(("thrust::system::detail::bad_alloc: " + std::string(e.what())).c_str());
@@ -153,11 +154,10 @@ void run_benchmark(
   }
 
 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::Benchmark*>& benchmarks, const std::string seed_type)
 {
   constexpr int entropy_reductions[] = {0, 2, 4200}; // 1.000, 0.544, 0.000;
-  std::vector<benchmark::internal::Benchmark*> bs;
+  std::vector<benchmark::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::internal::Benchmark*> benchmarks;
+  std::vector<benchmark::Benchmark*> benchmarks;
   add_benchmarks<_if>("if", benchmarks, seed_type);
 
   // Use manual timing

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