Nvshmem Enable Tma

NVIDIAOfficialCreated by NVIDIA NVSHMEM Team

Skill
Tool integration
Software development
Apache-2.0
Medium risk

From the author

Prepare or review NVSHMEM CUDA kernels for TMA SMEM registration and direct-SMEM transfers. Do not use for unrelated CUDA tuning.

— the description by NVIDIA NVSHMEM Team

AutoGineer summary

Written by AutoGineer from the skill’s own files, not by the author

This skill provides a structured process for assessing, patching, and reviewing NVSHMEM CUDA kernels to leverage Tensor Memory Accelerator (TMA) capabilities. It focuses on enabling efficient direct-SMEM transfers or NVSHMEM-managed GMEM staging while ensuring strict adherence to scope, safety, and performance requirements. The skill guides users through verifying prerequisites like NVSHMEM versioning, GPU architecture (SM90+), and peer-reachability, ensuring that changes are surgical and do not disrupt existing application logic.

What it does

Assesses CUDA kernels for TMA-backed NVSHMEM communication compatibility.
Provides a workflow for patching kernels to use direct-SMEM puts or managed GMEM staging.
Enforces strict scope boundaries to prevent unintended modifications to non-target kernels.
Validates prerequisites including NVSHMEM version, GPU architecture, and network topology.
Generates a structured report detailing the outcome, policy preconditions, and validation status.
Ensures existing application error handling and completion semantics are preserved.

When to use it

When you need to optimize NVSHMEM point-to-point transfers using TMA on SM90+ hardware.
When you want to transition from standard NVSHMEM transfers to direct-SMEM puts for proven-redundant staging buffers.
When performing a code review of NVSHMEM-enabled kernels to ensure correct TMA registration and launch configuration.
When diagnosing performance or configuration issues related to TMA routing in NVSHMEM applications.

Advantages

Minimizes risk by enforcing surgical, scope-limited code changes.
Ensures correctness by requiring proof of redundancy before removing staging buffers.
Provides clear, standardized reporting for validation and troubleshooting.
Maintains compatibility with regular NVSHMEM paths for ineligible operations or older hardware.

Usage examples

“Use nvshmem-enable-tma to prepare app.cu for NVSHMEM TMA on H100 NVLink while leaving non-NVSHMEM kernels unchanged.”: The agent will inspect app.cu, identify the target kernel, apply necessary registration and TMA-enabled launch parameters, and confirm that only the selected kernel was modified.
“Use nvshmem-enable-tma to inspect a shared-tile-to-GMEM-to-put sequence and, if the GMEM copy is provably redundant, pass the existing shared tile directly.”: The agent will verify if the GMEM buffer is redundant and, upon confirmation, provide a patch to remove the intermediate copy and use a direct-SMEM put with the required async proxy fence.

Install

Download the skill as a ZIP:
https://mp-staging.autogineer.com/agent-skills/nvidia/nvshmem-enable-tma/download

Unzip it into the cross-client skills folder of your project:

mkdir -p .agents/skills && unzip nvshmem-enable-tma.zip -d .agents/skills

Documentation by NVIDIA NVSHMEM Team

Enable TMA in NVSHMEM Applications

Purpose

Prepare an application for NVSHMEM's TMA-backed device-side transfers without changing its put/get APIs or unrelated kernels. Assess the target, choose the correct shared-memory path, make the smallest safe change, and explain launch-time requirements.

Prerequisites

Require readable CUDA/C++ source for a patch or code review. A static assessment needs no GPU, installed NVSHMEM runtime, or cluster access.
Require an NVSHMEM 3.7.0-or-newer development environment and CUDA toolchain only when compiling.
Require SM90-or-newer GPUs and peer GPU load/store reachability such as NVLink only when validating TMA execution or performance.
CFT requires TMA.

Reference Routing

Before changing code, read registration-and-launch.md completely and then every reference selected below. Use the references as shapes, not blind replacements: preserve the application's error handling, launch abstraction, completion contract, and shared-memory layout.

Situation	Required additional reference
Global-memory operands, or staging copy/buffer is not proven redundant	gmem-staging.md
Explicitly opted-in shared-memory put source with a proven-redundant GMEM copy	direct-smem-put.md
Compile, run, diagnose, or report validation	validation-and-troubleshooting.md

For a read-only assessment or review, read only the references needed to evaluate the proposed path. Do not make a code change until the complete registration reference and every selected path reference have been read.

Inputs

