### NVSHMEM4Py Quick Start

Source: https://docs.nvidia.com/nvshmem/api/api/language_bindings/python/index.html

A quick start guide demonstrating how to initialize NVSHMEM, allocate memory, perform communication operations, and finalize the NVSHMEM environment using NVSHMEM4Py.

```APIDOC
## NVSHMEM4Py Quick Start

To use NVSHMEM4Py in your Python application:

```python
import nvshmem.core as nvshmem
from mpi4py import MPI
from cuda.core.experimental import Device

dev = Device()
dev.set_current()
stream = dev.create_stream()

# Initialize MPI
comm = MPI.COMM_WORLD

# Initialize NVSHMEM with MPI
nvshmem.init(dev, mpi_comm=comm, init_method="mpi")

# Get information about the current PE
my_pe = nvshmem.my_pe()
n_pes = nvshmem.n_pes()

# Allocate symmetric memory
# array() returns a CuPy NDArray object
x = nvshmem.array((1024,), dtype="float32")
y = nvshmem.array((1024,), dtype="float32")

if my_pe == 0:
    y[:] = 1.0

# Perform communication operations
# Put y from PE 0 into x on PE 1
if my_pe == 0:
    nvshmem.put(x, y, pe=1, stream=stream)

# Synchronize PEs
stream.sync()

# Clean up
nvshmem.free_array(x)
nvshmem.free_array(y)
nvshmem.finalize()
```
```

--------------------------------

### NVSHMEM Python Quick Start and Key Features

Source: https://docs.nvidia.com/nvshmem/api/api/language_bindings/index.html

Provides a quick start guide and highlights the key features of the NVSHMEM Python bindings.

```APIDOC
## NVSHMEM Python Quick Start and Key Features

### Description
This section offers a starting point for using NVSHMEM with Python and outlines the main features of the NVSHMEM4Py library.

### Quick Start
[Quick Start](https://docs.nvidia.com/nvshmem/api/api/language_bindings/python/index.html.md#quick-start)

### Key Features
[Key Features](https://docs.nvidia.com/nvshmem/api/api/language_bindings/python/index.html.md#key-features)
```

--------------------------------

### NVSHMEM Installation Guide

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/index.html

Provides comprehensive instructions for installing NVSHMEM.

```APIDOC
## NVSHMEM Installation Guide

### Description
This document outlines the steps required to install the NVIDIA NVSHMEM library.

### Endpoint
/websites/nvidia_nvshmem/release-notes-install-guide/install-guide/abstract.html.md
```

--------------------------------

### PyTorch and Triton Interoperability Setup

Source: https://docs.nvidia.com/nvshmem/api/examples/language_bindings/python/index.html

This example demonstrates initializing NVSHMEM4Py using a Torch.distributed ProcessGroup, suitable for use with `torchrun`. It sets up the environment for a Triton kernel to interoperate with NVSHME4Py.

```python
"""
This example initializes NVSHMEM4Py with the `torchrun`
launcher and torch.distributed

It runs a kernel expressed with Triton

Run this program with `torchrun --nproc-per-node <NGPUs> torch_triton_interop.py`
"""

import torch.distributed as dist
import torch
import triton
import triton.language as tl
import nvshmem.core
import os
from cuda.core.experimental import Device, system

###
#  Helper code from https://github.com/NVIDIA/cuda-python/blob/main/cuda_core/examples/pytorch_example.py
#  Used to extract PyTorch Stream into a cuda.core.Stream for NVSHMEM APIs
###

# Create a wrapper class that implements __cuda_stream__

```

--------------------------------

### NVSHMEM4Py 'Hello World' Example

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem4py-install-proc.html

A 'Hello World' example demonstrating NVSHMEM4Py initialization, device configuration, and a basic AllReduce collective operation. This snippet shows how to create NVSHMEM arrays and perform a sum reduction.

