Skip to content

Commit

Permalink
6.1 bulk update from develop branch 2024-1-16 (#510)
Browse files Browse the repository at this point in the history
* StreamHPC 2023-10-18 (#480)

* Excessive shared memory usage in block_shuffle fix

* remove block_sort_algorithm template param from block_sort_kernel_impl and block_sort_impl

* fixed compile errors

* Updated ChangeLog.md

* remove unnecessary code

* fixed CHANGELOG.md to not be so verbose about non public api changes

* Add dynamic dispatch and autotuning to device_adjacent_difference

* Fix device_adjacent_difference storage type

* ci: remove autotune dependency from build:benchmark

The workaround needed to make this work is has major disadvantages,
and our current workflow does not make use of this dependency anyway
(Currently the generated configs are checked into the repository, so
the CI would run the benchmarks on them on the next push to the
merge-request).

When we improve automation around autotuning this could be implemented
with conditional jobs, but lets just drop the dependency for now.

* test: fix indexing error test_type_helper<custom_16aligned>::get_random_data

Indexing was 4 based when the type has 3 variables, therefore it was
overflowing. Caught with address sanitizer.

* fixes for compilation in debug for radix_sort

- Add force inline to onesweep kernel, to avoid too much shared memory
  errors
- Declare `block_radix_sort::radix_bits_per_pass` to fix linker errors

* fix: Detect DPP & DPP broadcast support with __GFX<GENERATION>__ macros

The amdgpu target in clang now provides the GFX generation as a
predefined macro, so we no longer need to explicitly list all targets,
which was bad for maintenance.

Also replace the use of the generic `ROCPRIM_NAVI` which signals navi support,
with `ROCPRIM_DETAIL_HAS_DPP_BROADCAST`, a macro that explicitly
states what we're after.

Also also makes sure that `ROCPRIM_DETAIL_USE_DPP` is always defined
(to 0 when DPP is disabled), previously it was undefined when
`ROCPRIM_DISABLE_DPP` was set.

* refactor: Use __GFX<GENERATION>__ to detect NAVI cards

* docs: Update CHANGELOG for DPP & ROCPRIM_NAVI fixes

* remove deprecated structs and functions

* rename scan_by_key_config_v2 to scan_by_key_config
remove the option to use custom implemented config for scan_by_key
update tests to not use custom implemented config for scan_by_key

* remove the option to use custom implemented config for histogram
update tests to not use custom implemented config for histogram

* update config compile time check to a different pattern

* update documentation comments for configs

* change documentation comments

* change documentation comments on device_radix_sort
rename radix_sort_config_v2 to radix_sort_config

* change documentation comment
add static_assert to check type for reduce_config

* update documentation comments
remove wrap_scan_config function
add static_assert to disallow custom scan_config type
rename scan_config_v2 to scan_config

* update documentation comments

* update documentation comments
make transform_config inherit from detail::transfomr_config_params
remove wrap_transform_config
add static assert to test for Config type in device_transform

* remove wrap_adjacent_difference_config function
add static_assert to test config type
create default ctor for adjacent_difference_config

* add missing transform_config ctor
rewrite adjacent_difference_config ctor to match other config structs

* fix binary search still using wrap_transform_config

* implement static_asset to make binary_search only use binary search configs, but also work with the underlying transform

* update changelog

* remove some *_v2s that went under the radar

* remove unnecessary default values

* Add binary search, lower_bound and upper_bound documentation

* host_warp_size() is replaced with two different versions with parameters.
the new versions use either a device id or a stream to figure out the warp size of the device

* comment out unused param names

* fix typos in the documentation

* move host_warp_size to config_type.hpp
changed host_warp_size signatures to fit other similar functions

* add error checks to host_warp_size calls in tests and benchmarks

* fix format

* add missing comment

* fix error handling in lookback_scan_state.hpp

* fix compilation error

* change block_radix_rank_match and block_histogram_atomic to use rocprim::match_any instead of implementing same functionality

* change radix_digit_count_helper to use rocprim::match_any instead of implementing same functionality
added predicate param to rocprim::match_any to set invalid lanes and added tests for this functionality

* add elect function to warp intrinsics
add test for elect
change block_histogram_atomic, block_radix_rank_match, device_histogram, device_radix_sort to use elect instead of copy-paste code

* update match_any to return 0 when predicate is false

* fix the bit check in elect function

* update changelog.md

* fix hard coded warps per block value to come from param in kernel

* remove unused variables

* fix review comments
minor name changes
update test
update comments

* update group_elect test
tests multiple groups per warp
doesn't check which exact thread is elected in a group, only that one is elected

* remove unnecessary comments

* remove expected from group_elect test
fix compile error

* fix overindexing

* fix review comments
update group_elect_test to have better coverage

* format

* fix review comments

* fix perf regression

* undo group_elect in block_histogram_atomic.hpp, because of perf impact

* fix bad func name in CHANGELOG.md

* fix merge errors

* Fix reduce_by_key algorithm so keys[0] is not flagged as a new run when is nan

* make device_radix_sort compatible with compiler provided __int128_t and __uint128_t

* add ifdefs to only compile int128 parts on clang/gcc

* update changelog

* fix for int128 to_string labdas

* add test for block_radix_sort int128 support

* Implement block run length decode

* Fix reduce_by_key algorithm so out of bounds items are not flagged as new runs for NaNs

* Add reduce_by_key test to check that flagging is correct when keys are all different

* Fix performance regression observed during tuning for gfx1030 and gfx1102

* Block Runlength Decode: Fix incorrect offsets and improve test

* Remove duplicate key from .clang-format

* Remove additional duplicates from clang-format

* Fix binary_search upper/lower_bound config tuning

Use specialized configurations for upper, lower, and binary search
algorithms when preforming tuning

* unify language around config params in documentation

* Make the autotune build job run nightly

* remove radix_sort_onesweep autotuning workaround

* Resolve doxygen warnings for upstream PR

* Enable get_device_from_stream for Windows

* Use _ENABLE_EXTENDED_ALIGNED_STORAGE for windows build in rmake.py

* Bump unreleased ROCm version

---------

Co-authored-by: Ivan Siutsou <ivan@streamhpc.com>
Co-authored-by: Bence Parajdi <bence@streamhpc.com>
Co-authored-by: Bálint Soproni <balint@streamhpc.com>
Co-authored-by: Gergely Meszaros <gergely@streamhpc.com>
Co-authored-by: Beatriz Navidad Vilches <beatriz@streamhpc.com>
Co-authored-by: Mátyás Aradi <matyas@streamhpc.com>

* StreamHPC 2023-11-17 (batch memcpy) (#485)

* Implemented batch memcpy algorithm and relevant tests and benchmarks

* Optimize match_any by using arithmetic shifts

The compiler seems to see through these much better than the conditional,
generating bit-field extract instructions, and recognizing that the loop
is a reduction.

* Pedantic / consistency changes for batch memcpy

* Improve interface and implementation of align_(up|down)

- Use the alignment of the destination type instead of its size
- Rename to emphasize that this does a form of reinterpret_cast
- Use the same type as the return type and template parameter, to
  match the interface of built-in casts
- Pedantic: use uintptr_t instead of size_t for the numerical value
  of a pointer
- Use clangs __builtin_align_(up|down) when available

* Take parameters as explicit const-ref in test_utils::bit_equal

Because these are templates this already works for non-copyable types,
(as `T` will be deduced to `Type&`) but its confusing, and wouldn't work
for r-values. Because we are comparing object representations taking a copy
isn't okay as that only guarantees that the value representation is copied.
(I.e. padding bytes are not required to be copied when taking a parameter
 by copy)

* Actually make custom_non(copyable|moveable)_type non (copy|move)-able

* Allow passing rocprim::default_config to batch_memcpy

As all the other device functions do too.

* Fix typo in cast_align_down documentation

* Fixup accidentally deleted constructor of custom_non_moveable_type

This was accidentally deleted, it was meant to be defaulted.
Currently no test calls this as batch-memcpy tests only use this type
at the device side.

* Improve error message of test_rocprim_package

The error message of the package test wasn't very nice, improve it
for easier debugging in the future.

Before:
```console
❯ ./a.out
98
```

After:
```console
❯ ./a.out
Error hipErrorInvalidDeviceFunction(98): invalid device function in main at test_rocprim_package.cpp:90
```

* Refactor test_utils::get_random_data into generate_random_data_n

- Writes the output into an output iterator instead of creating &
  returning a vector. This allows greater flexibility for users
  i.e. writing random values with differing options into the same
  container.
- Accepts a generator instead of a seed. This is more efficient, because
  creating an instance of an rng engine might be costly. It's also
  more consistent with how the standard library operates.
- The naming and interface tries to mirror the stl (i.e. `std::generate_n`)
- Backwards compatibility is maintained by adding test_utils::get_random_data
  that uses `generate_random_data_n` internally.

* Refactor get_random_data into generate_random_data_n in benchmark_utils

This mirrors the test changes in the previous commit

* Unify segmnented generation from test generate_random_data_n overloads

* Add missing include for iterator traits to benchmark_utils

* ci: use build instead rocm-build tag

This allows the build job to be performed by any runner configured
for building, instead of the ROCm-specialized builder. As the
target architectures are specified ahead of time, the GPU is not
needed during the build process, and may be performed by any builder.

* fix: Fixed doxygen warning in device_memcpy_config.hpp

* Speed up / Improve data-generation in test_device_batch_memcpy

Do bulk data-generation instead of individual calls, especially of
individual bytes for the data to copy.
Also changes the verification to do bulk memcmp instead of item-wise
test_utils::bit_equals for each buffer.
Overall this reduces the time it takes to run the test to ~1s from
around 3s.

* Refactor & Speedup benchmark_device_batch_memcpy

- Share the data generation between the naive and uut benchmarks
- Make the data-generation be bulk using a fast random number engine
  (mt19937) to significantly speed it up.

The overall runtime of the benchmark decreased from 14 minutes (!) to
around 2 minutes.

* Fix explanation comment in batch_memcpy test/benchmark

* fix include order in benchmark_device_batch_memcpy

* doc: add batch memcpy to changelog

---------

Co-authored-by: Gergely Meszaros <gergely@streamhpc.com>
Co-authored-by: Robin Voetter <robin@streamhpc.com>

* Add unit testing to verify that algorithms work with hipGraphs (#478)

* Basic hipGraph tests

* Add basic tests for graph creation, instantiation, and execution using:
  * stream capture
  * manual construction

* hipGraph test for device_reduce algorithms
* Added new unit tests for device_reduce, device_reduce_by_key algorithms
to verify basic support for hipGraphs (no synchronous API functions are
called within the algorithms).
* Fixed up CMakeLists compile issue for tests in the test/hipgraph folder
* Updated code documentation

* Add hipGraph unit tests for device level algorithms

* Added unit tests that run the following algorithms inside of a graph
(in isolation):
  - device_adjacent_difference
  - device_binary_search
  - device_histogram
  - device_merge
  - device_merge_sort
  - device_partition
  - device_radix_sort
  - device_scan
  - device_segmented_reduce
  - device_segmented_scan
  - device_select
  - device_transform

* Updated existing tests for:
  - device_reduce
  - device_reduce_by_key

* Moved graph test helper functions to a separate file

* Add hipGraph unit tests

* Added remaining device level hipGraph unit tests

* Note: currently, there are two device level algorithms that
do no work with hipGraphs because they contain synchronization
barriers. No hipGraph unit tests have been added for these
algorithms:
  * device_run_length_encode
  * device_segmented_radix_sort

* Added a functional integration test for hipGraphs, which
runs several algorithms back-to-back within a graph.

* Refactored test helper code to remove unnecessary parameter

* Set hipgraph test pointers to nullptr

* Set key_type device pointers to nullptr when they are declared, for
  safety.

* Several minor fixes for hipGraph tests
* Fixed up spelling error in comments
* Moved call to hipGetLastError to a more appropriate position
* Removed old commented test code

* Minor fixes for hipgraph unit tests
* Moved several synchronization barriers so they are now outside of graph capture blocks
  in the test_device_partition source
* Changed several loop counters to unsigned type
* Updatedpgraph  cmake files - removed test/hipgraph
  directory's CMakeLists.txt

* Additional test and bugfix for hipgraph tests
* Removed syncrhonization barrier in test_device_scan
* Added basic test to exercise atomic function within a hipgraph
* Rebased and resolved merge conflicts

* readme and changelog updates (#486)

* Skip device_adjacent_difference hipGraph test on Windows for Navi3x (#490)

* Currently, the LargeIndices hipGraphs test for gfx1030 on Windows is skipped
* This change causes this test case to also get skiped on gfx1100, gfx1101, gfx1102 on Windows
* The reason this test fails on Navi on Windows appears to be related to
  the check_output class (used by OutputIterator in the test).
  * this may be releated to using atomics inside of graphs, but further
    investigation is needed

* Bump cryptography from 41.0.4 to 41.0.6 in /docs/.sphinx (#488)

Bumps [cryptography](https://github.com/pyca/cryptography) from 41.0.4 to 41.0.6.
- [Changelog](https://github.com/pyca/cryptography/blob/main/CHANGELOG.rst)
- [Commits](pyca/cryptography@41.0.4...41.0.6)

---
updated-dependencies:
- dependency-name: cryptography
  dependency-type: indirect
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

* Bump rocm-docs-core from 0.27.0 to 0.30.0 in /docs/.sphinx (#489)

Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.27.0 to 0.30.0.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/RadeonOpenCompute/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](ROCm/rocm-docs-core@v0.27.0...v0.30.0)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

* Lookback state fixes (#491)

* Do not call fence in the wait loop

* Use __hip_atomic_load/store instead of atomicExch/atomicAdd

atomicExch is compiled to global_atomic_swap even when the results is not
used.

* Use faster fences in lookback algorithms on gfx94*

This version is specific for devices with slow __threadfence ("agent" fence which does
L2 cache flushing and invalidation).
Fences with "workgroup" scope are used instead to ensure ordering only but not coherence,
they do not flush and invalidate cache.
Global coherence of prefixes_*_values is ensured by atomic_load/atomic_store that bypass
cache.

* Rename ROCPRIM_DETAIL_LOOKBACK_SCAN_STATE_WITHOUT_SLOW_FENCES

from ROCPRIM_LOOKBACK_WITHOUT_SLOW_FENCES.
This is more verbose to communicates that it is implementation detail

It uses 0 and 1 instead of the presence of the macro now, and won't
be overriden if set by a developer on the command line.

* Add WITHOUT_SLOW_FENCES version to lookback_scan_state::get_complete_value

* refactor: lookback_scan_state WITHOUT_SLOW_FENCES misc changes

- use sizeof(variable)
- use auto* and const auto* instead of just auto
- use void* instead of char* to avoid yet another cast
- make the atomic order fence a separate function and add docs &
  warning

* fix: Restore removed interfaces of lookback_scan_state

Even though these are in the detail namespace and as such explicitly
not meant for usage by users, some projects did start depending on them.

The interfaces for these are slightly broken and rocPRIM developers
discourage any users from using them (or the newer interfaces for that
matter) because they are implementation details. No further guarantees
are provided for these APIs.

In the future a public interface is planned for lookback_scan_state
as we have recognized that this is a useful primitive, and it's
unreasonable to expect users to implement for themselves.

* refactor: rename __builtin_amdgcn_fence as atomic_fence_acquire_order_only

---------

Co-authored-by: Anton Gorenko <anton@streamhpc.com>

* Bump rocm-docs-core from 0.30.0 to 0.30.3 in /docs/.sphinx (#496)

Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.30.0 to 0.30.3.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/RadeonOpenCompute/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](ROCm/rocm-docs-core@v0.30.0...v0.30.3)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-type: direct:production
  update-type: version-update:semver-patch
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

* 6.0 final mergeback to develop (#498)

* Fix cpp-check reported issues

Fixed a number of issues that static the analysis tool picked up:
  - Made some functions const since they don't modify member state
  - Made some parameters const, since they're never modified
  - Made some functions static (for performance), since they don't require access to the class instance
  - Fixes for several benchmark/test functions
    - Removed unused variable declarations
    - Added missing input data transfer from host to device
    - Added default return value for one overlooked execution path
    - Added some member variables to constructor initializer list
    - Added override keyword in several places
    - Fixed up item placeholders in some printf statements

* Separate gfx942 specific code (#468)

Co-authored-by: Stanley Tsang <stanley.tsang@amd.com>

* Fix cpp-check reported issues
* Removed host to data transfer from memcpy benchmark.
Since this benchmark only tests memcpy performance between device buffers,
we don't really need to copy data into these from the host.

* Remove Unnecessary Newline & Re-trigger Performance Checks

* Update comment for extra clarification

* Updated comment in memcpy benchmark to make the purpose of the code a little clearer.

* Update googlebenchmark version (#477)

* 6.0 cherry pick for changelog and version update (#483)

* Fix changelog for 6.0

* Fix version

* Fix up changelog

---------

Co-authored-by: Wayne Franz <wayfranz@amd.com>
Co-authored-by: Eiden Yoshida <47196116+eidenyoshida@users.noreply.github.com>
Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com>

* Add CODEOWNERS file (#504)

* Standardize documentation for ReadtheDocs (#497)

* Bump jinja2 from 3.1.2 to 3.1.3 in /docs/sphinx (#506)

Bumps [jinja2](https://github.com/pallets/jinja) from 3.1.2 to 3.1.3.
- [Release notes](https://github.com/pallets/jinja/releases)
- [Changelog](https://github.com/pallets/jinja/blob/main/CHANGES.rst)
- [Commits](pallets/jinja@3.1.2...3.1.3)

---
updated-dependencies:
- dependency-name: jinja2
  dependency-type: indirect
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

* Bump gitpython from 3.1.37 to 3.1.41 in /docs/sphinx (#508)

Bumps [gitpython](https://github.com/gitpython-developers/GitPython) from 3.1.37 to 3.1.41.
- [Release notes](https://github.com/gitpython-developers/GitPython/releases)
- [Changelog](https://github.com/gitpython-developers/GitPython/blob/main/CHANGES)
- [Commits](gitpython-developers/GitPython@3.1.37...3.1.41)

---
updated-dependencies:
- dependency-name: gitpython
  dependency-type: indirect
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

* Bump rocm-docs-core from 0.30.3 to 0.31.0 in /docs/sphinx (#507)

Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.30.3 to 0.31.0.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/RadeonOpenCompute/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](ROCm/rocm-docs-core@v0.30.3...v0.31.0)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

* Update links in README.md

- Update links to other ROCm repositories.

* Update package version

---------

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: Nara <nara@streamhpc.com>
Co-authored-by: Ivan Siutsou <ivan@streamhpc.com>
Co-authored-by: Bence Parajdi <bence@streamhpc.com>
Co-authored-by: Bálint Soproni <balint@streamhpc.com>
Co-authored-by: Gergely Meszaros <gergely@streamhpc.com>
Co-authored-by: Beatriz Navidad Vilches <beatriz@streamhpc.com>
Co-authored-by: Mátyás Aradi <matyas@streamhpc.com>
Co-authored-by: Robin Voetter <robin@streamhpc.com>
Co-authored-by: Wayne Franz <wayfranz@amd.com>
Co-authored-by: Lisa <lisajdelaney@gmail.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
Co-authored-by: Anton Gorenko <anton@streamhpc.com>
Co-authored-by: Eiden Yoshida <47196116+eidenyoshida@users.noreply.github.com>
Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com>
Co-authored-by: Sam Wu <sam.wu2@amd.com>
Co-authored-by: David Galiffi <dgaliffi@amd.com>
  • Loading branch information
17 people authored Jan 17, 2024
1 parent 06c7860 commit 8b86fed
Show file tree
Hide file tree
Showing 48 changed files with 173 additions and 169 deletions.
1 change: 1 addition & 0 deletions .github/CODEOWNERS
Validating CODEOWNERS rules …
Original file line number Diff line number Diff line change
@@ -0,0 +1 @@
* @stanleytsang-amd @umfranzw @RobsonRLemos @lawruble13
8 changes: 7 additions & 1 deletion .github/dependabot.yml
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,13 @@
version: 2
updates:
- package-ecosystem: "pip" # See documentation for possible values
directory: "/docs/.sphinx" # Location of package manifests
directory: "/docs/sphinx" # Location of package manifests
open-pull-requests-limit: 10
schedule:
interval: "daily"
labels:
- "documentation"
- "dependencies"
- "ci:docs-only"
reviewers:
- "samjwu"
6 changes: 0 additions & 6 deletions .gitignore
Original file line number Diff line number Diff line change
Expand Up @@ -4,12 +4,6 @@ build*/
### clangd. ###
/.cache

### Docs dirs ###
doc/html/
doc/xml/
doc/latex/
doc/*.tag

# Created by https://www.gitignore.io/api/c++,cmake

### C++ ###
Expand Down
2 changes: 1 addition & 1 deletion .readthedocs.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@ formats: [htmlzip, pdf, epub]

python:
install:
- requirements: docs/.sphinx/requirements.txt
- requirements: docs/sphinx/requirements.txt

build:
os: ubuntu-22.04
Expand Down
50 changes: 15 additions & 35 deletions CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@
Documentation for rocPRIM is available at
[https://rocm.docs.amd.com/projects/rocPRIM/en/latest/](https://rocm.docs.amd.com/projects/rocPRIM/en/latest/).

## [Unreleased rocPRIM-3.0.0 for ROCm 6.1.0]
## Unreleased rocPRIM-3.1.0 for ROCm 6.1.0

### Additions

Expand Down Expand Up @@ -35,48 +35,28 @@ Documentation for rocPRIM is available at
* 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-2.13.1 for ROCm 5.7.0
## rocPRIM-3.0.0 for ROCm 6.0.0

### Additions
- `block_sort::sort()` overload for keys and values with a dynamic size, for all block sort algorithms. Additionally, all `block_sort::sort()` overloads with a dynamic size are now supported for `block_sort_algorithm::merge_sort` and `block_sort_algorithm::bitonic_sort`.
- New two-way partition primitive `partition_two_way` which can write to two separate iterators.

* `block_sort::sort()` overload for keys and values with a dynamic size, for all block sort algorithms
* All `block_sort::sort()` overloads with a dynamic size are now supported for
`block_sort_algorithm::merge_sort` and `block_sort_algorithm::bitonic_sort`
* New two-way partition primitive (`partition_two_way`) that can write to two separate iterators
* Added config tuning and dynamic dispatch to the `device_adjacent_difference` algorithm
* New `rocprim::group_elect` warp intrinsic, which chooses one lane from the lanes enabled by a mask

### Changes
### Optimizations
- Improved the performance of `partition`.

* Removed erroneous implementation of device-level `inclusive_scan` and `exclusive_scan` (the prior
default implementation that uses `lookback-scan` is now the only available implementation)
* The benchmark metric indicating the bytes processed for `exclusive_scan_by_key` and
`inclusive_scan_by_key` has been changed to incorporate the key type; the benchmark log has been
changed so that these algorithms are reported as `scan` and `scan_by_key` instead of
`scan_exclusive` and `scan_inclusive`
* Improved the performance of `partition`
* `merge_sort_block_sort` always uses stable merge sort because it's faster than the fallback
implementation
* The `rocprim::match_any` interface has a new parameter (`valid`) to enable and disable lanes; the
default value is true, so it doesn't change the previous behavior
### Fixes
- Fixed `rocprim::MatchAny` for devices with 64-bit warp size. The function `rocprim::MatchAny` is deprecated and `rocprim::match_any` is preferred instead.

### Deprecations
## rocPRIM-2.13.1 for ROCm 5.7.0

* `radix_sort_config` for device-level Radix Sort (it no longer matches the algorithm's parameters); use
`radix_sort_config_v2` instead
* `scan_config` and `scan_by_key_config` for device-level scans (they no longer match the algorithm's
parameters); use`scan_config_v2` and `scan_by_key_config_v2` instead)
### Changes
- Deprecated configuration `radix_sort_config` for device-level radix sort as it no longer matches the algorithm's parameters. New configuration `radix_sort_config_v2` is preferred instead.
- Removed erroneous implementation of device-level `inclusive_scan` and `exclusive_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` and `inclusive_scan_by_key` has been changed to incorporate the key type. Furthermore, the benchmark log has been changed such that these algorithms are reported as `scan` and `scan_by_key` instead of `scan_exclusive` and `scan_inclusive`.
- Deprecated configurations `scan_config` and `scan_by_key_config` for device-level scans, as they no longer match the algorithm's parameters. New configurations `scan_config_v2` and `scan_by_key_config_v2` are preferred instead.

### Fixes

* Build issue caused by a missing header in `thread/thread_search.hpp`
* `rocprim::MatchAny` for devices with 64-bit warp size (`rocprim::MatchAny` is deprecated and is
replaced with `rocprim::match_any`)
* Fixed `device_adjacent_difference` using more shared memory than required
* Fixed a compilation error when `ROCPRIM_DISABLE_DPP` is defined
* Improved robustness for detecting GPU architecture features
* Explicitly listing each architecture is no longer required
* Fixed compilation failures when targeting devices are not known by rocPRIM
- Fixed build issue caused by missing header in `thread/thread_search.hpp`.

## rocPRIM-2.13.0 for ROCm 5.5.0

Expand Down
2 changes: 1 addition & 1 deletion CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -114,7 +114,7 @@ if(USE_HIP_CPU)
endif()

# Setup VERSION
set(VERSION_STRING "2.13.1")
set(VERSION_STRING "3.1.0")
rocm_setup_version(VERSION ${VERSION_STRING})

# Print configuration summary
Expand Down
16 changes: 8 additions & 8 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@ develop performant GPU-accelerated code on AMD ROCm platforms.
* CMake (3.16 or later)
* AMD [ROCm](https://rocm.docs.amd.com/en/latest/) platform (1.8.2 or later)
* Including
[HIP-clang](https://github.com/ROCm-Developer-Tools/HIP/blob/master/INSTALL.md#hip-clang)
[HIP-clang](https://github.com/ROCm/HIP/blob/master/INSTALL.md#hip-clang)
compiler
* C++14
* Python 3.6 or higher (HIP on Windows only, required only for install script)
Expand Down Expand Up @@ -37,7 +37,7 @@ To build our documentation locally, use the following code:
cd rocPRIM; cd docs

# Install Python dependencies
python3 -m pip install -r .sphinx/requirements.txt
python3 -m pip install -r sphinx/requirements.txt

# Build the documentation
python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html
Expand All @@ -54,7 +54,7 @@ You can build and install rocPRIM on Linux or Windows.
* Linux:

```shell
git clone https://github.com/ROCmSoftwarePlatform/rocPRIM.git
git clone https://github.com/ROCm/rocPRIM.git

# Go to rocPRIM directory, create and go to the build directory.
cd rocPRIM; mkdir build; cd build
Expand Down Expand Up @@ -101,7 +101,7 @@ You can build and install rocPRIM on Linux or Windows.
We've added initial support for HIP on Windows; to install, use the provided `rmake.py` python script:
```shell
git clone https://github.com/ROCmSoftwarePlatform/rocPRIM.git
git clone https://github.com/ROCm/rocPRIM.git
cd rocPRIM
# the -i option will install rocPRIM to C:\hipSDK by default
Expand Down Expand Up @@ -289,17 +289,17 @@ algorithms optimized for gfx803 GCN version, or to `900` for gfx900.

## hipCUB

[hipCUB](https://github.com/ROCmSoftwarePlatform/hipCUB/) is a thin wrapper library on top of
[rocPRIM](https://github.com/ROCmSoftwarePlatform/rocPRIM) or
[hipCUB](https://github.com/ROCm/hipCUB/) is a thin wrapper library on top of
[rocPRIM](https://github.com/ROCm/rocPRIM) or
[CUB](https://github.com/NVlabs/cub). You can use it to port projects that use the CUB library to the
[HIP](https://github.com/ROCm-Developer-Tools/HIP) layer and run them on AMD hardware. In the
[HIP](https://github.com/ROCm/HIP) layer and run them on AMD hardware. In the
[ROCm](https://rocm.docs.amd.com/en/latest/) environment, hipCUB uses the rocPRIM library as a
backend; on CUDA platforms, it uses CUB as a backend.

## Support

You can report bugs and feature requests through our GitHub
[issue tracker](https://github.com/ROCmSoftwarePlatform/rocPRIM/issues).
[issue tracker](https://github.com/ROCm/rocPRIM/issues).

## Contributions and license

Expand Down
8 changes: 4 additions & 4 deletions benchmark/benchmark_block_sort.parallel.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -187,11 +187,11 @@ struct block_sort_benchmark : public config_autotune_interface
static constexpr unsigned int warmup_size = 5;
static constexpr bool debug_synchronous = false;

auto dispatch_block_sort(std::false_type /*stable_sort*/,
static auto dispatch_block_sort(std::false_type /*stable_sort*/,
size_t size,
const hipStream_t stream,
KeyType* d_input,
KeyType* d_output) const
KeyType* d_output)
{
hipLaunchKernelGGL(
HIP_KERNEL_NAME(
Expand All @@ -204,11 +204,11 @@ struct block_sort_benchmark : public config_autotune_interface
d_output);
}

auto dispatch_block_sort(std::true_type /*stable_sort*/,
static auto dispatch_block_sort(std::true_type /*stable_sort*/,
size_t size,
const hipStream_t stream,
KeyType* d_input,
KeyType* d_output) const
KeyType* d_output)
{
hipLaunchKernelGGL(HIP_KERNEL_NAME(stable_sort_kernel<KeyType,
ValueType,
Expand Down
23 changes: 7 additions & 16 deletions benchmark/benchmark_device_memory.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -65,7 +65,7 @@ template<class T, unsigned int ItemsPerThread, unsigned int BlockSize>
struct operation<no_operation, T, ItemsPerThread, BlockSize>
{
ROCPRIM_HOST_DEVICE inline
void operator()(T (&)[ItemsPerThread], void* = nullptr, unsigned int = 0, T* = nullptr)
void operator()(T (&)[ItemsPerThread], void* = nullptr, unsigned int = 0, T* = nullptr) const
{
// No operation
}
Expand All @@ -80,7 +80,7 @@ struct operation<custom_operation, T, ItemsPerThread, BlockSize>
ROCPRIM_HOST_DEVICE inline
void operator()(T (&input)[ItemsPerThread],
void* shared_storage = nullptr, unsigned int shared_storage_size = 0,
T* global_mem_output = nullptr)
T* global_mem_output = nullptr) const
{
(void) shared_storage;
(void) shared_storage_size;
Expand All @@ -105,7 +105,7 @@ struct operation<block_scan, T, ItemsPerThread, BlockSize>
ROCPRIM_HOST_DEVICE inline
void operator()(T (&input)[ItemsPerThread],
void* shared_storage = nullptr, unsigned int shared_storage_size = 0,
T* global_mem_output = nullptr)
T* global_mem_output = nullptr) const
{
(void) global_mem_output;
using block_scan_type = typename rocprim::block_scan<
Expand Down Expand Up @@ -419,19 +419,10 @@ void run_benchmark_memcpy(benchmark::State& state,
size_t size,
const hipStream_t stream)
{
std::vector<T> input;
if(std::is_floating_point<T>::value)
{
input = get_random_data<T>(size, (T)-1000, (T)+1000);
}
else
{
input = get_random_data<T>(
size,
std::numeric_limits<T>::min(),
std::numeric_limits<T>::max()
);
}
// Allocate device buffers
// Note: since this benchmark only tests performance by memcpying between device buffers,
// we don't really need to transfer data into these from the host - whatever happens
// to be in device memory will do.
T * d_input;
T * d_output;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&d_input), size * sizeof(T)));
Expand Down
4 changes: 3 additions & 1 deletion benchmark/benchmark_device_radix_sort_onesweep.parallel.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -49,6 +49,8 @@ constexpr const char* radix_rank_algorithm_name(rp::block_radix_rank_algorithm a
return "block_radix_rank_algorithm::basic_memoize";
case rp::block_radix_rank_algorithm::match: return "block_radix_rank_algorithm::match";
}

return ""; // unknown algorithm
}

template<typename Config>
Expand Down Expand Up @@ -403,7 +405,7 @@ struct device_radix_sort_onesweep_benchmark_generator
RadixRankAlgorithm,
std::enable_if_t<(!is_buildable<ItemsPerThread, RadixRankAlgorithm>())>>
{
void operator()(std::vector<std::unique_ptr<config_autotune_interface>>&) {}
void operator()(std::vector<std::unique_ptr<config_autotune_interface>>&) const {}
};

template<rocprim::block_radix_rank_algorithm RadixRankAlgorithm>
Expand Down
3 changes: 0 additions & 3 deletions benchmark/benchmark_device_select.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -51,7 +51,6 @@ void run_flagged_benchmark(benchmark::State& state,
{
std::vector<T> input;
std::vector<FlagType> flags = get_random_data01<FlagType>(size, true_probability);
std::vector<unsigned int> selected_count_output(1);
if(std::is_floating_point<T>::value)
{
input = get_random_data<T>(size, T(-1000), T(1000));
Expand Down Expand Up @@ -181,7 +180,6 @@ void run_selectop_benchmark(benchmark::State& state,
float true_probability)
{
std::vector<T> input = get_random_data<T>(size, T(0), T(1000));
std::vector<unsigned int> selected_count_output(1);

auto select_op = [true_probability] __device__ (const T& value) -> bool
{
Expand Down Expand Up @@ -308,7 +306,6 @@ void run_unique_benchmark(benchmark::State& state,
input[i] = op(acc, input01[i]);
}
}
std::vector<unsigned int> selected_count_output(1);
auto equality_op = rocprim::equal_to<T>();

T * d_input;
Expand Down
2 changes: 1 addition & 1 deletion benchmark/benchmark_utils.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -467,7 +467,7 @@ struct bench_naming
static format storage = human;
return storage;
}
static void set_format(std::string argument)
static void set_format(const std::string& argument)
{
format result = human;
if(argument == "json")
Expand Down
10 changes: 5 additions & 5 deletions benchmark/benchmark_warp_exchange.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,7 @@ struct BlockedToStripedOp
ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE
void operator()(warp_exchange_type warp_exchange,
T (&items)[ItemsPerThread],
typename warp_exchange_type::storage_type& storage)
typename warp_exchange_type::storage_type& storage) const
{
warp_exchange.blocked_to_striped(items, items, storage);
}
Expand All @@ -68,7 +68,7 @@ struct StripedToBlockedOp
ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE
void operator()(warp_exchange_type warp_exchange,
T (&items)[ItemsPerThread],
typename warp_exchange_type::storage_type& storage)
typename warp_exchange_type::storage_type& storage) const
{
warp_exchange.striped_to_blocked(items, items, storage);
}
Expand All @@ -84,7 +84,7 @@ struct BlockedToStripedShuffleOp
ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE
void operator()(warp_exchange_type warp_exchange,
T (&items)[ItemsPerThread],
typename warp_exchange_type::storage_type& /*storage*/)
typename warp_exchange_type::storage_type& /*storage*/) const
{
warp_exchange.blocked_to_striped_shuffle(items, items);
}
Expand All @@ -100,7 +100,7 @@ struct StripedToBlockedShuffleOp
ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE
void operator()(warp_exchange_type warp_exchange,
T (&items)[ItemsPerThread],
typename warp_exchange_type::storage_type& /*storage*/)
typename warp_exchange_type::storage_type& /*storage*/) const
{
warp_exchange.striped_to_blocked_shuffle(items, items);
}
Expand All @@ -118,7 +118,7 @@ struct ScatterToStripedOp
void operator()(warp_exchange_type warp_exchange,
T (&thread_data)[ItemsPerThread],
const OffsetT (&ranks)[ItemsPerThread],
typename warp_exchange_type::storage_type& storage)
typename warp_exchange_type::storage_type& storage) const
{
warp_exchange.scatter_to_striped(thread_data, thread_data, ranks, storage);
}
Expand Down
Loading

0 comments on commit 8b86fed

Please sign in to comment.