edge-fm-x

NCU-guided CUDA/CUTLASS kernel optimization workflow for edge-fm-x, including short standalone/repo-in-place tuning and long Humanize + KernelPilot loops. Use when optimizing kernels under src/, tuning standalone CUDA/CUTLASS reproductions, or when the user asks to speed up an Edge-FM CUDA kernel and mentions Nsight Compute, Nsight Systems, NCU, benchmark plateaus, Humanize, KernelPilot, long-running kernel search, source provenance, ledgers, CUTLASS, latency regressions, or iterative optimization.

Query NVIDIA PTX ISA 9.1, CUDA Runtime API 13.1, Driver API 13.1, Programming Guide v13.1, Best Practices Guide, Nsight Compute, Nsight Systems local documentation. Debug and optimize GPU kernels with nsys/ncu/compute-sanitizer workflows. Includes hardware quick guides for T4, A100, H100, A800, RTX 3060, and Jetson Orin. Use when writing, debugging, or optimizing CUDA code, GPU kernels, PTX instructions, inline PTX, TensorCore operations (WMMA, WGMMA, TMA, tcgen05), or when the user mentions CUDA API functions, error codes, device properties, memory management, profiling, GPU performance, compute capabilities, CUDA Graphs, Cooperative Groups, Unified Memory, dynamic parallelism, or CUDA programming model concepts.

Develop, debug, and optimize CUTLASS C++, CuTe C++, and CuTe DSL kernels using the local third_party/cutlass repository in this workspace. Use when the user mentions CUTLASS, CuTe, CuTe DSL, cute::Layout, cute::Tensor, TiledMMA, TiledCopy, CollectiveBuilder, CollectiveMainloop, CollectiveEpilogue, GEMM kernel, grouped GEMM, sparse GEMM, FMHA, StreamK, warp specialization, TMA, WGMMA, tcgen05, blockscaled GEMM, or asks about writing or understanding high-performance CUDA kernels with CUTLASS/CuTe.

Generate detailed benchmark reports inside this edge-fm repo, covering configurable Qwen2.5 LLM/VLM model matrices on CUDA. The default fair comparison remains Transformers vs EdgeFM(cuda-graph) vs TRT-Edge-LLM for LLM, while VLM currently uses Transformers vs EdgeFM(cuda-graph). Use when the user asks for benchmark指标、性能对比、详细报告、cuda-graph vs TRT-Edge-LLM analysis, VLM benchmark reports, or wants the existing tests/engine/test_qwen2_generate.py benchmark flow wrapped into a repeatable local workflow.

Run an end-to-end Humanize Kernel Agent Loop for GPU kernel optimization: plan, refine, create a clean standalone repo, use kernel-knowledge, benchmark, profile with Nsight Compute, maintain lineage/ledgers, and start RLCR.

Use when optimizing GPU kernels with Humanize or Codex and the agent should consult the local kernel knowledge pack before planning, choosing optimization directions, or doing plateau research. Covers code-first framework/topic lookup, source provenance, baseline-derived candidates, and user-directed kernel stacks.

Use when a GPU kernel optimization loop needs to turn Nsight Compute (ncu) metrics into a Profile Evidence Digest, classify bottlenecks, explain regressions, plateaus, or surprising wins, and choose one concrete next kernel edit.

Build edge-fm for Jetson Orin inside a mounted Docker workspace using `nvcr.io/nvidia/l4t-jetpack:r36.4.0`. Use when Codex needs to validate Orin compilation, reproduce a JetPack 6.1 build environment, or produce a repeatable Docker-based Orin build command without relying on Thor.

Add or extend operators in edge-fm using the current layer + operators + operator_impl_table design. Use when Codex needs to add a new operator implementation (`impl_id`) to an existing op kind, introduce a new `op_kind` and layer pair, update the 最优算子表 / `operator_impl_table`, wire operator selection by `model_name`, `hw_profile`, `layer_role`, `op_name`, `stage`, or `shape_sig`, update CMake or pybind touchpoints, or keep Horizon `graph_tuning` aligned with the operator changes.