```python
import cupy
import nvshmem.core as nvshmem
from cuda.core.experimental import Device, system
import mpi4py

local_rank_per_node = MPI.COMM_WORLD.Get_rank() % system.num_devices
dev = Device(local_rank_per_node)
dev.set_current()
stream = dev.create_stream()
vshmem.core.init(device=dev, mpi_comm=MPI.COMM_WORLD, initializer_method="mpi")

arr_src = nvshmem.core.array((2,2), dtype="float32")
arr_dst = nvshmem.core.array((2,2), dtype="float32")
arr_dst[:] = 0
arr_src[:] = local_rank_per_node + 1

# Perform a sum reduction from arr_src to arr_dst across all PEs in TEAM_WORLD (an AllReduce)
vshmem.core.reduce(nvshmem.core.Teams.TEAM_WORLD, arr_dst, arr_src, "sum", stream=stream)
stream.sync()

# Print dst, src after
print(f"Dest after collective from PE {nvshmem.core.my_pe()}:", arr_dst)
print(f"Src after collective from PE {nvshmem.core.my_pe()}:", arr_src)
```

--------------------------------

### Install NVSHMEM Network Repository Key on Ubuntu

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

Installs the downloaded NVSHMEM network repository key using dpkg. This command should be run after downloading the key.

```bash
sudo dpkg -i cuda-keyring_1.1-1_all.deb
```

--------------------------------

### Install NVSHMEM Runtime and Development Packages on Ubuntu

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

Installs the NVSHMEM runtime library and development headers using APT. This command installs the latest available versions.

```bash
sudo apt install libnvshmem3-cuda-12 libnvshmem3-dev-cuda-12
```

--------------------------------

### Synchronize and Broadcast Example

Source: https://docs.nvidia.com/nvshmem/api/api/language_bindings/python/device/numba/collectives.html

This example demonstrates synchronizing all PEs with a barrier and then broadcasting data from the root PE. Only thread 0 performs the broadcast for demonstration.

```python
# This example synchronizes all PEs with barrier_all and then broadcasts the value in buf[0] from the root PE (PE 0) to all PEs.
# Only thread 0 performs the broadcast for demonstration purposes. In practice, you can have multiple threads participate as needed.
# For more details, see the Numba test suite and the NVSHMEM4Py documentation.
```

--------------------------------

### Install Local NVSHMEM Repository

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

Install the NVSHMEM repository from a local RPM file.

```bash
sudo rpm -i nvshmem-repo-<version>.rpm
```

--------------------------------

### Simple P2P Kernel Example

Source: https://docs.nvidia.com/nvshmem/api/index.html

A basic example demonstrating point-to-point communication within a kernel using NVSHMEM. This snippet likely involves sending and receiving data between ranks.

```python
import numba
import nvshmem.core.rma as nr

@numba.cuda.jit
def kernel(send_buf, recv_buf):
    rank = nr.my_rank()
    if rank == 0:
        nr.put(send_buf, 0, recv_buf, 0, 1, 1) # Put to rank 1
    elif rank == 1:
        # Receive operation would typically be here or implicitly handled
        pass

# Example usage:
# send_data = np.array([42], dtype=np.int32)
# recv_data = np.zeros(1, dtype=np.int32)
# kernel[1, 1](send_data, recv_data)
```

--------------------------------

### Installing NVSHMEM Language Bindings for Python

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/index.html

Instructions for installing NVSHMEM Python bindings.

```APIDOC
## Installing NVSHMEM Language Bindings for Python

### Description
This guide details the process of installing the NVSHMEM language bindings specifically for Python.

### Endpoint
/websites/nvidia_nvshmem/release-notes-install-guide/install-guide/nvshmem4py-install-proc.html.md
```

--------------------------------

### NVSHMEM Installation Procedures

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/index.html

Detailed procedures for installing NVSHMEM.

```APIDOC
## NVSHMEM Installation Procedures

### Description
This section provides detailed procedures for installing NVSHMEM on your system.

### Endpoint
/websites/nvidia_nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html.md
```

--------------------------------

### Install NVSHMEM Local Repository on Ubuntu

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

Installs NVSHMEM from a local .deb package. Ensure to follow prompts for installing the local key.

```bash
sudo dpkg -i nvshmem-repo-<version>.deb
```

--------------------------------

### Install NVSHMEM Packages

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

Install NVSHMEM libraries and development headers using YUM.

```bash
sudo yum install libnvshmem3 libnvshmem3-devel libnvshmem3-static
```

```bash
sudo yum install libnvshmem3-cuda-11-3.1.7-1 libnvshmem3-devel-cuda-11-3.1.7-1 libnvshmem3-static-cuda-11-3.1.7-1
```

