NVSHMEM APIs

• Overview of the APIs
  • Unsupported OpenSHMEM 1.3 APIs
  • OpenSHMEM 1.3 APIs Not Supported over Remote Network Transports
  • Supported OpenSHMEM APIs (OpenSHMEM 1.4 and 1.5)
  • NVSHMEM API Extensions for CPU Threads
  • NVSHMEM API Extensions for GPU Threads
  • Tile-Granular Collective APIs
• Library Setup, Exit, and Query
  • NVSHMEM_INIT
    • nvshmem_init()
  • NVSHMEMX_INIT_ATTR
    • nvshmemx_init_attr()
  • NVSHMEMX_HOSTLIB_INIT_ATTR
    • nvshmemx_hostlib_init_attr()
  • NVSHMEMX_HOSTLIB_FINALIZE
    • nvshmemx_hostlib_finalize()
  • NVSHMEMX_GET_UNIQUE_ID
    • nvshmemx_get_uniqueid()
  • NVSHMEMX_SET_ATTR_UNIQUEID_ARGS
    • nvshmemx_set_attr_uniqueid_args()
  • NVSHMEMX_CUMODULE_INIT
    • nvshmemx_cumodule_init()
  • NVSHMEMX_INIT_STATUS
    • nvshmemx_init_status()
  • NVSHMEM_MY_PE
    • nvshmem_my_pe()
  • NVSHMEM_N_PES
    • nvshmem_n_pes()
  • NVSHMEM_FINALIZE
    • nvshmem_finalize()
  • NVSHMEM_GLOBAL_EXIT
    • nvshmem_global_exit()
  • NVSHMEM_PTR
    • nvshmem_ptr()
  • NVSHMEMX_MC_PTR
    • nvshmemx_mc_ptr()
  • NVSHMEM_INFO_GET_VERSION
    • nvshmem_info_get_version()
  • NVSHMEM_INFO_GET_NAME
    • nvshmem_info_get_name()
  • NVSHMEMX_VENDOR_GET_VERSION_INFO
    • nvshmemx_vendor_get_version_info()
• Thread Support
  • NVSHMEM_INIT_THREAD
    • nvshmem_init_thread()
  • NVSHMEM_QUERY_THREAD
    • nvshmem_query_thread()
• Kernel Launch Routines
  • NVSHMEMX_COLLECTIVE_LAUNCH
    • nvshmemx_collective_launch()
  • NVSHMEMX_COLLECTIVE_LAUNCH_QUERY_GRIDSIZE
    • nvshmemx_collective_launch_query_gridsize()
• Memory Management
  • NVSHMEM_MALLOC, NVSHMEM_FREE, NVSHMEM_ALIGN
    • nvshmem_malloc()
    • nvshmem_free()
    • nvshmem_align()
  • NVSHMEM_CALLOC
    • nvshmem_calloc()
  • Memory Registration
    • NVSHMEMX_BUFFER_REGISTER
      • nvshmemx_buffer_register()
    • NVSHMEMX_BUFFER_UNREGISTER
      • nvshmemx_buffer_unregister()
    • NVSHMEMX_BUFFER_UNREGISTER_ALL
      • nvshmemx_buffer_unregister_all()
    • NVSHMEMX_BUFFER_REGISTER_SYMMETRIC
      • nvshmemx_buffer_register_symmetric()
    • NVSHMEMX_BUFFER_REGISTER_SYMMETRIC_AT_PREFERRED_ADDRESS
      • nvshmemx_buffer_register_symmetric_at_preferred_address()
    • NVSHMEMX_BUFFER_UNREGISTER_SYMMETRIC
      • nvshmemx_buffer_unregister_symmetric()
• Queue Pair (QP) Specific APIs
  • Overview
  • QP Handle
  • Special QP Handle Values
  • Special PE Values for Synchronization
  • Special QP Values for Synchronization
  • API Categories
  • Device-Side Only
  • Ordering Semantics
  • Thread Safety
  • Transport Support and Compatibility
    • NVSHMEMX_QP_CREATE
      • nvshmemx_qp_create()
  • QP Remote Memory Access
    • NVSHMEMX_QP_PUT
    • NVSHMEMX_QP_P
    • NVSHMEMX_QP_GET
    • NVSHMEMX_QP_G
    • NVSHMEMX_QP_PUT_NBI
    • NVSHMEMX_QP_GET_NBI
  • QP Signaling Operations
    • NVSHMEMX_QP_SIGNAL_OP
    • NVSHMEMX_QP_PUT_SIGNAL
    • NVSHMEMX_QP_PUT_SIGNAL_NBI
  • QP Memory Ordering
    • NVSHMEMX_QP_FENCE
    • NVSHMEMX_QP_QUIET
