Cute Dsl Kernel

Write a CuTe DSL GPU Kernel

Goal

Use the bundled CuTe DSL API snapshots in this skill and the workspace CUTLASS checkout to design, implement, debug, and integrate CuTe DSL GPU kernels in a way that is reusable across projects, including cache-dit.

When to Use

Use this skill when you need to:

• write or modify a CuTe DSL GPU kernel in Python
• study CuTe DSL types, runtime helpers, architecture APIs, or pipeline abstractions
• port or rewrite an existing CUDA or C++ operator into CuTe DSL
• use CuTe DSL examples from the workspace CUTLASS checkout as reference material
• debug CuTe DSL compilation, runtime behavior, layout issues, or integration problems

Do not use this skill for:

• CUTLASS or CuTe C++ template work as the primary task; use cutlass-cpp-kernel
• generic CUDA/PTX documentation lookup with no CuTe DSL angle; use cuda-cpp-kernel
• repository integration plumbing by itself; pair with operator-migration

Core Rule

Read the relevant API reference files before writing kernel code.

Do not guess CuTe DSL APIs or architecture helpers from memory when the bundled docs or workspace CUTLASS examples can answer the question precisely.

Reference Style Rule

Use Copilot-friendly sibling-file references for bundled docs in this skill, for example:

• cute.md
• cute_runtime.md
• utils.md
• cute_nvgpu_tcgen05.md
• pipeline.md

Use workspace-relative paths for CUTLASS sources, for example:

• vipshop/cutlass/python/CuTeDSL/
• vipshop/cutlass/examples/python/CuTeDSL/
• vipshop/cutlass/python/pycute/
• vipshop/cutlass/include/cute/
• vipshop/cutlass/media/docs/pythonDSL/

Do not use agent-specific skill paths or placeholder-driven argument text in the final skill content.

Read These Files First

Core API references:

• cute.md — core CuTe DSL types and tensor or layout operations
• cute_runtime.md — runtime helpers and data interop
• utils.md — helper utilities and hardware info

Architecture-specific references:

• cute_nvgpu.md — architecture API index
• cute_nvgpu_warp.md — warp-level APIs for SM80 to SM89
• cute_nvgpu_warpgroup.md — warpgroup APIs for SM90
• cute_nvgpu_tcgen05.md — tcgen05 and SM100+ APIs
• cute_nvgpu_cpasync.md — async-copy APIs
• cute_arch.md — low-level architecture primitives
• utils_sm90.md and utils_sm100.md — architecture helpers

Pipeline and overview:

• pipeline.md
• intro.md

Additional workflow and concept references from the workspace CUTLASS docs:

• vipshop/cutlass/media/docs/pythonDSL/overview.rst — high-level positioning of CUTLASS DSLs and how CuTe DSL relates to CUTLASS C++
• vipshop/cutlass/media/docs/pythonDSL/quick_start.rst — environment, install, and setup assumptions
• vipshop/cutlass/media/docs/pythonDSL/functionality.rst — supported dtypes, architectures, and current feature scope
• vipshop/cutlass/media/docs/pythonDSL/limitations.rst — current CuTe DSL limitations and unsupported cases
• vipshop/cutlass/media/docs/pythonDSL/faqs.rst — common issues and expected behavior
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl.rst — CuTe DSL workflow overview
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_api.rst — API documentation entrypoint
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/dsl_introduction.rst — DSL programming model and mental model
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/dsl_control_flow.rst — control-flow semantics and restrictions
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/dsl_dynamic_layout.rst — static vs dynamic layout handling
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/dsl_jit_arg_generation.rst — JIT argument typing and signature generation
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/dsl_jit_caching.rst — JIT cache behavior
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/dsl_jit_compilation_options.rst — compilation flags and debugging options
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/framework_integration.rst — framework interop patterns
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/dsl_ahead_of_time_compilation.rst — AOT compilation and export flow
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/debugging.rst — debugging workflow and generated-artifact inspection
• vipshop/cutlass/media/docs/pythonDSL/cute_dsl_general/autotuning_gemm.rst — autotuning guidance for GEMM kernels

These workspace docs are especially valuable when the bundled API snapshots are too terse for workflow, compilation, debugging, or integration questions.

CUDA architecture and profiling references bundled in this skill:

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

Use these files when interpreting nsys and ncu results for generated CuTe DSL kernels on different GPU families.

Workspace Source Map

Use the workspace CUTLASS checkout for source examples and implementation patterns.

Key locations:

• vipshop/cutlass/python/CuTeDSL/ — CuTe DSL implementation sources
• vipshop/cutlass/examples/python/CuTeDSL/ — CuTe DSL examples by architecture and topic
• vipshop/cutlass/python/pycute/ — pycute helpers and layout utilities
• vipshop/cutlass/include/cute/ — CuTe C++ headers for semantic grounding

Use the shell path /workspace/dev/vipshop/cutlass only when you need a literal command path.

Architecture-Specific Profiling Guidance

CuTe DSL kernels often need architecture-aware profiling because the generated kernel structure can look similar while the best bottleneck diagnosis differs by GPU generation.

Use the bundled optimization guides as follows:

1. On sm89 and sm120, prioritize memory throughput, L2 hit rate, occupancy, and fusion opportunity; these targets do not have TMA, TMEM, or cluster features.
2. On sm90, inspect whether TMA-style overlap, warpgroup execution, and shared-memory staging are actually visible in the timeline and counters.
3. On sm100 and sm103, inspect whether tcgen05 or WGMMA, TMEM, TMA v2, and cluster-capable execution are being used effectively.

Recommended profiling order:

1. Read the relevant smXX-optimization-guide.md file.
2. Use nsys to identify launch gaps, missing overlap, copy or compute imbalance, and end-to-end bottlenecks.
3. Use ncu to inspect occupancy, memory throughput, L2 hit rate, register pressure, shared-memory pressure, tensor core utilization, and stall reasons.
4. Only then decide whether to change tiling, pipelining, copy strategy, or fusion structure.