--------------------------------

### Install NVSHMEM4Py via PyPI

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem4py-install-proc.html

Commands to create a virtual environment and install the package using pip.

```bash
virtualenv nvshmem4py-env # Create a virtual environment
source nvshmem4py-env/bin/activate # Activate the virtual environment
pip install nvshmem4py-cu12 # Install NVSHMEM4Py
```

--------------------------------

### nvshmem.core.init_status()

Source: https://docs.nvidia.com/nvshmem/api/api/language_bindings/python/initialization.html

Get the current initialization status of the NVSHMEM runtime.

```APIDOC
## GET /nvshmem/core/init_status

### Description
Get the current initialization status of the NVSHMEM runtime.

### Method
GET

### Endpoint
/nvshmem/core/init_status

### Response
#### Success Response (200)
- **status** (nvshmem.bindings.nvshmem.Init_status) - An enum representing the status of initialization.

#### Response Example
```json
{
  "status": "NVSHMEM_INITED"
}
```
```

--------------------------------

### Example of nvshmem_fence in C

Source: https://docs.nvidia.com/nvshmem/api/gen/api/ordering.html

This C example demonstrates the usage of nvshmem_fence. It ensures that 'Put1' operations are ordered before 'put3', and 'put2' operations are ordered before 'put4' in terms of delivery.

```c
#include <nvshmem.h>

int main() {
    // ... setup code ...

    // Example of Put1 and put3 ordering
    nvshmem_put(&data_on_pe1, &local_data, 1, PE1);
    nvshmem_fence();
    nvshmem_put(&data_on_pe1, &local_data_2, 1, PE1);

    // Example of put2 and put4 ordering
    nvshmem_put(&data_on_pe2, &local_data_3, 1, PE2);
    nvshmem_fence();
    nvshmem_put(&data_on_pe2, &local_data_4, 1, PE2);

    // ... other code ...

    return 0;
}
```

--------------------------------

### Run nvshmem_hello_world.py on Two Hosts (1 GPU Each)

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem4py-install-proc.html

Use this command to run the `nvshmem_hello_world.py` sample across two hosts, each with one GPU, connected via InfiniBand. Replace `hostname1` and `hostname2` with actual hostnames.

```bash
$HYDRA_HOME/bin/nvshmrun -n 2 -ppn 1 --hosts hostname1,hostname2 python3 nvshmem_hello_world.py
```

--------------------------------

### NVSHMEM4Py Build Error Log

Source: https://docs.nvidia.com/nvshmem/api/faq.html

Example of a compilation failure when the CUDA runtime headers are missing during the installation of NVSHMEM4Py.

```text
  building 'nvshmem.bindings.nvshmem' extension
  creating build/temp.linux-x86_64-cpython-310/nvshmem/bindings
  x86_64-linux-gnu-g++ -Wno-unused-result -Wsign-compare -DNDEBUG -g -fwrapv -O2 -Wall -g -fstack-protector-strong -Wformat -Werror=format-security -g -fwrapv -O2 -fPIC -I/home/benjaming/repo/nvshmem_venv/include -I/usr/include/python3.10 -c nvshmem/bindings/nvshmem.cpp -o build/temp.linux-x86_64-cpython-310/nvshmem/bindings/nvshmem.o
  nvshmem/bindings/nvshmem.cpp:1137:10: fatal error: cuda_runtime_api.h: No such file or directory
  1137 | #include "cuda_runtime_api.h"
      |          ^~~~~~~~~~~~~~~~~~~~
  compilation terminated.
  error: command '/bin/x86_64-linux-gnu-g++' failed with exit code 1
  [end of output]

note: This error originates from a subprocess, and is likely not a problem with pip.
ERROR: Failed building wheel for nvshmem4py-cu12
Failed to build nvshmem4py-cu12
ERROR: Could not build wheels for nvshmem4py-cu12, which is required to install pyproject.toml-based projects
```

--------------------------------

### Perform Put and Get RMA Operations in NVSHMEM4Py

Source: https://docs.nvidia.com/nvshmem/api/api/language_bindings/python/rma.html

Examples demonstrating data transfer between processing elements using Put and Get operations. Explicit memory management via nvshmem.free is required for symmetric memory allocations.