• Team Management
  • Predefined and Application-Defined Teams
  • Team Handles
  • Thread Safety
  • Collective Ordering
  • Team Creation
    • Team Splitting
    • Arbitrary Team Initialization
  • NVSHMEM_TEAM_MY_PE
    • nvshmem_team_my_pe()
  • NVSHMEM_TEAM_N_PES
    • nvshmem_team_n_pes()
  • NVSHMEM_TEAM_CONFIG_T
  • NVSHMEM_TEAM_GET_CONFIG
    • nvshmem_team_get_config()
  • NVSHMEM_TEAM_TRANSLATE_PE
    • nvshmem_team_translate_pe()
  • NVSHMEM_TEAM_SPLIT_STRIDED
    • nvshmem_team_split_strided()
  • NVSHMEM_TEAM_SPLIT_2D
    • nvshmem_team_split_2d()
  • NVSHMEM_TEAM_DESTROY
    • nvshmem_team_destroy()
  • NVSHMEMX_TEAM_INIT
    • nvshmemx_team_init()
  • NVSHMEMX_TEAM_GET_UNIQUEID
    • nvshmemx_team_get_uniqueid()
• Remote Memory Access
  • Blocking RMA
    • NVSHMEM_PUT
      • nvshmem_TYPENAME_put()
      • nvshmemx_TYPENAME_put_on_stream()
      • nvshmem_putSIZE()
      • nvshmemx_putSIZE_on_stream()
      • nvshmem_putmem()
      • nvshmemx_putmem_on_stream()
    • NVSHMEM_P
      • nvshmem_TYPENAME_p()
    • NVSHMEM_IPUT
      • nvshmem_TYPENAME_iput()
      • nvshmemx_TYPENAME_iput_on_stream()
      • nvshmem_iputSIZE()
      • nvshmemx_iputSIZE_on_stream()
    • NVSHMEM_GET
      • nvshmem_TYPENAME_get()
      • nvshmemx_TYPENAME_get_on_stream()
      • nvshmem_getSIZE()
      • nvshmemx_getSIZE_on_stream()
      • nvshmem_getmem()
      • nvshmemx_getmem_on_stream()
    • NVSHMEM_G
      • nvshmem_TYPENAME_g()
    • NVSHMEM_IGET
      • nvshmem_TYPENAME_iget()
      • nvshmemx_TYPENAME_iget_on_stream()
      • nvshmem_igetSIZE()
      • nvshmemx_igetSIZE_on_stream()
  • Nonblocking RMA
    • NVSHMEM_PUT_NBI
      • nvshmem_TYPENAME_put_nbi()
      • nvshmemx_TYPENAME_put_nbi_on_stream()
      • nvshmem_putSIZE_nbi()
      • nvshmemx_putSIZE_nbi_on_stream()
      • nvshmem_putmem_nbi()
      • nvshmemx_putmem_nbi_on_stream()
    • NVSHMEM_GET_NBI
      • nvshmem_TYPENAME_get_nbi()
      • nvshmemx_TYPENAME_get_nbi_on_stream()
      • nvshmem_getSIZE_nbi()
      • nvshmemx_getSIZE_nbi_on_stream()
      • nvshmem_getmem_nbi()
      • nvshmemx_getmem_nbi_on_stream()
  • Tile RMA
    • TILE_PUT
    • TILE_GET
• Atomic Memory Operations
  • NVSHMEM_ATOMIC_FETCH
  • NVSHMEM_ATOMIC_SET
  • NVSHMEM_ATOMIC_COMPARE_SWAP
  • NVSHMEM_ATOMIC_SWAP
  • NVSHMEM_ATOMIC_FETCH_INC
  • NVSHMEM_ATOMIC_INC
  • NVSHMEM_ATOMIC_FETCH_ADD
  • NVSHMEM_ATOMIC_ADD
  • NVSHMEM_ATOMIC_FETCH_AND
  • NVSHMEM_ATOMIC_AND
  • NVSHMEM_ATOMIC_FETCH_OR
  • NVSHMEM_ATOMIC_OR
  • NVSHMEM_ATOMIC_FETCH_XOR
  • NVSHMEM_ATOMIC_XOR
