DOCSEARCHbyAlgolia
Submit Docs

rocSHMEM

repository·develop_deprecated·Indexed 41 days ago

https://github.com/rocm/rocshmem

A GPU-centric networking runtime providing an OpenSHMEM-like interface for intra-kernel communication. It enables fine-grained overlap of communication and computation using a single symmetric heap on GPU memories, supporting IPC, Reverse Offload (RO), and GPU Direct Async (GDA) backends.

Tokens
16.8K
Snippets
62
Records
87
Agent score
46%

What's inside rocSHMEM

  1. What is rocSHMEM?

    develop_deprecated

    rocSHMEM (ROCm OpenSHMEM) is an intra-kernel networking library designed for AMD GPUs on the ROCm platform. It provides a GPU-centric networking interface similar to the OpenSHMEM specification.

    Key features include:

    • Simplifies application complexity by enabling communication directly from GPU kernels.
    • Enables finer overlap between communication and computation compared to traditional host-driven networking.
    • Uses a single symmetric heap allocated on GPU memories.

  2. Overview of rocSHMEM

    develop_deprecated
    rocSHMEM (ROCm OpenSHMEM) is an intra-kernel networking library designed for GPU-centric networking. It provides an OpenSHMEM-like interface that simplifies application code complexity and enables finer overlap between communication and computation compared to traditional host-driven networking. A key architectural feature is that rocSHMEM uses a single symmetric heap allocated directly to GPU memories.

  3. Overview of ROCm OpenSHMEM (rocSHMEM)

    develop_deprecated
    rocSHMEM is a GPU-centric networking runtime that provides an OpenSHMEM-like interface. It is designed to simplify application code complexity and enable fine-grained communication/computation overlap by using a single symmetric heap allocated on GPU memories. Instead of traditional host-driven networking, rocSHMEM allows for intra-kernel networking directly from the device.

  4. Use Workgroup APIs with the `_wg` suffix or `_wg_` infix

    develop_deprecated

    Workgroup APIs (identified by the _wg suffix or _wg_ infix) are designed for collective communication within a GPU workgroup.

    Constraints and Behavior:

    • Uniformity: Every thread in the workgroup must call the routine with the exact same parameters. Calling these routines with different parameters results in undefined behavior.
    • Blocking: These APIs will block until the entire calling workgroup has completed the operation.

  5. Use Wavefront APIs with the `_wave` suffix

    develop_deprecated

    Wavefront APIs (identified by the _wave suffix) are designed for collective communication within a GPU wavefront.

    Constraints and Behavior:

    • Uniformity: Every thread in the wavefront must call the routine with the exact same parameters. Calling these routines with different parameters results in undefined behavior.
    • Blocking: These APIs will block until the entire calling wavefront has completed the operation.

  6. Understand rocSHMEM RMA routine naming conventions

    develop_deprecated

    rocSHMEM Remote Memory Access (RMA) routines follow specific naming suffixes that dictate their execution behavior and synchronization requirements:

    • _nbi (Non-Blocking): These routines return as soon as the request is posted. They do not wait for the transfer to complete.
    • Blocking (No suffix): Routines without the _nbi substring will block the calling thread until the operation completes locally.
    • _wave: These routines require all threads in a wavefront to call the routine with the same parameters.
    • _wg: These routines require all threads in a workgroup to call the routine with the same parameters.
    • _ctx_: These versions of the routines allow specifying a rocshmem_ctx_t context for the operation.

    Note: All dest and source addresses used in RMA operations must be located on the symmetric heap.

  7. How the rocSHMEM programming model works

    develop_deprecated

    rocSHMEM is designed to interface with the HIP programming model. It provides two main categories of APIs:

    1. __host__ APIs: Intended for host-side code.
    2. __device__ APIs: Intended for use within GPU kernels.

    Calling Device APIs

    When using __device__ APIs, the calling context determines how they must be invoked:

    • Standard Device APIs: APIs without special suffixes (like _wg or _wave) must be called by a single thread.
    • GPU-specific APIs: APIs with _wg or _wave suffixes are designed to be called by multiple threads (a workgroup or a wavefront) and will block until the calling scope completes.

    Note: While these APIs can be called in divergent code paths, this is not recommended.

  8. Understand rocSHMEM backends: IPC, RO, and GDA

    develop_deprecated

    rocSHMEM supports three different backends that implement intra-kernel networking differently:

    1. IPC (Inter-Process Communication): Implements communication primitives using load/store operations issued directly from the GPU.
    2. Reverse Offload (RO): The GPU runtime forwards networking operations to a host-side runtime, which then calls into a traditional MPI or OpenSHMEM implementation. This process is transparent to the programmer.
    3. GPU Direct Async (GDA): Allows the device-side code to issue communication operations directly to the NIC without a CPU proxy. During initialization, network resources are prepared using vendor-appropriate Direct Verbs APIs. Device threads construct Work Queue Entries (WQEs) and post them to the NIC's send queues, while Completion Queues (CQs) are polled from the device-side code.

  9. Use Atomic Memory Operations (AMO) in rocSHMEM

    develop_deprecated

    rocSHMEM provides a suite of Atomic Memory Operations (AMO) that can be called from divergent control paths at the per-thread level. These operations allow for atomic updates to memory locations on a remote Processing Element (PE).

    Key Requirements

    • Symmetric Heap: The dest (destination address) must be an address located on the symmetric heap.
    • Blocking Behavior: Most AMO functions in rocSHMEM are blocking operations.
    • Contexts: Many functions provide two versions: a standard version and a context-aware version (rocshmem_ctx_TYPENAME_...) that accepts a rocshmem_ctx_t ctx parameter.

    Operation Categories

    • Fetch/Set: atomic_fetch, atomic_set, atomic_swap.
    • Arithmetic: atomic_inc, atomic_add (and their fetch variants).
    • Bitwise: atomic_and, atomic_or, atomic_xor (and their fetch variants).
    • Comparison: atomic_compare_swap.

  10. Understand rocSHMEM network backends

    develop_deprecated

    rocSHMEM supports three primary communication backends that can be activated in the same build:

    • IPC (Inter-Process Communication): Enables fast communication between GPUs on the same host using ROCm inter-process mechanisms. It does not support inter-node communication.
    • RO (Reverse Offload): Enables communication between GPUs on different nodes through a NIC using a host-based proxy. It is built on an MPI-RMA compatibility layer. Requires ROCm-aware Open MPI and UCX.
    • GDA (GPU Direct Async): Enables communication between GPUs on different nodes through a NIC where the GPU directly interacts with the NIC without CPU involvement in the critical path.

    Backend Interaction Rules

    • If RO + IPC are both active: All atomic operations use the RO backend, even for intra-node communication.
    • If GDA + IPC are both active: All atomic operations use the GDA backend, even for intra-node communication.

  11. Build PDF documentation on macOS

    develop_deprecated

    To build PDF documentation on macOS, you must have a LaTeX installation on your machine. After installing LaTeX, install the Sphinx dependencies and use sphinx-build with the latexpdf builder to generate the PDF in the _build/latex directory.

    pip3.10 install -r ./sphinx/requirements.txt
sphinx-build -M latexpdf . _build
open _build/latex/rocshmem.pdf

  12. Get a unique ID for group processes

    develop_deprecated

    To facilitate initialization using unique IDs, follow these steps:

    1. Generate a Unique ID: Call rocshmem_get_uniqueid() to obtain a unique ID handle.
    2. Configure Attributes: Use rocshmem_set_attr_uniqueid_args() to initialize a rocshmem_init_attr_t structure with the rank, number of PEs, and the generated unique ID.
    3. Initialize: Pass the configured attributes to rocshmem_init_attr() using the ROCSHMEM_INIT_WITH_UNIQUEID flag.
    rocshmem_uniqueid_t uid;
rocshmem_get_uniqueid(&uid);

rocshmem_init_attr_t attr;
rocshmem_set_attr_uniqueid_args(my_rank, total_ranks, &uid, &attr);

rocshmem_init_attr(ROCSHMEM_INIT_WITH_UNIQUEID, &attr);