Required: Application source or source path, and the requested outcome: assessment, patch, review, or debugging.
Optional: Derive NVSHMEM version, CUDA architecture, topology, kernel launch sites, and launch configuration from the project or environment when possible.
Conditional for execution: Obtain build command, launch command, and target hardware access only when the user asks to compile or run validation.
Ask only for facts that cannot be inspected and materially affect the result. Do not require topology details merely to prepare code; mark performance benefit unverified when topology is unknown.

Scope

Identify exactly one user-opted-in CUDA translation unit and one NVSHMEM communication kernel before editing. Its required launch code is the complete change boundary.
Derive the target from the user's named file/kernel or attachment metadata, not from files discovered by globbing. A singular request does not authorize editing every eligible kernel.
Preserve all other kernels, host-only NVSHMEM use, and unrelated launch code. Never modify adjacent candidates merely because they contain NVSHMEM calls.
Expand beyond one kernel only when the user explicitly names multiple targets. Apply this scope check to each target independently.

Before editing, answer every question below. Stop without editing if a required answer is no or unverified.

Explicit scope: Is this the selected file and kernel?
Eligible operation: Does the selected kernel's device call graph reach a TMA-relevant NVSHMEM operation?
Necessary edit: Is every planned change needed for registration, lifetime, direct-SMEM dataflow, or launch sizing?
Proven redundancy: If removing a staging copy or buffer, is it proven to have no producer, consumer, alias, synchronization, completion, or lifetime role? If the answer is unverified because the GMEM buffer is a parameter, escapes the selected translation unit, or has uninspectable consumers, ask the user whether it is externally observable before editing. Do not infer redundancy.
No speculation: Does the selected target require application-side TMA registration for the operation being prepared?
Surgical launch change: Can the existing grid, block, stream, arguments, and error handling remain intact apart from required registration parameters and dynamic-SMEM bytes?

After editing, inspect the changed-file list and diff. Remove only your own out-of-scope edits before reporting completion; the changed application file set must contain only the opted-in target.

Instructions

1. Establish the Boundary

Treat an assessment or review as read-only.
Treat a request to enable, prepare, update, or fix TMA as permission to edit only the selected kernel and in-scope launch configuration.
Inspect the selected kernel's device call graph for device-side point-to-point put, get, put-with-signal, and their typed/thread/warp/block variants.

2. Decide Whether TMA Is Useful and Available

State these conclusions before or alongside a patch:

Require NVSHMEM 3.7.0 or newer. Stop a requested conversion for an older confirmed version and recommend upgrading.
Require SM90 or newer for TMA execution. With NVSHMEM_TMA_POLICY=ENABLE, older GPUs preserve correctness through the regular path but receive no TMA benefit.
TMA helps only peer-reachable GPU memory paths through the GPU load/store fabric, normally NVLink. Do not recommend it for IB, RoCE, EFA, or another network transport.
It is promising for large point-to-point transfers, low communication-thread counts, fused kernels, or tiles already in shared memory. Never promise a speedup without measurement.

When version, architecture, or topology is unknown, prepare code only if requested and label the corresponding conclusion unverified.

3. Classify Each Operation

Inspect local operand spaces, application shared memory, staging-copy dataflow, dynamic/static shared memory, alignment, early returns, CTA-uniform control flow, existing fences/flushes/quiet calls, and possible concurrent staged callers in a CTA.

Situation	Path
Global-memory application operands, or no proven-redundant staging	NVSHMEM-managed GMEM staging
Shared-memory put source with a proven-redundant SMEM-to-GMEM copy	Offer direct shared-memory put as an explicit opt-in
Shared-memory get destination	Keep the documented staging path

Default to NVSHMEM-managed staging. Never introduce a new application shared-memory layout solely to force a direct-operand path.

4. Apply the Selected References

Follow the complete registration/lifetime protocol in the shared registration reference. Then apply the selected path reference. Preserve existing completion semantics and do not hard-code the current recommended shared-memory byte count.

5. Validate and Report

Read the validation reference for all compile, run, diagnosis, and validation-report work. Lead the final report with one of:

TMA is useful and the application is prepared
The application is prepared, but TMA benefit is unverified
TMA is not useful or not supported for this target

Use this report template:

Outcome: <one required outcome line above>
Changed scope: <the selected file, kernel, and launch site; confirm no adjacent files changed>
Selected path: <NVSHMEM-managed GMEM staging | explicitly opted-in direct-SMEM put>
Policy precondition: Set NVSHMEM_TMA_POLICY=ENABLE before nvshmem_init*; <where to set it>
Eligibility: NVSHMEM <version, must be >= 3.7.0>; <SM90+ evidence>; <NVLink/peer load-store evidence>
Operation constraints: <alignment, byte count, scope/concurrency, completion, shared-memory sizing>
Fallback: ENABLE preserves the regular path for ineligible operations; registration alone does not guarantee TMA routing
Unverified: <topology, occupancy, or none>
Validation: <checks run and exact remaining test commands>