• Signaling Operations
  • Atomicity Guarantees for Signaling Operations
  • Available Signal Operators
  • NVSHMEM_PUT_SIGNAL
    • nvshmem_TYPENAME_put_signal()
    • nvshmemx_TYPENAME_put_signal_on_stream()
    • nvshmem_putSIZE_signal()
    • nvshmemx_putSIZE_signal_on_stream()
    • nvshmem_putmem_signal()
    • nvshmemx_putmem_signal_on_stream()
  • NVSHMEM_PUT_SIGNAL_NBI
    • nvshmem_TYPENAME_put_signal_nbi()
    • nvshmemx_TYPENAME_put_signal_nbi_on_stream()
    • nvshmem_putSIZE_signal_nbi()
    • nvshmemx_putSIZE_signal_nbi_on_stream()
    • nvshmem_putmem_signal_nbi()
    • nvshmemx_putmem_signal_nbi_on_stream()
  • NVSHMEM_SIGNAL_FETCH
  • NVSHMEMX_SIGNAL
  • NVSHMEMX_SIGNAL_OP
• Collective Communication
  • Team-based collectives
  • Implicit team collectives
  • Tile-based Collectives
    • Tile helper functions
    • Tile collective algorithms
  • Error codes returned from team-based collectives
  • Collective operations scopes and active sets
    • nvshmem_barrier_all()
    • nvshmemx_barrier_all_on_stream()
  • NVSHMEM_BARRIER
    • nvshmem_barrier()
    • nvshmemx_barrier_on_stream()
  • NVSHMEM_SYNC
    • nvshmem_sync()
    • nvshmemx_sync_on_stream()
    • nvshmem_team_sync()
    • nvshmemx_team_sync_on_stream()
  • NVSHMEM_SYNC_ALL
    • nvshmem_sync_all()
    • nvshmemx_sync_all_on_stream()
  • NVSHMEM_ALLTOALL
    • nvshmem_TYPENAME_alltoall()
    • nvshmemx_TYPENAME_alltoall_on_stream()
    • nvshmem_alltoallmem()
    • nvshmemx_alltoallmem_on_stream()
  • NVSHMEM_BROADCAST
    • nvshmem_TYPENAME_broadcast()
    • nvshmemx_TYPENAME_broadcast_on_stream()
    • nvshmem_broadcastmem()
    • nvshmemx_broadcastmem_on_stream()
  • NVSHMEM_FCOLLECT
    • nvshmem_TYPENAME_fcollect()
    • nvshmemx_TYPENAME_fcollect_on_stream()
    • nvshmem_fcollectmem()
    • nvshmemx_fcollectmem_on_stream()
  • NVSHMEM_REDUCTIONS
    • AND
      • nvshmem_TYPENAME_and_reduce()
      • nvshmemx_TYPENAME_and_reduce_on_stream()
      • nvshmem_TYPENAME_and_reducescatter()
      • nvshmemx_TYPENAME_and_reducescatter_on_stream()
    • OR
      • nvshmem_TYPENAME_or_reduce()
      • nvshmemx_TYPENAME_or_reduce_on_stream()
      • nvshmem_TYPENAME_or_reducescatter()
      • nvshmemx_TYPENAME_or_reducescatter_on_stream()
    • XOR
      • nvshmem_TYPENAME_xor_reduce()
      • nvshmemx_TYPENAME_xor_reduce_on_stream()
      • nvshmem_TYPENAME_xor_reducescatter()
      • nvshmemx_TYPENAME_xor_reducescatter_on_stream()
    • MAX
      • nvshmem_TYPENAME_max_reduce()
      • nvshmemx_TYPENAME_max_reduce_on_stream()
      • nvshmem_TYPENAME_max_reducescatter()
      • nvshmemx_TYPENAME_max_reducescatter_on_stream()
    • MIN
      • nvshmem_TYPENAME_min_reduce()
      • nvshmemx_TYPENAME_min_reduce_on_stream()
      • nvshmem_TYPENAME_min_reducescatter()
      • nvshmemx_TYPENAME_min_reducescatter_on_stream()
    • SUM
      • nvshmem_TYPENAME_sum_reduce()
      • nvshmemx_TYPENAME_sum_reduce_on_stream()
      • nvshmem_TYPENAME_sum_reducescatter()
      • nvshmemx_TYPENAME_sum_reducescatter_on_stream()
    • PROD
      • nvshmem_TYPENAME_prod_reduce()
      • nvshmemx_TYPENAME_prod_reduce_on_stream()
      • nvshmem_TYPENAME_prod_reducescatter()
      • nvshmemx_TYPENAME_prod_reducescatter_on_stream()
  • TILE_REDUCTIONS
  • TILE_ALLGATHER
  • TILE_BROADCAST
  • TILE_WAIT
