Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add changelog entry about hipGraph support #412

Merged
merged 1 commit into from
Jun 19, 2024
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
1 change: 1 addition & 0 deletions CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@ Documentation for rocThrust available at

### Known issues
* `thrust::reduce_by_key` outputs are not bit-wise reproducible, as run-to-run results for pseudo-associative reduction operators (e.g. floating-point arithmetic operators) are not deterministic on the same device.
* Note that currently, rocThrust memory allocation is performed in such a way that most algorithmic API functions cannot be called from within hipGraphs.

## rocThrust 3.0.0 for ROCm 6.0

Expand Down
24 changes: 24 additions & 0 deletions docs/hipgraph-support.rst
Original file line number Diff line number Diff line change
@@ -0,0 +1,24 @@
.. meta::
:description: rocThrust documentation and API reference
:keywords: rocThrust, ROCm, API, reference, hipGraph

.. _hipgraph-support:

******************************************
hipGraph Support
******************************************
Currently, rocThrust does not support the use of ``hipGraphs``. ``hipGraphs`` are not allowed to contain any synchronous
function calls or barriers. Thrust API functions are blocking (synchronous with respect to the host) by default.

Thrust does provide asynchronous versions of a number of algorithms. These are contained in the ``thrust::async`` namespace
(see the headers in ``rocThrust/thrust/async/``). However, these algorithms operate asynchronously by returning futures.
This approach is different from the form of asynchronous execution required within ``hipGraphs``, which must be achieved by
issuing calls into a user-defined ``hipStream``.

While it is possible to create execution policies that encourage Thrust API algorithms to execute within a user-defined stream,
(eg. ``thrust::hip_rocprim::par.on(stream)``), this is not enough to guarentee that synchronization will not occur within
a given algorithm. This is because some Thrust functions require execution policies to be passed in at compile time (via template
arguments) rather than at runtime. Since streams must be created at runtime, there is no way to pass these functions a stream.
Adding a stream argument to such functions breaks compatibility with the Thrust API.

For these reasons, we recommend that you do not use hipGraphs together with rocThrust code.
1 change: 1 addition & 0 deletions docs/index.rst
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,7 @@ The documentation is structured as follows:

* :ref:`data-type-support`
* :ref:`bitwise-repro`
* :ref:`hipgraph-support`
* :ref:`api-reference`
* :ref:`genindex`

Expand Down