Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion projects/hiprand/CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@
Documentation for hipRAND is available at
[https://rocm.docs.amd.com/projects/hipRAND/en/latest/](https://rocm.docs.amd.com/projects/hipRAND/en/latest/).

## hipRAND 2.13.0 for ROCm 7.0
## hipRAND 3.0.0 for ROCm 7.0

### Added

Expand Down
2 changes: 1 addition & 1 deletion projects/hiprand/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -41,7 +41,7 @@ endif()
# hipRAND project
#
project(hipRAND CXX)
set(hipRAND_VERSION "2.13.0")
set(hipRAND_VERSION "3.0.0")

# Build options
option(BUILD_ADDRESS_SANITIZER "Build with address sanitizer enabled" OFF)
Expand Down
142 changes: 55 additions & 87 deletions projects/rocprim/CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -15,47 +15,45 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
* Added new optimization to the backend for `device_transform` when the input and output are pointers.
* Added `LoadType` to `transform_config`, which is used for the `device_transform` when the input and output are pointers.
* Added `rocprim:device_transform` for n-ary transform operations API with as input `n` number of iterators inside a `rocprim::tuple`.
* Added gfx950 support.
* Added `rocprim::key_value_pair::operator==`.
* Added the `rocprim::unrolled_copy` thread function to copy multiple items inside a thread.
* Added the `rocprim::unrolled_thread_load` function to load multiple items inside a thread using `rocprim::thread_load`.
* Added `rocprim::int128_t` and `rocprim::uint128_t` to benchmarks for improved performance evaluation on 128-bit integers.
* Added `rocprim::int128_t` to the supported autotuning types to improve performance for 128-bit integers.
* Added the `rocprim::merge_inplace` function for merging in-place.
* Added initial value support for warp- and block-level inclusive scan.
* Added support for building tests with device-side random data generation, making them finish faster. This requires rocRAND, and is enabled with the `WITH_ROCRAND=ON` build flag.
* Added tests and documentation to `lookback_scan_state`. It is still in the `detail` namespace.

### Changed

* The default scan accumulator types for device-level scan algorithms have changed. This is a breaking change.
The previous default accumulator types could lead to situations in which unexpected overflow occured, such as
when the input or inital type was smaller than the output type.

This is a complete list of affected functions and how their default accumulator types are changing:
* `rocprim::inclusive_scan`
* past default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* new default: `class AccType = rocprim::invoke_result_binary_op_t<typename std::iterator_traits<InputIterator>::value_type, BinaryFunction>`
* `rocprim::deterministic_inclusive_scan`
* past default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* new default: `class AccType = rocprim::invoke_result_binary_op_t<typename std::iterator_traits<InputIterator>::value_type, BinaryFunction>`
* `rocprim::exclusive_scan`
* past default: `class AccType = detail::input_type_t<InitValueType>>`
* new default: `class AccType = rocprim::invoke_result_binary_op_t<rocprim::detail::input_type_t<InitValueType>, BinaryFunction>`
* `rocprim::deterministic_exclusive_scan`
* past default: `class AccType = detail::input_type_t<InitValueType>>`
* new default: `class AccType = rocprim::invoke_result_binary_op_t<rocprim::detail::input_type_t<InitValueType>, BinaryFunction>`

* Changed the parameters `long_radix_bits` and `LongRadixBits` from `segmented_radix_sort` to `radix_bits` and `RadixBits` respectively.
* Marked the initialisation constructor of `rocprim::reverse_iterator<Iter>` `explicit`, use `rocprim::make_reverse_iterator`.
* Merged `radix_key_codec` into type_traits system.
* Renamed `type_traits_interface.hpp` to `type_traits.hpp`, rename the original `type_traits.hpp` to `type_traits_functions.hpp`.
* Changed the default accumulator type for various device-level scan algorithms:
* `rocprim::inclusive_scan`
* Previous default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, typename std::iterator_traits<InputIterator>::value_type>`
* `rocprim::deterministic_inclusive_scan`
* Previous default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, typename std::iterator_traits<InputIterator>::value_type>`
* `rocprim::exclusive_scan`
* Previous default: `class AccType = detail::input_type_t<InitValueType>>`
* Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, rocprim::detail::input_type_t<InitValueType>>`
* `rocprim::deterministic_exclusive_scan`
* Previous default: `class AccType = detail::input_type_t<InitValueType>>`
* Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, rocprim::detail::input_type_t<InitValueType>>`
* The default scan accumulator types for device-level scan algorithms have changed. This is a breaking change.
The previous default accumulator types could lead to situations in which unexpected overflow occured, such as
when the input or inital type was smaller than the output type.
* This is a complete list of affected functions and how their default accumulator types are changing:
* `rocprim::inclusive_scan`
* Previous default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, typename std::iterator_traits<InputIterator>::value_type>`
* `rocprim::deterministic_inclusive_scan`
* Previous default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, typename std::iterator_traits<InputIterator>::value_type>`
* `rocprim::exclusive_scan`
* Previous default: `class AccType = detail::input_type_t<InitValueType>>`
* Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, rocprim::detail::input_type_t<InitValueType>>`
* `rocprim::deterministic_exclusive_scan`
* Previous default: `class AccType = detail::input_type_t<InitValueType>>`
* Current default: `class AccType = rocprim::accumulator_t<BinaryFunction, rocprim::detail::input_type_t<InitValueType>>`
* Undeprecated internal `detail::raw_storage`.
* A new version of `rocprim::thread_load` and `rocprim::thread_store` replace the deprecated `rocprim::thread_load` and `rocprim::thread_store` functions. The versions avoid inline assembly where possible, and don't hinder the optimizer as much as a result.
* Renamed `rocprim::load_cs` to `rocprim::load_nontemporal` and `rocprim::store_cs` to `rocprim::store_nontemporal` to express the intent of these load and store methods better.
* All kernels now have hidden symbol visibility. All symbols now have inline namespaces that include the library version, for example, `rocprim::ROCPRIM_300400_NS::symbol` instead of `rocPRIM::symbol`, letting the user link multiple libraries built with different versions of rocPRIM.

### Deprecations
### Upcoming changes

* `rocprim::invoke_result_binary_op` and `rocprim::invoke_result_binary_op_t` are deprecated. Use `rocprim::accumulator_t` now.

Expand Down Expand Up @@ -87,79 +85,49 @@ This is a complete list of affected functions and how their default accumulator
* Use `rocprim::arch::wavefront::min_size()` or `rocprim::arch::wavefront::max_size()` instead.
* `__AMDGCN_WAVEFRONT_SIZE`
* This was a fallback define for the compiler's removed symbol, having the same name.
* This release removes support for custom builds on gfx940 and gfx941.

### Resolved issues

* Fixed an issue where `device_batch_memcpy` reported benchmarking throughput being 2x lower than it was in reality.
* Fixed an issue where `device_segmented_reduce` reported autotuning throughput being 5x lower than it was in reality.

## rocPRIM 3.5.0 for ROCm 6.5.0

### Removed
* This release removes support for custom builds on gfx940 and gfx941.

### Added

* Added gfx950 support.
* Added `rocprim::key_value_pair::operator==`.
* Added the `rocprim::unrolled_copy` thread function to copy multiple items inside a thread.
* Added the `rocprim::unrolled_thread_load` function to load multiple items inside a thread using `rocprim::thread_load`.
* Added `rocprim::int128_t` and `rocprim::uint128_t` to benchmarks for improved performance evaluation on 128-bit integers.
* Added `rocprim::int128_t` to the supported autotuning types to improve performance for 128-bit integers.
* Added the `rocprim::merge_inplace` function for merging in-place.
* Added initial value support for warp- and block-level inclusive scan.
* Added support for building tests with device-side random data generation, making them finish faster. This requires rocRAND, and is enabled with the `WITH_ROCRAND=ON` build flag.
* Added tests and documentation to `lookback_scan_state`. It is still in the `detail` namespace.

### Changed

* A new version of `rocprim::thread_load` and `rocprim::thread_store` replace the deprecated `rocprim::thread_load` and `rocprim::thread_store` functions. The versions avoid inline assembly where possible, and don't hinder the optimizer as much as a result.
* Renamed `rocprim::load_cs` to `rocprim::load_nontemporal` and `rocprim::store_cs` to `rocprim::store_nontemporal` to express the intent of these load and store methods better.
* All kernels now have hidden symbol visibility. All symbols now have inline namespaces that include the library version, for example, `rocprim::ROCPRIM_300400_NS::symbol` instead of `rocPRIM::symbol`, letting the user link multiple libraries built with different versions of rocPRIM.

### Resolved issues

* Fixed device radix sort not returning the correct required temporary storage when a double buffer contains `nullptr`.
* Fixed constness of equality operators (`==` and `!=`) in `rocprim::key_value_pair`.

### Upcoming changes

* The next major release may change the template parameters of warp and block algorithms.

* The default scan accumulator types for device-level scan algorithms will be changed. This is a breaking change.

Previously, the default accumulator type was set to the input type for inclusive scans and to the initial value type for exclusive scans. These default types could cause unexpected overflow in situations where the input or initial type is smaller than the output type when the user doesn't explicitly set an accumulator type using the `AccType` template parameter.

The new default types will be set to the type that results when the input or initial value type is applied to the scan operator.

The following is the complete list of affected functions and how their default accumulator types are changing:
* `rocprim::inclusive_scan`
* current default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* future default: `class AccType = rocprim::invoke_result_binary_op_t<typename std::iterator_traits<InputIterator>::value_type, BinaryFunction>`
* `rocprim::deterministic_inclusive_scan`
* current default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* future default: `class AccType = rocprim::invoke_result_binary_op_t<typename std::iterator_traits<InputIterator>::value_type, BinaryFunction>`
* `rocprim::exclusive_scan`
* current default: `class AccType = detail::input_type_t<InitValueType>>`
* future default: `class AccType = rocprim::invoke_result_binary_op_t<rocprim::detail::input_type_t<InitValueType>, BinaryFunction>`
* `rocprim::deterministic_exclusive_scan`
* current default: `class AccType = detail::input_type_t<InitValueType>>`
* future default: `class AccType = rocprim::invoke_result_binary_op_t<rocprim::detail::input_type_t<InitValueType>, BinaryFunction>`
### Known issues
* When using `rocprim::deterministic_inclusive_scan_by_key` and `rocprim::deterministic_exclusive_scan_by_key` the intermediate values can change order on Navi3x
* However if a commutative scan operator is used then the final scan value (output array) will still always be consistent between runs

### Deprecations
## rocPRIM 3.4.1 for ROCm 6.4.2

* `rocprim::load_cs` and `rocprim::store_cs` are deprecated. Use `rocprim::load_nontemporal` and `rocprim::store_nontemporal` now.
### Upcoming changes
* Changes to the template parameters of warp and block algorithms will be made in an upcoming release.

* Due to an upcoming compiler change the following warp size-related symbols will be removed in the next major release and are thus marked as deprecated:
* `rocprim::device_warp_size()`
* For compile-time constants, this is replaced with `rocprim::arch::wavefront::min_size()` and `rocprim::arch::wavefront::max_size()`. Use this when allocating global or shared memory.
* For run-time constants, this is replaced with `rocprim::arch::wavefront::size().`
* `rocprim::warp_size()`
* `ROCPRIM_WAVEFRONT_SIZE`

* The default scan accumulator types for device-level scan algorithms will be changed in an upcoming release, resulting in a breaking change. Previously, the default accumulator type was set to the input type for the inclusive scans and to the initial value type for the exclusive scans. This could lead to unexpected overflow if the input or initial type was smaller than the output type when the accumulator type was't explicitly set using the `AccType` template parameter. The new default accumulator types will be set to the type that results when the input or initial value type is applied to the scan operator.

### Known issues
* When using `rocprim::deterministic_inclusive_scan_by_key` and `rocprim::deterministic_exclusive_scan_by_key` the intermediate values can change order on Navi3x
* However if a commutative scan operator is used then the final scan value (output array) will still always be consistent between runs
The following is the complete list of affected functions and how their default accumulator types are changing:

* `rocprim::inclusive_scan`
* current default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* future default: `class AccType = rocprim::invoke_result_binary_op_t<typename std::iterator_traits<InputIterator>::value_type, BinaryFunction>`
* `rocprim::deterministic_inclusive_scan`
* current default: `class AccType = typename std::iterator_traits<InputIterator>::value_type>`
* future default: `class AccType = rocprim::invoke_result_binary_op_t<typename std::iterator_traits<InputIterator>::value_type, BinaryFunction>`
* `rocprim::exclusive_scan`
* current default: `class AccType = detail::input_type_t<InitValueType>>`
* future default: `class AccType = rocprim::invoke_result_binary_op_t<rocprim::detail::input_type_t<InitValueType>, BinaryFunction>`
* `rocprim::deterministic_exclusive_scan`
* current default: `class AccType = detail::input_type_t<InitValueType>>`
* future default: `class AccType = rocprim::invoke_result_binary_op_t<rocprim::detail::input_type_t<InitValueType>, BinaryFunction>`

* `rocprim::load_cs` and `rocprim::store_cs` are deprecated and will be removed in an upcoming release. Alternatively, you can use `rocprim::load_nontemporal` and `rocprim::store_nontemporal` to load and store values in specific conditions (like bypassing the cache) for `rocprim::thread_load` and `rocprim::thread_store`.

## rocPRIM 3.4.0 for ROCm 6.4.0

Expand Down
41 changes: 24 additions & 17 deletions projects/rocrand/CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -3,48 +3,55 @@
Documentation for rocRAND is available at
[https://rocm.docs.amd.com/projects/rocRAND/en/latest/](https://rocm.docs.amd.com/projects/rocRAND/en/latest/)

## (Unreleased) rocRAND-4.x.x for ROCm 7.0.0

### Changed

* Changed return type for `rocrand_generate_poisson` for `SOBOL64` and `SCRAMBLED_SOBOL64` engines
* Changed unnecessarily large 64-bit data type of constants used for skipping in `MRG32K3A` to 32-bit data type

## rocRAND 3.4.0 for ROCm 6.5
## rocRAND 4.0.0 for ROCm 7.0

### Added

* gfx950 support
* Additional unit tests for `test_log_normal_distribution.cpp`
* Additional unit tests for `test_normal_distribution.cpp`
* Additional unit tests for `test_rocrand_mtgp32_prng.cpp`
* Additional unit tests for `test_rocrand_scrambled_sobol32_qrng.cpp`
* Additional unit tests for `test_rocrand_scrambled_sobol64_qrng.cpp`
* Additional unit tests for `test_rocrand_sobol32_qrng.cpp`
* Additional unit tests for `test_rocrand_sobol64_qrng.cpp`
* Additional unit tests for `test_rocrand_threefry2x32_20_prng.cpp`
* Additional unit tests for `test_rocrand_threefry2x64_20_prng.cpp`
* Additional unit tests for `test_rocrand_threefry4x32_20_prng.cpp`
* Additional unit tests for `test_rocrand_threefry4x64_20_prng.cpp`
* Additional unit tests for `test_uniform_distribution.cpp`
* New unit tests for `include/rocrand/rocrand_discrete.h` in `test_rocrand_discrete.cpp`
* New unit tests for `include/rocrand/rocrand_mrg31k3p.h` in `test_rocrand_mrg31k3p_prng.cpp`
* New unit tests for `include/rocrand/rocrand_mrg32k3a.h` in `test_rocrand_mrg32k3a_prng.cpp`
* New unit tests for `include/rocrand/rocrand_poisson.h` in `test_rocrand_poisson.cpp`

### Changed

* Changed the return type for `rocrand_generate_poisson` for the `SOBOL64` and `SCRAMBLED_SOBOL64` engines.
* Changed the unnecessarily large 64-bit data type for constants used for skipping in `MRG32K3A` to the 32-bit data type.
* Updated several `gfx942` auto tuning parameters.
* Modified error handling and expanded the error information for the case of double-deallocation of the (scrambled) sobol32 and sobol64 constants and direction vectors.

### Removed

* Removed inline assembly and the `ENABLE_INLINE_ASM` CMake option. Inline assembly was used to optimizate of multiplications in the Mrg32k3a and Philox 4x32-10 generators. It is no longer needed because the current HIP compiler is able to produce code with the same or better performance.
* Removed instances of the deprecated clang definition `__AMDGCN_WAVEFRONT_SIZE`.

### Upcoming changes

* Deprecated the rocRAND Fortran API in favor of hipfort.
* Deprecated C++14 and set the default target to C++17. C++14 will be removed in the next major release.
* Directly accessing the (scrambled) sobol32 and sobol64 constants and direction vectors is deprecated and will be removed in the next major release. For:
* Removed C++14 support. Beginning with this release, only C++17 is supported.
* Directly accessing the (scrambled) sobol32 and sobol64 constants and direction vectors is no longer supported. For:
* `h_scrambled_sobol32_constants`, use `rocrand_get_scramble_constants32` instead.
* `h_scrambled_sobol64_constants`, use `rocrand_get_scramble_constants64` instead.
* `rocrand_h_sobol32_direction_vectors`, use `rocrand_get_direction_vectors32` instead.
* `rocrand_h_sobol64_direction_vectors`, use `rocrand_get_direction_vectors64` instead.
* `rocrand_h_scrambled_sobol32_direction_vectors`, use `rocrand_get_direction_vectors32` instead.
* `rocrand_h_scrambled_sobol64_direction_vectors`, use `rocrand_get_direction_vectors64` instead.

### Fixed
### Resolved issues

* Fixed an issue where `mt19937.hpp` would cause kernel errors during auto tuning.

### Removed
### Upcoming changes

* Removed C++14 support, only C++17 is supported.
* Deprecated the rocRAND Fortran API in favor of hipfort.

## rocRAND 3.3.0 for ROCm 6.4

Expand Down
2 changes: 1 addition & 1 deletion projects/rocrand/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -191,7 +191,7 @@ if(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY AND NOT WIN32)
endif()

# Set version variables
rocm_setup_version( VERSION "3.4.0" )
rocm_setup_version( VERSION "4.0.0" )
set ( rocrand_VERSION ${rocRAND_VERSION} )
# Old-style version number used within the library's API. rocrand_get_version should be modified.
math(EXPR rocrand_VERSION_NUMBER "${rocRAND_VERSION_MAJOR} * 100000 + ${rocRAND_VERSION_MINOR} * 100 + ${rocRAND_VERSION_PATCH}")
Expand Down
Loading