```python
import nvshmem.core as nvshmem
import cupy as cp
from cuda.core.experimental import Device

# Initialize NVSHMEM (initialization code not shown)

# Get current device and create a stream
device = Device()
stream = device.create_stream()

# Allocate symmetric memory
size = 10
src_array = nvshmem.array((size,), dtype=cp.float32)
dest_array = nvshmem.array((size,), dtype=cp.float32)

# Set values on local PE
my_pe = nvshmem.my_pe()
n_pes = nvshmem.n_pes()

# Fill source array with PE ID
src_array[:] = cp.ones(size, dtype=cp.float32) * my_pe

# Target PE (circular next PE)
target_pe = (my_pe + 1) % n_pes

# Put data to the target PE
nvshmem.put(dest_array, src_array, size, target_pe, stream=stream)

# Ensure operation is complete
stream.synchronize()

# Clean up - explicit free is required
nvshmem.free(src_array)
nvshmem.free(dest_array)
```

```python
import nvshmem.core as nvshmem
import cupy as cp
from cuda.core.experimental import Device

# Initialize NVSHMEM (initialization code not shown)

# Get current device and create a stream
device = Device()
stream = device.create_stream()

# Allocate symmetric memory
size = 10
src_array = nvshmem.array((size,), dtype=cp.float32)
dest_array = nvshmem.array((size,), dtype=cp.float32)

# Set values on each PE
my_pe = nvshmem.my_pe()
n_pes = nvshmem.n_pes()

# Fill source array with PE ID
src_array[:] = cp.ones(size, dtype=cp.float32) * my_pe

# Target PE to get data from (circular previous PE)
target_pe = (my_pe - 1 + n_pes) % n_pes

# Get data from the target PE
nvshmem.get(dest_array, src_array, size, target_pe, stream=stream)

# Ensure operation is complete
stream.synchronize()
```

--------------------------------

### Install Hydra Process Manager

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

Use this bash script to install the Hydra Process Manager. Provide the download and install locations as arguments.

```bash
./install_hydra.sh <download_path> <install_path>
```

--------------------------------

### Run nvshmem_hello_world.py on One Host (2 GPUs)

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem4py-install-proc.html

Use this command to run the `nvshmem_hello_world.py` sample on a single host with two GPUs. Ensure `$HYDRA_HOME` is set correctly.

```bash
$HYDRA_HOME/bin/nvshmrun -n 2 -ppn 2 python3 nvshmem_hello_world.py
```

--------------------------------

### Initialize and Execute NVSHMEM Main Routine

Source: https://docs.nvidia.com/nvshmem/api/examples.html

Main entry point for initializing NVSHMEM, validating hardware requirements, and executing kernels.

```cpp
int main(int argc, char const **args) {
    // initialize nvshmem
    nvshmem_init();
    mype = nvshmem_my_pe();
    npes = nvshmem_n_pes();
    CUDA_CHECK(cudaSetDevice(mype));
    printf(" Executing PE: %d out of %d\n", mype, npes);

    // CUTLASS must be compiled with CUDA 12.0 Toolkit to run this example
    // and must have compute capability at least 100a.

    if (__CUDACC_VER_MAJOR__ < 12 ||
       (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ < 8)) {
        std::cerr << "This example requires CUDA 12.8 or newer."
                  << std::endl;
        // Returning zero so this test passes on older Toolkits.
        // Its actions are no-op.
        return 0;
    }

    cudaDeviceProp props;
    int current_device_id;
    CUDA_CHECK(cudaGetDevice(&current_device_id));
    CUDA_CHECK(cudaGetDeviceProperties(&props, current_device_id));
    cudaError_t error = cudaGetDeviceProperties(&props, 0);
    if (props.major != 10 || props.minor != 0) { std::cerr
        << "This example requires a GPU with compute capability 100a)."
        << std::endl;
        return 0;
    }

    //
    // Parse options
    //

    Options options;

    options.parse(argc, args);

    if (options.help) {
        options.print_usage(std::cout) << std::endl;
        return 0;
    }

    //
    // Evaluate CUTLASS kernels
    //

#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
    run<Gemm>(options);
#endif  // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)

    nvshmem_barrier_all();

    for (int i = 0; i < num_teams; ++i) {
        nvshmem_team_destroy(teams[i]);
    }
    nvshmem_barrier_all();

    block_D.free();
    block_D_red.free();
    block_ref_D.free();
    free(teams);
    CUDA_CHECK(cudaFree(teams_dev));
    nvshmem_finalize();
    return 0;
}
```