Mention any direct-SMEM candidate deliberately not applied because it was outside scope or its redundancy was not proven.

Limitations

Do not convert host-side RMA, collectives without an applicable TMA path, network-transport operations, or remote shared-memory targets.
Do not claim registration guarantees TMA routing; policy, architecture, topology, operand space, alignment, size, scope, and runtime limits determine each transfer.

Troubleshooting

For configuration failures, incorrect results, hangs near shared-memory release, or an unselected TMA path, read validation-and-troubleshooting.md. Treat regular-path fallback under NVSHMEM_TMA_POLICY=ENABLE as valid behavior; do not diagnose it as a correctness failure.

Examples

Use $nvshmem-enable-tma to prepare app.cu for NVSHMEM TMA on H100 NVLink while leaving non-NVSHMEM kernels unchanged.

Use $nvshmem-enable-tma to inspect a shared-tile-to-GMEM-to-put sequence and, if the GMEM copy is provably redundant, pass the existing shared tile directly.

Use $nvshmem-enable-tma to assess an NVSHMEM 3.6 application running on A100 peers connected only through InfiniBand.

Files

ATTRIBUTION.mdAdded by AutoGineer925 B
BENCHMARK.md7.5 KiB
LICENSELicence file added by AutoGineer11.2 KiB
SKILL.md9.0 KiB
evals/evals.json6.0 KiB
evals/files/cuda-transpose.cu628 B
evals/files/direct-smem-put.cu908 B
evals/files/gmem-staging.cu872 B
evals/files/multi-node-put.cu527 B
references/direct-smem-put.md2.0 KiB
references/gmem-staging.md1.2 KiB
references/registration-and-launch.md6.7 KiB
references/validation-and-troubleshooting.md2.5 KiB
skill-card.md4.2 KiB
skill.oms.sig6.4 KiB

references/registration-and-launch.md

TMA Shared-Memory Registration and Launch

Read this reference completely for every TMA code change. Read the selected path reference as well.

Canonical NVSHMEM API (3.7.X and later)

Use only the symbols below. They are declared by <nvshmemx.h>; do not redeclare them, invent aliases, or substitute CUDA TMA APIs. If the target installation does not declare them, stop and report that its headers do not match the requested 3.7.0-or-newer TMA API.

The NVSHMEM 3.7.0, 3.7.1, and 3.7.2 TMA documentation contains an incorrect single-thread requirement for nvshmemx_give_smem and nvshmemx_release_smem. Both calls are CTA-wide: every thread in each participating CTA must call both functions on the same CTA-uniform path. Every thread must pass the same shared-memory pointer and size to nvshmemx_give_smem. Never guard either call with threadIdx.x == 0, an elected lane, or another leader-only condition. This requirement overrides the archived 3.7.0 through 3.7.2 wording.

Pinned NVIDIA sources:

typedef enum {
    NVSHMEMX_SMEM_RECOMMENDED,
    NVSHMEMX_SMEM_MINIMUM,
    NVSHMEMX_SMEM_BARRIERS_ONLY
} nvshmemx_smem_amount_t;

__host__ __device__ int nvshmemx_ask_smem(
    nvshmemx_smem_amount_t flag);
__device__ void nvshmemx_give_smem(void *smem, size_t size);
__device__ void nvshmemx_release_smem(void);

Treat NVSHMEM_TMA_POLICY as an environment variable, not a C/C++ symbol. Set NVSHMEM_TMA_POLICY=ENABLE before any nvshmem_init* call.
Keep the existing NVSHMEM put/get API. NVSHMEM exposes no TMA-specific RMA API.

Select the SMEM Amount

Select the nvshmemx_ask_smem flag from the selected operations' data path:

Use NVSHMEMX_SMEM_RECOMMENDED when any selected operation may require NVSHMEM GMEM staging.
Use NVSHMEMX_SMEM_BARRIERS_ONLY only when every selected operation is a direct application-owned shared-memory put source and needs only NVSHMEM's internal TMA state.
Use NVSHMEMX_SMEM_MINIMUM only after the user explicitly accepts its occupancy/capability tradeoff and the inspected operations support the smaller allocation.

Query the selected amount on the host, pass the returned size to the kernel, and keep it identical across every participating CTA in a launch. Never hard-code a byte count from a particular release.

Mandatory Lifetime Protocol

Reject an incomplete patch. Every participating CTA must follow this protocol on a CTA-uniform path:

