diff --git a/projects/rocprim/test/rocprim/test_device_segmented_topk.cpp b/projects/rocprim/test/rocprim/test_device_segmented_topk.cpp index 774a3808100c..8ea85172595c 100644 --- a/projects/rocprim/test/rocprim/test_device_segmented_topk.cpp +++ b/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 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 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; diff --git a/projects/rocprim/test/rocprim/test_device_topk.cpp b/projects/rocprim/test/rocprim/test_device_topk.cpp index 3ec91b809a61..f885ff9f2291 100644 --- a/projects/rocprim/test/rocprim/test_device_topk.cpp +++ b/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 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 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; diff --git a/projects/rocthrust/.clang-format b/projects/rocthrust/.clang-format index 676f9822da5f..ba9b78224095 100644 --- a/projects/rocthrust/.clang-format +++ b/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', diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index 1e3caeb39ac1..f77dc955cca4 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/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 diff --git a/projects/rocthrust/CMakeLists.txt b/projects/rocthrust/CMakeLists.txt index 9bdb858bfd21..467bb97f7549 100644 --- a/projects/rocthrust/CMakeLists.txt +++ b/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) diff --git a/projects/rocthrust/benchmark/bench/adjacent_difference/basic.cu b/projects/rocthrust/benchmark/bench/adjacent_difference/basic.cu index c0fcecdf658f..33f8f023449b 100644 --- a/projects/rocthrust/benchmark/bench/adjacent_difference/basic.cu +++ b/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 -void add_benchmarks(const std::string& name, std::vector& benchmarks, const std::string seed_type) +void add_benchmarks( + const std::string& name, std::vector& benchmarks, const std::string seed_type) { - std::vector bs; + std::vector 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 benchmarks; + std::vector benchmarks; add_benchmarks("basic", benchmarks, seed_type); // Use manual timing diff --git a/projects/rocthrust/benchmark/bench/adjacent_difference/custom.cu b/projects/rocthrust/benchmark/bench/adjacent_difference/custom.cu index 8626c923135e..2be6dc85c80e 100644 --- a/projects/rocthrust/benchmark/bench/adjacent_difference/custom.cu +++ b/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