Cuda Cpp Kernel

CUDA C++ and PTX Kernel Development

Goal

Use the bundled CUDA, PTX, and profiling references in this skill to implement, debug, and optimize CUDA kernels without relying on agent-specific install paths or ad hoc web searches.

When to Use

Use this skill when you need to:

• write or review CUDA C++ kernels or supporting host code
• reason about PTX instructions, inline PTX, Tensor Core instructions, or memory model details
• debug CUDA Runtime API or Driver API failures
• profile or optimize a kernel with Nsight Systems, Nsight Compute, compute-sanitizer, or cuda-gdb
• investigate shared memory bank conflicts, memory coalescing, occupancy, register pressure, async copy, TMA, or cluster behavior
• compare a custom operator or kernel against a PyTorch baseline for correctness or performance

Do not use this skill for:

• CUTLASS or CuTe template design as the primary task; use cutlass-cpp-kernel
• CuTe DSL Python kernel authoring as the primary task; use cute-dsl-kernel
• cache-dit operator registration, packaging, or public API migration work by itself; pair with operator-migration

Reference Style Rule

Use skill-local relative paths for bundled references, for example:

• references/ptx-docs/
• references/cuda-runtime-docs/
• references/ncu-docs/ProfilingGuide.md

Do not write agent-specific install paths into follow-up notes or generated docs.

Bundled Reference Map

The following directories are bundled under references/ inside this skill:

• references/ptx-docs/ — full PTX ISA reference
• references/ptx-simple/ — condensed PTX quick reference
• references/cuda-runtime-docs/ — CUDA Runtime API reference
• references/cuda-driver-docs/ — CUDA Driver API reference
• references/cuda-guide/ — CUDA Programming Guide
• references/best-practices-guide/ — CUDA C++ Best Practices Guide
• references/ncu-docs/ — Nsight Compute docs
• references/nsys-docs/ — Nsight Systems docs
• references/debugging-tools.md — debugging workflow notes
• references/performance-traps.md — common optimization traps

The following architecture and kernel reference files are also bundled at the top level of this skill:

• sm89-optimization-guide.md — Ada-specific optimization and profiling guidance
• sm90-optimization-guide.md — Hopper-specific optimization and profiling guidance
• sm100-optimization-guide.md — Blackwell datacenter optimization and profiling guidance
• sm103-optimization-guide.md — Blackwell Ultra optimization and profiling guidance
• sm120-optimization-guide.md — Blackwell desktop or workstation optimization and profiling guidance
• kernel-templates.md — low-level CUDA kernel templates and implementation patterns
• troubleshooting.md — debugging, compute-sanitizer, and profiling troubleshooting notes

How to Search the Bundle

Prefer narrow text search over loading large reference files into context.

Suggested workflow:

1. Identify the exact instruction, API, metric, or concept.
2. Search the narrowest relevant subdirectory first.
3. Read only the matching file or a short relevant range.
4. Translate the documentation into the specific kernel or operator constraint you are implementing.

Typical search targets:

• PTX instruction syntax: references/ptx-docs/9-instruction-set/
• Quick PTX lookup: references/ptx-simple/
• CUDA Runtime APIs: references/cuda-runtime-docs/modules/
• CUDA Driver APIs: references/cuda-driver-docs/modules/
• architecture and programming-model behavior: references/cuda-guide/
• profiling metrics and sections: references/ncu-docs/ProfilingGuide.md
• timeline and launch behavior: references/nsys-docs/UserGuide.md

For architecture-specific tuning and profiling, read the matching optimization guide first:

• sm89-optimization-guide.md
• sm90-optimization-guide.md
• sm100-optimization-guide.md
• sm103-optimization-guide.md
• sm120-optimization-guide.md

Architecture-Specific Profiling Workflow

When using Nsight Systems or Nsight Compute, interpret the results in the context of the target architecture instead of treating all GPUs the same.

Use the bundled architecture guides as the first reference for cross-architecture analysis:

1. For sm89, focus on memory throughput, L2 hit rate, kernel fusion opportunity, and the lack of TMA or cluster features.
2. For sm90, focus on TMA overlap, warpgroup behavior, shared-memory staging, and whether the timeline shows good load or compute overlap.
3. For sm100 and sm103, focus on tcgen05 usage, TMEM behavior, TMA v2 overlap, cluster behavior, and whether the kernel actually benefits from Blackwell datacenter features.
4. For sm120, treat it closer to Ada than to datacenter Blackwell for profiling purposes: watch memory throughput, L2 hit rate, shared-memory limits, and the lack of TMEM or cluster features while deciding explicitly whether TMA or cp.async is the better staging path.

Recommended order:

1. Read the matching smXX-optimization-guide.md file.
2. Use nsys to identify launch gaps, overlap, copy or compute concurrency, and end-to-end bottlenecks.
3. Use ncu to inspect architecture-specific limits such as occupancy, memory throughput, L2 hit rate, tensor core utilization, stall reasons, register pressure, or shared-memory pressure.
4. Compare the observations against the architecture guide before changing tile shapes, pipelines, or memory movement.