[ ] Reserve a 16-byte-aligned prefix at least as large as the queried amount.
[ ] Have every CTA thread call nvshmemx_give_smem(prefix, prefix_size) once per registration lifetime on the same CTA-uniform path. All threads must pass the same pointer and size. Then execute __syncthreads() before any CTA thread enters a TMA-eligible operation.
[ ] Drain work before release: use the scope-matched nvshmemx_flush, nvshmemx_flush_warp, or nvshmemx_flush_block only when source reuse is sufficient; use nvshmem_quiet for remote completion or full ordering. nvshmem_fence is not a completion drain.
[ ] Execute __syncthreads() after all issuing threads reach the drain and before release.
[ ] Have every CTA thread call nvshmemx_release_smem() on the same CTA-uniform path before the CTA returns. No thread that calls give_smem may bypass its matching release. Synchronize again before reusing the prefix after release.

Use this device-side shape, adapting only the RMA call and completion strength:

nvshmemx_give_smem(nvshmem_prefix, nvshmem_prefix_size);
__syncthreads();

// Existing NVSHMEM put/get/put-with-signal call(s).
nvshmem_quiet();  // Or the correct scope-matched flush when sufficient.

__syncthreads();
nvshmemx_release_smem();

Restructure early returns only as needed to guarantee that every CTA thread reaches matching registration and release calls. Never place nvshmemx_give_smem, nvshmemx_release_smem, or __syncthreads() in CTA-divergent control flow.

Default GMEM Registration and Launch

For GMEM staging, keep the operands and RMA call unchanged; donate an NVSHMEM-owned dynamic-SMEM region. Adapt the application's error handling and launch abstraction rather than replacing them:

__global__ void put_kernel(float *remote_dst, const float *local_src,
                           size_t nelems, int peer,
                           size_t nvshmem_smem_bytes) {
    extern __shared__ __align__(16) unsigned char dynamic_smem[];

    nvshmemx_give_smem(dynamic_smem, nvshmem_smem_bytes);
    __syncthreads();

    nvshmemx_float_put_nbi_block(remote_dst, local_src, nelems, peer);
    nvshmem_quiet();  // Keep or adapt the application's required completion.

    __syncthreads();
    nvshmemx_release_smem();
}

size_t nvshmem_smem_bytes =
    (size_t)nvshmemx_ask_smem(NVSHMEMX_SMEM_RECOMMENDED);
size_t total_dynamic_smem = nvshmem_smem_bytes;

cudaFuncSetAttribute(
    put_kernel,
    cudaFuncAttributeMaxDynamicSharedMemorySize,
    (int)total_dynamic_smem);

put_kernel<<<grid, block, total_dynamic_smem, stream>>>(
    remote_dst, local_src, nelems, peer, nvshmem_smem_bytes);

Fit cudaFuncSetAttribute into the application's CUDA error-checking convention. Check the device's opt-in shared-memory limit before claiming the launch is valid.

Coexisting Application Shared Memory

Keep application shared memory outside the NVSHMEM prefix and preserve its alignment. A kernel may retain static application shared memory and add a separate dynamic NVSHMEM region; account for static plus dynamic use when checking the per-block limit and occupancy.

static size_t align_up(size_t value, size_t alignment) {
    return (value + alignment - 1) / alignment * alignment;
}

size_t nvshmem_smem_bytes =
    (size_t)nvshmemx_ask_smem(NVSHMEMX_SMEM_RECOMMENDED);
size_t app_offset = align_up(nvshmem_smem_bytes, app_alignment);
size_t total_dynamic_smem = app_offset + app_smem_bytes;

kernel<<<grid, block, total_dynamic_smem, stream>>>(
    args..., nvshmem_smem_bytes, app_offset);

__global__ void kernel(/* application args */,
                       size_t nvshmem_smem_bytes,
                       size_t app_offset) {
    extern __shared__ __align__(16) unsigned char dynamic_smem[];
    void *nvshmem_smem = dynamic_smem;
    float *application_tile =
        reinterpret_cast<float *>(dynamic_smem + app_offset);

    nvshmemx_give_smem(nvshmem_smem, nvshmem_smem_bytes);
    __syncthreads();

    // Existing application work and eligible NVSHMEM operations.
    nvshmem_quiet();  // Or the correct scope-matched flush when sufficient.

    __syncthreads();
    nvshmemx_release_smem();
}

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.

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-enable-tma · commit fce3cca

Licence

Apache License 2.0 (redistribution allowed)

LICENSE (added by AutoGineer from LICENSE-APACHE)
ATTRIBUTION.md (written by AutoGineer)

Synced

2026-11-08

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-enable-tma.zip