Documentation for rocPRIM is available at https://rocm.docs.amd.com/projects/rocPRIM/en/latest/.
- Add --test smoke option in rtest.py. It will run a subset of tests such that the total test time is in 5 minutes. Use python3 ./rtest.py --test smoke or python3 ./rtest.py -t smoke to execute smoke test.
- Option
--seed
to benchmarks to specify a seed for the generation of random inputs. The default behavior is to keep using a random seed per benchmark measurement. - Added configuration autotuning to device partition (
rocprim::partition
,rocprim::partition_two_way
, androcprim::partition_three_way
), device select (rocprim::select
,rocprim::unique
, androcprim::unique_by_key
), and device reduce by key (rocprim::reduce_by_key
) for improved performance on selected architectures. - Added
rocprim::uninitialized_array
which provides uninitialized storage in local memory for user-defined types. - Added large segment support for
rocprim:segmented_reduce
. - Added a parallel
nth_element
device function similar tostd::nth_element
, this function rearranges elements smaller than the n-th before and bigger than the n-th after the n-th element. - Added deterministic (bitwise reproducible) algorithm variants
rocprim::deterministic_inclusive_scan
,rocprim::deterministic_exclusive_scan
,rocprim::deterministic_inclusive_scan_by_key
,rocprim::deterministic_exclusive_scan_by_key
, androcprim::deterministic_reduce_by_key
. These provide run-to-run stable results with non-associative operators such as float operations, at the cost of reduced performance. - Added a parallel
partial_sort
andpartial_sort_copy
device function similar tostd::partial_sort
andstd::partial_sort_copy
, these functions rearranges elements such that the elements are the same as a sorted list up to and including the middle index.
- Modified the input size in device adjacent difference benchmarks. Observed performance with these benchmarks might be different.
- Changed the default seed for
device_benchmark_segmented_reduce
.
- Fixed an issue in rtest.py where if the build folder was made without release or debug directory it would crash the program
- Fixed an issue where while running rtest.py on windows and passing in an absolute path to
--install_dir
causes aFileNotFound
error. - rocPRIM functions are no longer forcefully inlined on Windows, significantly reducing the build time in debug builds.
block_load
,block_store
,block_shuffle
,block_exchange
andwarp_exchange
now use placementnew
instead of copy assignment (operator=
) when writing to local memory. This fixes the behavior of custom types with non-trivial copy assignments.- Fixed a bug in the generation of input data for benchmarks, which caused incorrect performance to be reported in specific cases. It may affect the reported performance for one-byte types (
uint8_t
andint8_t
) and instantiations ofcustom_type
. Specifically, device binary search, device histogram, device merge and warp sort are affected. - Fixed a bug for
rocprim::merge_path_search
where usingunsigned
offsets would output wrong results. - Fixed a bug for
rocprim::thread_load
androcprim::thread_store
wherefloat
anddouble
were not casted to the correct type resulting in wrong results. - Fix tests failing when compiling with
-D_GLIBCXX_ASSERTIONS=ON
.
rocprim::thread_load
androcprim::thread_store
, use dereference instead. Not all of those functions are available on every device architecture, and their usage can hurt performance, because inline assembly inhibits optimizations.
- Improved performance of block_reduce_warp_reduce when warp size == block size.
- New overloads for
warp_scan::exclusive_scan
that take no initial value. These new overloads will write an unspecified result to the first value of each warp. - The internal accumulator type of
inclusive_scan(_by_key)
andexclusive_scan(_by_key)
is now exposed as an optional type parameter.- The default accumulator type is still the value type of the input iterator (inclusive scan) or the initial value's type (exclusive scan). This is the same behaviour as before this change.
- New overload for
device_adjacent_difference_inplace
that allows separate input and output iterators, but allows them to point to the same element. - New public API for deriving resulting type on device-only functions:
rocprim::invoke_result
rocprim::invoke_result_t
rocprim::invoke_result_binary_op
rocprim::invoke_result_binary_op_t
- New
rocprim::batch_copy
function added. Similar torocprim::batch_memcpy
, but copies by element, not with memcpy. - Added more test cases, to better cover supported data types.
- Updated some tests to work with supported data types.
- An optional
decomposer
argument for all member functions ofrocprim::block_radix_sort
and all functions ofdevice_radix_sort
. To sort keys of an user-defined type, a decomposer functor should be passed. The decomposer should produce arocprim::tuple
of references to arithmetic types from the key. - New
rocprim::predicate_iterator
which acts as a proxy for an underlying iterator based on a predicate. It iterates over proxies that holds the references to the underlying values, but only allow reading and writing if the predicate istrue
. It can be instantiated with:rocprim::make_predicate_iterator
rocprim::make_mask_iterator
- Added custom radix sizes as the last parameter for
block_radix_sort
. The default value is 4, it can be a number between 0 and 32. - New
rocprim::radix_key_codec
, which allows the encoding/decoding of keys for radix-based sorts. For user-defined key types, a decomposer functor should be passed.
- Improved the performance of
warp_sort_shuffle
andblock_sort_bitonic
. - Created an optimized version of the
warp_exchange
functionsblocked_to_striped_shuffle
andstriped_to_blocked_shuffle
when the warpsize is equal to the items per thread. - Improved the performance of
device_transform
.
- Fixed incorrect results of
warp_exchange::blocked_to_striped_shuffle
andwarp_exchange::striped_to_blocked_shuffle
when the block size is larger than the logical warp size. The test suite has been updated with such cases. - Fixed incorrect results returned when calling device
unique_by_key
with overlappingvalues_input
andvalues_output
. - Fixed incorrect output type used in
device_adjacent_difference
. - Hotfix for incorrect results on the GFX10 (Navi 10/RDNA1, Navi 20/RDNA2) ISA and GFX11 ISA (Navi 30 GPUs) on device scan algorithms
rocprim::inclusive_scan(_by_key)
androcprim::exclusive_scan(_by_key)
with large input types. device_adjacent_difference
now considers both the input and the output type for selecting the appropriate kernel launch config. Previously only the input type was considered, which could result in compilation errors due to excessive shared memory usage.- Fixed incorrect data being loaded with
rocprim::thread_load
when compiling with-O0
. - Fixed a compilation failure in the host compiler when instantiating various block and device algorithms with block sizes not divisible by 64.
- The internal header
detail/match_result_type.hpp
has been deprecated. TwiddleIn
andTwiddleOut
have been deprecated in favor ofradix_key_codec
.- The internal
::rocprim::detail::radix_key_codec
has been deprecated in favor of the new public utility with the same name.
- New primitive:
block_run_length_decode
- New primitive:
batch_memcpy
- Renamed:
scan_config_v2
toscan_config
scan_by_key_config_v2
toscan_by_key_config
radix_sort_config_v2
toradix_sort_config
reduce_by_key_config_v2
toreduce_by_key_config
radix_sort_config_v2
toradix_sort_config
- Removed support for custom config types for device algorithms
host_warp_size()
was moved intorocprim/device/config_types.hpp
; it now uses eitherdevice_id
or astream
parameter to query the proper device and adevice_id
out parameter- The return type is
hipError_t
- The return type is
- Added support for
__int128_t
indevice_radix_sort
andblock_radix_sort
- Improved the performance of
match_any
, andblock_histogram
which uses it
- Removed
reduce_by_key_config
,MatchAny
,scan_config
,scan_by_key_config
, andradix_sort_config
- Build issues with
rmake.py
on Windows when using VS 2017 15.8 or later (due to a breaking fix with extended aligned storage) - Fix tests for
block_histogram
,block_exchange
,device_histogram
anddevice_reduce_by_key
for various types
device_run_length_encode
,warp_exchange
andwarp_load
tests fail withrocprim::half
block_sort::sort()
overload for keys and values with a dynamic size, for all block sort algorithms. Additionally, allblock_sort::sort()
overloads with a dynamic size are now supported forblock_sort_algorithm::merge_sort
andblock_sort_algorithm::bitonic_sort
.- New two-way partition primitive
partition_two_way
which can write to two separate iterators.
- Improved the performance of
partition
.
- Fixed
rocprim::MatchAny
for devices with 64-bit warp size. The functionrocprim::MatchAny
is deprecated androcprim::match_any
is preferred instead.
- Deprecated configuration
radix_sort_config
for device-level radix sort as it no longer matches the algorithm's parameters. New configurationradix_sort_config_v2
is preferred instead. - Removed erroneous implementation of device-level
inclusive_scan
andexclusive_scan
. The prior default implementation using lookback-scan now is the only available implementation. - The benchmark metric indicating the bytes processed for
exclusive_scan_by_key
andinclusive_scan_by_key
has been changed to incorporate the key type. Furthermore, the benchmark log has been changed such that these algorithms are reported asscan
andscan_by_key
instead ofscan_exclusive
andscan_inclusive
. - Deprecated configurations
scan_config
andscan_by_key_config
for device-level scans, as they no longer match the algorithm's parameters. New configurationsscan_config_v2
andscan_by_key_config_v2
are preferred instead.
- Fixed build issue caused by missing header in
thread/thread_search.hpp
.
- New block level
radix_rank
primitive - New block level
radix_rank_match
primitive - Added a stable block sorting implementation, which can be used with
block_sort
by adding theblock_sort_algorithm::stable_merge_sort
algorithm
- Improved the performance of:
block_radix_sort
device_radix_sort
device_merge_sort
- Updated the
docs
directory structure to match the standard of rocm-docs-core
- Disabled GPU error messages relating to incorrect warp operation usage with Navi GPUs on
Windows (due to GPU
printf
performance issues on Windows) - When
ROCPRIM_DISABLE_LOOKBACK_SCAN
is set,device_scan
fails for input sizes larger thanscan_config::size_limit
, which defaults tostd::numeric_limits<unsigned int>::max()
device_partition
,device_unique
, anddevice_reduce_by_key
now support problem sizes larger than 2^32 items- Device algorithms now return
hipErrorInvalidValue
if the amount of passed temporary memory is insufficient - Lists of sizes for tests are unified, restored scan and reduce tests for
half
andbfloat16
values
block_sort::sort()
overload for keys and values with a dynamic size- This overload was documented but the implementation is missing; to avoid further confusion, the documentation is removed until a decision is made on implementing the function
- Fixed the compilation failure in
device_merge
when the two key iterators don't match
device_merge
doesn't correctly support different types forkeys_input1
andkeys_input2
(as of the 5.3.0 release)
- New functions
subtract_left
andsubtract_right
inblock_adjacent_difference
to apply functions on pairs of adjacent items distributed between threads in a block - New device-level
adjacent_difference
primitives - Experimental tooling for automatic kernel configuration tuning for various architectures
- Benchmarks collect and output more detailed system information
- CMake functionality improves build parallelism of the test suite that splits compilation units by function or by parameters
- Reverse iterator
- Support for problem sizes over
UINT_MAX
in device functionsinclusive_scan_by_key
andexclusive_scan_by_key
- Improved the performance of warp primitives using the swizzle operation on Navi
- Improved build parallelism of the test suite by splitting up large compilation units
device_select
now supports problem sizes larger than 2^32 itemsdevice_segmented_radix_sort
now partitions segments to group small, medium, and large segments- Each segment group can be sorted by specialized kernels to improve throughput
- Improved histogram performance for the case of highly uneven sample distribution
- Packages for tests and benchmark executables on all supported operating systems using CPack
- Added file and folder reorganization changes with backward compatibility support using wrapper headers
- Fixed Radix Sort
int64_t
bug introduced in version 2.10.11
- Future value
- Device
partition_three_way
to partition input to three output iterators based on two predicates
- The reduce/scan algorithm precision issues in the tests has been resolved for half types
- The device Radix Sort algorithm supports indexing with 64-bit unsigned integers
- The indexer type is chosen based on the type argument of parameter
size
- If
sizeof(size)
is not larger than 4 bytes, the indexer type is 32-bit unsigned int, otherwise, the indexer type is 64-bit unsigned int - The maximum problem size is based on the compile time configuration of the algorithm according to the following formula:
max_problem_size = (UINT_MAX + 1) * config::scan::block_size * config::scan::items_per_thread
- The indexer type is chosen based on the type argument of parameter
- Flags API of
block_adjacent_difference
device_segmented_radix_sort
unit test is failing for HIP on Windows
- Enable bfloat16 tests and reduce threshold for bfloat16
- Fix device scan
limit_size
feature - Non-optimized builds no longer trigger local memory limit errors
- Scan size limit feature
- Reduce size limit feature
- Transform size limit feature
block_load_striped
andblock_store_striped
gather_to_blocked
to gather values from other threads into a blocked arrangement- The block sizes for device merge sorts initial block sort and its merge steps are now separate in its
kernel config
- The block sort step supports multiple items per thread
- you can now set the
size_limit
for scan, reduce, and transform in the config struct instead of using a parameter device_scan
anddevice_segmented_scan
:inclusive_scan
now uses theinput-type
asaccumulator-type
;exclusive_scan
usesinitial-value-type
- This changes the behavior of small-size input types with large-size output types (e.g.,
short
input,int
output) and low-res input with high-res output (e.g.,float
input,double
output)
- This changes the behavior of small-size input types with large-size output types (e.g.,
- Revert an old Fiji workaround because they solved the issue at the compiler side
- Update README CMake minimum version number
- Added block sort support for multiple items per thread
- Currently only powers of two block sizes, and items per threads are supported and only for full blocks
- Bumped the minimum required version of CMake to 3.16
device_segmented_radix_sort
anddevice_scan
unit tests failing for HIP on WindowsReduceEmptyInput
causes random failure with bfloat16
- Initial HIP on Windows support
- bfloat16 support added
- Packaging has been split into a runtime package (
rocprim
) and a development package (rocprim-devel
): The development package depends on the runtime package. When installing the runtime package, the package manager will suggest the installation of the development package to aid users transitioning from the previous version's combined package. This suggestion by package manager is for all supported operating systems (except CentOS 7) to aid in the transition. Thesuggestion
feature in the runtime package is introduced as a deprecated feature and will be removed in a future ROCm release.- Because rocPRIM is a header-only library, the runtime package is an empty placeholder used to aid in the transition. This package is also a deprecated feature and will be removed in a future rocm release.
- Unit tests may soft hang on MI200 when running in
hipMallocManaged
mode
- Code coverage tools build option
- AddressSanitizer build option
- gfx1030 support added
- Experimental HIP-CPU support; build using GCC/Clang/MSVC on Windows and Linux (this is work in progress and many algorithms are known to fail)
- Added single tile Radix Sort for smaller sizes
- Improved performance for Radix Sort for larger element sizes
- Bug fix and minor performance improvement for
merge_sort
when input and output storage are the same
- gfx90a support added
warp_size()
function; usehost_warp_size()
anddevice_warp_size()
for host and device references, respectively
- Size zero inputs are now properly handled with newer ROCm builds that no longer allow zero-size kernel grid and block dimensions
- Minimum CMake version required is now 3.10.2
- Device scan unit test is currently failing due to an LLVM bug
- Texture cache iteration support has been re-enabled
- Benchmark builds have been re-enabled
- Unique operator is no longer called on invalid elements
- Device scan unit test is currently failing because of an LLVM bug
- No new features
- Updates to DPP instructions for warp shuffle
- Benchmark builds are disabled due to compiler bug
- HIP CMake dependency
- Updates to warp shuffle for gfx10
- Disabled DPP functions on gfx10++
- Benchmark builds are disabled due to compiler bug
- Fix for rocPRIM texture cache iterator
- Package dependency correct to
hip-rocclr
- rocPRIM texture cache iterator functionality is broken in the runtime (this will be fixed in the next release); you can use the prior release if calling this function
- No new features
- Point release with compilation fix
- Improved tests with fixed and random seeds for test data
- Network interface improvements with API v3
- Switched to HIP-Clang as the default compiler
- CMake searches for rocPRIM locally first; if t's not found, CMake downloads it from GitHub