skills.rest
vipshop
vipshopOFFICIAL
Explore all skills by this author →
WebsiteGitHub

cuda-cpp-kernel

Implements, debugs, and optimizes CUDA C++ and PTX kernels across NVIDIA GPU architectures.

1.3k|80|Updated Jun 12, 2025
Search Tags:#CUDA#gpu profiling#kernel optimization#ptx#blackwell#tensor cores#Nsight Compute
One-click install
npx skills add https://github.com/vipshop/cache-dit --skill cuda-cpp-kernel
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: cuda-cpp-kernel
Source: https://github.com/vipshop/cache-dit/tree/main/.github/skills/cuda-cpp-kernel
Command: npx skills add https://github.com/vipshop/cache-dit --skill cuda-cpp-kernel

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill includes references (resource) components.

What problem does it solve?

Writing and tuning CUDA kernels requires deep knowledge of PTX instructions, memory hierarchies, Tensor Core behavior, and architecture-specific limits across Ampere, Hopper, and Blackwell GPUs, and this Skill bundles the references and workflows needed to do that work without ad hoc web searches.

Core Features & Use Cases

  • Kernel Implementation and Review: Write or review CUDA C++ kernels and host code using bundled PTX ISA, CUDA Runtime, Driver API, and programming guide references.
  • Architecture-Specific Optimization: Apply dedicated optimization guides for sm89 (Ada), sm90 (Hopper), sm100/sm103 (Blackwell datacenter), and sm120 (Blackwell desktop), covering TMA, WGMMA/tcgen05, TMEM, clusters, shared memory, and occupancy.
  • Profiling and Debugging Workflows: Diagnose failures with compute-sanitizer and cuda-gdb, and profile with Nsight Systems and Nsight Compute using architecture-aware metric interpretation.
  • Use Case: When porting an attention kernel from H100 to B200, consult the sm100 guide to enlarge tiles for 228 KB shared memory, replace wgmma with tcgen05, and validate accuracy and speed against a PyTorch baseline.

Quick Start

Ask the assistant to implement or optimize a CUDA kernel for a specific GPU architecture, providing the shape, dtype, and current bottleneck.

Frequently Asked Questions about cuda-cpp-kernel

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I optimize a CUDA kernel for a specific GPU architecture?▼

Start with the matching smXX optimization guide (sm89, sm90, sm100, sm103, or sm120), then profile with Nsight Systems for end-to-end bottlenecks and Nsight Compute for per-kernel root causes. Change one dimension at a time, such as tile shape or memory movement, and re-measure against the baseline.

How do I debug illegal memory access or race conditions in CUDA kernels?▼

Reproduce the failure with the smallest failing input, then use compute-sanitizer or cuda-gdb for correctness problems. For kernels using shared memory or cp.async, first suspect missing barriers or premature shared-memory slot reuse before assuming the math is wrong.

What is the difference between Hopper wgmma and Blackwell tcgen05?▼

Hopper uses wgmma.mma_async with register-fragment operands, while Blackwell replaces it with tcgen05 tensor-core instructions backed by Tensor Memory (TMEM). TMEM keeps tensor operands out of the 64K register file, reducing register pressure and enabling higher occupancy for GEMM kernels.

Can I port a Blackwell datacenter kernel to RTX 5090 (sm120)?▼

Yes, but sm120 lacks clusters, TMEM, and WGMMA, so you must replace tcgen05 with WMMA, shrink tiles to fit 128 KB shared memory, and rely on cp.async or TMA without multicast. Kernel fusion becomes critical because GDDR7 bandwidth is far lower than HBM3e.

When should I not use this CUDA kernel skill?▼

Do not use it when the primary task is CUTLASS or CuTe template design, or CuTe DSL Python kernel authoring, since dedicated skills cover those. For cache-dit operator registration and packaging work, pair it with the operator-migration skill instead.

How do I avoid shared memory bank conflicts in CUDA?▼

Shared memory has 32 banks of 4 bytes each, so consecutive thread access to consecutive addresses avoids conflicts while 32-stride access causes them. Padding arrays, such as declaring float data[32][33] instead of [32][32], shifts accesses across banks.

Related Skills

Explore More Software Engineering Skills→
moirai-internalmoirai-internal

diffusion-kernel

Integrate optimized CUDA and Triton kernels for diffusion model inference.

Official
Advanced
danielfleischerdanielfleischer
1

cuda-kernels

Benchmark optimized CUDA kernels for diffusers and transformers on NVIDIA GPUs.

Community
Intermediate
vipshopvipshop
1.2k

cutlass-cpp-kernel

Implement and optimize CUTLASS and CuTe C++ kernels with validation.

Official
Advanced
hung-phanhung-phan
1

unsloth-sft

Fine-tune large language models on single GPUs with custom Triton kernels.

Community
Advanced