--------------------------------

### Library Setup, Exit, and Query APIs

Source: https://docs.nvidia.com/nvshmem/api/index.html

APIs for initializing, finalizing, and querying the NVSHMEM library, as well as managing process information.

```APIDOC
## NVSHMEM Initialization and Finalization APIs

### Description
Functions to initialize and finalize the NVSHMEM library, query process information, and manage unique IDs.

### NVSHMEM_INIT
Initializes the NVSHMEM library.

### NVSHMEMX_INIT_ATTR
Initializes the NVSHMEM library with attributes.

### NVSHMEMX_HOSTLIB_INIT_ATTR
Initializes the host library with attributes.

### NVSHMEMX_HOSTLIB_FINALIZE
Finalizes the host library.

### NVSHMEMX_GET_UNIQUE_ID
Retrieves a unique ID for the NVSHMEM process.

### NVSHMEMX_SET_ATTR_UNIQUEID_ARGS
Sets arguments for unique ID attributes.

### NVSHMEMX_CUMODULE_INIT
Initializes the CUDA module for NVSHMEM.

### NVSHMEMX_INIT_STATUS
Initializes the NVSHMEM library and returns its status.

### NVSHMEM_MY_PE
Returns the process ID of the calling process.

### NVSHMEM_N_PES
Returns the total number of processes in the job.

### NVSHMEM_FINALIZE
Finalizes the NVSHMEM library.

### NVSHMEM_GLOBAL_EXIT
Exits all processes in the job.

### NVSHMEM_PTR
Returns a remote pointer to a local memory address.

### NVSHMEMX_MC_PTR
Returns a remote pointer to a memory address.

### NVSHMEM_INFO_GET_VERSION
Retrieves the version information of the NVSHMEM library.

### NVSHMEM_INFO_GET_NAME
Retrieves the name of the NVSHMEM library.

### NVSHMEMX_VENDOR_GET_VERSION_INFO
Retrieves vendor-specific version information.
```

--------------------------------

### Extract NVSHMEM Tarball

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

Navigate to the installation directory to extract the OS-agnostic installer.

```bash
# cd /usr/local
```

--------------------------------

### NVSHMEM Best Practice Guide

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/index.html

Guide on best practices for using NVSHMEM.

```APIDOC
## NVSHMEM Best Practice Guide

### Description
This document provides best practices for effectively using the NVIDIA NVSHMEM library.

### Endpoint
/websites/nvidia_nvshmem/release-notes-install-guide/best-practice-guide/index.html.md
```

--------------------------------

### Library Setup and Query APIs

Source: https://docs.nvidia.com/nvshmem/api/api.html

Core functions for initializing the NVSHMEM library, querying environment information, and finalizing the library state.

```APIDOC
## NVSHMEM_INIT

### Description
Initializes the NVSHMEM library for the calling process.

## NVSHMEM_MY_PE

### Description
Returns the processing element (PE) number of the calling PE.

## NVSHMEM_N_PES

### Description
Returns the total number of PEs in the current job.

## NVSHMEM_FINALIZE

### Description
Releases resources and shuts down the NVSHMEM library.
```

--------------------------------

### Attribute-Based Initialization Example

Source: https://docs.nvidia.com/nvshmem/api/index.html

Shows how to initialize NVSHMEM4Py using attributes, which might include parameters like network transport or specific device configurations.

```python
import nvshmem.core as nc

# Initialize NVSHMEM with specific attributes
# nc.init(network='verbs', device_type='gpu')

# ... rest of the NVSHMEM operations ...

# nc.finalize()
```

--------------------------------

### NVSHMEM Initialization Best Practices

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/index.html

Best practices for initializing NVSHMEM.

```APIDOC
## NVSHMEM Initialization

### Description
This section covers best practices for initializing the NVSHMEM library.

### Endpoint
/websites/nvidia_nvshmem/release-notes-install-guide/best-practice-guide/init.html.md
```

