Nvshmem Get Started
NVIDIAOfficialCreated by NVIDIA NVSHMEM Team
- Skill
- Tool integration
- Software development
- Apache-2.0
- Medium risk
AutoGineer summary
Written by AutoGineer from the skill’s own files, not by the author
This skill provides a structured onboarding path for developers new to NVSHMEM. It focuses on teaching the core mental model, including symmetric memory and one-sided communication, before guiding users through the lifecycle of C/C++ or Python NVSHMEM programs. It emphasizes best practices for initialization, memory management, and synchronization, while providing adaptable templates for compilation and execution. The skill acts as a central hub, routing complex installation, version-specific, or troubleshooting requests to specialized documentation and support tools.
What it does
- Assesses whether NVSHMEM is a suitable communication model for a given application.
- Explains the NVSHMEM mental model, including PEs, symmetric memory, and one-sided operations.
- Provides standard patterns for initializing, allocating symmetric memory, and finalizing programs.
- Offers adaptable command templates for compiling and linking C/C++ applications.
- Guides users through the launch workflow for multi-PE programs.
- Routes technical failures and version-specific queries to appropriate documentation or support tools.
When to use it
- When you are new to NVSHMEM and need a conceptual overview.
- When you need to determine if NVSHMEM fits your multi-GPU communication requirements.
- When you are ready to design, compile, or run your first NVSHMEM C/C++ or Python program.
- When you need a standard template for building or launching an NVSHMEM-enabled application.
Advantages
- Provides a clear, step-by-step learning path for beginners.
- Focuses on conceptual understanding before introducing API-specific details.
- Supplies reusable, adaptable command shapes for common development tasks.
- Prevents common pitfalls by emphasizing correct symmetric memory and synchronization patterns.
- Maintains a clean separation between conceptual guidance and environment-specific troubleshooting.
Usage examples
- “Does NVSHMEM fit my multi-GPU stencil application?”
- The agent will ask for specific communication details and provide an assessment of whether NVSHMEM is a strong fit, a possible fit, or if another model like MPI or NCCL is simpler.
- “Help me design my first C++ NVSHMEM program.”
- The agent will explain the mental model, outline the program lifecycle, and provide a code structure for initialization, symmetric allocation, and communication.
- “How do I compile a C++ NVSHMEM application?”
- The agent will provide a standard nvcc command template, explaining the necessary flags for relocatable device code and library linking.
- “How do I run my first NVSHMEM4Py program?”
- The agent will collect environment details, provide a launch command shape, and explain the requirements for running the program across multiple PEs.
Install
Download the skill as a ZIP:
https://autogineer.com/agent-skills/nvidia/nvshmem-get-started/downloadUnzip it into the cross-client skills folder of your project:
mkdir -p .agents/skills && unzip nvshmem-get-started.zip -d .agents/skills
Documentation by NVIDIA NVSHMEM Team
Getting Started with NVSHMEM
Purpose
Guide a beginner to the next useful action. Teach the stable mental model, program structure, build pattern, and launch pattern directly. Route installation, version-sensitive variants, advanced integrations, and failures to the appropriate specialist skill.
Prerequisites
Require no credentials, installed software, or GPU access for fit assessment and conceptual guidance.
Keep the bundled mental model, basic API calls, and NVSHMEM4Py program guide available for the teaching paths that require them.
Keep
$nvshmem-install,$nvshmem-docs, and$nvshmem-troubleshoot-and-report-bugsavailable for complete installation, version-specific, and failure handoffs. If a specialist skill is unavailable, follow the documented fallback instead of inventing current details.
Inputs
Required: Obtain the user's goal or requested learning stage from the explicit user prompt.
Optional: Obtain application communication needs, language, NVSHMEM and NVSHMEM4Py versions, source or script path, Python interpreter/environment, installation prefix, target GPU architecture, launcher, PE and node counts, and target environment only when the selected path needs them.
Source precedence: Prefer the current explicit user prompt, then explicit invocation arguments, then non-conflicting agent context, then answers to focused follow-up questions. Report conflicting values instead of choosing silently.
Instructions
Execution Boundary
Default to explaining the workflow and printing commands for the user to run. Do not compile, launch a program, submit a scheduler job, allocate resources, or modify the environment unless the user explicitly asks for execution. Treat requests to learn, get started, or receive help running a program as requests for guidance, not execution approval.
Choose a Starting Point
When the requested stage is clear, skip the menu and handle that stage. Otherwise ask what the user wants to do and present all options in one response:
Decide whether NVSHMEM fits my use case.
Learn the NVSHMEM mental model.
Install NVSHMEM.
Run a precompiled NVSHMEM example.
Compile or link a C/C++ application.
Write an NVSHMEM C/C++ or NVSHMEM4Py program.
Recommend the next step.
For option 7, first determine whether NVSHMEM fits the intended communication pattern. Then ask whether NVSHMEM is installed, whether an example has run successfully, and whether source code already exists. Recommend the mental model when the user is completely new, installation when NVSHMEM is absent, a precompiled example after installation, program design when no source exists, and compilation when source is ready.
Decide When to Use NVSHMEM
Start from the application's communication needs rather than from a list of APIs. Describe NVSHMEM as a strong fit when several of these conditions apply:
GPU threads should initiate communication directly from CUDA kernels without returning control to the CPU for every exchange.
The application needs fine-grained, one-sided put, get, atomic, signaling, or synchronization operations between GPUs.
Long-running or fused kernels should combine computation and communication while reducing kernel-launch and CPU-GPU synchronization overhead.
Communication should overlap with computation to improve strong scaling as the amount of work per GPU becomes smaller.
A partitioned global address space with collectively allocated symmetric GPU memory fits the application's data layout.
The application already uses OpenSHMEM concepts or needs NVSHMEM alongside MPI or OpenSHMEM for GPU-resident communication.
Give representative workload patterns without claiming that NVSHMEM is mandatory for them:
neighbor or halo exchanges in multi-GPU stencils, solvers, and simulations;
irregular or fine-grained communication in graph and data-analytics workloads;
persistent GPU kernels that communicate repeatedly during computation;
custom GPU-initiated reductions, broadcasts, queues, and producer-consumer protocols;
applications whose strong scaling is limited by repeated CPU orchestration of GPU communication.
Explain when another model may be simpler:
Use ordinary CUDA for a single GPU or when no inter-GPU communication is required.
Prefer NCCL when the application primarily makes standard bulk-collective calls—especially host-issued all-reduce, all-gather, reduce-scatter, or broadcast operations. NCCL is generally better suited than NVSHMEM for this pattern.
Prefer MPI when communication is coarse-grained, CPU-driven send and receive is sufficient, or symmetric allocation would complicate the design.
Reconsider NVSHMEM when the target cannot provide the required GPU, launcher, bootstrap, or transport environment.
First determine whether the prompt already contains enough communication detail to assess fit. If it does, immediately conclude with one of strong fit, possible fit, or likely simpler with another model and explain the deciding factors; do not defer the recommendation behind a checklist of questions. Otherwise ask only for the missing decision-critical facts: where communication is initiated, whether it is fine- or coarse-grained, which operations dominate, whether communication must occur inside kernels, and what GPU/node topology is targeted. Then give the same explicit conclusion. Link to the official NVSHMEM Introduction and Advantages.
Teach the Mental Model
Read references/mental-model.md completely before teaching the conceptual path. Explain the mental model without API-name clutter and use its illustrative ring shift when a concrete example would help. Before introducing C/C++ primitives, read references/basic-api-calls.md completely and present only the API groups relevant to the user's goal. For a Python request, also read references/nvshmem4py-programs.md before outlining the program lifecycle or launch workflow. Keep examples framed as teaching aids rather than environment-specific recipes.
Resolve the Hands-On Path
Collect only the context required by the selected path. Do not ask for every environment detail up front.
Install NVSHMEM
Before invoking $nvshmem-install, follow this intake order: resolve the requested version or latest, ask whether the user needs C/C++, Python, or both, then at step 3 ask whether the user is an administrator on the target system with root or sudo rights. On a managed cluster, also ask whether the user is authorized to change the compute-node image or shared software stack. Do not infer either answer from shell access, package-manager availability, or target type.
Invoke $nvshmem-install. Pass the administrator and target-scope authority status together with the requested API or language, target environment, preferred install method, version policy, and whether the user needs runtime files, development headers and libraries, or static linking. For an unprivileged cluster user, do not frame DEB, RPM, APT, DNF, YUM, or another system package path as directly usable; prefer a site module, approved container, compatible user-owned prefix, or a clearly labeled administrator handoff.
Design a C/C++ Program
Read references/basic-api-calls.md before explaining API-family choices or recommending specific primitives.
Teach this stable structure directly:
Choose initialization. Use
nvshmem_initfor the basic standalone pattern. Usenvshmemx_init_attrwhen the application must integrate an MPI, OpenSHMEM, or other bootstrap context.Initialize, query a node-local PE ID, and select the CUDA device before allocating symmetric memory or launching kernels.
Allocate symmetric objects collectively. Require every PE to call allocation and deallocation routines in matching order with identical arguments.
Choose the operation family:
Use put or get for remote data movement.
Use atomics for remote updates.
Use signals and waits for point-to-point coordination.
Use team collectives for group communication.
Choose the invocation context: host API, device API inside a kernel, or on-stream API composed with CUDA stream work.
Identify remote data with a symmetric address and destination PE. Never send a symmetric pointer value to another PE for it to reuse as a local pointer.
Add the required ordering and completion. Use
nvshmem_fenceto order later operations after earlier operations to the same PE,nvshmem_quietor its on-stream form for completion, and barriers, signals and waits, or collectives to coordinate PEs.Free symmetric allocations collectively and finalize.
Use the ring-shift example in references/mental-model.md as the first concrete illustration; read that reference before presenting the example.
Design or Run an NVSHMEM4Py Program
Read references/nvshmem4py-programs.md completely. Ask only for the Python script path, Python interpreter or environment, NVSHMEM4Py version if known, initialization choice (standalone or MPI), launcher, PE and node counts, and whether the run is local or scheduled.
Teach the same PE, symmetric-memory, one-sided-operation, and completion model as C/C++, then map it to the Python lifecycle: choose and set the local CUDA device, initialize NVSHMEM4Py, collectively create symmetric arrays, perform the required communication and synchronization, free every symmetric array collectively, and finalize. Do not present C/C++ API names as Python APIs.
Invoke $nvshmem-docs with the requested NVSHMEM version or latest before giving Python API names, a runnable code sample, package-version compatibility, framework interoperability, initialization arguments, or exact launch commands. Pass the script path and all known Python, launcher, PE, node, and environment details. Use the documentation result to give a version-matched, user-runnable command shape. Confirm that the selected Python environment, the script, CUDA, NVSHMEM runtime libraries, and the NVSHMEM4Py package are available on every participating node.
When the user supplies a simple NVSHMEM4Py script, use the example in references/nvshmem4py-programs.md to trace its device selection, initialization, symmetric allocation, communication, stream completion, collective cleanup, and launch requirements. Identify missing lifecycle steps or a mismatch between the script's initialization and launcher; do not execute or rewrite the script unless requested.
Start with two PEs on one node and one PE per GPU by default. Require the launcher to propagate the selected Python environment and CUDA/NVSHMEM environment to every PE. Explain the expected per-PE output and validate the small run before suggesting a larger or multi-node run.
Compile and Link C/C++
Ask for the source path, NVSHMEM installation prefix, target GPU architecture, and whether the project uses a direct nvcc command or CMake. For a standard CUDA source file, explain that the application must enable relocatable device code, include the NVSHMEM headers, search the installed library directory, and link both the host and device libraries.
Provide this command shape after replacing or clearly labeling its variables:
nvcc -rdc=true \
-ccbin g++ \
-gencode="$NVCC_GENCODE" \
-I"$NVSHMEM_HOME/include" \
application.cu \
-L"$NVSHMEM_HOME/lib" \
-lnvshmem_host -lnvshmem_device \
-o applicationExplain that NVSHMEM_HOME names the installed prefix and NVCC_GENCODE must select the target GPU architecture. For CMake, prefer the config files installed with NVSHMEM and use find_package in config mode; inspect the installed package before naming an imported target.
After the build command, state that the resulting NVSHMEM binary is still a multi-PE program and must be launched with nvshmrun, mpirun, or srun at a chosen PE count, typically one PE per GPU for the first run. Point the user to the launcher command shapes in Run a Precompiled Program; do not leave nvcc or the CMake build as the final step.
Run a Precompiled Program
Ask for the executable path, launcher, PE and node counts, and whether the run is local or scheduled. Confirm that the executable and NVSHMEM runtime libraries are visible on every participating node.
Start with two PEs on one node, use one PE per GPU by default, and expand only after the small run succeeds. Present the command shape that matches the available launcher:
mpirun -n <number-of-pes> <program> [arguments]srun -n <number-of-pes> <program> [arguments]nvshmrun -n <number-of-pes> <program> [arguments]Require the launcher to propagate the CUDA and NVSHMEM environment to every PE. Explain the expected per-PE output and verify it before recommending a larger or multi-node run.
Invoke Documentation Only for Variants
Invoke $nvshmem-docs with a complete requested version or latest when the request needs any of the following:
exact package-specific example paths or binary names;
PMI, PMI-2, PMIx, Hydra installation, or non-default bootstrap wiring;
multi-node transport setup or transport-specific link dependencies;
static linking, shared-library composition, symbol visibility, or non-
nvcctoolchains;exact CMake package or imported-target names;
release-specific API availability, flags, requirements, limitations, or commands.
Limitations
Provide introductory guidance, not exhaustive API, installation, launcher, transport, or performance coverage.
Treat command blocks as adaptable shapes. Do not claim that they match an installed release, package layout, cluster policy, or target architecture without version-specific documentation and target inspection.
Cover first-program and launch guidance for C/C++ CUDA and NVSHMEM4Py. Route Python API details, installation, and version compatibility to
$nvshmem-docsrather than inferring them.Assess architectural fit from the stated communication pattern; do not promise a speedup without representative measurements on the target topology.
Keep execution outside the default scope. A request for explanation, onboarding, or help running a program does not authorize environment changes or workload execution.
Troubleshooting
| Symptom | Likely cause class | Response |
|---|---|---|
| Compilation or linking fails | Source, flags, library order, package layout, or version mismatch | Invoke $nvshmem-troubleshoot-and-report-bugs with the full command and diagnostics. |
| A program errors or hangs | Launcher, bootstrap, transport, topology, runtime configuration, or program behavior | Invoke $nvshmem-troubleshoot-and-report-bugs with the complete launch stanza and output from every PE. |
| A specialist skill is unavailable | The normal diagnostic or version-specific route cannot be used | Invoke $nvshmem-docs for version-matched troubleshooting, FAQ, and support routes; do not invent a workaround. |
For an execution failure, pass or collect the complete command or batch-script launch stanza and complete stdout and stderr from every PE, including output before the first error and per-PE logs. For a compile or link failure, pass or collect the complete compiler or linker command, full diagnostic output, undefined symbols, and library order. In both cases, include the NVSHMEM version, relevant CUDA and launcher versions, and target environment details when available.
Keep installation and program-writing requests on their normal routes unless the user explicitly reports a failure in one of the two stages above.
Examples
Fit assessment:
Use $nvshmem-get-started to assess whether a multi-GPU stencil with GPU-resident halo exchanges is a strong NVSHMEM fit or is simpler with MPI.Ask only for communication facts that affect the decision, then conclude with strong fit, possible fit, or likely simpler with another model and explain the deciding factors.
First-program guidance:
Use $nvshmem-get-started to help me design, compile, and launch my first C++ NVSHMEM ring program without executing anything.Teach the mental model before concrete primitives, present adaptable compile and launch shapes, and keep execution outside scope unless the user separately requests it.
NVSHMEM4Py first-run guidance:
Use $nvshmem-get-started to help me run my first NVSHMEM4Py program with two GPUs, without executing anything.Read the NVSHMEM4Py program guide, collect the Python environment and launcher context, obtain current version-matched Python API and launch details from $nvshmem-docs, and provide a user-runnable small-run plan.
<!-- ci trigger -->
Files
- ATTRIBUTION.mdAdded by AutoGineer928 B
- BENCHMARK.md8.0 KiB
- LICENSELicence file added by AutoGineer11.2 KiB
- SKILL.md17.1 KiB
- evals/evals.json5.4 KiB
- references/basic-api-calls.md12.3 KiB
- references/mental-model.md4.0 KiB
- references/nvshmem4py-programs.md5.3 KiB
- skill-card.md4.5 KiB
- skill.oms.sig5.2 KiB
references/basic-api-calls.md
Basic NVSHMEM API Calls
Use this reference after the conceptual mental model is understood or when a user directly asks which C/C++ NVSHMEM primitives to use. Present only the API groups relevant to the user's goal, and treat signatures containing TYPENAME, TYPE, SIZE, or OP as naming patterns rather than literal function names.s
Contents
Setup, Exit, and Query
Symmetric Memory
Remote Memory Access
Invocation Contexts
Blocking and Nonblocking RMA
Signaling and Atomics
Ordering and Completion
Point-to-Point Synchronization
Teams and Collectives
Collective Kernel Launch
Setup, Exit, and Query
Start with the core and extension headers and the job-lifecycle calls:
#include <nvshmem.h>
#include <nvshmemx.h>
void nvshmem_init(void);
void nvshmem_finalize(void);
int nvshmem_my_pe(void);
int nvshmem_n_pes(void);
void *nvshmem_ptr(const void *dest, int pe);
void nvshmem_info_get_version(int *major, int *minor);Use
nvshmem_initandnvshmem_finalizecollectively to establish and release the PE job.Use
nvshmem_my_peandnvshmem_n_pesto query the world-team rank and size.Use
nvshmem_ptronly as an optional direct-access optimization. It returns a directly usable pointer when the peer object is accessible and null otherwise; the result is not a symmetric address for use in other NVSHMEM calls.
For an application that already has an MPI communicator, initialize MPI first and use attribute-based initialization instead of nvshmem_init:
nvshmemx_init_attr_t attr = NVSHMEMX_INIT_ATTR_INITIALIZER;
MPI_Comm comm = MPI_COMM_WORLD;
attr.mpi_comm = &comm;
nvshmemx_init_attr(NVSHMEMX_INIT_WITH_MPI_COMM, &attr);Teach both supported device-selection patterns:
If a launcher-provided node-local rank is available before initialization, call
cudaSetDevice(local_rank)before the initialization routine.Otherwise use two-stage initialization: call the initialization routine, query a node-local PE with
nvshmem_team_my_pe(NVSHMEMX_TEAM_NODE), callcudaSetDevice, and only then call an operation such as symmetric allocation that completes device initialization.
Do not call general NVSHMEM operations between two-stage initialization and device selection; only use the documented setup and PE/team query routines needed to select the device.
Symmetric Memory
Introduce the collectively managed symmetric heap:
void *nvshmem_malloc(size_t size);
void *nvshmem_calloc(size_t count, size_t size);
void *nvshmem_align(size_t alignment, size_t size);
void nvshmem_free(void *ptr);Have every PE call allocation and deallocation in the same order with identical allocation arguments. Use the returned local symmetric address together with a destination PE to name the corresponding remote object. Require the remote operand of RMA, AMO, signal, wait, and collective operations to be symmetric or provided through another explicitly documented symmetric-memory mechanism; ordinary cudaMalloc memory is not automatically a remote symmetric object.
Mention NVSHMEM_SYMMETRIC_SIZE only for configurations that use a statically sized symmetric heap, and retrieve the exact release-specific behavior before recommending a value.
Remote Memory Access
Explain the three common naming families: TYPENAME for typed elements such as int or float, SIZE for fixed-width transfers, and mem for byte counts.
void nvshmem_TYPENAME_put(TYPE *dest, const TYPE *source,
size_t nelems, int pe);
void nvshmem_TYPENAME_get(TYPE *dest, const TYPE *source,
size_t nelems, int pe);
void nvshmem_putSIZE(void *dest, const void *source,
size_t nelems, int pe);
void nvshmem_putmem(void *dest, const void *source,
size_t nbytes, int pe);
void nvshmem_getmem(void *dest, const void *source,
size_t nbytes, int pe);
void nvshmem_TYPENAME_p(TYPE *dest, TYPE value, int pe);
TYPE nvshmem_TYPENAME_g(const TYPE *source, int pe);
void nvshmem_TYPENAME_iput(TYPE *dest, const TYPE *source,
ptrdiff_t dst_stride, ptrdiff_t src_stride,
size_t nelems, int pe);
void nvshmem_TYPENAME_iget(TYPE *dest, const TYPE *source,
ptrdiff_t dst_stride, ptrdiff_t src_stride,
size_t nelems, int pe);Use put to copy local data to a remote symmetric destination, get to copy from a remote symmetric source into local storage, p/g for a scalar, and iput/iget for strided data.
Invocation Contexts
Distinguish where an operation is issued:
| Form | Invocation context | Meaning |
|---|---|---|
nvshmem_* host form | CPU code | Issue the documented host operation directly. |
nvshmem_* device form | CUDA kernel | Let a GPU thread initiate the operation. |
nvshmemx_*_on_stream | CPU code with a CUDA stream | Enqueue a supported operation in stream order; it is asynchronous with respect to the host. |
nvshmemx_*_warp / nvshmemx_*_block | CUDA kernel | Have all threads in the warp or block cooperate on one operation. |
Require every participating thread in a cooperative warp/block call to use identical arguments. Recommend cooperative forms for sufficiently large kernel-initiated transfers, but verify availability for the chosen operation and release.
Blocking and Nonblocking RMA
A blocking put makes its local source reusable when the call returns but does not by itself synchronize a remote consumer with the update.
A blocking get or fetching operation returns its requested result to the caller, subject to NVSHMEM ordering rules.
An
_nbioperation initiates work and requires a later completion operation before relying on completion or reusing affected local buffers.
nvshmem_putmem_nbi(dest, source, nbytes, pe);
/* issue more independent work */
nvshmem_quiet();Use _nbi variants to pipeline independent transfers, then complete the batch with quiet in the same issuing context.
Signaling and Atomics
Use put-with-signal for a producer/consumer transfer whose data delivery is paired with an atomic update to a symmetric uint64_t signal object:
void nvshmem_TYPENAME_put_signal(
TYPE *dest, const TYPE *source, size_t nelems,
uint64_t *sig_addr, uint64_t signal, int sig_op, int pe);
__device__ uint64_t nvshmem_signal_wait_until(
uint64_t *sig_addr, int cmp, uint64_t cmp_value);Use NVSHMEM_SIGNAL_SET to replace the signal value or NVSHMEM_SIGNAL_ADD to add to it. The signal establishes delivery for the data coupled to that put; do not assume it completes unrelated earlier operations without the required ordering operation.
Use AMOs for remote counters, work queues, ownership, or dynamic scheduling on symmetric objects:
TYPE nvshmem_TYPENAME_atomic_fetch_add(TYPE *dest, TYPE value, int pe);
void nvshmem_TYPENAME_atomic_add(TYPE *dest, TYPE value, int pe);
TYPE nvshmem_TYPENAME_atomic_compare_swap(
TYPE *dest, TYPE cond, TYPE value, int pe);
TYPE nvshmem_TYPENAME_atomic_fetch_inc(TYPE *dest, int pe);Explain fetching versus non-fetching forms and verify datatype, invocation-context, and transport support for the target release.
Ordering and Completion
Keep the guarantees separate:
| Call | Guarantee |
|---|---|
nvshmem_fence() | Order earlier operations before later operations to the same destination PE; do not wait for completion. |
nvshmem_quiet() | Complete previously issued operations on symmetric objects from the calling PE in the matching host or device context; do not synchronize with another PE. |
nvshmemx_quiet_on_stream(stream) | Enqueue completion of prior device-side operations in stream order. |
Use fence when only per-destination ordering is needed. Use quiet when completion or safe local-buffer reuse is required. Follow either with the appropriate signal/wait or collective operation when a consumer must coordinate with the producer.
Point-to-Point Synchronization
Use wait/test operations on a local symmetric object that another PE updates:
void nvshmem_TYPENAME_wait_until(
TYPE *ivar, int cmp, TYPE cmp_value);
int nvshmem_TYPENAME_test(
TYPE *ivar, int cmp, TYPE cmp_value);Use comparisons such as NVSHMEM_CMP_EQ, NVSHMEM_CMP_NE, NVSHMEM_CMP_GT, NVSHMEM_CMP_GE, NVSHMEM_CMP_LT, and NVSHMEM_CMP_LE. Explain that wait blocks while test polls, and mention the documented _all, _any, _some, and vector families only when the user needs array-wide conditions. Pair signal objects with nvshmem_signal_wait_until or its documented on-stream form.
Teams and Collectives
Describe nvshmem_team_t as a PE subset and collective-ordering domain. Use NVSHMEM_TEAM_WORLD for all PEs and NVSHMEMX_TEAM_NODE for the predefined node-local team. Query, create, translate, and destroy application teams with:
int nvshmem_team_my_pe(nvshmem_team_t team);
int nvshmem_team_n_pes(nvshmem_team_t team);
int nvshmem_team_translate_pe(
nvshmem_team_t src_team, int src_pe, nvshmem_team_t dest_team);
int nvshmem_team_split_strided(/* parent and split arguments */);
int nvshmem_team_split_2d(/* parent and split arguments */);
void nvshmem_team_destroy(nvshmem_team_t team);Treat team creation as collective on the parent team and require team collectives to occur in the same program order on every member.
Introduce synchronization and representative data collectives:
void nvshmem_barrier(nvshmem_team_t team);
void nvshmem_barrier_all(void);
int nvshmem_sync(nvshmem_team_t team);
void nvshmem_sync_all(void);
int nvshmem_TYPENAME_broadcast(
nvshmem_team_t team, TYPE *dest, const TYPE *source,
size_t nelems, int root);
int nvshmem_TYPENAME_fcollect(
nvshmem_team_t team, TYPE *dest, const TYPE *source,
size_t nelems);
int nvshmem_TYPENAME_alltoall(
nvshmem_team_t team, TYPE *dest, const TYPE *source,
size_t nelems);
int nvshmem_TYPENAME_OP_reduce(
nvshmem_team_t team, TYPE *dest, const TYPE *source,
size_t nelems);Explain that barrier is a rendezvous that also completes prior remote updates from participants, whereas sync is a rendezvous that does not complete prior NVSHMEM RMA or AMO updates. Present fcollect as an all-gather-like operation and OP_reduce for operations such as sum, product, minimum, maximum, and bitwise reductions.
Although these calls exist, explicitly prefer NCCL when an application is dominated by standard bulk-collective calls, especially host-issued all-reduce, all-gather, reduce-scatter, or broadcast operations.
Collective Kernel Launch
Require collective launch when a CUDA kernel itself invokes NVSHMEM synchronization or collective APIs such as wait, barrier, or a collective operation:
void *args[] = {&arg0, &arg1};
nvshmemx_collective_launch(
(const void *)kernel, grid, block, args,
dynamic_shared_bytes, stream);Explain that this launch is collective across PEs and uses cooperative launch to provide forward progress. Check the supported grid size with nvshmemx_collective_launch_query_gridsize. Ordinary kernel launch remains valid for a kernel that only issues non-synchronizing NVSHMEM operations, such as the scalar put in the ring-shift example below.
Use the official API index, setup and query reference, memory management reference, RMA reference, signaling reference, point-to-point synchronization reference, memory ordering reference, collectives reference, teams reference, and kernel launch reference when exact signatures or version-specific availability matter.
Requirements and compatibility
Works with: Claude Code, Claude apps, OpenAI Codex, GitHub Copilot, Cursor, Gemini CLI, Kiro, OpenCode, Goose, Amp, Roo Code, JetBrains Junie, Factory Droid, Devin (ex-Windsurf), Mistral Vibe, OpenHands, Letta
“Works with” means the tool reads the open Agent Skills format; a skill written for one tool may use features others lack.
Needs: Needs Node.js
Security
What our scanner noticed· Medium risk
- Large base64-encoded blobs
A static scan is not a guarantee. Skills are third-party content: review a skill before you let an agent use it.
Attribution and licence
- Created by
- NVIDIA NVSHMEM Team
- Publisher
- NVIDIA (Official)
- Source
- NVIDIA/skills · skills/nvshmem-get-started · commit fce3cca
- Licence
- Apache License 2.0 (redistribution allowed)
- LICENSE (added by AutoGineer from LICENSE-APACHE)
- ATTRIBUTION.md (written by AutoGineer)
- Synced
AutoGineer is not affiliated with or endorsed by NVIDIA NVSHMEM Team or NVIDIA.
Download ZIP
The ZIP holds the skill’s files, its licence files and an ATTRIBUTION.md. Download nvshmem-get-started.zip