• Point-To-Point Synchronization
  • NVSHMEM_WAIT_UNTIL
    • nvshmemx_TYPENAME_wait_until_on_stream()
    • nvshmemx_TYPENAME_wait_on_stream()
  • NVSHMEM_WAIT_UNTIL_ALL
  • NVSHMEM_WAIT_UNTIL_ANY
  • NVSHMEM_WAIT_UNTIL_SOME
  • NVSHMEM_WAIT_UNTIL_ALL_VECTOR
  • NVSHMEM_WAIT_UNTIL_ANY_VECTOR
  • NVSHMEM_WAIT_UNTIL_SOME_VECTOR
  • NVSHMEM_TEST
  • NVSHMEM_TEST_ALL
  • NVSHMEM_TEST_ANY
  • NVSHMEM_TEST_SOME
  • NVSHMEM_TEST_ALL_VECTOR
  • NVSHMEM_TEST_ANY_VECTOR
  • NVSHMEM_TEST_SOME_VECTOR
  • NVSHMEM_SIGNAL_WAIT_UNTIL
    • nvshmemx_signal_wait_until_on_stream()
• Memory Ordering
  • NVSHMEM_FENCE
    • nvshmem_fence()
  • NVSHMEM_QUIET
    • nvshmem_quiet()
    • nvshmemx_quiet_on_stream()
  • NVSHMEMX_FLUSH
    • nvshmemx_flush_on_stream()
• Language Bindings
  • Python Bindings (NVSHMEM4Py)
    • NVSHMEM4Py Overview
      • Current API Support
      • Key Features
      • Usage Model
      • Limitations
      • Compatibility Guide
    • Initialization and Finalization
      • Examples
      • API Reference
      • Teams
        • Teams
          • TEAM_WORLD
          • TEAM_SHARED
          • TEAM_NODE
          • TEAM_SAME_MYPE_NODE
          • TEAM_SAME_GPU
          • TEAM_GPU_LEADERS
      • Initialization Methods
      • Querying Initialization Status
        • InitStatus
          • STATUS_NOT_INITIALIZED
          • STATUS_IS_BOOTSTRAPPED
          • STATUS_IS_INITIALIZED
          • STATUS_LIMITED_MPG
          • STATUS_FULL_MPG
          • STATUS_INVALID
      • Finalization
      • Retrieving Version Information
      • Retrieving a Unique ID
    • Memory Management
      • Symmetric Memory in Python
      • Memory Lifecycle
      • Memory API Reference
    • Interoperability
      • Interoperability with PyTorch and CuPy
        • Creating NVSHMEM-Backed Arrays
        • NVSHMEM Operations with PyTorch and CuPy
        • Memory Management
      • Interoperability API Reference
        • Interoperability with PyTorch
        • Interoperability with CuPy
    • Collective Operations
      • Supported Collective Operations
        • Reduce
        • Broadcast
        • ReduceScatter
        • FCollect
        • AlltoAll
      • Stream Requirement for Collectives
      • Collective Examples
      • API Reference
    • Remote Memory Access (RMA)
      • RMA Operations in NVSHMEM4Py
      • Supported RMA Operations
      • Stream Requirement for RMA
      • Memory Management for RMA
      • RMA Examples
      • RMA API Reference
        • SignalOp
          • SIGNAL_SET
          • SIGNAL_ADD
    • Utility Functions for NVSHMEM4Py
    • Python Device APIs
      • Python Device APIs for Numba-CUDA DSL
        • NVSHMEM Device Collectives with Numba-CUDA DSL
          • Using barrier and broadcast in a Numba-CUDA Kernel
        • NVSHMEM Device Remote Memory Access (RMA) with Numba-CUDA DSL
          • Using put and get in a Numba-CUDA Kernel
        • NVSHMEM4Py Memory Management with Numba-CUDA DSL
        • NVSHMEM Device Atomic Memory Operations with Numba-CUDA DSL
          • Using atomic_add in a Numba-CUDA Kernel
        • Available APIs
      • Python Device APIs for CuTe DSL
        • NVSHMEM Device Collectives with CuTe DSL
          • Using barrier_all and broadcast in a CuTe Kernel
        • NVSHMEM Device Remote Memory Access (RMA) with CuTe DSL
          • Using put and get in a CuTe Kernel
        • NVSHMEM4Py Memory Management with CuTe DSL
        • NVSHMEM Device Atomic Memory Operations with CuTe DSL
          • Using atomic_add in a CuTe Kernel
        • Available APIs
    • Quick Start
    • Key Features