--------------------------------

### Install NVSHMEM4Py via Conda-Forge

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem4py-install-proc.html

Command to install the package from the NVIDIA channel on Conda-Forge.

```bash
conda install -c nvidia nvshmem4py # Install NVSHMEM4Py
```

--------------------------------

### Initialization Methods

Source: https://docs.nvidia.com/nvshmem/api/api.html

Methods for initializing and finalizing the NVSHMEM environment, querying its status, and retrieving version information.

```APIDOC
## Initialization Methods

### Description
Provides functions to initialize and finalize the NVSHMEM environment, check its status, and retrieve version details.

### Methods
- `nvshmem.init()`: Initializes the NVSHMEM environment.
- `nvshmem.finalize()`: Finalizes the NVSHMEM environment.
- `nvshmem.is_initialized()`: Checks if NVSHMEM is initialized.
- `nvshmem.version()`: Returns the NVSHMEM version information.
- `nvshmem.my_pe()`: Returns the process identifier (PE) of the current process.
- `nvshmem.n_pes()`: Returns the total number of processes in the communication domain.
```

--------------------------------

### NVSHMEM Best Practices Guide Abstract

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/index.html

Abstract of the NVSHMEM Best Practices Guide.

```APIDOC
## NVSHMEM Best Practices Guide Abstract

### Description
An abstract summarizing the content of the NVSHMEM Best Practices Guide.

### Endpoint
/websites/nvidia_nvshmem/release-notes-install-guide/best-practice-guide/abstract.html.md
```

--------------------------------

### Install and Use NVSHMRun with Hydra

Source: https://docs.nvidia.com/nvshmem/api/using.html

Install the Hydra Process Manager using the provided script to enable standalone NVSHMEM application development. The installed launcher is nvshmrun.hydra, with a symbolic link nvshmrun for convenience.

```bash
scripts/install_hydra.sh
```

--------------------------------

### NVSHMEM Get Operations

Source: https://docs.nvidia.com/nvshmem/api/gen/api/rma.html

Overview of the NVSHMEM get API functions for remote memory access.

```APIDOC
## NVSHMEM_GET

### Description
The get routines provide a method for copying a contiguous symmetric data object from a different PE to a contiguous data object on the local PE. The routines return after the data has been delivered to the dest array on the local PE.

### Parameters
- **dest** (void* / TYPE*) - Required - Symmetric address or host/device address of the data object to be updated.
- **source** (void* / TYPE*) - Required - Symmetric address of the source data object.
- **nelems** (size_t) - Required - Number of elements in the dest and source arrays. For nvshmem_getmem, elements are bytes.
- **pe** (int) - Required - PE number of the remote PE.
- **stream** (cudaStream_t) - Optional - CUDA stream for asynchronous operations (used in _on_stream variants).

### Functions
- nvshmem_TYPENAME_get
- nvshmemx_TYPENAME_get_on_stream
- nvshmem_TYPENAME_get (device)
- nvshmemx_TYPENAME_get_block (device)
- nvshmemx_TYPENAME_get_warp (device)
- nvshmem_getSIZE
- nvshmemx_getSIZE_on_stream
- nvshmem_getSIZE (device)
- nvshmemx_getSIZE_block (device)
- nvshmemx_getSIZE_warp (device)
- nvshmem_getmem
- nvshmemx_getmem_on_stream
- nvshmem_getmem (device)
- nvshmemx_getmem_block (device)
- nvshmemx_getmem_warp (device)

### Returns
- None
```

--------------------------------

### NVSHMEM RMA Get Immediate

Source: https://docs.nvidia.com/nvshmem/api/api/language_bindings/python/device/numba/rma.html

Gets immediate data from a source on a remote PE to a local destination. This is a device-initiated operation.

```APIDOC
## `nvshmem.core.device.numba.rma.g`

### Description
Get immediate data from `src` on PE `pe` to local `dst`. Device initiated.
`src` must be a scalar value, passed as a symmetric Array of size 1.

### Method
`nvshmem.core.device.numba.rma.g`

### Parameters
- **dst** (Array) - Local destination array on this PE.
- **src** (Array) - Source symmetric array (size-1) on remote PE.
- **pe** (int) - PE to get from.

### Returns
Scalar value retrieved from `src` on PE `pe`.
```

