Anatolii Rozanov
d0c8380650
* Add functional test for barrier_all_on_stream
* Add rocshmem_barrier_all_on_stream support for GDA and RO backends
Implements rocshmem_barrier_all_on_stream operation for
GPU Direct Access and Reverse Offload backends.
Previously, rocshmem_barrier_all_on_stream was only supported for IPC backend.
* Add functional test for rocshmem_broadcastmem_on_stream
* Add host-side rocshmem_broadcastmem_on_stream API
Implement stream-based broadcast collective operation
- Add rocshmem_broadcastmem_on_stream host API and kernel implementation
- Add functional test TeamBroadcastmemOnStreamTester with multi-stream
support and correctness verification
- Use per-workgroup contexts to avoid contention across parallel streams
API:
rocshmem_broadcastmem_on_stream(team, dest, source, nelems, pe_root, stream)
* Add functional test for rocshmem_getmem_on_stream
* Add host-side rocshmem_getmem_on_stream API
Implement stream-based point-to-point RMA get operation
- Add rocshmem_getmem_on_stream host API and kernel implementation
- Support for asynchronous getmem operations on HIP streams
- Add backend support for GDA, RO, and IPC contexts
- Use work-group collective getmem for efficient memory transfer
API:
rocshmem_getmem_on_stream(dest, source, nelems, pe, stream)
(AI Assist)
* Add host-side rocshmem_putmem_on_stream API
- Add rocshmem_putmem_on_stream for asynchronous remote writes
- Support for concurrent RMA operations on HIP streams
- Add backend support for GDA, RO, and IPC contexts
- Use work-group device collective operation
API:
rocshmem_putmem_on_stream(dest, source, bytes, pe, stream)
(AI Assist)
* Add functional test for rocshmem_putmem_on_stream
* Add host-side rocshmem_putmem_signal_on_stream API
Enables asynchronous putmem operations with signaling on HIP streams.
The implementation includes:
- Kernel wrapper rocshmem_putmem_signal_kernel
- Host interface putmem_signal_on_stream method
- Context layer support across all backends (IPC, GDA, RO)
- Public API
Function signature:
void rocshmem_putmem_signal_on_stream(void *dest, const void *source,
size_t bytes, uint64_t *sig_addr,
uint64_t signal, int sig_op,
int pe, hipStream_t stream);
* Add functional test for rocshmem_putmem_signal_on_stream
* Add host-side rocshmem_signal_wait_until_on_stream API
Enables asynchronous signal wait operations on HIP streams.
The implementation includes:
- Kernel wrapper rocshmem_signal_wait_until_kernel
- Host interface signal_wait_until_on_stream method
- Context layer support across all backends (IPC, GDA, RO)
- Native uint64_t support in wait_until API (generated from P2P_SYNC.py)
Function signature:
void rocshmem_signal_wait_until_on_stream(uint64_t *sig_addr, int cmp,
uint64_t cmp_value,
hipStream_t stream);
(AI Assist)
* Add functional test for rocshmem_signal_wait_until_on_stream
* Add documentation for stream API functions
This commit adds API documentation for the following host-side
stream functions:
- rocshmem_barrier_all_on_stream (collective routines)
- rocshmem_broadcastmem_on_stream (collective routines)
- rocshmem_getmem_on_stream (RMA operations)
- rocshmem_putmem_on_stream (RMA operations)
- rocshmem_putmem_signal_on_stream (signaling operations)
- rocshmem_signal_wait_until_on_stream (point-to-point sync)
The documentation includes function signatures, parameter descriptions,
and detailed explanations of asynchronous behavior and stream handling.
(AI Assist)
* Rename "bytes" -> "nelems"
* Add "_TEST_" to the variables used in tests
* Remove incorrect hipStreamDefault usage
hipStreamDefault is not a default stream. This is a flag.
If stream == nullptr, then just pass it to kernel. It will launch the kernel on the default stream
143 líneas
5.3 KiB
ReStructuredText
143 líneas
5.3 KiB
ReStructuredText
.. meta::
|
|
:description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform.
|
|
:keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication
|
|
|
|
.. _rocshmem-api-pt2pt-sync:
|
|
|
|
-----------------------------------------
|
|
Point-to-point synchronization routines
|
|
-----------------------------------------
|
|
|
|
ROCSHMEM_WAIT_UNTIL
|
|
-------------------
|
|
|
|
.. cpp:function:: __device__ void rocshmem_TYPENAME_wait_until(TYPE *ivars, int cmp, TYPE val)
|
|
|
|
:param ivars: Pointer to memory on the symmetric heap to wait for.
|
|
:param cmp: Operation for the comparison.
|
|
:param val: Value to compare the memory at ``ivars`` to.
|
|
:returns: None.
|
|
|
|
**Description:**
|
|
This routine blocks the caller until the condition ``(*ivars cmp val)`` is true.
|
|
|
|
Valid ``cmp`` values are listed in :ref:`CMP_VALUES`.
|
|
|
|
Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`STANDARD_AMO_TYPES`.
|
|
|
|
ROCSHMEM_WAIT_UNTIL_ALL
|
|
-----------------------
|
|
|
|
.. cpp:function:: __device__ void rocshmem_TYPENAME_wait_until_all(TYPE *ivars, size_t nelems, const int* status, int cmp, TYPE val)
|
|
|
|
:param ivars: Pointer to memory on the symmetric heap to wait for.
|
|
:param nelems: Number of elements in the ``ivars`` array.
|
|
:param status: Array of length ``nelems`` to exclude elements from the wait.
|
|
:param cmp: Operation for the comparison.
|
|
:param val: Value to compare.
|
|
:returns: None.
|
|
|
|
**Description:**
|
|
This routine blocks the caller until the condition ``(ivars[i] cmp val)`` is true for all ``ivars``.
|
|
|
|
Valid ``cmp`` values are listed in :ref:`CMP_VALUES`.
|
|
|
|
Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`STANDARD_AMO_TYPES`.
|
|
|
|
ROCSHMEM_WAIT_UNTIL_ANY
|
|
-----------------------
|
|
.. cpp:function:: __device__ size_t rocshmem_TYPENAME_wait_until_any(TYPE *ivars, size_t nelems, const int* status, int cmp, TYPE val)
|
|
|
|
:param ivars: Pointer to memory on the symmetric heap to wait for.
|
|
:param nelems: Number of elements in the ``ivars`` array.
|
|
:param status: Array of length ``nelems`` to exclude elements from the wait.
|
|
:param cmp: Operation for the comparison.
|
|
:param val: Value to compare.
|
|
:returns: The index of an element in the ``ivars`` array that satisfies the wait condition. If the wait set is empty, this routine returns ``SIZE_MAX``.
|
|
|
|
**Description:**
|
|
This routine blocks the caller until any of the condition ``(ivars[i] cmp val)`` is true.
|
|
|
|
Valid ``cmp`` values are listed in :ref:`CMP_VALUES`.
|
|
|
|
Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`STANDARD_AMO_TYPES`.
|
|
|
|
ROCSHMEM_WAIT_UNTIL_SOME
|
|
------------------------
|
|
|
|
.. cpp:function:: __device__ size_t rocshmem_TYPENAME_wait_until_some(TYPE *ivars, size_t nelems, size_t* indices, const int* status, int cmp, TYPE val)
|
|
|
|
:param ivars: Pointer to memory on the symmetric heap to wait for.
|
|
:param nelems: Number of elements in the ``ivars`` array.
|
|
:param indices: List of indices with a length of at least ``nelems``.
|
|
:param status: Array of length ``nelems`` to exclude elements from the wait.
|
|
:param cmp: Operation for the comparison.
|
|
:param val: Value to compare.
|
|
:returns: The number of indices returned in the indices array. If the wait set is empty, this routine returns ``0``.
|
|
|
|
**Description:**
|
|
This routine blocks the caller until any of the conditions ``(ivars[i] cmp val)`` is true.
|
|
|
|
Valid ``cmp`` values are listed in :ref:`CMP_VALUES`.
|
|
|
|
Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`STANDARD_AMO_TYPES`.
|
|
|
|
ROCSHMEM_TEST
|
|
-------------
|
|
|
|
.. cpp:function:: __device__ int rocshmem_TYPENAME_test(TYPE *ivars, int cmp, TYPE val)
|
|
|
|
:param ivars: Pointer to memory on the symmetric heap to wait for.
|
|
:param cmp: Operation for the comparison.
|
|
:param val: Value to compare the memory at ``ivars`` to.
|
|
|
|
:returns: ``1`` if the evaluation is true. ``0`` otherwise.
|
|
|
|
**Description:**
|
|
This routine tests if the condition ``(*ivars cmp val)`` is true.
|
|
|
|
ROCSHMEM_SIGNAL_WAIT_UNTIL_ON_STREAM
|
|
-------------------------------------
|
|
|
|
.. cpp:function:: __host__ void rocshmem_signal_wait_until_on_stream(uint64_t *sig_addr, int cmp, uint64_t cmp_value, hipStream_t stream)
|
|
|
|
:param sig_addr: Address of the signal variable on the symmetric heap.
|
|
:param cmp: Comparison operator (e.g., ROCSHMEM_CMP_EQ, ROCSHMEM_CMP_GE, etc.).
|
|
:param cmp_value: Value to compare against.
|
|
:param stream: HIP stream on which to enqueue the operation.
|
|
:returns: None.
|
|
|
|
**Description:**
|
|
This routine enqueues a wait operation on a HIP stream. The function blocks the calling thread
|
|
until the signal variable at ``sig_addr`` satisfies the comparison condition ``(*sig_addr cmp cmp_value)``.
|
|
The wait operation is executed asynchronously on the specified stream. The caller must synchronize
|
|
the stream (e.g., using ``hipStreamSynchronize``) to ensure the wait condition has been satisfied.
|
|
|
|
Valid ``cmp`` values are listed in :ref:`CMP_VALUES`.
|
|
|
|
.. _CMP_VALUES:
|
|
|
|
Supported comparisons
|
|
---------------------
|
|
|
|
The following table lists the point-to-point comparison constants:
|
|
|
|
.. list-table:: Point-to-Point Comparison Constants
|
|
:widths: 20 20
|
|
:header-rows: 1
|
|
|
|
* - Constant
|
|
- Description
|
|
* - ROCSHMEM_CMP_EQ
|
|
- Equal
|
|
* - ROCSHMEM_CMP_NE
|
|
- Not equal
|
|
* - ROCSHMEM_CMP_GT
|
|
- Greater than
|
|
* - ROCSHMEM_CMP_GE
|
|
- Greater than or equal to
|
|
* - ROCSHMEM_CMP_LT
|
|
- Less than
|
|
* - ROCSHMEM_CMP_LE
|
|
- Less than or equal to
|