rocPRIM 3.3.0 for ROCm 6.3.1

rocPRIM code for ROCm 6.3.1 did not change. The library was rebuilt for the updated ROCm 6.3.1 stack.

rocPRIM 3.3.0 for ROCm 6.3.0

Added

• • Changed the default value of rmake.py -a to default_gpus. This is equivalent to gfx906:xnack-,gfx1030,gfx1100,gfx1101,gfx1102,gfx1151,gfx1200,gfx1201.
• The --test smoke option has been added to rtest.py. When rtest.py is called with this option it runs a subset of tests such that the total test time is 5 minutes. Use python3 ./rtest.py --test smoke or python3 ./rtest.py -t smoke to run the smoke test.
• The --seed option has been added to run_benchmarks.py. The --seed option specifies a seed for the generation of random inputs. When the option is omitted, the default behavior is to use a random seed for each benchmark measurement.
• Added configuration autotuning to device partition (rocprim::partition, rocprim::partition_two_way, and rocprim::partition_three_way), to device select (rocprim::select, rocprim::unique, and rocprim::unique_by_key), and to device reduce by key (rocprim::reduce_by_key) to improve performance on selected architectures.
• Added rocprim::uninitialized_array to provide 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 to std::nth_element. nth_element places elements that are smaller than the nth element before the nth element, and elements that are bigger than the nth element after the nth 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, and rocprim::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 and partial_sort_copy device functions similar to std::partial_sort and std::partial_sort_copy. partial_sort and partial_sort_copy arrange elements such that the elements are in the same order as a sorted list up to and including the middle index.

Changed

• 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.

Removed

• rocprim::thread_load() and rocprim::thread_store() have been deprecated. Use dereference() instead.

Resolved issues

• Fixed an issue in rmake.py where the list storing cmake options would contain individual characters instead of a full string of options.
• Resolved an issue in rtest.py where it crashed if the build folder was created without release or debug subdirectories.
• Resolved an issue with rtest.py on Windows where passing an absolute path to --install_dir caused a FileNotFound error.
• rocPRIM functions are no longer forcefully inlined on Windows. This significantly reduces the build
  time of debug builds.
• block_load, block_store, block_shuffle, block_exchange, and warp_exchange now use placement new 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 and int8_t) and instantiations of custom_type. Specifically, device binary search, device histogram, device merge and warp sort are affected.
• Fixed a bug for rocprim::merge_path_search where using unsigned offsets would produce incorrect results.
• Fixed a bug for rocprim::thread_load and rocprim::thread_store where float and double were not cast to the correct type, resulting in incorrect results.
• Resolved an issue where tests where failing when they were compiled with -D_GLIBCXX_ASSERTIONS=ON.
• Resolved an issue where algorithms that used an internal serial merge routine caused a memory access fault that resulted in potential performance drops when using block sort, device merge sort (block merge), device merge, device partial sort, and device sort (merge sort).
• Fixed memory leaks in unit tests due to missing calls to hipFree() and the incorrect use of hipGraphs.
• Fixed an issue where certain inputs to block_sort_merge(), device_merge_sort_merge_path(), device_merge(), and warp_sort_stable() caused an assertion error during the call to serial_merge().

rocPRIM 3.2.2 for ROCm 6.2.4

Additions

• GFX1151 Support

rocPRIM 3.2.1 for ROCm 6.2.2

rocPRIM code for ROCm 6.2.2 did not change. The library was rebuilt for the updated ROCm 6.2.2 stack.

rocPRIM 3.2.1 for ROCm 6.2.1

Optimizations

• Improved performance of block_reduce_warp_reduce when warp size == block size.

rocPRIM 3.2.0 for ROCm 6.2.0

Additions

• 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) and exclusive_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 to rocprim::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 of rocprim::block_radix_sort and all functions of device_radix_sort.
  To sort keys of an user-defined type, a decomposer functor should be passed. The decomposer should produce a rocprim::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 is true.
  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.

Optimizations

• Improved the performance of warp_sort_shuffle and block_sort_bitonic.
• Created an optimized version of the warp_exchange functions blocked_to_striped_shuffle and striped_to_blocked_shuffle when the warpsize is equal to the items per thread.

Fixes

• Fixed incorrect results of warp_exchange::blocked_to_striped_shuffle and warp_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 overlapping values_input and values_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) and rocprim::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.

Deprecations

• The internal header detail/match_result_type.hpp has been deprecated.
• TwiddleIn and TwiddleOut have been deprecated in favor of radix_key_codec.
• The internal ::rocprim::detail::radix_key_codec has been deprecated in favor of the new public utility with the same name.

rocPRIM 3.1.0 for ROCm 6.1.2

rocPRIM code for ROCm 6.1.2 did not change. The library was rebuilt for the updated ROCm 6.1.2 stack.

rocPRIM 3.1.0 for ROCm 6.1.1

rocPRIM code for ROCm 6.1.1 did not change. The library was rebuilt for the updated ROCm 6.1.1 stack.

rocPRIM 3.1.0 for ROCm 6.1.0

Additions

• New primitive: block_run_length_decode
• New primitive: batch_memcpy

Changes

• Renamed:
  • scan_config_v2 to scan_config
  • scan_by_key_config_v2 to scan_by_key_config
  • radix_sort_config_v2 to radix_sort_config
  • reduce_by_key_config_v2 to reduce_by_key_config
  • radix_sort_config_v2 to radix_sort_config
• Removed support for custom config types for device algorithms
• host_warp_size() was moved into rocprim/device/config_types.hpp; it now uses either device_id or
  a stream parameter to query the proper device and a device_id out parameter
  • The return type is hipError_t
• Added support for __int128_t in device_radix_sort and block_radix_sort
• Improved the performance of match_any, and block_histogram which uses it

Deprecations

• Removed reduce_by_key_config, MatchAny, scan_config, scan_by_key_config, and
  radix_sort_config

Fixes

• Build issues with rmake.py on Windows when using VS 2017 15.8 or later (due to a breaking fix with
  extended aligned storage)

rocPRIM 3.0.0 for ROCm 6.0.2

rocPRIM code for ROCm 6.0.2 did not change. The library was rebuilt for the updated ROCm 6.0.2 stack.