--------------------------------

### Performing Put with Signal

Source: https://docs.nvidia.com/nvshmem/api/api/language_bindings/python/rma.html

Demonstrates initializing symmetric memory, performing a put operation with a signal, and waiting for the signal to complete.

```python
import nvshmem.core as nvshmem
import cupy as cp
from cuda.core.experimental import Device

# Initialize NVSHMEM (initialization code not shown)

# Get current device and create a stream
device = Device()
stream = device.create_stream()

# Allocate symmetric memory
size = 10
src_array = nvshmem.array((size,), dtype=cp.float32)
dest_array = nvshmem.array((size,), dtype=cp.float32)

# Set values on each PE
my_pe = nvshmem.my_pe()
n_pes = nvshmem.n_pes()

# Fill source array with PE ID
src_array[:] = cp.ones(size, dtype=cp.float32) * my_pe

# Target PE to get data from (circular previous PE)
target_pe = (my_pe - 1 + n_pes) % n_pes
# Source PE (the opposite of the target PE)
src_pe = (my_pe + 1) % n_pes

# A signal is always
signal = nvshmem.core.array((1,), dtype="uint64")
signal [:] = 0
buf_sig, sz, type = nvshmem.core.array_get_buffer(signal)

# Put with signal to the target PE
nvshmem.core.put_signal(dest_array, src_array, buf_sig, 1, nvshmem.core.SignalOp.SIGNAL_SET, remote_pe=target_pe, stream=stream)

# Wait on signal from the source PE
nvshmem.core.signal_wait(buf_sig, 1, nvshmem.core.SignalOp.SIGNAL_SET, remote_pe=target_pe, stream=stream)

# Ensure operation is complete
stream.synchronize()

# Now dest_array contains data from the target PE

# Clean up - explicit free is required
nvshmem.core.free_array(src_array)
nvshmem.core.free_array(dest_array)
```

--------------------------------

### NVSHMEM Non-Blocking Get Operations

Source: https://docs.nvidia.com/nvshmem/api/gen/api/rma.html

API definitions for non-blocking get routines including standard, stream-based, block-level, and warp-level variants.

```APIDOC
## NVSHMEM Non-Blocking Get (nvshmem_TYPENAME_get_nbi)

### Description
Initiates a non-blocking copy of a contiguous symmetric data object from a remote PE to the local PE. The operation is completed after a call to nvshmem_quiet.

### Parameters
- **dest** (TYPE *) - Required - Local symmetric address or registered buffer to be updated.
- **source** (const TYPE *) - Required - Symmetric address of the remote source data object.
- **nelems** (size_t) - Required - Number of elements to copy.
- **pe** (int) - Required - PE number of the remote PE.
- **stream** (cudaStream_t) - Optional - CUDA stream for stream-based variants.

### Returns
- None
```

--------------------------------

### nvshmem.core.init()

Source: https://docs.nvidia.com/nvshmem/api/api/language_bindings/python/initialization.html

Initializes the NVSHMEM runtime with either MPI or UID-based bootstrapping. Supports multiple bootstrap methods to initialize the runtime.

```APIDOC
## POST /nvshmem/core/init

### Description
Initializes the NVSHMEM runtime with either MPI or UID-based bootstrapping. Supports multiple bootstrap methods to initialize the runtime.

### Method
POST

### Endpoint
/nvshmem/core/init

### Parameters
#### Query Parameters
- **initializer_method** (str) - Required - Specifies the initialization method. Must be either "mpi" or "uid".

#### Request Body
- **device** (cuda.core.Device) - Required - A Device() that will be bound to this process. All NVSHMEM operations on this process will use this Device.
- **uid** (nvshmem.UniqueID) - Optional - A unique identifier used for UID-based initialization. Must be provided if initializer_method is “uid”.
- **rank** (int) - Optional - Rank of the calling process in the NVSHMEM job. Required for UID-based init.
- **nranks** (int) - Optional - Total number of NVSHMEM ranks in the job. Required for UID-based init.
- **mpi_comm** (None) - Optional - MPI communicator to use for MPI-based initialization. Defaults to `MPI.COMM_WORLD` if `None` and `initializer_method` is “mpi”.

### Request Example
```json
{
  "device": "cuda:0",
  "initializer_method": "mpi",
  "mpi_comm": "MPI_COMM_WORLD"
}
```
```json
{
  "device": "cuda:0",
  "uid": "some_unique_id",
  "rank": 0,
  "nranks": 1,
  "initializer_method": "uid"
}
```

### Response
#### Success Response (200)
- **status** (None) - Indicates successful initialization.

#### Response Example
```json
{
  "status": null
}
```

### Errors
- **NvshmemInvalid**: If an invalid initialization method is provided, or required arguments for the selected method are missing or incorrect.
- **NvshmemError**: If NVSHMEM fails to initialize using the specified method.
```

--------------------------------

### Disable NVSHMEM Perftests and Examples with CMake

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

Control the build of perftests and examples by setting NVSHMEM_BUILD_TESTS and NVSHMEM_BUILD_EXAMPLES to 0 when configuring with CMake.

```bash
cmake -DNVSHMEM_BUILD_TESTS=0 -DNVSHMEM_BUILD_EXAMPLES=0
```

--------------------------------

### Well-defined nvshmemx_barrier_all_block with collective launch

Source: https://docs.nvidia.com/nvshmem/api/gen/api/collectives.html?highlight=allreduce

This example demonstrates a well-defined use of `nvshmemx_barrier_all_block` within a kernel launched using `nvshmemx_collective_launch`. It ensures that only one PE calls the barrier at a time, preventing issues.

```c++
__global__ void kernel() {
    if(blockIdx.x == nvshmem_my_pe()) {
        nvshmemx_barrier_all_block();
    }
}

gridDims grid(2, 1, 1);
blockDims(32, 1, 1);
vshmemx_collective_launch(kernel, grid, block, ...); // On 2 PEs

```

--------------------------------

### Install Specific NVSHMEM Version on Ubuntu

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

Installs a specific version of NVSHMEM runtime and development packages. Replace 'cuda-11_3.1.7-1' with the desired version.

```bash
sudo apt install libnvshmem3=cuda-11_3.1.7-1 libnvshmem3-dev=cuda-11_3.1.7-1
```

--------------------------------

### NVSHMEM Info Utility Usage

Source: https://docs.nvidia.com/nvshmem/api/faq.html

Demonstrates how to use the nvshmem-info utility to display help information. This utility provides details about NVSHMEM build and runtime configurations.

```bash
./build/src/bin/nvshmem-info -h
Print information about NVSHMEM
Usage: nvshmem-info [options]
Options:
    -h  This help message
    -a  Print all output
    -n  Print version number
    -b  Print build information
    -e  Print environment variables
    -d  Include hidden environment variables in output
    -r  RST format environment variable output
```

--------------------------------

### On-stream Kernels Example

Source: https://docs.nvidia.com/nvshmem/api/examples/language_bindings/python/index.html

Demonstrates GPU-centric communication using NVSHMEM on CUDA streams for reductions and collective operations. This is the Python version of the on-stream.cu example.

```python
"""
This is the Python version of the NVSHMEM on_stream.cu example.
This example demonstrates distributed GPU computing using NVSHMEM with CUDA streams,
showing how to perform parallel reductions and collective operations across multiple GPUs.
"""

import cupy
from numba import cuda
from cuda.core.experimental import Device, system

import nvshmem.core

from mpi4py import MPI

```

--------------------------------

### UID-Based Initialization Example

Source: https://docs.nvidia.com/nvshmem/api/index.html

Demonstrates initializing NVSHMEM4Py using a unique identifier (UID). This method is suitable when NVSHMEM is launched independently.

```python
import nvshmem.core as nc

# Initialize NVSHMEM using a UID
# nc.init_uid(uid='my_nvshmem_app')

# ... rest of the NVSHMEM operations ...

# nc.finalize()
```

--------------------------------

### Run NVSHMEM Job with nvshmrun

Source: https://docs.nvidia.com/nvshmem/release-notes-install-guide/install-guide/nvshmem-install-proc.html

After installing Hydra, use the 'nvshmrun' launcher to run your NVSHMEM job. This launcher is located in the 'bin/' directory of the Hydra installation path.

```bash
nvshmrun
```