{"owner":"NVIDIA","repo":"TensorRT-LLM","hasSkills":true,"totalSkillsCount":391,"totalTokensCount":155792,"categories":["anthropic-skill","root-instruction","claude-rule","codex-rule","subagent-persona","command-prompt","plugin-manifest"],"hasMcp":true,"mcpConfig":{"mcpServers":{"TensorRT-LLM":{"command":"npx","args":["-y","@modelcontextprotocol/server-TensorRT-LLM"]}}},"found":[".claude/skills/exec-local-slurm/SKILL.md",".claude/skills/exec-local-compile/SKILL.md",".claude/skills/exec-env-check/SKILL.md",".claude/skills/exec-slurm-compile/SKILL.md",".claude/skills/exec-local-docker/SKILL.md",".claude/skills/exec-remote-slurm/SKILL.md",".claude/skills/kernel-cute-writing/SKILL.md",".claude/skills/perf-host-analysis/SKILL.md",".claude/skills/perf-analysis/SKILL.md",".claude/skills/perf-host-optimization/SKILL.md",".claude/skills/kernel-triton-writing/SKILL.md",".claude/skills/kernel-tileir-optimization/SKILL.md",".claude/skills/perf-instability-cookbook/SKILL.md",".claude/skills/perf-nsight-compute-analysis/SKILL.md",".claude/skills/perf-nsight-systems/SKILL.md",".claude/skills/perf-optimization-casebook/SKILL.md",".claude/skills/perf-optimization/SKILL.md",".claude/skills/perf-regression-cookbook/SKILL.md",".claude/skills/perf-torch-cuda-graphs/SKILL.md",".claude/skills/perf-torch-sync-free/SKILL.md",".claude/skills/perf-workload-profiling/SKILL.md",".claude/skills/trtllm-case-executor/SKILL.md",".claude/skills/trtllm-code-contribution/SKILL.md",".claude/skills/trtllm-codebase-exploration/SKILL.md",".claude/skills/trtllm-flashinfer-upgrade/SKILL.md",".claude/skills/trtllm-model-onboard-multimodal/SKILL.md",".claude/skills/trtllm-moe-develop/SKILL.md",".claude/skills/trtllm-serve-config-guide/SKILL.md",".claude/skills/trtllm-test-script-builder/SKILL.md",".claude/skills/trtllm-test-specialist/SKILL.md",".claude/skills/trtllm-visualgen-bench/SKILL.md",".claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md",".codex/skills/visual-gen-component-test/SKILL.md","agent-flow/.claude/skills/perf-analyze/SKILL.md","agent-flow/.claude/skills/perf-optimize/SKILL.md","AGENTS.md","CLAUDE.md",".codex/AGENTS.md",".codex/config.toml","agent-flow/AGENTS.md","agent-flow/CLAUDE.md","tests/AGENTS.md","tests/CLAUDE.md",".claude/agents/exec-compile-specialist.md",".claude/agents/exec-local-slurm.md",".claude/agents/exec-remote-slurm.md",".claude/agents/kernel-cuda-specialist.md",".claude/agents/kernel-cute-specialist.md",".claude/agents/kernel-tileir-specialist.md",".claude/agents/kernel-triton-specialist.md",".claude/agents/modeling-test-specialist.md",".claude/agents/perf-profiling-specialist.md",".claude/agents/perf-test-sync.md",".claude/agents/perf-torch-cuda-graph-specialist.md",".codex/agents/onboard_update_reviewer.toml","triton_backend/tools/dataset/mini_cnn_eval.json","triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json",".claude/skills/exec-remote-slurm/scripts/build.sh",".claude/skills/exec-remote-slurm/scripts/build.slurm",".claude/skills/exec-slurm-compile/scripts/compile.sh",".claude/skills/exec-slurm-compile/scripts/compile.slurm",".claude/skills/exec-slurm-compile/scripts/enroot-import",".claude/skills/exec-slurm-compile/scripts/submit_compile.sh",".claude/skills/kernel-cute-writing/references/api-arch.md",".claude/skills/kernel-cute-writing/references/api-core.md",".claude/skills/kernel-cute-writing/references/api-nvgpu.md",".claude/skills/kernel-cute-writing/references/api-runtime-utils.md",".claude/skills/kernel-cute-writing/references/concepts-architecture.md",".claude/skills/kernel-cute-writing/references/concepts-layouts.md",".claude/skills/kernel-cute-writing/references/concepts-mma.md",".claude/skills/kernel-cute-writing/references/concepts-tensors.md",".claude/skills/kernel-cute-writing/references/patterns-compilation.md",".claude/skills/kernel-cute-writing/references/patterns-elementwise.md",".claude/skills/kernel-cute-writing/references/patterns-gemm.md",".claude/skills/kernel-cute-writing/references/patterns-getting-started.md",".claude/skills/kernel-cute-writing/references/patterns-memory.md",".claude/skills/kernel-cute-writing/references/patterns-pipeline.md",".claude/skills/kernel-cute-writing/references/patterns-reduction.md",".claude/skills/kernel-cute-writing/references/troubleshooting.md",".claude/skills/kernel-cute-writing/scripts/__init__.py",".claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py",".claude/skills/kernel-cute-writing/scripts/verify_kernel.py",".claude/skills/kernel-tileir-optimization/references/config-templates.md",".claude/skills/kernel-tileir-optimization/references/tma-conversion.md",".claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py",".claude/skills/kernel-tileir-optimization/scripts/tileir_check.py",".claude/skills/kernel-triton-writing/references/api-core.md",".claude/skills/kernel-triton-writing/references/api-language.md",".claude/skills/kernel-triton-writing/references/concepts-semantics.md",".claude/skills/kernel-triton-writing/references/operator-routing.md",".claude/skills/kernel-triton-writing/references/patterns-advanced.md",".claude/skills/kernel-triton-writing/references/patterns-basic.md",".claude/skills/kernel-triton-writing/references/patterns-fusion.md",".claude/skills/kernel-triton-writing/references/patterns-gemm.md",".claude/skills/kernel-triton-writing/references/troubleshooting.md",".claude/skills/kernel-triton-writing/scripts/__init__.py",".claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py",".claude/skills/kernel-triton-writing/scripts/verify_kernel.py",".claude/skills/perf-host-analysis/references/examples.md",".claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md",".claude/skills/perf-host-analysis/references/kernel-level-analysis.md",".claude/skills/perf-host-analysis/references/metrics.md",".claude/skills/perf-host-analysis/references/nsys-schema.md",".claude/skills/perf-host-analysis/references/output-format.md",".claude/skills/perf-host-analysis/references/phase-classification.md",".claude/skills/perf-host-analysis/references/thresholds.md",".claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md",".claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py",".claude/skills/perf-host-analysis/scripts/detect_host_overhead.py",".claude/skills/perf-host-optimization/references/communication-patterns.md",".claude/skills/perf-host-optimization/references/examples.md",".claude/skills/perf-host-optimization/references/hot-path-files.md",".claude/skills/perf-host-optimization/references/hotspot-classification.md",".claude/skills/perf-host-optimization/references/optimization-patterns.md",".claude/skills/perf-host-optimization/references/optimization-strategy.md",".claude/skills/perf-instability-cookbook/data/aliases.yaml",".claude/skills/perf-instability-cookbook/data/modules.yaml",".claude/skills/perf-instability-cookbook/data/patterns.yaml",".claude/skills/perf-instability-cookbook/data/pending.yaml",".claude/skills/perf-instability-cookbook/data/removed.yaml",".claude/skills/perf-instability-cookbook/data/tags.yaml",".claude/skills/perf-instability-cookbook/references/case-template.md",".claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md",".claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md",".claude/skills/perf-nsight-compute-analysis/references/cli-reference.md",".claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md",".claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md",".claude/skills/perf-nsight-compute-analysis/references/python-report-api.md",".claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md",".claude/skills/perf-nsight-compute-analysis/references/sections-guide.md",".claude/skills/perf-nsight-systems/references/app-preparation.md",".claude/skills/perf-nsight-systems/references/cli-post-collection.md",".claude/skills/perf-nsight-systems/references/cli-profiling.md",".claude/skills/perf-nsight-systems/references/expert-systems.md",".claude/skills/perf-nsight-systems/references/nvtx-analysis.md",".claude/skills/perf-nsight-systems/references/recipes-dl.md",".claude/skills/perf-nsight-systems/references/stats-reports.md",".claude/skills/perf-optimization-casebook/data/aliases.yaml",".claude/skills/perf-optimization-casebook/data/patterns.yaml",".claude/skills/perf-optimization-casebook/data/tags.yaml",".claude/skills/perf-optimization-casebook/references/case-template.md",".claude/skills/perf-regression-cookbook/data/aliases.yaml",".claude/skills/perf-regression-cookbook/data/modules.yaml",".claude/skills/perf-regression-cookbook/data/patterns.yaml",".claude/skills/perf-regression-cookbook/data/pending.yaml",".claude/skills/perf-regression-cookbook/data/removed.yaml",".claude/skills/perf-regression-cookbook/data/tags.yaml",".claude/skills/perf-regression-cookbook/references/case-template.md",".claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md",".claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md",".claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md",".claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md",".claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md",".claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py",".claude/skills/perf-torch-sync-free/references/sync-patterns.md",".claude/skills/perf-torch-sync-free/scripts/verify_workload.py",".claude/skills/perf-workload-profiling/references/benchmarking-patterns.md",".claude/skills/perf-workload-profiling/references/nvtx-api.md",".claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md",".claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh",".claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md",".claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md",".claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md",".claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py",".claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh",".claude/skills/trtllm-test-specialist/references/agg_config_template.yaml",".claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml",".claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml",".claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml",".claude/skills/trtllm-test-specialist/references/test_config_template.yaml",".claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md",".claude/skills/trtllm-test-specialist/scripts/build_test_command.py",".claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py",".claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py",".claude/skills/trtllm-test-specialist/scripts/generate_report.py",".claude/skills/trtllm-test-specialist/scripts/parse_config.py",".claude/skills/trtllm-visualgen-bench/media/cat_piano.png",".claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png",".claude/skills/trtllm-visualgen-bench/media/landscape.png",".claude/skills/trtllm-visualgen-bench/media/robot_153.jpg",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json",".claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json",".claude/skills/trtllm-visualgen-bench/scripts/serve_check.py",".claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py",".claude/skills/trtllm-visualgen-bench/serve/README.md",".claude/skills/trtllm-visualgen-bench/tests/test_discovery.py",".claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py",".claude/skills/trtllm-visualgen-bench/tests/test_workloads.py",".claude/skills/trtllm-visualgen-bench/workloads/README.md",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml",".claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml",".claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml",".claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml",".claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml",".claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml",".claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml",".claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml",".claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml",".claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml",".claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml",".claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml","docs/source/commands/trtllm-serve/render-endpoints.md","docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md","tensorrt_llm/tools/profiler/host_profile_tools/README.md","tensorrt_llm/tools/profiler/nsys_profile_tools/README.md","tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json","tests/unittest/api_stability/references/batched_logits_processor.yaml","tests/unittest/api_stability/references/calib_config.yaml","tests/unittest/api_stability/references/completion_output.yaml","tests/unittest/api_stability/references/conversation_params.yaml","tests/unittest/api_stability/references/decoding_base_config.yaml","tests/unittest/api_stability/references/guided_decoding_params.yaml","tests/unittest/api_stability/references/llm.yaml","tests/unittest/api_stability/references/logits_processor.yaml","tests/unittest/api_stability/references/quant_config.yaml","tests/unittest/api_stability/references/request_output.yaml","tests/unittest/api_stability/references/sampling_params.yaml","tests/unittest/api_stability/references/trtllm_serve_api.yaml","tests/unittest/api_stability/references/trtllm_serve_cli.yaml",".claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md",".claude/skills/perf-host-optimization/references/patterns/gpu-graph.md",".claude/skills/perf-host-optimization/references/patterns/loop-iteration.md",".claude/skills/perf-host-optimization/references/patterns/python-overhead.md",".claude/skills/perf-host-optimization/references/patterns/sync-alloc.md",".claude/skills/perf-host-optimization/references/patterns/system.md",".claude/skills/perf-instability-cookbook/references/communication/index.md",".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md",".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md",".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md",".claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md",".claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md",".claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md",".claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md",".claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md",".claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md",".claude/skills/perf-instability-cookbook/references/moe/index.md",".claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md",".claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md",".claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md",".claude/skills/perf-optimization-casebook/references/communication/deepep.md",".claude/skills/perf-optimization-casebook/references/communication/index.md",".claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md",".claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md",".claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md",".claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/index.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md",".claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md",".claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md",".claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md",".claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md",".claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md",".claude/skills/perf-regression-cookbook/references/attention-fmha/index.md",".claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md",".claude/skills/perf-regression-cookbook/references/autotuner/index.md",".claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md",".claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md",".claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md",".claude/skills/perf-regression-cookbook/references/communication/index.md",".claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md",".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md",".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md",".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md",".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md",".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md",".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md",".claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md",".claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md",".claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md",".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md",".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md",".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md",".claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md",".claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md",".claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md",".claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md",".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md",".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md",".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md",".claude/skills/perf-regression-cookbook/references/model-definition/index.md",".claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md",".claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md",".claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md",".claude/skills/perf-regression-cookbook/references/moe/index.md",".claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md",".claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md",".claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md",".claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md",".claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md",".claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md",".claude/skills/perf-regression-cookbook/references/sampler/index.md",".claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md",".claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md",".claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md",".claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md",".claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md",".claude/skills/perf-regression-cookbook/references/spec-decode/index.md",".claude/skills/perf-regression-cookbook/references/tokenization/index.md",".claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md","jenkins/scripts/cbts/cbts/rules/README.md","tests/integration/defs/accuracy/references/acceptance_length.yaml","tests/integration/defs/accuracy/references/cnn_dailymail.yaml","tests/integration/defs/accuracy/references/gpqa_diamond.yaml","tests/integration/defs/accuracy/references/gsm8k.yaml","tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml","tests/integration/defs/accuracy/references/humaneval.yaml","tests/integration/defs/accuracy/references/json_mode_eval.yaml","tests/integration/defs/accuracy/references/longbench_v1.yaml","tests/integration/defs/accuracy/references/longbench_v2.yaml","tests/integration/defs/accuracy/references/mmlu.yaml","tests/integration/defs/accuracy/references/mmmu.yaml","tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml","tests/integration/defs/accuracy/references/videomme.yaml","tests/integration/defs/accuracy/references/voxpopuli.yaml","tests/integration/defs/accuracy/references/zero_scrolls.yaml"],"skills":{".claude/skills/exec-local-slurm/SKILL.md":"---\nname: exec-local-slurm\ndescription: >-\n  Submit and monitor a Slurm job on a local cluster. Supports two modes:\n  (1) Persistent allocation (default) — allocates nodes once via nohup salloc,\n  imports the container once, installs once, and reuses across runs by setting\n  SLURM env vars and running the sbatch script via bash. (2) One-shot sbatch —\n  submits a fully-generated Slurm script via sbatch, polls job status, reads\n  logs on completion, and reports results. Workflow-agnostic — handles pytest,\n  eval, benchmark, and custom scripts identically.\ntags: [slurm, execution, infrastructure]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Local Slurm Executor\n\nSubmit a Slurm job locally, monitor it, and report results. Uses persistent\nallocation by default to eliminate queue wait, container import, and install\noverhead on repeated runs.\n\n## Input (from orchestrator prompt)\n\nThe orchestrator passes these fields in the skill prompt:\n\n| Field | Description |\n|-------|-------------|\n| `script_path` | Full local path to the generated `.slurm` script |\n| `work_dir` | Local work directory for logs and artifacts |\n| `model_name` | Short model name for reporting |\n| `workflow_type` | `pytest`, `eval`, `custom`, or `benchmark` — for output parsing hints |\n| `success_patterns` | Comma-separated patterns indicating success |\n| `failure_patterns` | Comma-separated patterns indicating failure |\n| `log_file_pattern` | Log filename pattern with `%j` placeholder (e.g., `llama_auto_test_%j.out`) |\n| `persistent_mode` | Default: `true` for all local slurm workflows. Set to `false` to force one-shot sbatch (opt-out). |\n| `release_allocation` | `true` to release the current allocation and stop. Only set when the user explicitly says no more jobs are needed. Default: `false`. Never auto-release. |\n| `alloc_time_limit` | Walltime for the persistent allocation. Default: `04:00:00`. |\n| `docker_image` | Container image for persistent container import. From `job_spec.json`. |\n| `container_name` | Container name used in the `.slurm` script (e.g., `llama_auto_test`). Must match exactly. From `job_spec.json`. |\n| `container_mounts` | Comma-separated mount mappings. From `job_spec.json`. |\n| `repo_root` | Repo root path (for locating state file and project path). |\n| `slurm_params` | Slurm parameters object: `partition`, `account`, `nodes`, `ntasks`, `ntasks_per_node`, `gpus_per_node`. From `job_spec.json`. |\n\n## Procedure\n\n### Pre-step: Resolve Image and Build (when `build_project=true`)\n\nThis executor owns image selection and the build for the local SLURM cluster. Skip the entire pre-step when `build_project=false`.\n\n1. **Read the cluster GPU type** from `device_type` in `job_spec.json` (resolved upstream by env-check / internal-env-info). If absent, probe from any compute node via `srun -p <partition> --ntasks-per-node=1 nvidia-smi --query-gpu=name --format=csv,noheader | head -1`. If `skills/internal-env-info/` is not installed, skip the upstream lookup silently and rely on the srun probe — do not report the missing skill as an error.\n2. **Determine the partition's CPU arch.** Look up `job_spec.slurm_env.partitions[]` for the entry where `name == partition` and read its `arch` (`x86_64` / `aarch64`). Case-executor's Step 2.5 already detected this — do **not** run `scontrol show node` here. If `slurm_env` is absent or the entry's `arch` is null (case-executor ran without SLURM tools), fall back to inferring from `device_type`: Grace-based parts (`GB*` / `GH*`) → `aarch64`; everything else → `x86_64`.\n3. **Resolve the container image.** If `docker_image` is already set in `job_spec.json`, use it directly. Otherwise read `<repo_root>/jenkins/current_image_tags.properties` and pick the tag matching the partition's CPU arch.\n4. **Map GPU → build arch (`-a` flag):** `H100`/`H200` → `90-real`; `B200`/`GB200`/`B300`/`GB300` → `100-real`; `A100` → `80-real`; `L40S` → `89-real`. Default `100-real`.\n5. **Compile.** Invoke the `exec-slurm-compile` skill with `repo_dir=<repo_root>`, `partition`, `account`, `container_image=<resolved tag>`, `user_root_dir`, `arch`. Wait for completion.\n6. **On failure**, do not allocate or run the workload. Report `BUILD_FAILED` with the last 100 lines of the build log.\n\n`build_project`, `device_type`, `repo_root`, `partition`, `account`, `user_root_dir`, and (optionally) `docker_image` come from `job_spec.json`.\n\n### Step 0: Allocation Management\n\nThis step runs before any execution. It determines whether to reuse an existing\npersistent allocation, create a new one, or fall through to one-shot sbatch.\n\n#### Step 0A — Release mode\n\nIf `release_allocation=true`:\n\n1. Read `<repo_root>/work_dirs/.slurm_alloc.json`\n2. If file exists and `job_id` is present:\n   ```bash\n   scancel <job_id>\n   ```\n3. Delete the state file and `.salloc.log`\n4. Report: \"Allocation <job_id> released.\" Stop.\n\n#### Step 0B — One-shot mode\n\nIf `persistent_mode=false`, skip entirely to Step 1 (One-Shot Path).\n\n#### Step 0C — Validate existing allocation\n\nRead `<repo_root>/work_dirs/.slurm_alloc.json`.\n\n- If state file **does not exist** → go to Step 0D.\n- If state file exists, validate the allocation:\n  ```bash\n  squeue -j <job_id> -h -o \"%T %L\"\n  ```\n  - **RUNNING + remaining > 5 min + params compatible** → **reuse** (skip to Step 0F)\n  - **RUNNING + remaining <= 5 min** → warn user (\"Allocation expiring soon\"), `scancel <job_id>`, delete state file, go to Step 0D\n  - **RUNNING + params incompatible** → `scancel <job_id>`, delete state file, go to Step 0D\n  - **Empty output / error** (job gone) → stale state file, delete it, go to Step 0D\n\n**Params compatibility check:**\n- `partition` must match\n- `nodes` in state must be **>=** requested nodes (a 2-node allocation can serve 1-node jobs)\n- `docker_image` must match\n\n#### Step 0D — Justify and prepare allocation\n\nBefore allocating, ensure no orphaned allocations exist and log the reason:\n\n1. Check for existing persistent jobs:\n   ```bash\n   squeue -u $(whoami) -h -o \"%i %T %j\" --name=<account>-trtllm.persistent\n   ```\n2. If found → orphan (state file was missing/corrupt). Cancel it:\n   ```bash\n   scancel <orphan_job_id>\n   ```\n   Log: \"Released orphaned allocation <orphan_job_id> — state file was missing.\"\n3. Log justification: \"No reusable allocation found. Allocating <nodes> node(s) on partition <partition> for <alloc_time_limit>.\"\n\n#### Step 0E — Allocate new\n\nConvert `alloc_time_limit` from `HH:MM:SS` to seconds (e.g., `04:00:00` → `14400`).\n\n```bash\nmkdir -p <repo_root>/work_dirs\nnohup salloc --partition=<partition> --account=<account> \\\n  --nodes=<nodes> --time=<alloc_time_limit> \\\n  --job-name=<account>-trtllm.persistent \\\n  sleep <alloc_time_limit_seconds> \\\n  > <repo_root>/work_dirs/.salloc.log 2>&1 &\n```\n\nRetrieve job ID:\n```bash\nsqueue -u $(whoami) -h -o \"%i\" --name=<account>-trtllm.persistent\n```\n\nIf no job appears after 10 seconds, read `<repo_root>/work_dirs/.salloc.log`\nfor errors (bad partition, invalid account, etc.) and report to user.\n\nOtherwise, poll until RUNNING (every 10s, max 60 polls). Once RUNNING, get\nthe nodelist:\n```bash\nsqueue -j <JOB_ID> -h -o \"%N\"\n```\n\n**Import container** on all nodes (using the **same container name** as the\n`.slurm` script — critical for pyxis reuse):\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-image=<docker_image> \\\n  --container-name=<container_name> true\n```\n\n**Warm up filesystem mounts** on all nodes — `ls` each mounted path so\nLustre/NFS metadata is cached for later use:\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-name=<container_name> \\\n  --container-mounts=<container_mounts> \\\n  bash -c 'for p in <mount_target_1> <mount_target_2> ...; do ls \"$p\" > /dev/null 2>&1; done'\n```\nParse mount targets from `container_mounts` — the right-hand side of each\n`host:container` pair.\n\n**Check GPU status** on all nodes (no container needed — `nvidia-smi` is on\nthe host):\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  bash -c 'echo \"=== $(hostname) ===\" && nvidia-smi --query-compute-apps=pid,name,used_memory --format=csv,noheader'\n```\n- If output shows no processes → GPUs are clean, proceed\n- If unexpected processes found → warn user: \"GPU processes found on\n  <node>: <process_list>. These may interfere with the job.\"\n\n**Run install** on all nodes with `--container-writable` (so packages persist\nin the named container across srun calls). Install all common requirements\nupfront so that subsequent jobs of any workflow type can skip install:\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-name=<container_name> --container-writable \\\n  --container-mounts=<container_mounts> \\\n  bash -c 'cd <repo_root> && pip install -e . && pip install -r requirements-dev.txt && \\\n    if [ -f examples/trtllm-eval/requirements.txt ]; then pip install -r examples/trtllm-eval/requirements.txt; fi'\n```\nFor custom workflow with `skip_install=true`, skip this step entirely.\n\n**Write state file** `<repo_root>/work_dirs/.slurm_alloc.json`:\n```json\n{\n  \"job_id\": \"<JOB_ID>\",\n  \"container_name\": \"<container_name>\",\n  \"nodelist\": \"<NODELIST>\",\n  \"partition\": \"<partition>\",\n  \"account\": \"<account>\",\n  \"nodes\": <nodes>,\n  \"gpus_per_node\": <gpus_per_node>,\n  \"docker_image\": \"<docker_image>\",\n  \"container_mounts\": \"<container_mounts>\",\n  \"allocated_at\": \"<ISO timestamp>\",\n  \"time_limit\": \"<alloc_time_limit>\",\n  \"installed\": true\n}\n```\n\nProceed to Step 1 (Persistent Path).\n\n#### Step 0F — Reuse path validation\n\nAllocation is valid. Check if container name matches current request:\n\n- If `container_name` in state file **matches** `container_name` from\n  `job_spec.json` → proceed to Step 1 (Persistent Path). Container and install\n  are already set up.\n- If **different** (model changed) → import the new container, warm up mounts,\n  and install:\n  ```bash\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-image=<docker_image> \\\n    --container-name=<new_container_name> true\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-name=<new_container_name> \\\n    --container-mounts=<container_mounts> \\\n    bash -c 'for p in <mount_targets>; do ls \"$p\" > /dev/null 2>&1; done'\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-name=<new_container_name> --container-writable \\\n    --container-mounts=<container_mounts> \\\n    bash -c 'cd <repo_root> && pip install -e . && pip install -r requirements-dev.txt'\n  ```\n  Update `container_name` in state file. Proceed to Step 1 (Persistent Path).\n\n---\n\n### Step 1: Execute\n\nTwo execution paths depending on `persistent_mode`.\n\n#### Persistent Path (persistent_mode=true)\n\n**Pre-flight time check** — verify the allocation has enough remaining time:\n```bash\nsqueue -j <job_id> -h -o \"%L\"\n```\nIf remaining time < job's `time_limit` from `slurm_params`, warn user:\n\"Allocation has <remaining> left but job expects <time_limit>. The job may\nbe killed early.\"\n\n**Pre-flight GPU check** — verify GPUs are not occupied by leftover processes:\n```bash\nsrun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n  bash -c 'procs=$(nvidia-smi --query-compute-apps=pid,name,used_memory --format=csv,noheader 2>/dev/null); [ -n \"$procs\" ] && echo \"WARNING: GPU processes on $(hostname): $procs\"'\n```\nIf unexpected processes found, warn user before proceeding.\n\n**Run the sbatch script** with SLURM env vars — the `#SBATCH` directives are\ncomments when run via `bash`; the inner `srun` inherits the env vars and uses\nthe persistent allocation:\n```bash\nexport SLURM_JOB_ID=<job_id>\nexport SLURM_JOB_NUM_NODES=<nodes>\nexport SLURM_NNODES=<nodes>\nexport SLURM_NTASKS=<ntasks>\nexport SLURM_NTASKS_PER_NODE=<ntasks_per_node>\nexport SLURM_NODELIST=<nodelist>\nexport SLURM_JOB_NODELIST=<nodelist>\nbash <script_path> 2>&1 | tee <work_dir>/<model_name>_<workflow_type>.log\n```\n\nRun with `run_in_background`. Container reuse is automatic (pyxis skips\nre-import when `--container-name` already exists on the node). The install\nstep (Step 1 srun) always runs to keep the container up to date.\n\n**Step cancellation:** If a job hangs or the user wants to abort, cancel just\nthe srun step without killing the allocation:\n```bash\n# List active steps:\nsqueue -s -j <job_id>\n# Cancel a specific step (e.g., step 0):\nscancel <job_id>.0\n```\nThe allocation stays RUNNING — new jobs can run immediately. The job name\n`<account>-trtllm.persistent` uniquely identifies allocations created by this\nskill via `squeue --name=<account>-trtllm.persistent`.\n\n#### One-Shot Path (persistent_mode=false)\n\n```bash\nsbatch <script_path>\n```\n\nParse the job ID from output: `Submitted batch job <JOB_ID>`.\n\nIf sbatch fails, report the error immediately and stop.\n\n### Step 2: Report Submission\n\n**One-shot mode:**\n```\nJob ID: <JOB_ID>\nScript: <script_path>\nWork directory: <work_dir>\nLog files: <work_dir>/<log_file_pattern with %j replaced>\nMonitor: squeue -j <JOB_ID>\n```\n\n**Persistent mode:**\n```\nAllocation: <job_id> (persistent, remaining: <time>)\nContainer: <container_name>\nScript: <script_path>\nWork directory: <work_dir>\nLog file: <work_dir>/<model_name>_<workflow_type>.log\n```\n\n### Step 3: Poll / Wait for Completion (with Hang Monitoring)\n\nBoth polling modes must also monitor the log file for hang indicators. Every\n60 seconds, check the log file for the pattern `hang detected`\n(case-insensitive). If found, terminate the job and skip to Step 5 with\n`HANG_DETECTED` status.\n\n**One-shot mode:** Poll every 30 seconds, max 60 polls (30 minutes):\n\n```bash\nsqueue -j <JOB_ID> -h -o \"%T %M %R\"\n```\n\n- Report state transitions: `PENDING` → `RUNNING` → `COMPLETING` → done\n- If `PENDING` for > 5 minutes, check the reason:\n  ```bash\n  squeue -j <JOB_ID> -o \"%i %T %r %S\"\n  ```\n- The job is done when `squeue -j <JOB_ID>` returns empty (no rows)\n- If max polls exceeded, report the job is still running and provide manual\n  check commands\n- **Hang check** (every other poll, i.e., every 60s): Once the job is RUNNING,\n  resolve the log file path (`<work_dir>/<log_file_pattern>` with `%j` replaced\n  by `<JOB_ID>`) and check:\n  ```bash\n  grep -qi \"hang detected\" <log_file> 2>/dev/null && echo \"HANG_DETECTED\"\n  ```\n  If `HANG_DETECTED` is printed, immediately cancel the job:\n  ```bash\n  scancel <JOB_ID>\n  ```\n  Then read the last 200 lines of the log file and skip to Step 5.\n\n**Persistent mode:** While waiting for `run_in_background` completion\nnotification, launch a hang monitoring loop using `run_in_background`:\n\n```bash\nwhile true; do\n  sleep 60\n  LOG_FILE=\"<work_dir>/<model_name>_<workflow_type>.log\"\n  if [ -f \"$LOG_FILE\" ] && grep -qi \"hang detected\" \"$LOG_FILE\"; then\n    echo \"HANG_DETECTED: Found 'hang detected' in log file\"\n    squeue -s -j <job_id> -h -o \"%i\" | head -1 | xargs -r scancel 2>/dev/null\n    exit 1\n  fi\ndone\n```\n\nWhen the main process completes normally, kill the monitoring loop. If the\nmonitor fires first, it cancels the active Slurm step (not the allocation\nitself — the allocation stays for future runs).\n\n### Step 4: Collect Results\n\nOn completion:\n\n1. **Get exit code:**\n   - **One-shot:** `sacct -j <JOB_ID> --format=JobID,State,ExitCode,Elapsed,MaxRSS -n`\n   - **Persistent:** From the `run_in_background` result exit code\n\n2. **Resolve log file paths:**\n   - **One-shot:** Replace `%j` in `log_file_pattern` with `<JOB_ID>`:\n     - Output: `<work_dir>/<MODEL_NAME>_..._<JOB_ID>.out`\n     - Error: `<work_dir>/<MODEL_NAME>_..._<JOB_ID>.err`\n   - **Persistent:** Single log at `<work_dir>/<model_name>_<workflow_type>.log`\n\n3. **Read log files:**\n   - Read the last 100 lines of the output log\n   - If job failed, also read the first 50 lines to catch early errors\n   - **One-shot only:** If job failed, read the last 50 lines of the `.err` file\n\n4. **Search for patterns:**\n   - Grep the log file for each `success_patterns` entry\n   - Grep the log file for each `failure_patterns` entry\n\n### Step 5: Report\n\nReturn a structured result:\n\n```\nJob ID: <JOB_ID> (or allocation <job_id> for persistent mode)\nStatus: COMPLETED | FAILED | TIMEOUT | CANCELLED | OUT_OF_MEMORY | HANG_DETECTED\nExit code: <exit code>\nElapsed: <time>\nLog files:\n  stdout: <path>\n  stderr: <path> (one-shot only)\nWork directory: <work_dir>\nSummary: <relevant output — pytest summary, accuracy score, throughput, or last lines>\nErrors: <if failed, relevant error output>\n```\n\n### Output Parsing by Workflow Type\n\n- **pytest**: Look for `X passed, Y failed in Zs` summary line. Status maps:\n  `COMPLETED` + 0 exit → `PASSED`; non-zero → `FAILED`\n- **eval**: Look for `accuracy:` lines. If `--check_accuracy` was used,\n  assertion error means threshold not met\n- **benchmark**: Look for throughput/latency numbers. Check for\n  `8_done_*.txt` completion marker\n- **custom**: No specific patterns — report last 10 lines of output\n\n### Failure Diagnosis\n\n| Symptom | Likely Cause | Action |\n|---------|-------------|--------|\n| `TIMEOUT` | Job exceeded time limit | Report and suggest increasing time |\n| `OUT_OF_MEMORY` | OOM on node | Report and suggest reducing batch size |\n| Exit `1:0` + `FAILED` | Non-zero exit from srun | Read `.err` file for details |\n| `CANCELLED` | Preempted or user cancelled | Check `sacct --format=JobID,State,Reason` |\n| Empty output log | Container failed to start | Check `.err` / `.salloc.log` for mount or image errors |\n| `HANG_DETECTED` in log | Process/GPU hang (deadlock, NCCL timeout) | Job auto-cancelled; check last 200 log lines for root cause |\n| `srun: error: Unable to create step` | Allocation expired | Delete state file, re-run (will re-allocate) |\n| `srun: error: ... does not exist` | Named container missing | Delete state file, re-run (will re-import) |\n\n## Rules\n\n- Never use `squeue` with the `-uall` flag\n- Always use `sacct` (not `squeue`) for final job status in one-shot mode\n- Poll with `run_in_background` + `sleep 30 &&` to avoid blocking (one-shot)\n- Never read full log files — use Read with offset/limit\n- If the Slurm script path doesn't exist, report the error immediately\n- For cluster hardware info (GPU types, partitions, constraints, storage\n  paths), use the `trtllm-agent-toolkit:internal-env-info` skill\n  when it is installed. If the skill is absent, skip the lookup silently and\n  use the caller-supplied fields or the srun probe in Step 1 — do not report\n  the missing skill as an error.\n- Never auto-release an allocation after job completion — it should persist\n  for potential future runs until walltime expires or user explicitly releases\n- Must pass `-N <nodes>` to every direct `srun --jobid` call — without it,\n  srun defaults to 1 node even in a multi-node allocation\n- The container name used in setup steps must **exactly match** the\n  `--container-name` in the `.slurm` script, otherwise pyxis will re-import\n  the image and lose `--container-writable` state\n",".claude/skills/exec-local-compile/SKILL.md":"---\nname: exec-local-compile\ndescription: Compile TensorRT-LLM on a compute node inside a Docker container. Use this when already on a compute node with GPUs visible.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Compile TensorRT-LLM (Local / Compute Node)\n\nCompile TensorRT-LLM from source on a compute node inside a Docker container.\n\n## When to Use\n\n| Scenario | Use This Skill? |\n|----------|----------------|\n| On a compute node with GPUs visible (`nvidia-smi` works) | Yes |\n| On a SLURM login node (no GPUs) | No — use `exec-slurm-compile` instead |\n\n## Prerequisites\n\n- You are inside a Docker/enroot container on a compute node\n- `nvidia-smi` succeeds (GPUs visible)\n- `/usr/local/tensorrt` exists (TensorRT installation in the container)\n\n## Instructions\n\n### Step 1: Verify Environment\n\nRun `nvidia-smi` to confirm you are on a compute node with GPU access.\n\n### Step 2: Locate the Codebase\n\n`cd` to the TensorRT-LLM repository. If the path is not provided by the user, ask for it.\n\n### Step 3: (Optional) Checkout Branch\n\nIf the user specifies a branch (e.g., \"compile ToT\"), checkout and pull:\n```bash\ngit checkout main && git pull\n```\n\n### Step 4: Build\n\nRun the build command (**incremental by default** — omit `-c`/`--clean` unless explicitly requested or the incremental build fails):\n\n```bash\n./scripts/build_wheel.py --use_ccache -a \"<arch>\" -f --nvtx\n```\n\nReplace `<arch>` with the target GPU architecture (see Architecture Reference below). If not specified by the user, auto-detect from `nvidia-smi`.\n\n### Step 5: Install\n\n```bash\npip install -e .[devel]\n```\n\n### Step 6: Verify\n\n```bash\npython3 -c \"import tensorrt_llm; print(tensorrt_llm.__version__)\"\n```\n\n## Build Flags\n\n| Flag | Description |\n|------|-------------|\n| `-a \"<arch>\"` | Target GPU architecture(s) |\n| `--nvtx` | Enable NVTX markers for profiling |\n| `--use_ccache` | Use ccache for faster recompilation |\n| `-f` / `--fast_build` | Skip some kernels for faster dev compilation. **Always use for dev builds.** |\n| `-c` / `--clean` | Clean build directory before building. Only when needed (see below). |\n| `--skip_building_wheel` | Build in-place without creating a wheel file |\n| `--no-venv` | Skip virtual environment creation |\n\n## Architecture Reference\n\n| Value | GPU Family |\n|-------|-----------|\n| `\"100-real\"` | Blackwell (B200, GB200) |\n| `\"90-real\"` | Hopper (H100, H200) |\n| `\"89-real\"` | Ada Lovelace (L40S) |\n| `\"80-real\"` | Ampere (A100) |\n| `\"90;100-real\"` | Multiple architectures |\n\n## Incremental vs. Clean Builds\n\n**Default to incremental builds** — CMake only recompiles changed files, saving significant time.\n\nUse a **clean build** (`-c`) only when:\n- The user explicitly requests a clean/fresh build\n- An incremental build fails with linker errors, stale object files, or CMake cache issues\n- Major branch changes (e.g., rebasing across many commits) that may invalidate the build cache\n- Build system files changed (`CMakeLists.txt`, `*.cmake`)\n",".claude/skills/exec-env-check/SKILL.md":"---\nname: exec-env-check\ndescription: >-\n  Check the local execution environment for GPU availability, Docker support,\n  and Slurm access. Returns the execution scenario (`satisfied, local, docker`,\n  `satisfied, local, direct`, `satisfied, slurm, local`, or `not_satisfied`),\n  the number of available GPUs, and the GPU type. On Slurm login nodes\n  without local GPUs, the cluster is identified by delegating the hostname\n  to internal-env-info (hostname-based mode), which owns the\n  hostname → cluster_name patterns; GPU type and gpus_per_node then come\n  from that skill's reference files. If internal-env-info is not\n  installed, the scenario falls back to `not_satisfied` without probing\n  compute nodes via srun.\ntags: [infrastructure, slurm, environment]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# TensorRT-LLM Environment Check\n\nDetect whether the current machine can run a GPU workload locally (Docker) or via Slurm, and report hardware details.\n\n## Input\n\n| Field | Description | Required |\n|-------|-------------|----------|\n| `required_devices` | Minimum number of GPUs needed | Yes |\n| `account` | Slurm account (unused for GPU probing, kept for compatibility) | No |\n\n## Procedure\n\n### 1. Check local GPUs\n\n```bash\ntimeout 5 nvidia-smi --query-gpu=name,memory.total --format=csv,noheader 2>/dev/null\n```\n\nIf this succeeds:\n- Count the number of GPU lines → `available_gpus`\n- Extract the GPU name from the first line → `device_type` (e.g., `NVIDIA B200`, `NVIDIA H100 80GB HBM3`)\n- Normalize `device_type`: strip `NVIDIA ` prefix and trailing memory info to get the short name (e.g., `B200`, `H100`, `A100`, `L40S`, `RTX 6000`)\n\nIf `nvidia-smi` fails or returns no GPUs → `available_gpus = 0`, `device_type = null`.\n\n### 2. Check if local GPUs are sufficient\n\nIf `available_gpus >= required_devices`, continue to step 2a to determine whether Docker is available on this host.\n\n### 2a. Check Docker availability\n\n```bash\ncommand -v docker >/dev/null 2>&1 && timeout 3 docker info >/dev/null 2>&1\n```\n\n- If the command succeeds (Docker CLI exists **and** the daemon responds) → **Result**: `satisfied, local, docker`\n- Otherwise (Docker CLI missing, daemon not running, or permission denied) → **Result**: `satisfied, local, direct`\n\nInclude `available_gpus` in the result. Do NOT include `device_type` — local execution does not need it.\n\n### 3. Check Slurm availability\n\nIf local GPUs are insufficient (or as additional detection), check for Slurm:\n\n```bash\nwhich squeue 2>/dev/null && squeue --version 2>/dev/null\n```\n\nIf Slurm is NOT available → go to step 5.\n\n### 4. Resolve GPU type\n\n**Optional dependency.** Before doing anything else in this step, check whether `skills/internal-env-info/` exists in the toolkit. If it does **not**, skip this step entirely and return `scenario: not_satisfied` with `available_gpus = 0`, `device_type = null`, `gpus_per_node = null`, `cluster_name = null`, `default_models_repo = null`, and `default_user_root_dir = null`. Do **not** report this as an error — the skill is an internal-only dependency.\n\nWhen Slurm is available but local `nvidia-smi` returned no GPUs or `device_type` is null (login nodes typically have no GPUs), capture the hostname and delegate cluster identification to `internal-env-info`:\n\n```bash\nhostname -f 2>/dev/null || hostname\n```\n\n- Pass the captured hostname to `internal-env-info` in **hostname-based mode**. That skill owns the NVIDIA-internal login-host patterns and the hostname → `<cluster>` mapping; do **not** parse the hostname or hard-code any cluster identifier here.\n- It returns the standard output template (`cluster_name`, `device_type`, `gpus_per_node`, `default_models_repo`, `default_user_root_dir`) plus supplementary field (`mfa_style`).\n- Set `available_gpus` = `gpus_per_node` (if resolved).\n- Set `cluster_name` = the value returned (a placeholder `<cluster>` token in this skill). The orchestrator uses this to fetch per-cluster info (`mfa_style`, `default_models_repo`, `default_user_root_dir`, `gpus_per_node`) from `internal-env-info`; connection fields (`mounts`, `ssh_host`, `partition`, etc.) come from caller-supplied inputs in `job_spec.json` (with `internal-env-info` default values / ask-the-user fallbacks).\n- Set `default_user_root_dir` = the user root directory returned (with `<user_name>` substituted with the actual SLURM username). Set to `null` if not found.\n- If `internal-env-info` returns `null` for `cluster_name` (no pattern matched), set `device_type = null`, `cluster_name = null`, `default_user_root_dir = null` and fall through to Step 5 (`not_satisfied`).\n\n**Result**: `satisfied, slurm, local` with `device_type`, `gpus_per_node`, `cluster_name`, and `default_user_root_dir`.\n\n### 5. Not satisfied\n\nIf neither local GPUs nor Slurm is available:\n- **Result**: `not_satisfied`\n\n## Output\n\nReturn a single structured result:\n\n```\nscenario: <satisfied, local, docker | satisfied, local, direct | satisfied, slurm, local | not_satisfied>\navailable_gpus: <N>\ndevice_type: <short GPU/device name or null>\ngpus_per_node: <N or null>\ncluster_name: <cluster name or null>\ndefault_models_repo: <host path to llm-models directory, or null>\ndefault_user_root_dir: <per-user root directory on cluster storage, or null>\n```\n\n**Examples:**\n\n```\nscenario: satisfied, local, docker\navailable_gpus: 4\n```\n\n```\nscenario: satisfied, local, direct\navailable_gpus: 4\n```\n\n```\nscenario: satisfied, slurm, local\navailable_gpus: 4\ndevice_type: B200\ngpus_per_node: 4\ncluster_name: <cluster>\ndefault_user_root_dir: <default_user_root_dir>/<user_name>\n```\n\n```\nscenario: not_satisfied\navailable_gpus: 0\ndevice_type: null\ncluster_name: null\n```\n\n## Rules\n\n- Never install drivers or modify the system\n- If `nvidia-smi` hangs, use a 5-second timeout: `timeout 5 nvidia-smi ...`\n- GPU type is derived from the hostname by matching the cluster name via the `internal-env-info` skill when it is installed — no `srun` allocation needed. If that skill is absent, the scenario falls back to `not_satisfied` (see Step 4); do not error.\n- Report the GPU type exactly as found in the mapping (do not guess or fabricate)\n",".claude/skills/exec-slurm-compile/SKILL.md":"---\nname: exec-slurm-compile\ndescription: Compile TensorRT-LLM on a SLURM cluster. Covers submitting a batch job with a container image, monitoring the job, and verifying the build. Use when the user wants to compile TRT-LLM remotely via SLURM rather than on a local compute node.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Compile TensorRT-LLM on SLURM Cluster\n\nSubmit, monitor, and verify a TensorRT-LLM compilation job on a SLURM cluster using enroot containers.\n\n## When to Use\n\n| Scenario | Use This Skill? |\n|----------|----------------|\n| User wants to compile TRT-LLM on a SLURM cluster | Yes |\n| User is already on a compute node and wants to compile | No — use `exec-local-compile` skill instead |\n\n## Finding the Docker Image\n\nThe official Docker image tag for a given TensorRT-LLM version is recorded in the repo itself:\n\n```\n<repo_dir>/jenkins/current_image_tags.properties\n```\n\nRead this file to find the current image URL. The format is `<registry-host>/<namespace>/tensorrt-llm:<tag>`, for example `<your-registry>/<your-namespace>/tensorrt-llm:pytorch-25.12-py3-aarch64-ubuntu24.04-trt10.14.1.48-skip-tritondevel-202602011118-10901`. Substitute `<your-registry>` and `<your-namespace>` with your own container registry coordinates.\n\n\n## Pre-dumping the Container Image (enroot import)\n\nSLURM clusters using enroot/pyxis require a `.sqsh` container image. To avoid download overhead at compile time, **pre-dump the image in advance** using the `enroot-import` companion script:\n\n```bash\n# Basic usage — submits a SLURM job on a CPU partition to import the image\nenroot-import --partition cpu_datamover --debug <docker_image_url>\n```\n\nThe script submits an `sbatch` job that runs `enroot import docker://<image_url>` and produces a `.sqsh` file in the current directory. The output on stdout is the SLURM job ID.\n\n### enroot-import flags\n\n| Flag | Description |\n|------|-------------|\n| `-p, --partition` | SLURM partition for the import job (use a CPU partition like `cpu_datamover`) |\n| `-d, --debug` | Enable debug output and preserve the SLURM log (recommended) |\n| `-o, --output` | Custom output path for the `.sqsh` file |\n| `-A, --account` | SLURM account (defaults to user's first account) |\n| `-t, --time` | Time limit for the import job (default: 1 hour) |\n| `-n, --just-print` | Print the sbatch command without executing |\n| `-J, --job-name` | Custom job name |\n\n### enroot-import workflow\n\n1. Read the image tag from `jenkins/current_image_tags.properties` in the TRT-LLM repo.\n2. Run `enroot-import` to submit the import job:\n   ```bash\n   cd <directory_where_sqsh_should_be_stored>\n   <path_to>/enroot-import --partition cpu_datamover --debug <image_url>\n   ```\n   **IMPORTANT:** Replace the first `/` after the registry host with `#` to avoid credential issues. For example, `<your-registry>/<your-namespace>/tensorrt-llm:xxx` becomes `<your-registry>#<your-namespace>/tensorrt-llm:xxx`.\n3. Wait for the import job to complete (`squeue -j <job_id>`).\n4. The resulting `.sqsh` file is the `container_image` used in the compile step.\n\n\n## Prerequisites\n\nThe user must provide (or you must ask for) these values:\n\n| Parameter | Description | Example |\n|-----------|-------------|---------|\n| `container_image` | Path to `.sqsh` container image (see enroot import above) | `/path/to/pytorch.sqsh` |\n| `repo_dir` | Path to the TensorRT-LLM repository | `/path/to/TensorRT-LLM` |\n| `user_root_dir` | Top-level directory to bind-mount into the container | `/shared/users` |\n| `partition` | SLURM partition | `batch` |\n| `account` | SLURM account | `my_account` |\n\nOptional parameters:\n\n| Parameter | Description | Default |\n|-----------|-------------|---------|\n| `jobname` | SLURM job name | `trtllm-compile.<username>` |\n| `gpu_count` | Number of GPUs to request | `4` |\n| `time_limit` | Job time limit | `02:00:00` |\n| `arch` | GPU architecture(s) for `-a` flag | `100-real` |\n| `extra_build_args` | Extra flags for `build_wheel.py` | (none) |\n\n## Companion Scripts\n\nThis skill includes three companion scripts in `scripts/`:\n\n| Script | Purpose |\n|--------|---------|\n| `enroot-import` | Pre-dump a Docker image to `.sqsh` via a SLURM batch job |\n| `submit_compile.sh` | Template for submitting the SLURM job — copy and customize |\n| `compile.slurm` | SLURM batch script — launches the container and calls `compile.sh` |\n| `compile.sh` | Runs inside the container — executes `build_wheel.py` |\n\nScripts directory: `skills/exec-slurm-compile/scripts/`\n\n## Instructions\n\nFollow these steps in order:\n\n### Step 0: Resolve the Container Image (if needed)\n\nIf the user does not already have a `.sqsh` container image:\n\n1. Read the Docker image tag from `<repo_dir>/jenkins/current_image_tags.properties`.\n2. Use `enroot-import` to pre-dump it:\n   ```bash\n   cd <directory_for_sqsh_files>\n   <scripts_dir>/enroot-import --partition cpu_datamover --debug <image_url>\n   ```\n3. Monitor the import job with `squeue -j <job_id>`.\n4. Once complete, the `.sqsh` file path becomes the `container_image` parameter.\n\nIf the user already has a `.sqsh` file, skip this step.\n\n### Step 1: Gather Information\n\nAsk the user for any missing prerequisite values listed above. At minimum you need:\n- `container_image` (or the Docker image URL — then run Step 0 first)\n- `repo_dir`\n- `user_root_dir`\n- `partition` and `account`\n\nIf the user has used this workflow before, check if previous values are stored in memory files.\n\n### Step 2: Prepare the Scripts Directory\n\nThe compile scripts must be accessible from inside the container (i.e., under `user_root_dir`). Either:\n\n**Option A** — Copy companion scripts to a location under `user_root_dir`:\n```bash\nscripts_dir=<user_root_dir>/<username>/workspace/tensorrt_llm_scripts\nmkdir -p ${scripts_dir}/log\ncp skills/exec-slurm-compile/scripts/compile.sh ${scripts_dir}/\ncp skills/exec-slurm-compile/scripts/compile.slurm ${scripts_dir}/\nchmod +x ${scripts_dir}/compile.sh ${scripts_dir}/compile.slurm\n```\n\n**Option B** — If the user already has scripts at a known location, use those directly.\n\n### Step 3: Submit the Job\n\nRun `sbatch` from the login node (or a node with SLURM client access):\n\n```bash\nsbatch \\\n    --nodes=1 --ntasks=1 --ntasks-per-node=1 \\\n    --gres=gpu:<gpu_count> \\\n    --partition=<partition> \\\n    --account=<account> \\\n    --job-name=<jobname> \\\n    --time=<time_limit> \\\n    <scripts_dir>/compile.slurm \\\n    <container_image> <user_root_dir> <scripts_dir> <repo_dir>\n```\n\nCapture and report the job ID from the `sbatch` output.\n\n### Step 4: Monitor the Job (Proactive — Do NOT Wait for User)\n\n**You MUST actively poll the job until it completes.** Do not submit and walk away.\n\n```bash\n# Check job status (repeat every 30-60 seconds)\nsqueue -j <job_id> -o \"%.18i %.9P %.30j %.8u %.2t %.10M %.6D %R\"\n\n# Once running, periodically tail the log (do NOT use tail -f, use tail -30 instead)\ntail -30 <scripts_dir>/log/compile_<job_id>.srun.log\n```\n\n**Monitoring loop:**\n1. Poll `squeue -j <job_id>` to check state\n2. If `PD` (pending) — report the reason, keep polling every 30-60s\n3. If `R` (running) — tail the build log every 30-60s; look for `[XX%] Building`, errors, or completion\n4. If the job disappears from `squeue`, it has finished — proceed to Step 5\n5. If `F` (failed) — immediately read the full log and report the error\n\n**Progress indicators to look for in the log:**\n- `[XX%] Building CXX object...` — compilation progress\n- `Linking CXX...` — link phase\n- `FAILED:`, `error:`, `fatal error:` — build failure\n- `Successfully built` — success\n\n### Step 5: Verify the Build\n\nOnce the job completes, check for success:\n\n```bash\n# Check SLURM exit code\nsacct -j <job_id> --format=JobID,State,ExitCode,Elapsed\n\n# Check the build log for errors\ntail -50 <scripts_dir>/log/compile_<job_id>.srun.log\n```\n\nA successful build ends with a message like `Successfully built tensorrt_llm` or completes without error.\n\n## Common Build Flags Reference\n\n| Flag | Description |\n|------|-------------|\n| `-a \"100-real\"` | Target architecture — `100` for Blackwell, `90` for Hopper, etc. |\n| `--nvtx` | Enable NVTX markers for profiling |\n| `--no-venv` | Skip virtual environment creation |\n| `--use_ccache` | Use ccache to speed up recompilation |\n| `--skip_building_wheel` | Build in-place without creating a wheel file |\n| `-f` | Fast build — skip some kernels for faster dev compilation |\n| `-c` | Clean build — wipe build directory before building |\n\nCommon architecture values:\n- `\"100-real\"` — Blackwell (B200, GB200)\n- `\"90-real\"` — Hopper (H100, H200)\n- `\"89-real\"` — Ada Lovelace (L40S)\n- `\"80-real\"` — Ampere (A100)\n- `\"90;100-real\"` — Multiple architectures\n\n## Troubleshooting\n\n| Issue | Solution |\n|-------|----------|\n| `sbatch: error: invalid partition` | Verify partition name with `sinfo -s` |\n| `sbatch: error: invalid account` | Check available accounts with `sacctmgr show assoc user=$USER` |\n| Container image not found | Verify the `.sqsh` path exists and is readable |\n| Build OOM (out of memory) | Reduce parallelism with `-j <N>` flag to `build_wheel.py` |\n| `srun: error: Unable to create step` | The node may lack enroot/pyxis — check with cluster admin |\n| Job stuck in `PD` state | Check `squeue -j <id> -o %R` for the reason (e.g., resource limits, priority) |\n| `enroot import` fails with auth error | Check `~/.config/enroot/.credentials` has the correct registry credentials |\n| `enroot import` produces empty/corrupt `.sqsh` | Re-run with `--debug` and check the SLURM log; verify the image URL has no `https://` prefix |\n| Weird compile issues | Retry with a clean build (`-c` flag) |\n| `QOSGrpNodeLimit` shown in `NODELIST(REASON)` | Not a blocker, just wait for the job to get scheduled |\n\n## Example Interaction\n\n**User**: \"Compile TRT-LLM on the OCI cluster\"\n\n**Agent actions**:\n1. Ask for container image path, repo path, mount dir (if not known)\n2. Confirm partition/account for OCI cluster\n3. Copy scripts to accessible location under user_root_dir\n4. Submit with `sbatch`\n5. Report job ID\n6. Monitor with `squeue` until complete\n7. Check logs and report success/failure\n",".claude/skills/exec-local-docker/SKILL.md":"---\nname: exec-local-docker\ndescription: >-\n  Execute a TensorRT-LLM workload locally in Docker. Runs a fully-resolved\n  Docker command in background, monitors completion, reads logs, and reports\n  results. Workflow-agnostic — does not need to know if the workload is pytest,\n  eval, benchmark, or a custom script.\ntags: [docker, execution, infrastructure]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Local Docker Executor\n\nRun a Docker command locally, monitor it, and report results.\n\n## Input (from orchestrator prompt)\n\nThe orchestrator passes these fields in the skill prompt:\n\n| Field | Description |\n|-------|-------------|\n| `docker_cmd` | Complete `docker run` command string, ready to execute |\n| `work_dir` | Local work directory for logs and artifacts |\n| `log_file` | Full path to the log file (output redirected here) |\n| `model_name` | Short model name for reporting |\n| `workflow_type` | `pytest`, `eval`, `custom`, or `benchmark` — for output parsing hints |\n| `success_patterns` | Comma-separated patterns indicating success (e.g., `passed,accuracy:`) |\n| `failure_patterns` | Comma-separated patterns indicating failure (e.g., `FAILED,Error,AssertionError`) |\n\n## Procedure\n\n### Step 0: Resolve Image and Build (when `build_project=true`)\n\nThis executor owns image selection and the build for the local Docker target. Skip the entire step when `build_project=false`.\n\n1. **Detect the target GPU type.** Use `gpu_type` from `job_spec.json` if upstream env-check resolved it; otherwise probe locally with `nvidia-smi --query-gpu=name --format=csv,noheader | head -1`.\n2. **Detect the host CPU arch** with `uname -m` (`x86_64` or `aarch64`).\n3. **Resolve the container image.** Read `<repo_root>/jenkins/current_image_tags.properties` and pick the tag whose CPU-arch flavor matches the host. If the orchestrator already passed a `container_image` field in the job spec, use that and skip the lookup.\n4. **Map GPU → build arch (`-a` flag):** `H100`/`H200` → `90-real`; `B200`/`GB200`/`B300`/`GB300` → `100-real`; `A100` → `80-real`; `L40S` → `89-real`. Default `100-real` when the GPU is unknown.\n5. **Compile.** Invoke the `exec-local-compile` skill with `repo_dir=<repo_root>`, `image=<resolved tag>`, `arch=<arch>`. Wait for completion.\n6. **On failure**, do not launch the workload. Report `BUILD_FAILED` with the last 100 lines of the compile log.\n\n`build_project`, `gpu_type`, `repo_root`, and (optionally) `container_image` come from `job_spec.json`.\n\n### Step 1: Launch\n\nRun the Docker command in background using `run_in_background`:\n\n```bash\n<docker_cmd> 2>&1 | tee <log_file>\n```\n\nReport to the orchestrator: \"Launched locally, log at `<log_file>`\"\n\n### Step 2: Monitor for Hangs\n\nWhile waiting for the background process to complete, actively monitor the log\nfile for hang indicators. Launch a monitoring loop using `run_in_background`:\n\n```bash\nwhile true; do\n  sleep 60\n  if [ -f \"<log_file>\" ] && grep -qi \"hang detected\" \"<log_file>\"; then\n    echo \"HANG_DETECTED: Found 'hang detected' in log file\"\n    docker ps --filter \"ancestor=<image>\" -q | xargs -r docker kill 2>/dev/null\n    exit 1\n  fi\ndone\n```\n\n- If the monitor detects a hang, it kills the Docker container and exits with\n  code 1. The main background process will also terminate.\n- When the main process completes normally (background notification received),\n  kill the monitoring loop (it is no longer needed).\n- If a hang is detected, skip to Step 4 and report `HANG_DETECTED` status\n  instead of proceeding to normal result collection.\n\n### Step 3: Wait for Completion\n\nThe Bash tool's `run_in_background` will notify when the process finishes\n(either normally or because the container was killed by the hang monitor).\n\n### Step 4: Read Results\n\nOn completion:\n\n1. **Read exit code** from the background command result.\n2. **Read the last 100 lines** of `<log_file>` using the Read tool.\n3. **If exit code != 0**, also read the first 50 lines to catch early errors (import failures, setup crashes).\n4. **Search for patterns**:\n   - Grep `<log_file>` for each `success_patterns` entry\n   - Grep `<log_file>` for each `failure_patterns` entry\n\n### Step 5: Report\n\nReturn a structured result:\n\n```\nStatus: PASSED | FAILED | ERROR | HANG_DETECTED\nExit code: <N>\nLog file: <log_file>\nWork directory: <work_dir>\nSummary: <relevant output lines — pytest summary, accuracy score, or last few lines>\nErrors: <if failed, the relevant error output>\n```\n\n### Output Parsing by Workflow Type\n\n- **pytest**: Look for `X passed, Y failed in Zs` summary line\n- **eval**: Look for `accuracy:` or `score:` lines; check for `Expected accuracy >= X, but got Y` assertion\n- **custom**: No specific patterns — report last 10 lines of output\n- **benchmark**: Look for throughput/latency numbers\n\n## Rules\n\n- Never run the Docker command in foreground — always use `run_in_background`\n- Never `cat` the full log file — use Read with offset/limit or tail\n- If the Docker command fails immediately (exit code within seconds), check if the image exists locally\n- Report results even if the log file is empty (container may have failed to start)\n",".claude/skills/exec-remote-slurm/SKILL.md":"---\nname: exec-remote-slurm\ndescription: >-\n  Remote SLURM cluster development via SSH. Use when running jobs,\n  profiling, or developing on a remote SLURM cluster with pyxis/enroot\n  containers. Covers SSH connection management, srun/sbatch/salloc job\n  patterns, tmux-based allocation persistence, file transfer, and safe\n  remote file access. Works with any SLURM cluster accessible via SSH.\ntags: [infrastructure, remote, slurm]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Remote SLURM Development\n\nDevelop and execute on remote SLURM clusters via SSH commands. No filesystem\nsync needed — all remote operations go through SSH.\n\n## Step 1: Resolve Cluster\n\n**Optional dependency — `internal-env-info`.** Before doing\nanything else in Step 1, check whether `skills/internal-env-info/`\nexists in the toolkit. If it does **not**, skip every sub-step that invokes\nthe skill or reads its `references/` files and proceed using only the\ncluster fields the orchestrator wrote into `job_spec.json` (`ssh_host`,\n`slurm_user`, `remote_cwd`, `partition`, `account`, `container_image`,\n`mounts`, `gpus_per_node`, `mfa_style`). If a required field is also\nmissing, stop and ask the user to supply it — do **not** emit an error\nabout the missing skill.\n\nThe orchestrator supplies one of two input shapes; pick the right path:\n\n- **Explicit cluster** — `slurm_cluster` name is provided. Skip\n  auto-selection and go straight to **Load the per-cluster info**.\n- **Auto-select cluster** — only hardware constraints are provided\n  (`device_type`, `total_required_devices`, `required_devices_per_node`).\n  Resolve a cluster from the constraints, then load its per-cluster info.\n\n### 1a. Auto-select (only when no explicit cluster was provided)\n\nInvoke the `internal-env-info` skill in **constraint-based mode**, passing\n`device_type`, `required_devices_per_node`, and (optionally)\n`total_required_devices`. That skill owns the filter rules and either\nreturns a single resolved cluster entry or stops and asks the user. Do not\nimplement the filter here.\n\nUse `node_count` from `job_spec.json` directly as `--nodes` and\n`required_devices_per_node` as `--ntasks-per-node`. **Do not recompute\n`node_count`** — `trtllm-case-executor` is the single source of truth and\nwrites the resolved value into `job_spec.json` for every executor to read.\n\nThe returned entry's `cluster_name`, `mfa_style`, `default_user_root_dir`,\n`default_models_repo`, and `gpus_per_node` feed Step 1b.\n\n### 1b. Load the per-cluster info\n\nConnection fields are read **only** from `job_spec.json` (written by\n`trtllm-case-executor`): `ssh_host`, `slurm_user`, `remote_cwd`, `account`,\n`partition`, `mounts`, `container_image`, `gpus_per_node`, and `packages`.\nThere is no per-cluster connection-config file to parse — if a required\nfield is absent from `job_spec.json`, stop and ask the user to supply it.\nDo not silently default.\n\nInvoke the `internal-env-info` skill in single-cluster mode with\nthe resolved `cluster_name` to retrieve per-cluster info: partition limits,\ntopology-aware scheduling guidance, `default_models_repo`,\n`default_user_root_dir`, `gpus_per_node`, `mfa_style` (used by Step 2 to\nchoose the SSH path — see below), and shared constraint and tuning\npatterns. The runtime `container_image` always comes from `job_spec.json`,\nnever from this skill.\n\n## Step 2: SSH Preflight\n\nThe cluster's `mfa_style` (returned by `internal-env-info` in Step 1b) is\nthe **single source of truth for which SSH path to use**. Switch on the\nvalue — do **not** probe-and-fall-back on the SSH error string.\n\n| `mfa_style` | Path | Why |\n|-------------|------|-----|\n| `false` | Direct (Step 2a) | Cluster accepts key-based / passwordless SSH; no special handling needed. |\n| `true`  | MFA (Step 2b)    | Cluster requires interactive MFA (Duo / device-code / one-time PIN). |\n| `null`  | Unknown (Step 2c)| `mfa_style` not declared in the cluster reference (or `internal-env-info` is not installed). Probe direct first, fall back to MFA. When the reference exists, also ask the user to update it so future runs skip the probe; otherwise just continue silently. |\n\n### 2.0. Build `<ssh_host>` from `slurm_user` + `cluster_name`\n\nBefore any SSH attempt, derive the SSH destination. Priority order:\n\n1. **Caller-supplied `ssh_host`** (passed in through case-executor) → use\n   verbatim, skip everything below. This is the canonical escape hatch.\n2. **`internal-env-info` skill installed** → delegate for the\n   `ssh_host_pattern` key (defined in the skill's `Default Environment Values`\n   section). It returns a single-line snippet that concatenates `slurm_user`\n   and `cluster_name` into `SSH_HOST`. Do **not** re-derive or re-suffix\n   `slurm_user` here — `internal-env-info` already applied whatever its\n   `slurm_user` construction rule requires for that cluster.\n\n   Inputs to substitute:\n   - `slurm_user` — pre-resolved from `internal-env-info` (already in its\n     final, SSH-ready form for that cluster). Read it from\n     `job_spec.json`.\n   - `cluster_name` — from `internal-env-info` Step 1b.\n3. **Skill not installed** → stop and ask the user to supply `ssh_host`\n   explicitly. Do **not** hard-code a login-host pattern in this skill —\n   the NVIDIA-internal convention lives only in `internal-env-info`.\n\nThe resolved value is the canonical `<ssh_host>` placeholder used everywhere\nin the rest of this skill (Steps 2a/2b/2c, Step 3, recipes, file-transfer,\nhang detection).\n\n### 2a. Direct path — `mfa_style: false`\n\nVerify connectivity using the user's existing SSH config:\n\n```bash\nssh -o ConnectTimeout=5 <ssh_host> \"echo ok\"\n```\n\nIf this succeeds, **skip to Step 3**. Do not create a trtllm-managed\nControlMaster — reuse whatever the user already has.\n\nIf it fails, do **not** silently switch to the MFA path. Stop and ask the\nuser to verify their `ssh_host` value or their key-based SSH setup.\nA direct-path cluster failing here is a configuration issue, not an MFA\nprompt to interpret.\n\n### 2b. MFA path — `mfa_style: true`\n\nFirst check whether a ControlMaster is already live (the user may have set\none up earlier in this session):\n\n```bash\nssh -O check <ssh_host> 2>&1 || ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -O check <ssh_host> 2>&1\n```\n\nIf either succeeds, **skip to Step 3** — reuse it.\n\nOtherwise the user must complete an interactive MFA handshake. MFA uses\nbrowser-based device-code verification that needs an interactive SSH\nsession — the skill cannot complete this on the user's behalf because the\nsession must stay alive while the user authenticates in their browser and\npresses ENTER.\n\nTell the user to run the ControlMaster command:\n\n> Run this in your terminal:\n>\n>     ssh -fNM -o ControlMaster=auto -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -o ControlPersist=24h <ssh_host>\n>\n> If the cluster uses a device-code flow you will see a PIN and a URL like:\n>\n>     Authenticate with PIN XXXXXX at <your-identity-provider-device-url>\n>\n> Open that URL, enter the PIN, approve, then press ENTER. The SSH session\n> will background itself as a persistent ControlMaster. Let me know when done.\n\nAfter the user confirms, verify with\n`ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -O check <ssh_host>` and\nproceed to Step 3.\n\n### 2c. Unknown — `mfa_style: null`\n\nThe cluster reference (`references/<cluster>.md`) does not declare an\n`mfa_style` value. Try the direct path (Step 2a) first. If it fails with\n`Permission denied (keyboard-interactive)` — and only that — switch to\nthe MFA path (Step 2b). Any other failure stops here and is reported to\nthe user.\n\nAfter connecting successfully, **ask the user to update the cluster\nreference file** with the discovered value (`mfa_style: true` or\n`false`, in the `## Authentication` table) so that future runs skip\nthis probe and use the deterministic switch in Step 2a / 2b.\n\n## SSH Command Shorthand\n\nEvery SSH and SCP invocation in the rest of this skill uses the placeholders below. Both pin `ControlPath` to the trtllm-managed socket so they reuse the ControlMaster set up in Step 2 — without this, OpenSSH starts a fresh handshake on every call, which on MFA clusters re-prompts for keyboard-interactive auth (and defeats the whole point of multiplexing):\n\n| Placeholder | Expansion |\n|-------------|-----------|\n| `<ssh_cmd>` | `ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p` |\n| `<scp_cmd>` | `scp -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p` |\n\nWhen the user has their own ControlMaster configured in `~/.ssh/config` (no trtllm-managed socket needed), the explicit `-o ControlPath` is harmless: if the path has no live master, OpenSSH falls back to a fresh connection. The connectivity probe in Step 2 Check 2 (`ssh -o ConnectTimeout=5 <ssh_host> \"echo ok\"`) deliberately uses bare `ssh` to test un-multiplexed access; everywhere else use `<ssh_cmd>` / `<scp_cmd>`.\n\n## Step 3: Remote Environment Check\n\nVerify the remote environment is ready:\n\n```bash\n<ssh_cmd> <ssh_host> \"test -d <remote_cwd> && echo ok\"\n<ssh_cmd> <ssh_host> \"which tmux\"\n<ssh_cmd> <ssh_host> \"squeue --version\"\n```\n\nIf `remote_cwd` does not exist, create it or ask the user. If tmux or SLURM\ncommands are not available, stop and inform the user.\n\n## Step 4: Prepare Remote Environment\n\nThis step sets up the user's working directory and clones the TensorRT-LLM\nrepository on the remote cluster. It runs once before any job submission.\n\n> **MANDATORY: two-mechanism sync — both must run.**\n> - **Tracked changes** → temp branch (**4c–4f**): commit on a session-scoped\n>   branch, push to `origin`, fetch + checkout on remote, delete the ref both\n>   sides.\n> - **Untracked files** (`git ls-files --others --exclude-standard`) →\n>   `tar | ssh` pipe (**4e.1**), right after the remote checkout. Never touches\n>   `origin`.\n>\n> **Do not** shortcut with `git fetch origin <repo_branch> && git checkout <repo_branch>` —\n> `repo_branch` may not exist on `origin`, may be stale, or may lack uncommitted\n> edits. Always run 4c–4f **and** 4e.1.\n\n**Inputs received from `trtllm-case-executor`:**\n\n| Field | Description |\n|-------|-------------|\n| `default_user_root_dir` | Per-user root directory on the remote cluster (from `exec-env-check` / `internal-env-info`) |\n| `repo_url` | Git remote URL of the local TensorRT-LLM repo (from `git remote get-url origin`). **Used to clone the bare repository on the remote (Step 4b) and to push the temp branch to (Step 4d).** |\n| `repo_branch` | **Informational only** — the local current branch name (from `git rev-parse --abbrev-ref HEAD`). Used in Step 4c to remember which branch to restore locally after the temp branch is deleted. **Never used as the ref the remote checks out** — the remote always checks out the temp branch from Step 4c. |\n| `device_type` | Required short GPU name (e.g., `B200`, `B300`). Used by Step 1a auto-selection. Ignored when an explicit cluster was provided. |\n| `total_required_devices` | Total GPU processes across all nodes. Used by Step 1a auto-selection only. SLURM job sizing reads `node_count` from `job_spec.json` (precomputed by `trtllm-case-executor`) — do not recompute it here. |\n| `required_devices_per_node` | Minimum GPUs per node. Used by Step 1a auto-selection (cluster must have `gpus_per_node >= required_devices_per_node`) and as `--ntasks-per-node`. |\n\n**Procedure:**\n\n### 4a. Create user root directory\n\n```bash\n<ssh_cmd> <ssh_host> \"mkdir -p <default_user_root_dir>\"\n```\n\nIf `default_user_root_dir` is null or empty, skip this step and use `remote_cwd` directly.\n\n### 4b. Clone or update the repository\n\nDetermine the repo directory name from `repo_url` (strip `.git` suffix if present):\n\n```bash\nREPO_DIR_NAME=$(basename <repo_url> .git)\nREMOTE_REPO_PATH=\"<default_user_root_dir>/${REPO_DIR_NAME}\"\n```\n\nCheck if it already exists:\n\n```bash\n<ssh_cmd> <ssh_host> \"test -d <remote_repo_path>/.git && echo exists\"\n```\n\nIf the repository does **not** exist, clone it:\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <default_user_root_dir> && git clone <repo_url>\"\n```\n\nIf the repository **already exists**, refresh remote refs only — do **not**\ncheckout or reset to any branch here:\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <remote_repo_path> && git fetch origin --prune\"\n```\n\nThe actual ref the remote will check out is the temp branch created in 4c\nand pushed in 4d; the remote checkout happens in 4e. **Do not** add a\n`git checkout <repo_branch>` or `git reset --hard origin/<repo_branch>`\nhere as a shortcut — that path is forbidden (see the MANDATORY callout\nat the top of Step 4).\n\n### 4c. Create a temporary local branch capturing all local changes\n\nGenerate a unique temp branch name tied to this session. The branch name\n**must start with `<user_name>_<model_name>`** so it is easy to identify\nthe owner and the workload it was created for. A timestamp suffix keeps\nit unique across sessions:\n\n```bash\n# Resolve the user name (prefer the slurm_user input over whoami,\n# so the temp branch is attributed to the cluster user rather than the\n# host that happens to dispatch the job).\nUSER_NAME=\"${slurm_user:-$(whoami)}\"\n\n# Resolve the workload name. job_spec.json's `model_name` is authoritative;\n# fall back to the repo basename if it is missing.\nMODEL_NAME=\"${model_name:-$(basename \"<REPO_ROOT>\")}\"\n\n# Sanitize to a git-ref-safe identifier (alphanum + `_`/`-` only).\n# Any other character is replaced with `_`.\nsanitize() { printf '%s' \"$1\" | tr -c 'A-Za-z0-9_-' '_'; }\nUSER_NAME_SAFE=$(sanitize \"$USER_NAME\")\nMODEL_NAME_SAFE=$(sanitize \"$MODEL_NAME\")\n\nTEMP_BRANCH=\"${USER_NAME_SAFE}_${MODEL_NAME_SAFE}_$(date +%Y%m%d%H%M%S)\"\n```\n\nExample resolved value: `<user_name>_dwdp_deepseek_v3_lite_20260430145200`.\n\nStash the current branch name so it can be restored later:\n\n```bash\nORIGINAL_BRANCH=$(git -C \"<REPO_ROOT>\" rev-parse --abbrev-ref HEAD)\n```\n\nCreate the temp branch at HEAD and stage tracked changes only\n(`-u`, not `-A` — untracked files ship in 4e.1):\n\n```bash\ngit -C \"<REPO_ROOT>\" checkout -b \"$TEMP_BRANCH\"\ngit -C \"<REPO_ROOT>\" add -u\ngit -C \"<REPO_ROOT>\" commit --allow-empty -m \"sync: temp snapshot for remote execution\"\n```\n\n`--allow-empty` covers the no-tracked-changes case; the branch still exists\nat HEAD and 4e.1 still runs.\n\n### 4d. Push temp branch to the remote origin\n\nPush the temp branch so the remote cluster can fetch it:\n\n```bash\ngit -C \"<REPO_ROOT>\" push origin \"$TEMP_BRANCH\"\n```\n\n### 4e. Fetch and check out temp branch on the remote cluster\n\nOn the remote: stash prior-session leftovers, fetch, force-checkout the temp\nbranch, and re-materialize `cpp/` from HEAD (stale `cpp/` produces\nmissing-symbol `ImportError`s in Step 5):\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <remote_repo_path> && \\\n    git stash push -u -m 'auto-stash before remote sync' || true && \\\n    git fetch origin --prune && \\\n    git checkout -B $TEMP_BRANCH origin/$TEMP_BRANCH && \\\n    rm -rf cpp/ && \\\n    git checkout HEAD -- cpp/\"\n```\n\n- `stash -u` sweeps both tracked edits and untracked files left from a prior\n  session; `|| true` covers the no-op.\n- `checkout -B` force-resets even if a same-named branch already exists.\n- `rm -rf cpp/ && git checkout HEAD -- cpp/` always runs, even when checkout\n  reports no diffs.\n\n`remote_cwd` is updated to `<remote_repo_path>`. Untracked files for this\nsession arrive next in 4e.1.\n\n### 4e.1. Transfer untracked files via tar pipe over SSH\n\nShip untracked, non-gitignored files on top of the freshly checked-out tree:\n\n```bash\nUNTRACKED_COUNT=$(git -C \"<REPO_ROOT>\" ls-files --others --exclude-standard | wc -l)\nif [ \"$UNTRACKED_COUNT\" -gt 0 ]; then\n    git -C \"<REPO_ROOT>\" ls-files --others --exclude-standard -z | \\\n        tar -C \"<REPO_ROOT>\" --null -T - -cf - | \\\n        <ssh_cmd> <ssh_host> \"tar -C <remote_repo_path> -xf -\"\nfi\n```\n\n`-z` / `--null` keeps filenames with spaces or newlines intact; `tar -x`\nauto-creates parent dirs; the guard skips the round-trip when nothing to\nship. Do **not** substitute `rsync` or `sshfs` (not reliably available on HPC).\n\n### 4f. Delete the temp branch (local and remote)\n\nAfter the remote checkout succeeds, immediately delete the temp branch from\nthe local repo and the remote origin, then restore the original local branch:\n\n```bash\ngit -C \"<REPO_ROOT>\" checkout \"$ORIGINAL_BRANCH\"\ngit -C \"<REPO_ROOT>\" branch -D \"$TEMP_BRANCH\"\ngit -C \"<REPO_ROOT>\" push origin --delete \"$TEMP_BRANCH\"\n```\n\nThe remote cluster keeps the checked-out working tree from the deleted branch\n— only the ref is removed, not the files.\n\n### 4f.1 Create remote work_dir under the cloned repo\n\nThe remote work_dir must mirror the local layout (`<REPO_ROOT>/work_dirs/<name>`)\nso that all per-run artefacts — slurm script, stdout/stderr logs, intermediate\nfiles — live inside the cloned repo and are co-located with the code that\nproduced them.\n\n`remote_work_dir` is precomputed by `trtllm-case-executor` (Step 1) and read\nverbatim from `job_spec.json` — it is `<remote_repo_path>/work_dirs/<basename(local_work_dir)>`.\n**Do not** fall back to `<remote_cwd>/work_dirs/...` or any other root; that\nolder layout is deprecated.\n\n```bash\n<ssh_cmd> <ssh_host> \"mkdir -p <remote_work_dir>\"\n```\n\nAll subsequent steps (file uploads, sbatch submission, log paths) reference\n`<remote_work_dir>` from `job_spec.json`.\n\n### 4g. Copy build scripts to remote node\n\nThe skill ships two build scripts in `scripts/` alongside this SKILL.md.\nCopy them to a dedicated directory on the remote so they are accessible\nfrom inside the SLURM container:\n\n```bash\nREMOTE_SCRIPTS_DIR=\"<default_user_root_dir>/trtllm-build-scripts\"\nSKILL_SCRIPTS_DIR=\"<LOCAL_SKILL_DIR>/scripts\"   # local path to this skill's scripts/\n\n<ssh_cmd> <ssh_host> \"mkdir -p ${REMOTE_SCRIPTS_DIR}/log\"\n\n<ssh_cmd> <ssh_host> \"cat > ${REMOTE_SCRIPTS_DIR}/build.slurm\" < \"${SKILL_SCRIPTS_DIR}/build.slurm\"\n<ssh_cmd> <ssh_host> \"cat > ${REMOTE_SCRIPTS_DIR}/build.sh\"    < \"${SKILL_SCRIPTS_DIR}/build.sh\"\n<ssh_cmd> <ssh_host> \"chmod +x ${REMOTE_SCRIPTS_DIR}/build.slurm ${REMOTE_SCRIPTS_DIR}/build.sh\"\n```\n\n`<LOCAL_SKILL_DIR>` is the directory containing this SKILL.md file\n(`.claude/skills/exec-remote-slurm/`).\n\n---\n\n## Step 5: Build on Remote SLURM Node\n\n> **Mandatory unless `build_project: false`.** Step 4e wipes `cpp/`, so\n> the remote has no compiled C++ extensions until Step 5 succeeds (exit\n> code `0:0`). Submitting any test/bench job before then fails on the\n> first `import tensorrt_llm` with missing-symbol errors.\n>\n> Do **not** replace this with `pip install -e .` inside the user's\n> `.slurm` script — it does not reliably recompile `cpp/`, and when it\n> does it blocks the GPU allocation behind a 30+ minute compile. Step 5\n> runs the build as a 1-node side job and installs into the persistent\n> container volume.\n>\n> If the caller sets `build_project: false`, skip Step 5 (the caller is\n> asserting the install is current). Do not silently downgrade `true` →\n> `false` because Step 5 looks expensive.\n\nSubmit a SLURM batch job that compiles TensorRT-LLM on the remote cluster\ninside the official container. Wait for it to complete before proceeding to\njob submission.\n\n### 5a. Read the container image\n\nThe container image is resolved by `trtllm-case-executor` (Step 3) and\nforwarded in `job_spec.json` as `container_image`. Read it from there —\n**do not grep `current_image_tags.properties` and do not re-derive from\nthe cluster config**. case-executor is the single source of truth for the\nimage URI; this skill only consumes it.\n\nApply the enroot URI rewrite at use-time so the registry path is not\nmistaken for a Docker-style credential lookup. The rewrite recipe lives\nin the `internal-env-info` skill's `Default Environment Values` section\n(`enroot_uri_rewrite` key):\n\n- **Skill installed** → delegate to it, passing the canonical\n  `<container_image>` from `job_spec.json`. The skill returns the bash\n  snippet (an `sed` over the registry-host prefix); run it here to\n  produce `<IMAGE_PATH>`.\n- **Skill not installed** → use `<container_image>` verbatim as `<IMAGE_PATH>`.\n\nUse `<IMAGE_PATH>` as the value passed to `--container-image=` in the\nbuild sbatch below and in every `srun`/`sbatch` recipe later in this\nskill. The canonical (un-rewritten) URI in `job_spec.json` is unchanged;\nthe rewrite is purely a transport adaptation for enroot.\n\n### 5b. Derive build architecture\n\nMap `device_type` (from the env-check output) to the `build_wheel.py` `-a` flag:\n\n| `device_type` | `-a` value |\n|------------|------------|\n| B200, GB200, GB300, B300 | `100-real` |\n| H100, H200 | `90-real` |\n| L40S | `89-real` |\n| A100 | `80-real` |\n\nDefault to `100-real` if `device_type` is unknown.\n\n### 5c. Submit the build job\n\n```bash\n<ssh_cmd> <ssh_host> \"sbatch \\\n    --nodes=1 --ntasks=1 --ntasks-per-node=1 \\\n    --partition=<partition> \\\n    --account=<account> \\\n    --job-name=<account>-trtllm.build \\\n    --time=02:00:00 \\\n    ${REMOTE_SCRIPTS_DIR}/build.slurm \\\n    <IMAGE_PATH> <default_user_root_dir> ${REMOTE_SCRIPTS_DIR} <remote_repo_path> \\\n    -a \\\"<arch>\\\" --nvtx\"\n```\n\nCapture the job ID from `sbatch` stdout (format: `Submitted batch job <id>`).\n\nIf `gpus_per_node` is set in the cluster config, add\n`--gpus-per-node=<gpus_per_node>` to the `sbatch` flags.\n\n### 5d. Monitor the build job\n\nPoll every 30–60 seconds until the job leaves the queue:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <job_id> -o '%.18i %.9P %.30j %.8u %.2t %.10M %R'\"\n```\n\nOnce running, tail the log to track compile progress:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -30 ${REMOTE_SCRIPTS_DIR}/log/build_<job_id>.srun.log\"\n```\n\nProgress indicators in the log:\n- `[XX%] Building CXX object ...` — compilation in progress\n- `Linking CXX ...` — link phase\n- `FAILED:` / `error:` / `fatal error:` — build failure\n- `Successfully built` — success\n\nIf the job disappears from `squeue`, proceed to Step 5e.\n\n### 5e. Verify the build\n\n```bash\n<ssh_cmd> <ssh_host> \"sacct -j <job_id> --format=JobID,State,ExitCode,Elapsed -n\"\n<ssh_cmd> <ssh_host> \"tail -50 ${REMOTE_SCRIPTS_DIR}/log/build_<job_id>.srun.log\"\n```\n\nA successful build shows exit code `0:0` and ends with\n`Successfully built tensorrt_llm` or no error lines.\n\nIf the build fails, report the last 100 lines of the log to the user and\nstop — do not proceed to job submission.\n\n## Build Scripts Reference\n\n| Script | Location | Purpose |\n|--------|----------|---------|\n| `build.slurm` | `scripts/build.slurm` | SLURM batch script — launches container, calls `build.sh` |\n| `build.sh` | `scripts/build.sh` | Runs inside container — calls `build_wheel.py` with defaults |\n\nBoth scripts accept the same positional arguments:\n`<container_image> <user_root_dir> <scripts_dir> <repo_dir> [build_wheel_args...]`\n\n---\n\n## Recipes\n\nPick the recipe that matches your workflow. Each recipe is self-contained —\ncopy it, substitute the `<placeholders>` from config, and run.\n\n**Placeholders** (all from `job_spec.json` unless noted):\n\n| Placeholder | Source |\n|-------------|--------|\n| `<ssh_host>`, `<account>`, `<partition>`, `<mounts>` | `job_spec.json` |\n| `<image>` | `container_image` from `job_spec.json` (resolved by `trtllm-case-executor` Step 3). Apply enroot rewrite (`/` → `#`) at use-time per Step 5a — substitute the rewritten `<IMAGE_PATH>` here. Do not re-resolve from the cluster config or `current_image_tags.properties`. |\n| `<job_name>` | Pre-resolved by `trtllm-case-executor` and forwarded as `job_name` in `job_spec.json`. Substitute directly into `-J <job_name>`; do not reconstruct from `account` / `subproject` / `detail`. |\n| `<nodes>` | `node_count` from `job_spec.json` (computed by `trtllm-case-executor`). Do not recompute. |\n| `<gpus>` | `required_devices_per_node` from `job_spec.json` — probe only if absent (see bottom of this section) |\n| `<remote_repo_path>` | `remote_repo_path` from `job_spec.json` — `<default_user_root_dir>/<repo_dir_name>` precomputed by `trtllm-case-executor` Step 1. The cloned repo and `--container-workdir` for srun. Do not recompute. |\n| `<remote_work_dir>` | `remote_work_dir` from `job_spec.json` — `<remote_repo_path>/work_dirs/<basename(local_work_dir)>` precomputed by `trtllm-case-executor` Step 1. Used for `--output`/`--error` slurm log paths and any per-run artefact uploads. Do not recompute or fall back to `<remote_cwd>/work_dirs/...`. |\n| `<time>` | Walltime in HH:MM:SS |\n\n**Job naming:** Every recipe includes `-J <job_name>` using the\npre-resolved `job_name` value from `job_spec.json` (built by\n`trtllm-case-executor`). This is required — without `-J`, SLURM emits\nwarnings and jobs are hard to identify in `squeue`. Do not reconstruct\nthe name from `account` / `subproject` / `detail` here; that convention\nlives in `trtllm-case-executor`.\n\n**GPU allocation:** Cluster behavior varies:\n\n- **Exclusive-access clusters** (most NVIDIA internal): Nodes are allocated\n  exclusively with all GPUs available. No `--gpus-per-node` needed.\n- **Shared/QOS clusters**: Require `--gpus-per-node=<gpus_per_node>`. Without\n  it, you may get a CPU-only node (no GPUs at all), a \"Cannot find GPU\n  specification\" error, or a \"QOSMinGRES\" error (below minimum).\n\nIf `gpus_per_node` is set in the cluster config, add `--gpus-per-node=<gpus_per_node>`\nto every `srun`, `salloc`, and `sbatch` command. If not set, omit it.\n\nIf a job fails with GPU-related errors and `gpus_per_node` is not configured,\nprobe the node's GPU count, set `gpus_per_node` in the cluster config, and retry.\n\nDo not set `CUDA_VISIBLE_DEVICES` for MPI workloads — the MPI runtime handles\nGPU assignment.\n\n---\n\n### Recipe 1: Single-process command (srun)\n\nOne-off commands on a single node: profiling, inspecting, lightweight scripts.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N 1 --ntasks-per-node=1 \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <command>\"\n```\n\n---\n\n### Recipe 2: MPI workload (srun)\n\nMPI applications across nodes: NCCL tests, HPCX benchmarks. The binary must\nbe MPI-aware (e.g., `all_reduce_perf_mpi`, not `all_reduce_perf`).\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=<gpus> --mpi=pmix \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <command>\"\n```\n\n---\n\n### Recipe 3: Distributed training (srun)\n\nFramework-launched distributed training (PyTorch DDP, DeepSpeed). SLURM starts\n`<nodes> × <gpus>` processes. Pyxis/enroot sets `RANK`, `LOCAL_RANK`, and\n`WORLD_SIZE`. Each process runs the training script directly.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=<gpus> \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  python train.py\"\n```\n\n---\n\n### Recipe 4: Launcher-based training (srun)\n\nOne process per node, an in-container launcher forks locally (torchrun,\ndeepspeed launcher).\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=1 \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  torchrun --nproc-per-node=<gpus> train.py\"\n```\n\n---\n\n### Recipe 5: Batch job (sbatch)\n\nFire-and-forget background job:\n\n```bash\n<ssh_cmd> <ssh_host> \"sbatch -A <account> -p <partition> -N <nodes> \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <script-path>\"\n```\n\nCheck status:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <jobid>\"\n<ssh_cmd> <ssh_host> \"sacct -j <jobid> --format=JobID,State,ExitCode,Elapsed -n\"\n```\n\n---\n\n### Recipe 6: Multi-command session (salloc + srun)\n\nInteractive development with persistent container state. Use this when you need\nmultiple commands on the same allocation — installing packages, compiling,\nrunning experiments iteratively.\n\n> **Relationship to `persistent_mode`:** The `persistent_mode` field in\n> `job_spec.json` is consumed only by `exec-local-slurm`. This\n> skill does **not** read `persistent_mode`, `container_name`, or\n> `container_mounts` from `job_spec.json` for allocation reuse — remote\n> persistence is driven by this recipe (tmux + named `salloc` + named\n> container) and controlled by inputs to this skill directly (e.g.,\n> `release_allocation`, `alloc_time_limit`). If a caller forwards a\n> `persistent_mode=true` job spec into a `remote_slurm` run, treat the field\n> as informational only and use this recipe's flow.\n\n**Step 1: Allocate via tmux**\n\ntmux keeps the allocation alive on the frontend — invisible plumbing.\n\n```bash\n<ssh_cmd> <ssh_host> \"tmux new-session -d -s trtllm-<id> \\\n  'salloc -A <account> -p <partition> -N <nodes> \\\n  -J <job_name> -t <time>'\"\n```\n\n**Step 2: Check allocation status**\n\nWait ~30 seconds for the scheduler, then get job ID and status:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -u <user> -h -o '%i %T %S' --state=RUNNING,PENDING\"\n```\n\nIf RUNNING → proceed to Step 3.\n\nIf PENDING → follow the **walltime negotiation flow** (see Walltime strategy\nsection). Only proceed to Step 3 once the job is RUNNING.\n\n**Step 3: Create a persistent container**\n\nImport the image once. This creates a named container on each node that\npersists across `srun` calls (container persistence):\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-image=<image> --container-name=trtllm-<jobid> true\"\n```\n\n**Container packages:** If `packages` is defined in config, install them\nusing the exact names listed — do not guess or substitute package names:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> --container-writable \\\n  pip install <packages>\"\n```\n\nOnly install when the workload actually needs them — not for simple\ncommands like `nvidia-smi` or `hostname`.\n\n**Step 4: Run commands**\n\nUse `--container-name` (without `--container-image`) to reuse the persistent\ncontainer. All in-container state is preserved between commands.\n\nSingle-process command:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nPer-node operation (mkdir, pip install) — use `--ntasks-per-node=1`:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nMPI workload within the allocation:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --ntasks-per-node=<gpus> --mpi=pmix \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nDistributed training within the allocation:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --ntasks-per-node=<gpus> \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  python train.py\"\n```\n\nAdd `--container-writable` if the workload writes to container paths outside\nof mounted directories (e.g., `pip install` into site-packages).\n\n**Step 5: Release the allocation**\n\nClean up the named container, then kill tmux. If the allocation has already\nexpired (walltime exceeded), the `srun` will fail — that is fine, the container\nis already gone. Just kill the tmux session.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  bash -c 'enroot remove -f trtllm-<jobid> 2>/dev/null; true'\" 2>/dev/null\n<ssh_cmd> <ssh_host> \"tmux kill-session -t trtllm-<id>\"\n```\n\n---\n\n### Probing GPUs per node\n\nIf you do not know `<gpus>`, probe once and record in cluster notes:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> bash -c 'nvidia-smi -L | wc -l'\"\n```\n\n## Cluster-Specific Flags\n\nAfter substituting recipe placeholders, apply the cluster-specific flags\nreturned by the `internal-env-info` skill in Step 1:\n\n- **SHARP**: Add `--network=sharp` if the cluster supports it\n- **Topology constraints**: Add `--constraint=<feature>` for topology-aware\n  placement (e.g., `--constraint=nvlblk01` for a single rack)\n- **GPU clocks**: Run `sudo nvidia-smi -ac <mem>,<graphics>` as a separate\n  `srun --ntasks-per-node=1` pre-step in sbatch scripts\n- **Comments**: Add `--comment=metrics` for Prometheus/Grafana collection\n\n## File Transfer\n\nUse SSH pipes for file transfer — no extra tools needed, works on every cluster.\n\n### Local → Remote (writing files)\n\nPipe file content directly via SSH heredoc. Never write locally then scp:\n\n```bash\n<ssh_cmd> <ssh_host> \"cat > <remote_cwd>/script.py\" << 'PYEOF'\nimport torch\n# ... file content ...\nPYEOF\n```\n\nFor multiple files, repeat the pattern — each file is one SSH command.\n\n### Remote → Local (reading results)\n\nRead remote files via SSH stdout:\n\n```bash\n<ssh_cmd> <ssh_host> \"cat <remote_cwd>/output.json\"\n```\n\nFor large files, read only what you need:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/slurm-<jobid>.out\"\n<ssh_cmd> <ssh_host> \"head -n 50 <remote_cwd>/results.csv\"\n```\n\n### When to use scp\n\nFall back to `scp` only for binary files or large directories that cannot be\npiped through SSH:\n\n```bash\n<scp_cmd> <ssh_host>:<remote_cwd>/profile.nsys-rep .    # binary profile\n<scp_cmd> -r local_dir/ <ssh_host>:<remote_cwd>/        # whole directory\n```\n\n**Decision rule:** text files → SSH pipe. Binary files or directories → scp.\nNever use rsync or sshfs — not reliably available on HPC clusters.\n\n## Rules\n\n### Never read full remote files\n\nRemote output can be gigabytes. Always use `tail`, `head`, or `grep`:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/output.log\"\n<ssh_cmd> <ssh_host> \"head -n 50 <remote_cwd>/script.py\"\n<ssh_cmd> <ssh_host> \"grep -n 'error\\|Error\\|ERROR' <remote_cwd>/output.log\"\n```\n\nNever use `cat` on remote files unless you are certain the file is small.\n\n### Never run heavy computation on the frontend\n\nThe frontend node is shared. Only use it for SLURM commands, lightweight file\noperations, and tmux management.\n\n### Check before creating\n\nBefore creating resources, check if they already exist:\n- SSH ControlMaster: `ssh -O check <ssh_host> 2>&1 || <ssh_cmd> -O check <ssh_host>` (probe both the user's default socket and the trtllm-managed socket from Step 2)\n- tmux sessions: `<ssh_cmd> <ssh_host> \"tmux has-session -t trtllm-<id> 2>/dev/null && echo exists\"`\n- Existing allocations: `<ssh_cmd> <ssh_host> \"squeue -u <user> -h -o '%i %j' --state=RUNNING\"`\n\n### Clean up when done\n\n- Remove named containers: `enroot remove -f` (see Recipe 6, Step 5)\n- Kill tmux sessions: `<ssh_cmd> <ssh_host> \"tmux kill-session -t trtllm-<id>\"`\n- Cancel unneeded jobs: `<ssh_cmd> <ssh_host> \"scancel <jobid>\"`\n\n### Walltime strategy\n\n- **Production (sbatch, srun):** Tight walltime with ~20% margin — gets\n  scheduled sooner.\n- **Dev/debug (salloc):** Use **walltime negotiation** — submit with max\n  walltime, then reduce if the queue wait is too long:\n\n**Walltime negotiation flow:**\n\n1. Submit `salloc` with **max walltime** for the partition.\n2. Wait ~30 seconds for the scheduler to process, then check:\n   ```bash\n   <ssh_cmd> <ssh_host> \"squeue -j <jobid> -h -o '%i %T %S'\"\n   ```\n   - `%S` = estimated start time. Parse it to calculate wait.\n3. **If RUNNING** → proceed to next step.\n4. **If PENDING and wait ≤ 10 minutes** → wait for it, proceed.\n5. **If PENDING and wait > 10 minutes** (or no start time shown):\n   a. Infer how long you actually need from context (profiling → ~30 min,\n      iterative dev → ~2h, training run → match expected duration).\n      Round up to the next hour boundary.\n   b. Reduce walltime in-place:\n      ```bash\n      <ssh_cmd> <ssh_host> \"scontrol update JobId=<jobid> TimeLimit=<new_time>\"\n      ```\n   c. Tell the user: \"Reduced walltime from Xh to Yh for faster scheduling —\n      you'll have less time on the allocation.\"\n   d. Wait ~15 seconds, re-check estimated start time.\n   e. **If still > 10 minutes:** inform user with the ETA and suggest\n      alternatives — different partition (e.g., backfill), fewer nodes.\n\n## Hang Detection Monitoring\n\nWhen monitoring a running job (polling via `squeue`), also check the remote log\nfile every 60 seconds for the pattern `hang detected` (case-insensitive). This\ncatches GPU hangs, NCCL deadlocks, and other stalls that keep the job running\nindefinitely without making progress.\n\n**Check for hang (run alongside each `squeue` poll once job is RUNNING):**\n\n```bash\n<ssh_cmd> <ssh_host> \"grep -qi 'hang detected' <remote_work_dir>/slurm-<jobid>.out 2>/dev/null && echo HANG_DETECTED\"\n```\n\nAlso check the `.err` file if the output file has no match:\n\n```bash\n<ssh_cmd> <ssh_host> \"grep -qi 'hang detected' <remote_work_dir>/slurm-<jobid>.err 2>/dev/null && echo HANG_DETECTED\"\n```\n\n**If `HANG_DETECTED` is printed:**\n\n1. Immediately cancel the job:\n   ```bash\n   <ssh_cmd> <ssh_host> \"scancel <jobid>\"\n   ```\n2. Read the last 200 lines of the log file for context:\n   ```bash\n   <ssh_cmd> <ssh_host> \"tail -n 200 <remote_work_dir>/slurm-<jobid>.out\"\n   ```\n3. Report status as `HANG_DETECTED` with the relevant log lines. Do not\n   continue polling — the job is terminated.\n\n## Recovering from Failures\n\n**SSH connection drops:** Re-run SSH preflight (Step 2). tmux-held allocations\nsurvive disconnects.\n\n**Job stuck in PENDING:**\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <jobid> -o '%i %T %r %S'\"\n```\n\nCheck the reason (`%r`) and estimated start (`%S`). If wait > 10 minutes,\nfollow the **walltime negotiation flow** (see Walltime strategy). If already\nnegotiated and still stuck, suggest alternatives: backfill partition, fewer\nnodes, or a different cluster.\n\n**Job fails:** Read last lines, never dump full output:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/slurm-<jobid>.out\"\n<ssh_cmd> <ssh_host> \"tail -n 50 <remote_cwd>/slurm-<jobid>.err\"\n```\n\n**Hang detected:** If the log file contains `hang detected`, the job is\nauto-cancelled (see Hang Detection Monitoring section above). Report\n`HANG_DETECTED` with the last 200 lines of log context.\n",".claude/skills/kernel-cute-writing/SKILL.md":"---\nname: kernel-cute-writing\ndescription: >\n  Write and implement GPU kernels using NVIDIA CuTe DSL (CUTLASS 4.x Python\n  API) — NOT for Triton, CUDA C++, or conceptual explanations.\n  Trigger only when the user wants to write or implement a kernel, not when\n  asking questions about CuTe DSL concepts or layouts.\n  CuTe DSL uses cute.jit/cute.kernel decorators and cutlass.cute imports.\n  Covers element-wise kernels, GEMM patterns, reductions, memory hierarchy\n  (global/shared/register/TMA), MMA tensor core operations, software\n  pipelining, and framework integration.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# CuTe DSL\n\nCuTe DSL is a Python-based domain-specific language for GPU kernel development,\npart of CUTLASS 4.x. It provides Python abstractions over CUTLASS C++ templates\nwith JIT compilation to optimized CUDA kernels via MLIR and ptxas.\n\n## When to Use\n\n**Triggers:**\n- Writing CUDA kernels in Python (element-wise, GEMM, custom ops)\n- Optimizing GPU memory access patterns (vectorized loads, TMA, shared memory)\n- Building tensor core (MMA) kernels for Ampere/Hopper/Blackwell\n- Integrating custom GPU kernels with PyTorch or JAX\n- Prototyping high-performance kernels without C++ metaprogramming\n\n**Symptoms (wrong tool otherwise):**\n- Need shared memory coordination or tensor core MMA → use CuTe DSL (not Triton for complex patterns)\n- Need simple element-wise ops with no shared memory → CuTe DSL or Triton both work\n- Need to call existing CUTLASS C++ kernels → use CUTLASS C++ APIs instead\n- Need reductions, scans, or non-GEMM collective ops → consider CUB/Thrust\n\n**Keywords:** cute, cutlass, cute.jit, cute.kernel, from_dlpack, zipped_divide,\nTiledMMA, TiledCopy, TMA, WGMMA, tcgen05, pipeline, mbarrier\n\n## Requirements\n\n| Requirement | Detail |\n|-------------|--------|\n| Platform | Linux x86_64 only |\n| Python | 3.10–3.13 |\n| GPU | NVIDIA Ampere+ (SM80, SM90, SM100) |\n| CUDA Driver | ≥ 575.51.03 (Toolkit 12.9 compat) |\n| Install | `pip install nvidia-cutlass-dsl` |\n| Optional | `apache-tvm-ffi`, `torch-c-dlpack-ext` |\n\n## Workflows\n\n### Workflow 0: Starting from Examples (Recommended)\n\nFor any non-trivial kernel (GEMM, attention, pipelined, fused ops), start by\nfinding the most similar existing example to use as a **starting point** — study\nits structure, then rework it for your use case. Do not copy examples verbatim;\nthey target specific dtypes, architectures, and problem shapes that likely differ.\n\n1. **Pick the closest example** from the index below.\n   **Prefer examples matching the target GPU architecture** (check with\n   `torch.cuda.get_device_capability()`) when the operation is similar.\n\n   Fetch via `web_fetch` with base URL\n   `https://raw.githubusercontent.com/NVIDIA/cutlass/main/examples/python/CuTeDSL`\n\n   | Operation | Arch | Example path (append to base URL) |\n   |-----------|------|-----------------------------------|\n   | Element-wise add | SM80 | `ampere/elementwise_add.py` |\n   | Element-wise + autotune | SM80 | `ampere/elementwise_add_autotune.py` |\n   | Element-wise apply | SM80 | `ampere/elementwise_apply.py` |\n   | SGEMM (scalar) | SM80 | `ampere/sgemm.py` |\n   | Tensor-core GEMM | SM80 | `ampere/tensorop_gemm.py` |\n   | Flash Attention v2 | SM80 | `ampere/flash_attention_v2.py` |\n   | HSTU Attention | SM80 | `ampere/hstu_attention.py` |\n   | Shared memory allocator | SM80 | `ampere/smem_allocator.py` |\n   | CTA norm (LayerNorm) | SM90 | `hopper/cta_norm.py` |\n   | Dense GEMM | SM90 | `hopper/dense_gemm.py` |\n   | Dense GEMM persistent | SM90 | `hopper/dense_gemm_persistent.py` |\n   | Flash MHA | SM90 | `hopper/fmha.py` |\n   | Dense GEMM | SM100 | `blackwell/dense_gemm.py` |\n   | Dense GEMM persistent | SM100 | `blackwell/dense_gemm_persistent.py` |\n   | Dense GEMM + alpha/beta | SM100 | `blackwell/dense_gemm_alpha_beta_persistent.py` |\n   | RMSNorm | SM100 | `blackwell/rmsnorm.py` |\n   | Reduce | SM100 | `blackwell/reduce.py` |\n   | Flash MHA | SM100 | `blackwell/fmha.py` |\n   | Grouped GEMM | SM100 | `blackwell/grouped_gemm.py` |\n   | Mamba2 SSD | SM100 | `blackwell/mamba2_ssd/` |\n   | GEMM tutorial (notebook) | SM100 | `notebooks/tour_to_sol_gemm.ipynb` |\n\n   **Example:** To fetch the Hopper dense GEMM:\n   ```bash\n   web_fetch https://raw.githubusercontent.com/NVIDIA/cutlass/main/examples/python/CuTeDSL/hopper/dense_gemm.py\n   ```\n\n2. **Read reference materials first** — before diving into example code, read\n   the relevant `references/` docs to understand the patterns and APIs:\n   - For GEMM: `references/patterns-gemm.md` (3-level tiling, epilogue fusion,\n     `cute.compile` with `mark_layout_dynamic`, shared memory layouts)\n   - For reductions: `references/patterns-reduction.md` (warp reductions,\n     `cute.compile` cache pattern)\n   - For element-wise: `references/patterns-elementwise.md` (variations A–E)\n   - Always: `references/api-arch.md` (available APIs, arch-specific caveats)\n\n   This gives you the conceptual foundation so you can rework the example\n   intelligently rather than trying to copy-paste complex pipelines.\n\n3. **Fetch and study the example source** — read for structure, not to copy:\n   - Identify: decorators, tiling strategy, shared memory usage, mainloop flow\n   - Note which dtype/arch it targets (many examples are fp16/bf16-specific)\n   - Check if it uses APIs tied to a specific arch (TMA → SM90+, tcgen05 → SM100)\n\n4. **Rework for the user's workload** (do not copy-paste):\n   - Change shapes, data types, tile sizes to match requirements\n   - Replace compute logic (epilogue, activation fusion) as needed\n   - If dtype differs (e.g., example is fp16, need fp32), expect vectorization\n     and layout changes — the scalar-loop patterns in `references/` may be a\n     better starting point than adapting a vectorized example\n   - **Runtime wrapper must be lightweight**: `kernel_fn()` should only call\n     `from_dlpack()` + the compiled kernel. Never allocate intermediate tensors,\n     copy data, or re-compile per call — these belong in one-time setup\n   - Apply optimizations from this skill's reference docs\n\n   **⛔ Blackwell/Hopper GEMM + extra tensors — STOP:**\n   If the target GPU is SM90+ (Hopper/Blackwell) **and** the GEMM requires\n   extra tensors beyond A, B, C in the epilogue (e.g., bias vector, activation\n   inputs), **do not attempt it**. These examples use TMA descriptors for all\n   data movement — adding tensors requires modifying TMA descriptor setup,\n   which is prohibitively complex. Instead, tell the user this limitation and\n   suggest a **two-kernel approach**: run the GEMM kernel as-is, then apply\n   bias + activation in a separate element-wise kernel (Workflow 1).\n   Plain GEMM (just A×B→C with scalar alpha/beta) on Hopper/Blackwell is fine.\n\n5. **Validate and benchmark** using companion scripts:\n   ```bash\n   python scripts/verify_kernel.py kernel.py --rtol 1e-3 --atol 1e-3\n   python scripts/benchmark_kernel.py kernel.py\n   ```\n   The kernel file must export `kernel_fn`, `reference_fn`, and `get_inputs()`.\n\n**When to skip examples:** Pure element-wise operations (Workflow 1) have\ncomplete patterns in `references/patterns-elementwise.md` — no need to fetch\nexternal examples.\n\n**Reduction kernels** (softmax, layernorm, RMSNorm): Use\n`references/patterns-reduction.md` which provides complete, proven patterns\nfor float32 reductions using scalar loops + butterfly shuffle + shared memory.\n\n### Workflow 1: Element-wise Kernel\n\nFor unary/binary/in-place operations that map inputs to outputs 1:1.\n\n1. **Determine kernel structure**: inputs/outputs count, tensor rank, target arch\n2. **Select pattern** from `references/patterns-elementwise.md` (Variations A–E)\n3. **Write kernel** applying all four invariant principles:\n   - P1: `from_dlpack(tensor, assumed_align=16)` for vector loads\n   - P2: Derive `vec_size` from `element_type.width`\n   - P3: `cute.zipped_divide(mA, tiler)` for coalesced access\n   - P4: `cutlass.dynamic_expr(thread_idx < total)` for bounds\n4. **Critical rules**: No early return, no `a * 2` (use `a + a`), no `cute.math.sigmoid`\n5. **Pre-compile with `cute.compile()`**: Always pre-compile the kernel once\n   using `cute.compile()` so that `kernel_fn` calls the compiled object, not\n   `@cute.jit` directly.  Without pre-compilation, every call recompiles\n   (~20-50ms overhead).  Use `.mark_layout_dynamic()` so a single compiled\n   kernel handles arbitrary input shapes without recompilation:\n   ```python\n   # Compile once with dynamic layouts — works for any shape\n   fake_x = from_dlpack(torch.empty(1, 1, dtype=torch.float16, device=\"cuda\"),\n                         assumed_align=16).mark_layout_dynamic()\n   fake_out = from_dlpack(torch.empty(1, 1, dtype=torch.float16, device=\"cuda\"),\n                           assumed_align=16).mark_layout_dynamic()\n   compiled_kernel = cute.compile(host_fn, fake_x, fake_out)\n\n   def kernel_fn(x):\n       out = torch.empty_like(x)\n       compiled_kernel(from_dlpack(x, assumed_align=16).mark_layout_dynamic(),\n                       from_dlpack(out, assumed_align=16).mark_layout_dynamic())\n       return out\n   ```\n6. **Verify correctness** using companion script:\n   ```bash\n   python scripts/verify_kernel.py kernel.py --rtol 1e-3 --atol 1e-3\n   ```\n   The kernel file must export `kernel_fn`, `reference_fn`, and `get_inputs()`.\n7. **Benchmark** using companion script:\n   ```bash\n   python scripts/benchmark_kernel.py kernel.py\n   ```\n\n### Workflow 2: GEMM Kernel\n\nFor matrix multiplication with tiling, shared memory, and tensor cores.\n\n1. **Define problem**: shapes (M, N, K), data types, target architecture\n2. **Choose tiling**: CTA tile (bM, bN, bK), pipeline stages, cluster shape\n3. **Three-level partitioning** (see `references/patterns-gemm.md`):\n   - Level 1: CTA tiling with `local_tile()`\n   - Level 2: Copy partitioning (global → shared) with `TiledCopy`\n   - Level 3: Compute partitioning (shared → register) with `TiledMMA`\n4. **Shared memory**: Use swizzled layouts (`make_smem_layout_atom`) to avoid bank conflicts\n5. **Mainloop**: K-tile loop with copy → sync → MMA → sync\n6. **Pipeline**: Use `PipelineTmaAsync` (Hopper) or `PipelineTmaUmma` (Blackwell).\n   ⚠️ TMA-based pipelines manage data movement via TMA descriptors — adding\n   extra tensors (bias, activation inputs) to the epilogue requires modifying\n   descriptor setup, which is prohibitively complex. See the stop condition in\n   Workflow 0 step 4.\n7. **Epilogue**: Predicated store with alpha/beta scaling\n8. **Pre-compile with `cute.compile()`**: Always pre-compile the GEMM kernel\n   so `kernel_fn` calls the compiled object, not `@cute.jit` directly.\n   Without pre-compilation, every call recompiles (~20-50ms overhead).\n9. **Autotune**: Search over tile sizes, cluster shapes, pipeline depths\n\n### Workflow 3: Framework Integration\n\nFor wrapping CuTe DSL kernels as PyTorch/JAX custom operators.\n\n1. **Write kernel** using Workflow 1 or 2\n2. **Create wrapper**: Accept `torch.Tensor`, convert via `from_dlpack`, call host fn\n3. **For production**: Compile with TVM FFI for zero-overhead tensor passing:\n   ```python\n   compiled = cute.compile(host_fn, *fake_tensors, options=\"--enable-tvm-ffi\")\n   compiled(torch_a, torch_b)  # Direct torch.Tensor, no from_dlpack\n   ```\n4. **For deployment**: Use AOT compilation → export to `.o` → load at runtime\n\n### Workflow 4: Debugging & Profiling\n\n1. **Set environment**: `CUTE_DSL_PRINT_IR=1`, `CUTE_DSL_KEEP_PTX=1`\n2. **Use `cute.printf()`** for runtime values (not Python `print`)\n3. **Inspect generated code**: `compiled.__ptx__`, `compiled.__mlir__`\n4. **Profile**: Enable `CUTE_DSL_LINEINFO=1`, use Nsight Compute/Systems\n5. **Debug memory**: Run with `compute-sanitizer python script.py`\n\n## Output Formats\n\nA typical CuTe DSL kernel project:\n\n```\nkernel_dir/\n  kernel.py          # @cute.kernel + @cute.jit functions\n  test_kernel.py     # Correctness test vs PyTorch reference\n  bench_kernel.py    # Benchmark with cute.compile() setup\n```\n\n**Success indicators:**\n- Correctness test passes (`torch.testing.assert_close`)\n- Nsight shows vector loads (LDG.128/LDG.256), not scalar loads\n- For GEMM: tensor core utilization > 80% in Nsight Compute\n\n## Companion Script Contract\n\nKernel files used with `scripts/verify_kernel.py` and `scripts/benchmark_kernel.py`\nmust export three names:\n\n- `kernel_fn(*inputs)` — the CuTe DSL kernel wrapper (calls `cute.compile` + runs kernel)\n- `reference_fn(*inputs)` — PyTorch reference implementation (same signature)\n- `get_inputs()` — returns a list of CUDA tensors for testing\n\n```python\n# Example kernel.py contract\nimport torch\nimport cutlass.cute as cute\nfrom cutlass.cute.runtime import from_dlpack\n\ndef kernel_fn(x):\n    out = torch.empty_like(x)\n    # ... call compiled cute kernel ...\n    return out\n\ndef reference_fn(x):\n    return torch.nn.functional.gelu(x)\n\ndef get_inputs():\n    return [torch.randn(1024, 512, dtype=torch.float16, device=\"cuda\")]\n```\n\n## Examples\n\n### Example: 2D Unary Element-wise (ReLU)\n\n```python\nimport torch, cutlass, cutlass.cute as cute\nfrom cutlass.cute.runtime import from_dlpack\n\n@cute.kernel\ndef relu_kernel(gA: cute.Tensor, gC: cute.Tensor):\n    tidx, _, _ = cute.arch.thread_idx()\n    bidx, _, _ = cute.arch.block_idx()\n    bdim, _, _ = cute.arch.block_dim()\n    idx = bidx * bdim + tidx\n    m, n = gA.shape[1]\n    total = m * n\n    if cutlass.dynamic_expr(idx < total):\n        a = gA[(None, (idx // n, idx % n))].load()\n        gC[(None, (idx // n, idx % n))] = cute.where(a > 0, a, 0)\n\n@cute.jit\ndef relu_host(mA: cute.Tensor, mC: cute.Tensor):\n    vec = 16 // (mA.element_type.width // 8)\n    gA = cute.zipped_divide(mA, (1, vec))\n    gC = cute.zipped_divide(mC, (1, vec))\n    T = 256\n    N = cute.size(gA.shape[1])\n    relu_kernel(gA, gC).launch(grid=((N+T-1)//T,1,1), block=(T,1,1))\n\nx = torch.randn(1024, 512, dtype=torch.float16, device=\"cuda\")\nout = torch.empty_like(x)\nrelu_host(from_dlpack(x, assumed_align=16), from_dlpack(out, assumed_align=16))\n```\n\n## Error Handling\n\n| Error | Cause | Fix |\n|-------|-------|-----|\n| `MLIR function requires a Context` | Called @kernel from Python | Launch via @cute.jit host function |\n| `DSLAstPreprocessorError` on return | Early return in @kernel | Use `if cutlass.dynamic_expr(cond):` |\n| Type mismatch on store | `a * 2` promotes FP16→FP32 | Use `a + a` or `.to(cutlass.Float16)` |\n| `could not get source code` | Kernel in `exec()` context | Write to file and import |\n| Scalar loads in Nsight | Missing alignment hint | Add `assumed_align=16` to `from_dlpack` |\n| `Missing required argument` | Not all @jit params passed | Pass ALL declared parameters |\n| `AttributeError: sigmoid` | No `cute.math.sigmoid` | Use `1.0/(1.0+cute.math.exp(-x))` |\n\nSee `references/troubleshooting.md` for the full error table and limitations.\n\n**Debugging rule:** Never delete kernel.py during debugging. Use `backup_file`\nto save a checkpoint, then `edit_file` to iterate. If stuck, `revert_file` to\nrestore the backup. A partially-working kernel is always better than no kernel.\n\n## Finding More Information\n\n### Tier 1: This File (SKILL.md)\n\nWorkflows above cover element-wise kernels, GEMM, framework integration, and\ndebugging. Search this file first for procedural questions.\n\n### Tier 2: references/ Directory\n\nGrep for keywords across `references/`. Headers are grep-friendly.\n\n| File | Content |\n|------|---------|\n| `concepts-architecture.md` | Core abstractions, terminology, compilation pipeline |\n| `concepts-layouts.md` | Layout algebra: composition, complement, divide, swizzle |\n| `concepts-tensors.md` | Tensor types, partitioning, tiling, predication |\n| `concepts-mma.md` | MMA atoms, TiledMMA, per-architecture tensor core ops |\n| `patterns-getting-started.md` | Installation, decorators, first kernel walkthrough |\n| `patterns-elementwise.md` | Invariant principles, pattern variations, reference impl |\n| `patterns-gemm.md` | 3-level tiling, shared memory, pipelining, autotuning |\n| `patterns-memory.md` | from_dlpack, TMA, cp.async, TMEM, copy atoms |\n| `patterns-compilation.md` | Control flow, JIT caching, TVM FFI, AOT compilation |\n| `patterns-pipeline.md` | Producer-consumer, pipeline classes, barriers, warp specialization |\n| `api-core.md` | cute module: layouts, tensors, math, copy, gemm, printing |\n| `api-arch.md` | cute.arch: thread indexing, sync, atomics, memory ops |\n| `api-nvgpu.md` | cute.nvgpu: warp/warpgroup/cpasync/tcgen05 MMA and copy |\n| `api-runtime-utils.md` | Runtime: from_dlpack, fake tensors, utils, schedulers |\n| `troubleshooting.md` | Debugging, env vars, common errors, limitations, FAQ |\n\n**How to search:** Grep for your keyword across `references/`. Read only the\nfile and section that Grep points to.\n\n### Tier 3: Original Documentation\n\nIf Tiers 1–2 don't answer, consult the source:\n- **Web**: https://docs.nvidia.com/cutlass/latest/\n- **GitHub**: https://github.com/NVIDIA/cutlass\n- Fetch specific doc pages or search for \"CUTLASS CuTe DSL <topic>\"\n- Consider distilling the answer back into references/\n",".claude/skills/perf-host-analysis/SKILL.md":"---\nname: perf-host-analysis\ndescription: >\n  Analyze host/CPU overhead in TensorRT-LLM inference from nsys traces.\n  Detect whether host overhead is the bottleneck using GPU idle ratio, host prep\n  exposed ratio, and per-phase evidence. For regressions, isolate forward steps\n  via allreduce/NVTX patterns, compare host operation breakdowns across versions,\n  and identify scheduling or request-management overhead. Supports optional\n  inter-kernel gap, eager-vs-graph, pattern mapping, and multi-rank straggler\n  drill-down. Use standalone or within perf-analysis. Triggers: host overhead,\n  inter-step gap, scheduling overhead, forward step isolation, nsys iteration\n  analysis, NVTX breakdown, request management overhead, GPU idle, host bottleneck,\n  host prep exposed, inter-kernel gap, bubble analysis, graph coverage, eager kernel,\n  rank imbalance, straggler detection.\nlicense: Apache-2.0\ntags:\n  - analysis\n  - profiling\n  - host-overhead\n  - nsys\n  - inference\ndependencies:\n  - workload-instrumentation\n  - trace-interpretation\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Host Performance Analysis\n\nAnalyze host/CPU overhead in TensorRT-LLM inference workloads from nsys traces. This skill operates in two phases:\n\n| Phase | Question | Input | Output |\n|-------|----------|-------|--------|\n| **Detection** | Is host overhead the bottleneck? | Single nsys trace | YES/NO verdict with metric evidence |\n| **Root Cause** | What specifically regressed? | One or two nsys traces | NVTX per-step breakdown, regression sources, optional kernel-level drill-down |\n\n## When to Use\n\n- Before starting host optimization work -- confirms the bottleneck is real (Detection)\n- As a sub-step of `perf-analysis` for bottleneck classification (Detection)\n- When GPU utilization is suspiciously low and you need to know why (Detection)\n- When throughput regressed but GPU kernel execution times are unchanged (Root Cause)\n- When the gap between forward step iterations has increased (Root Cause)\n- To compare inter-iteration overhead between two versions of the inference engine (Root Cause)\n- When you need sub-operation granularity on inter-kernel gaps or graph coverage (Root Cause, kernel-level drill-down)\n- When piecewise CUDA graph coverage is unexpectedly low (Root Cause, kernel-level drill-down)\n- When multi-rank inference shows unexplained performance asymmetry (Root Cause, kernel-level drill-down)\n\nDo NOT use when:\n- The regression is in individual kernel performance (use `perf-nsight-compute-analysis`)\n- You need to profile a workload from scratch (use `workload-instrumentation` first)\n- The issue is NCCL communication (use distributed analysis)\n\n## Prerequisites\n\n- An nsys trace file (`.sqlite` or `.nsys-rep`) from a TRT-LLM benchmark run\n- For Root Cause comparison: two traces (baseline and target)\n- Python 3 with sqlite3 support\n\n---\n\n## Key Concepts\n\n### Host Overhead in LLM Inference\n\nIn an LLM inference loop, each iteration consists of:\n```\n[inter-step gap] -> [_forward_step] -> [inter-step gap] -> [_forward_step] -> ...\n```\n\nThe **forward step** includes GPU kernel execution (GEMM, attention, normalization, allreduce) plus host-side preparation. The **inter-step gap** includes host-side work between forward steps (scheduling, request fetching, broadcasting, sampling, response handling).\n\nSee [references/trtllm-nvtx-ranges.md](references/trtllm-nvtx-ranges.md) for the full per-operation breakdown and timing ranges.\n\n### Hidden vs Exposed Host Overhead\n\nHost overhead only hurts performance when it is **exposed** -- the GPU is idle waiting for work. When host prep overlaps with GPU execution, it is **hidden** and free. See [references/metrics.md](references/metrics.md) (M3 section) for diagrams and the exposed/hidden computation.\n\n### Forward Step Isolation\n\nIn TP configurations, forward steps are isolated via allreduce kernel grouping (deterministic count per transformer layer). For TP=1, NVTX `_forward_step` ranges are used directly. See [references/iteration-isolation-techniques.md](references/iteration-isolation-techniques.md) for the full algorithm.\n\n### Phase Classification (Context vs Generation)\n\nIterations are classified by NVTX marker text into **context** (eager, no CUDA graphs) and **generation** (CUDA graph replay). Per-phase analysis is critical because aggregate metrics can mask phase-specific bottlenecks. See [references/phase-classification.md](references/phase-classification.md).\n\n---\n\n## Phase 1: Detection (YES/NO Verdict)\n\nDetermine whether host overhead is the primary bottleneck.\n\n### Detection Metrics\n\nSix metrics in four categories. See [references/metrics.md](references/metrics.md) for full definitions, formulas, and SQL queries.\n\n| # | Metric | Threshold | What it answers |\n|---|--------|-----------|-----------------|\n| M1 | GPU idle ratio | > 0.30 | Is the GPU starved for work? |\n| M2 | Launch overhead ratio | > 0.10 | Is kernel launch itself expensive? |\n| M3a | Host prep exposed ratio | > 0.50 | How well is host prep pipelined? |\n| M3b | Host prep perf impact | > 0.05 | How much throughput does exposed prep cost? |\n| M3c | Host prep idle attribution | > 0.50 | Is host prep the main cause of GPU idle? |\n| M4 | GPU utilization | < 0.60 | Is GPU utilization too low? |\n| M5 | NCCL ratio (caveat) | > 0.20 | Is communication a confounding factor? |\n\n**Host prep confirmation rule**: Host prep is a confirmed bottleneck only when **both** M3b AND M3c cross their thresholds.\n\nThresholds are configurable with per-phase variants. See [references/thresholds.md](references/thresholds.md).\n\n### Detection Workflow\n\n#### Step 1: Input Validation\n\n```bash\n# Accept .sqlite or .nsys-rep\nls -la <trace_file>\n\n# If .nsys-rep, export to SQLite first\nnsys export -t sqlite -o <output.sqlite> <input.nsys-rep>\n```\n\n#### Step 2: Run Detection Script\n\n```bash\npython scripts/detect_host_overhead.py \\\n  --trace /path/to/trace.sqlite \\\n  --output /path/to/verdict.json\n```\n\nThe script computes M1, M2, M4, M5 from SQL, optionally M3 via range intersection, applies the verdict logic, and outputs structured JSON. See [references/output-format.md](references/output-format.md) for the output schema.\n\nFor manual metric extraction via SQL, see [references/nsys-schema.md](references/nsys-schema.md).\n\n#### Step 3: Interpret Verdict\n\n**Overall Verdict:**\n```\nif aggregate_verdict == YES or context_verdict == YES or generation_verdict == YES:\n    overall_verdict = YES\n```\n\nPer-phase analysis can **elevate** the verdict but never **demote** it.\n\nFormat using the template in [references/output-format.md](references/output-format.md).\n\n**Next Steps:**\n- **If YES** -> Proceed to Phase 2 (Root Cause) below, then use `perf-host-optimization` skill\n- **If NO** -> Use `perf-nsight-compute-analysis` for kernel SOL% or `perf-analysis` for full classification\n\n---\n\n## Phase 2: Root Cause Analysis\n\nIdentify which specific host operations regressed and by how much. Works with a single trace (breakdown) or two traces (comparison).\n\n### Principles\n\n1. **Isolate forward steps, not the full trace.** nsys traces contain warmup, JIT, model loading, and teardown.\n2. **Use structural kernel patterns for iteration detection.** Allreduce grouping is more robust than kernel density.\n3. **Compare steady-state iterations.** Filter to identical workload (same batch size, same ctx/gen mix).\n4. **Per-step metrics, not totals.** Always compare per-step averages.\n\n### Root Cause Workflow\n\n#### Step 1: Collect nsys Traces\n\nProfile both versions (if comparing) with identical settings:\n\n```bash\nnsys profile -o /path/to/trace \\\n  -t cuda,nvtx,osrt \\\n  --force-overwrite=true \\\n  --cuda-memory-usage=true \\\n  -w true \\\n  <benchmark_command> --num_requests 500\n```\n\n#### Step 2: Export to SQLite\n\n```bash\nnsys export --type=sqlite --force-overwrite=true -o trace.sqlite trace.nsys-rep\n```\n\n#### Step 3: Run Host Overhead Analysis\n\n```bash\n# Two-trace comparison\npython scripts/analyze_host_overhead.py \\\n  --baseline /path/to/baseline/trace.sqlite \\\n  --target /path/to/target/trace.sqlite \\\n  --baseline-label \"v1.1\" \\\n  --target-label \"main\" \\\n  --output /path/to/output/analysis.txt\n\n# Single-trace breakdown\npython scripts/analyze_host_overhead.py \\\n  --baseline /path/to/trace.sqlite \\\n  --baseline-label \"current\"\n```\n\n#### Step 4: Interpret Results\n\nThe script produces:\n1. **Allreduce-based iteration detection** -- confirms forward step boundaries\n2. **Per-step wall time comparison** -- quantifies the regression\n3. **NVTX per-step breakdown** -- identifies which host operations regressed\n4. **GPU kernel comparison** -- confirms GPU execution is unchanged\n5. **CUDA API comparison** -- detects kernel launch overhead changes\n\n### Reading the Output\n\n**Per-Step Wall Time:**\n```\nAvg wall time per step: 3,317 us (baseline) vs 3,978 us (target)  +19.9%\n```\nThis is the primary regression metric.\n\n**NVTX Breakdown:**\n```\nOperation           | baseline (us/step) | target (us/step) | Delta    | Status\n_fetch_new_requests |               36   |             270  | +234     | REGRESSION\nbroadcast_requests  |                -   |             250  | +250     | NEW\n_update_requests    |              413   |             723  | +310     | REGRESSION\n```\nFocus on operations with large absolute deltas.\n\n**GPU Kernel Comparison:**\n```\nKernels per step (launched): 6.2 (baseline) vs 21.9 (target)  +253%\n```\nMore individual launches = more host-side launch overhead.\n\n### Step 5: Kernel-Level Drill-Down (Optional)\n\nWhen the NVTX breakdown identifies a regressing operation but does not reveal *why* (the overhead is inside the GPU dispatch, not between NVTX ranges), drill below NVTX operations into individual GPU kernel launches.\n\nSee [references/kernel-level-analysis.md](references/kernel-level-analysis.md) for full technique details, SQL queries, and examples.\n\n**When to drill down:**\n- An operation has high wall time but the overhead is inside GPU dispatch, not between NVTX ranges\n- You need to understand how much of the forward pass is graph-captured vs eager\n- Per-layer overhead is significant and you need to map kernels to functional groups\n- Multi-rank inference shows unexplained performance asymmetry\n\n#### Kernel-Level Techniques\n\n| Technique | Question | Key Output |\n|-----------|----------|------------|\n| **Inter-Kernel Gap Analysis** | Where is the GPU idle between kernels? | Gap bucket distribution, top-N largest gaps with source mapping |\n| **Eager vs Graph Classification** | What fraction of kernels are graph-captured? | Graph coverage ratio, list of eager kernels with source attribution |\n| **Repeating-Pattern Mapping** | Which functional group within a layer has the most overhead? | Per-group gap totals, priority ranking |\n| **Straggler Detection** | Is one rank consistently slower? | Straggler rank ID, root cause (extra host work, queue depth feedback loop) |\n\n#### Workflow\n\n1. **Start with Inter-Kernel Gap Analysis** — bucket the gap distribution to understand the dominant overhead type (graph dispatch, Python interpreter, host-device sync)\n2. **If piecewise graph is in use**, run Eager vs Graph Classification to measure graph coverage and identify unnecessary eager kernels\n3. **For per-layer overhead**, use Repeating-Pattern Mapping to isolate the highest-overhead functional group within a single layer\n4. **For multi-rank setups**, run Straggler Detection if per-step wall time varies across ranks\n\n#### Kernel-Level Findings to Optimization Patterns\n\n| Finding | Optimization Pattern |\n|---------|---------------------|\n| Large gaps from Python tensor view chains | CUSTOM_OP — replace with C++ custom op |\n| Graph-capturable kernels running eagerly | GRAPH_EXPAND — fix partition poisoning |\n| Monolithic custom op blocking graph capture | GRAPH_SPLIT — split into capturable + eager parts |\n| Host-device sync (`.item()`) in per-layer code | SYNC (Pattern 1: pre-compute on CPU) + HOIST (Variant B: pass from step level) |\n| Per-layer buffer allocation | ALLOC — pre-allocate at init |\n| Straggler rank with extra host work | Apply targeted optimization to coordinator-only code paths |\n\n---\n\n## Common Patterns and Root Causes\n\n### Pattern 1: Request Management Refactor\n**Symptom**: `_fetch_new_requests` regressed 5-10x, new `broadcast_requests` operation.\n**Cause**: Request fetching refactored for multi-rank broadcasting in TP.\n**Mitigation**: Optimize broadcast path; batch request state updates.\n\n### Pattern 2: Increased Kernel Launch Count\n**Symptom**: 3-5x more `cudaLaunchKernel` calls per step, similar GPU time.\n**Cause**: Operations that were fused or graph-captured are now individual launches.\n**Mitigation**: Re-fuse kernels; extend CUDA graph capture scope.\n\n### Pattern 3: New Bookkeeping Operations\n**Symptom**: New NVTX ranges like `_write_finish_reasons`, `handle_additional_outputs`.\n**Cause**: New features added to the inference loop without overhead budgeting.\n**Mitigation**: Defer non-critical bookkeeping to async paths; batch updates.\n\n### Pattern 4: Flashinfer JIT Warmup Masquerading as Inference\n**Symptom**: Massive elementwise/reduce kernel counts in \"steady state\" analysis.\n**Cause**: Analysis window includes flashinfer JIT compilation phase.\n**Fix**: Use allreduce-based iteration isolation, not kernel density or time windows.\n\n### Pattern 5: Context-Only Bottleneck (Masked by Aggregate)\n**Symptom**: Aggregate metrics below threshold, but context iterations have 50% GPU idle.\n**Cause**: Generation iterations dilute the context-phase bottleneck.\n**Fix**: Per-phase analysis in Detection phase catches this.\n\n---\n\n## Pitfalls\n\n### 1. shortName is an Integer ID\nIn `CUPTI_ACTIVITY_KIND_KERNEL`, `shortName` is an integer referencing `StringIds.id`. Always join. See [references/nsys-schema.md](references/nsys-schema.md).\n\n### 2. NVTX textId vs text\nMost NVTX events have `textId` (integer) but NULL `text`. Join with StringIds. See [references/nsys-schema.md](references/nsys-schema.md).\n\n### 3. Duplicate NVTX Ranges from TP Ranks\nIn TP configurations, each rank reports NVTX ranges independently. De-duplicate by grouping entries within 100us of each other.\n\n### 4. Negative Inter-Step Gaps\nWhen TP ranks report overlapping NVTX ranges, `gap = next_start - prev_end` can be negative. Use the maximum end time when de-duplicating.\n\n### 5. Benchmark Window Selection\nThe allreduce-based window captures context+generation phases; steady-state NVTX filtering captures generation-only. Both are valid; use the appropriate one for your comparison goal.\n\n---\n\n## Handoff to Optimization\n\nWhen analysis is complete and the verdict is **YES**, hand off to the `perf-host-optimization` skill with:\n\n1. **Detection verdict and evidence**: Which metrics crossed thresholds (M1-M5), whether host prep was confirmed (M3b+M3c), and per-phase breakdown.\n\n2. **NVTX-based triage** (from Root Cause): Top regressing operations by absolute delta (us/step). Map NVTX range names to source functions -- see [references/trtllm-nvtx-ranges.md](references/trtllm-nvtx-ranges.md).\n\n3. **Handoff data block**: Include structured data from [references/output-format.md](references/output-format.md) (see \"Handoff to Optimization\" section).\n\n4. **Kernel-level findings** (from drill-down, if performed): Inter-kernel gap distribution, graph coverage ratio, per-group overhead map, and straggler rank identification. Map findings to optimization patterns using the table in the Root Cause kernel-level drill-down section above.\n\n---\n\n## Reference\n\n| File | Contents |\n|------|----------|\n| [references/metrics.md](references/metrics.md) | Full metric definitions, formulas, SQL queries, M3 sub-metric analysis |\n| [references/thresholds.md](references/thresholds.md) | Aggregate and per-phase threshold tables |\n| [references/phase-classification.md](references/phase-classification.md) | NVTX marker parsing, iteration classification, per-phase aggregation |\n| [references/output-format.md](references/output-format.md) | Report template and integration JSON schema |\n| [references/examples.md](references/examples.md) | Worked scenarios (aggregate, phase-specific, and case study) |\n| [references/iteration-isolation-techniques.md](references/iteration-isolation-techniques.md) | Allreduce, NVTX, and kernel-density iteration isolation techniques |\n| [references/trtllm-nvtx-ranges.md](references/trtllm-nvtx-ranges.md) | TRT-LLM NVTX range reference with per-operation timings |\n| [references/kernel-level-analysis.md](references/kernel-level-analysis.md) | Kernel-level drill-down techniques: gap analysis, graph classification, pattern mapping, straggler detection |\n| [references/nsys-schema.md](references/nsys-schema.md) | nsys SQLite schema reference and useful queries |\n| [scripts/analyze_host_overhead.py](scripts/analyze_host_overhead.py) | Python script for Phase 2 root cause analysis |\n| [scripts/detect_host_overhead.py](scripts/detect_host_overhead.py) | Python script for Phase 1 detection verdict |\n",".claude/skills/perf-analysis/SKILL.md":"---\nname: perf-analysis\ndescription: >\n  Performance analysis coordination workflow. Guides profiling delegation,\n  bottleneck classification (compute/memory/launch/communication/sync),\n  and structured report generation. Use when the user asks to analyze\n  performance, profile a workload, check MFU/SOL, or diagnose bottlenecks.\ntags:\n  - analysis\n  - profiling\n  - bottleneck\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Performance Analysis\n\n## Principles\n\n1. **Delegate profiling, own analysis.** You coordinate the analysis workflow\n   but do not run profiling tools directly. Delegate all profiling and\n   measurement tasks to **perf-profiling-specialist** or other domain specialists.\n2. **Metrics from tools, never invented.** All performance numbers must come\n   from profiling tool output. Never fabricate metrics.\n3. **Classify before recommending.** Identify the bottleneck type before\n   suggesting optimizations. The wrong classification leads to wasted effort.\n4. **Structured reports.** Every analysis produces a report with Summary,\n   Metrics, Findings, and Recommendations.\n\n## Key Performance Metrics\n\n- **Throughput**: samples/sec, tokens/sec, iterations/sec\n- **Latency**: end-to-end time, kernel time, communication time\n- **MFU (Model FLOPs Utilization)**: actual FLOPs / theoretical peak FLOPs\n- **% of SOL (Speed of Light)**: current perf / hardware peak perf\n- **GPU Utilization**: SM occupancy, tensor core usage\n- **Memory Bandwidth**: DRAM bandwidth utilization vs peak\n\n## Analysis Workflow\n\n1. **Understand**: Clarify what metrics the user needs (MFU, SOL, latency, etc.)\n2. **Plan**: Plan your profiling and analysis steps before starting\n3. **Profile**: Delegate to **perf-profiling-specialist** for actual measurements\n4. **Measure**: Extract requested metrics from profiling results\n5. **Classify**: If diagnosing issues, determine primary bottleneck type\n6. **Report**: Generate performance analysis report with findings\n\n## Bottleneck Classification\n\nWhen diagnosing performance issues, classify the primary bottleneck:\n\n| Type | Indicator | Description |\n|------|-----------|-------------|\n| **Compute-bound** | High GPU utilization, low memory bandwidth usage | Limited by compute capacity (FLOPs) |\n| **Memory-bound** | High memory bandwidth, low compute utilization | Limited by DRAM throughput |\n| **Launch-overhead** | Many small kernels, high CPU time | CPU becoming bottleneck from kernel launch overhead |\n| **Communication-bound** | Significant time in collective operations | Limited by inter-GPU or inter-node communication |\n| **Sync-bound** | Excessive CPU-GPU synchronization points | Stalls from unnecessary synchronization |\n\nAfter classifying, consult the **perf-optimization-casebook** skill for\ncandidate optimizations (decision precedents) matching the bottleneck class,\nmodel, and hardware. Use it to populate the Recommendations section with\nproven, adapted approaches rather than inventing them — but report only\nmeasured numbers, never a case's expected effect as an achieved result.\n\n## Delegation Guidelines\n\nWhen delegating to specialists, describe the **desired outcome** -- not the\ntool methodology.\n\n**DO include:**\n- The workload: file path, code snippet, or command to profile\n- Problem context: dimensions, dtypes, FLOPs calculations, batch sizes\n- Desired metrics: SOL%, MFU, throughput, occupancy, bottleneck classification\n- Any constraints: specific kernel to target, profiling region markers\n\n**DO NOT include:**\n- Specific tool flags or command patterns (e.g., `--set=full`, `--section SpeedOfLight`)\n- Step-by-step tool usage instructions\n- Fallback strategies for tool failures\n- Example commands\n- Output file paths or artifact locations (specialists create their own workspace artifacts)\n\nSpecialists have their own skills that encode best practices for tool usage\nand their own workspace artifacts for output. Prescribing commands in the\ndelegation overrides their skills and may lead to suboptimal profiling\nstrategies (e.g., collecting 8000+ metrics with `--set=full` when a targeted\nsection analysis would be faster and more surgical).\n\n### Good Example\n\n```\nProfile the batched GEMM kernel in bmm_workload.py with NCU.\nThe workload uses cudaProfilerStart/Stop markers to isolate the region of interest.\nCollect kernel-level metrics: SOL%, compute/memory throughput, DRAM bandwidth,\ntensor core utilization, occupancy, warp stall reasons, and roofline classification.\nThe batched GEMM performs 68.72 GFLOP per call (B=32, M=512, N=1024, K=2048, FP16).\nCalculate MFU against the GPU's peak FP16 tensor core TFLOP/s.\n```\n\n### Bad Example\n\n```\nRun NCU with --set=full --profile-from-start off --target-processes all.\nIf --set=full fails, try --set=detailed. Parse the CSV output for\nsm__throughput.avg.pct_of_peak_sustained_elapsed.\nSave raw NCU output to ncu_output.txt.\n```\n\n### Remote Profiling\n\nWhen profiling on a remote SLURM cluster, include the\n**Remote Execution Context** block in the delegation prompt with the SSH+srun\nwrapper for the target cluster. The perf-profiling-specialist will prefix its\ncommands (nsys, ncu, nvidia-smi) with this wrapper.\n\nThe perf-profiling-specialist does not need the `exec-remote-slurm` skill — the\ncontext block provides everything it needs to execute remotely.\n\n## Available Specialists\n\nDelegate profiling and domain-specific analysis to these specialists:\n\n- **perf-profiling-specialist**: Runs nvidia-smi, nsys, ncu, torch.profiler. Use for ALL profiling tasks.\n- **perf-torch-cuda-graph-specialist**: Analyzes CUDA Graph compatibility and applies capture workflows\n\n## Report Format\n\nStructure every analysis report with these four sections:\n\n1. **Summary**: High-level performance status or bottleneck classification\n2. **Metrics**: Key performance numbers from profiling\n3. **Findings**: Detailed observations with evidence\n4. **Recommendations**: Prioritized list of optimizations (if applicable)\n\n### Example Report\n\n```\n## Summary\nTraining at 42% MFU, memory-bound due to large attention tensors.\n\n## Metrics\n- Throughput: 1,247 samples/sec\n- MFU: 42% (vs 65% theoretical for this model)\n- % of SOL: 58% (room for 1.7x improvement)\n- GPU Utilization: 45%\n- Memory Bandwidth: 850 GB/s (89% of peak)\n- Kernel Count: 1,247 per iteration\n\n## Findings\n1. Self-attention consumes 60% of memory bandwidth\n2. Optimizer step has 3 unnecessary synchronizations\n3. Batch size could be increased by 2x\n\n## Recommendations\n1. Enable FlashAttention (expected: +15% MFU)\n2. Remove synchronizations in optimizer (expected: +5% throughput)\n3. Increase batch size to improve GPU utilization\n```\n",".claude/skills/perf-host-optimization/SKILL.md":"---\nname: perf-host-optimization\ndescription: Profiles and optimizes TensorRT-LLM host/CPU overhead using line_profiler (with nsys support planned). Runs iterative profile-analyze-optimize-validate rounds. Use when GPU utilization is low or optimizing PyExecutor throughput.\nlicense: Apache-2.0\ntags:\n  - optimization\n  - profiling\n  - host-overhead\n  - line-profiler\n  - inference\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Host Performance Optimization Skill\n\nAutomates detection and optimization of host-side (CPU) overhead in TensorRT-LLM's PyTorch backend.\n\n## When to Use\n\n- GPU utilization is low during inference (CPU bottleneck suspected)\n- User asks to reduce host overhead or CPU latency\n- Optimizing PyExecutor throughput (requests/sec)\n- Need line-by-line profiling of specific Python functions\n\n### Confirming the Bottleneck\n\nline_profiler measures *where* CPU time is spent but not *whether* CPU is the bottleneck.\nIf you need to confirm CPU is the limiting factor, run the `perf-host-analysis` skill first -- it provides a YES/NO verdict with metric evidence.\n\nAs a rough heuristic without nsys: if doubling the batch size does not proportionally increase GPU utilization or throughput, CPU overhead is likely the bottleneck.\n\n### Using Analysis Skill Results\n\nIf the `perf-host-analysis` skill has already been run, use its output to skip the confirmation step and prioritize targets:\n\n1. **Detection verdict**: If YES with host_prep_confirmed, start with `_prepare_tp_inputs`.\n2. **NVTX triage** (from Root Cause): The `top_regressing_ops` in the handoff data block maps NVTX range names to source functions. Profile the function with the largest absolute delta first.\n3. **Cross-function triage**: When the top NVTX regression is NOT in `_prepare_tp_inputs` (e.g., `_fetch_new_requests`, `broadcast_requests`, `_update_requests`), target that function's source file directly instead of defaulting to `_prepare_tp_inputs`. See [references/trtllm-nvtx-ranges.md](../perf-host-analysis/references/trtllm-nvtx-ranges.md) for the NVTX-to-source mapping.\n\n---\n\n## Profiling Setup\n\n### line_profiler (Primary Method)\n\n**Environment Variables:**\n- `TLLM_LINE_PROFILER_ENABLED=True` — Enable the profiler\n- `TLLM_LINE_PROFILER_PATH` — Output file path\n- `TLLM_LINE_PROFILER_FUNCTIONS` — Additional functions to profile (comma-separated)\n\n**Function specification format:**\n```bash\n# Class methods: module.path.ClassName.method_name\nTLLM_LINE_PROFILER_FUNCTIONS=\"tensorrt_llm._torch.pyexecutor.model_engine.PyTorchModelEngine._prepare_tp_inputs\"\n\n# Standalone functions: module.path::function_name\nTLLM_LINE_PROFILER_FUNCTIONS=\"tensorrt_llm._torch.pyexecutor.sampler::_group_requests_by_strategy_key\"\n\n# Multiple functions (comma-separated)\nTLLM_LINE_PROFILER_FUNCTIONS=\"module.Class.method1,module.Class.method2\"\n```\n\n### CPU Affinity (Environment Factor)\n\nCPU core affinity can significantly affect host overhead measurements,\nespecially on multi-socket systems (e.g., B300). Pinning processes to cores\nnear the GPU's NUMA node reduces cross-socket memory access latency.\n\n- Check current affinity: `taskset -p <pid>` or `numactl --show`\n- Pin to local NUMA node: `numactl --cpunodebind=<node> --membind=<node>`\n- **Impact**: Up to 2x difference in host overhead on B300 systems\n\nWhen comparing profiling results across runs, ensure CPU affinity is consistent.\nDo not externally modify the affinity, unless user requires to do this to examine the affects of this part.\nDocument the affinity setting in each round's report if it varies.\n\n### Workspace & Suffix Management\n\nEach profiling run should have a unique suffix to track progress across rounds:\n```bash\nEXTRA_SUFFIX=round0_baseline bash profile.sh\nEXTRA_SUFFIX=round1_eliminate_redundant_iter bash profile.sh\n```\n\n---\n\n## Autonomous Optimization Loop\n\nBefore starting the loop, review [references/optimization-strategy.md](references/optimization-strategy.md) for strategic guidance on ordering (zero-risk-first), measurement traps, and overhead scoping.\n\n**Key insight**: Optimizations are NOT independent. Fixing a 50ms bottleneck may reveal a 30ms bottleneck that was previously masked (hidden behind the larger one). Always re-profile after each significant change — the bottleneck landscape shifts.\n\n**Ordering principle**: Within each round, prefer zero-risk optimizations (caching, pre-allocation, hoisting invariants) over medium/high-risk ones (graph partition changes, algorithm fusion). Zero-risk changes provide free gains and make subsequent profiling cleaner.\n\nRun N rounds (default 3) of the following cycle:\n\n```\nFOR round = 1 to MAX_ROUNDS:\n\n  1. PROFILE (with Drill-Down)\n  2. ANALYZE (Multi-Option)\n  3. OPTIMIZE (Apply Change — prefer zero-risk first)\n  4. TEST (Unit Test Validation)\n  5. VALIDATE (Re-Profile — expect bottleneck landscape to shift)\n  6. REPORT\n\nEND FOR → FINAL SUMMARY\n```\n\n### Phase 1: PROFILE (with Drill-Down)\n\n- Run workload with profiler enabled\n- Parse output: identify functions with highest Total time and lines with highest % Time\n- **CRITICAL: Drill down into sub-functions that are not yet profiled** (see below)\n\n#### Drill-Down Profiling\n\nThe default profiler covers top-level executor functions but **not all sub-functions**. When a profiled function shows most time in a single sub-call, you must drill down.\n\n**When:** A single line consumes >80% of a function's time calling an unprofiled sub-function:\n```\nLine #      Hits         Time    Per Hit   % Time  Line Contents\n==============================================================\n  2848      4100  59200000000.0  14439024.4   98.7      output = self.model_engine.forward(...)\n```\n\n**How:**\n1. Identify the sub-function's full qualified path (e.g., `tensorrt_llm._torch.pyexecutor.model_engine.PyTorchModelEngine._prepare_tp_inputs`)\n2. Add it to `TLLM_LINE_PROFILER_FUNCTIONS`\n3. Re-profile to get line-level data inside it\n4. Now analyze the **inner** hotspots\n\nFor common drill-down targets, see [references/hot-path-files.md](references/hot-path-files.md).\n\n### Phase 2: ANALYZE (Multi-Option)\n\nFor the chosen hotspot:\n\n1. **Identify** the top hotspots by **absolute time** (not just %) within the target function\n2. **Classify** each hotspot by type. Summary table:\n\n| Type | Indicators | Severity |\n|------|------------|----------|\n| **HOST_SYNC** | `.item()`, `.cpu()` in per-layer forward path | Critical |\n| **SYNC** | `.item()`, `.cpu()`, `synchronize()` in step-level code | Critical |\n| **CUSTOM_OP** | Chain of Python tensor ops (view/slice/cast) before kernel launch | Critical |\n| **GRAPH_BREAK** | Op that prevents CUDA graph capture of surrounding code (fix via GRAPH_EXPAND / GRAPH_SPLIT) | High |\n| **ALLOC** | `torch.zeros/empty/tensor()` in loops, `.clone()` | High |\n| **HOIST** | Per-layer recomputation of step-invariant values | High |\n| **PYLOOP** | `for x in collection:` with many iterations | High |\n| **REDUNDANT_ITER** | Multiple passes over the same collection | High |\n| **DEAD_WORK** | Object construction whose results are always discarded | High |\n| **CONTAINER** | Dict/set lookups in hot loops | Medium |\n| **FUNCALL** | Repeated method/property calls | Medium |\n| **COMM** | `dist.all_reduce`, `dist.barrier`, NCCL overhead in TP/PP paths | Medium |\n| **GIL** | Lock/queue contention | Medium |\n| **SERIALIZE** | `pickle.dumps/loads`, `json.dumps/loads` in request processing | Medium |\n| **GC** | Periodic latency spikes, non-deterministic pauses (tail latency) | Low |\n| **COMPUTE** | Actual computation (may not be optimizable) | Low |\n\nFor detailed classification with code examples, see [references/hotspot-classification.md](references/hotspot-classification.md).\n\n3. **Propose 2-4 optimization options** in a table:\n\n| Option | Description | Estimated Savings | Risk | Complexity |\n|--------|-------------|-------------------|------|------------|\n| A | ... | ... | Low/Med/High | ... |\n| B | ... | ... | ... | ... |\n\n4. **Select the best option** and explain reasoning (prefer high-savings + low-risk; follow zero-risk-first ordering from [references/optimization-strategy.md](references/optimization-strategy.md))\n\nFor optimization patterns by type, see [references/optimization-patterns.md](references/optimization-patterns.md) (index) — it links to the relevant sub-file for each hotspot type. For GPU-specific patterns (CUSTOM_OP, GRAPH_SPLIT, GRAPH_EXPAND), see [references/patterns/gpu-graph.md](references/patterns/gpu-graph.md).\n\n### Phase 3: OPTIMIZE (Apply Change)\n\n- Apply the selected code change with Edit tool\n- **One optimization per round** — keep changes minimal and targeted\n- Record the exact change (file, line range, before/after) for potential rollback\n\n### Phase 4: TEST (Unit Test Validation)\n\n**Mandatory** after each optimization. Find and run related UTs to verify correctness.\n\n**Finding related tests:**\n```bash\n# Search by modified file name\ngrep -rl \"model_engine\\|PyTorchModelEngine\" tests/unittest/_torch/executor/\n\n# Search by modified function name\ngrep -rl \"_prepare_tp_inputs\\|prepare_inputs\" tests/\n```\n\n**Running tests:**\n```bash\n# Run specific test file with stop-on-first-failure\npytest tests/unittest/_torch/executor/test_pytorch_model_engine.py -v -x --timeout=120\n\n# Run specific test method\npytest tests/unittest/_torch/executor/test_pytorch_model_engine.py::PyTorchModelEngineTestCase::test_position_id_preparation -v -x\n```\n\nFor the full UT-to-file mapping, see [references/hot-path-files.md](references/hot-path-files.md).\n\n**If tests fail:**\n1. Read the failure message\n2. Rollback immediately (`git checkout -- <file>`)\n3. Analyze why the optimization broke correctness\n4. Try the next-best option from Phase 2\n\n### Phase 5: VALIDATE (Re-Profile)\n\n- Re-run profiler with identical workload, using suffix `round<N>_<description>`\n- Compare three things:\n  1. Did the **target hotspot** time decrease?\n  2. Did the **overall function** Total time decrease?\n  3. Did **benchmark metrics** (TPOT, throughput) improve?\n\n**If regression detected** (function time increased or metrics worsened):\n- The \"optimization\" may have triggered a CPython pitfall — see [references/patterns/compound-pitfalls.md](references/patterns/compound-pitfalls.md) (CPython Pitfalls section)\n- Rollback and try the next-best option from Phase 2\n\n### Phase 6: REPORT\n\nLog for this round:\n- Round number\n- Hotspot location (file:line) and classification\n- Optimization applied (with before/after code summary)\n- Time delta: function Total time before → after\n- Benchmark delta: TPOT, throughput before → after\n\n---\n\n## Reading Profile Output\n\n```\nTimer unit: 1e-06 s\nTotal time: 1.234 s\nFile: /path/to/file.py\nFunction: my_function at line 100\n\nLine #      Hits         Time  Per Hit   % Time  Line Contents\n==============================================================\n   100                                           def my_function(self):\n   101       500    890000.0   1780.0     72.1       result = tensor.item()\n   102       500    234567.0    469.1     19.0       return result\n```\n\n**How to read effectively:**\n1. Start with **Total time** for each function — this is the overall budget\n2. Sort lines mentally by **absolute Time**, not just % Time (3% of a 60s function = 1.8s)\n3. Check **Hits count** to understand iteration patterns:\n   - Hits = 2 × expected count → `for x in range(1):` loop overhead (2 hits = enter + exit check)\n   - Hits ≫ expected → the line is inside a nested loop\n4. Look for **repeated patterns**: if 10 lines each take 3% in a loop body, the loop itself costs 30%\n\n---\n\n## Stopping Criteria\n\nStop the optimization loop when:\n1. **Iteration limit reached**: Completed N rounds (default 3)\n2. **No actionable hotspots**: Top hotspots are pure GPU compute (COMPUTE type)\n3. **Diminishing returns**: < 5% improvement in last 2 rounds\n4. **Risk threshold**: Further optimizations require architectural changes (e.g., Cython, struct-of-arrays)\n5. **Test failures**: Cannot find an optimization that passes UTs\n\n**Primary success metric**: Benchmark throughput (requests/sec or tokens/sec) as measured by the profiling script. line_profiler time reductions are leading indicators, but throughput is the ground truth — a function-level speedup that doesn't improve throughput is not a real win.\n\n---\n\n## Final Summary Output\n\nThe final report should include:\n- **Rounds executed**: Number of profile-optimize cycles completed\n- **Cumulative improvement**: Total host time reduction (percentage and absolute)\n- **Benchmark metrics**: Before/after comparison table (TPOT, throughput, ITL, E2EL)\n- **Optimizations applied**: List of changes with file:line locations and classification\n- **Failed attempts**: Any optimizations tried and reverted (with why)\n- **Remaining hotspots**: Top bottlenecks that couldn't be optimized (with classification)\n- **Recommendations**: Suggested follow-up for architectural changes if needed\n\nFor a concrete multi-round example, see [references/examples.md](references/examples.md).\n\n---\n\n## Reference Files\n\n| File | Contents |\n|------|----------|\n| [references/optimization-patterns.md](references/optimization-patterns.md) | Pattern index — links to 6 sub-files: [sync-alloc](references/patterns/sync-alloc.md), [loop-iteration](references/patterns/loop-iteration.md), [python-overhead](references/patterns/python-overhead.md), [gpu-graph](references/patterns/gpu-graph.md), [system](references/patterns/system.md), [compound-pitfalls](references/patterns/compound-pitfalls.md) |\n| [references/optimization-strategy.md](references/optimization-strategy.md) | Zero-risk-first ordering, metric traps, three scopes of host overhead, pattern selection guide |\n| [references/hotspot-classification.md](references/hotspot-classification.md) | Extended per-type indicators and code examples (including CUSTOM_OP, GRAPH_BREAK, HOST_SYNC) |\n| [references/communication-patterns.md](references/communication-patterns.md) | Communication overhead patterns (NCCL batching, barrier removal, async overlap, reduce_scatter) |\n| [references/hot-path-files.md](references/hot-path-files.md) | Key file tables, drill-down targets, UT mapping |\n| [references/examples.md](references/examples.md) | Usage examples and multi-round walkthrough |\n| [trtllm-nvtx-ranges.md](../perf-host-analysis/references/trtllm-nvtx-ranges.md) | TRT-LLM NVTX range reference (from analysis skill) — maps range names to source functions |\n",".claude/skills/kernel-triton-writing/SKILL.md":"---\nname: kernel-triton-writing\ntags: [triton]\ndescription: >\n  ONLY for OpenAI Triton (@triton.jit) kernel development. NEVER use for\n  CUDA C++ kernels, CuTe DSL kernels, TileIR, or profiling tools (ncu, nsys).\n  The user's request must involve Triton explicitly. Covers Triton-specific\n  patterns: fused elementwise, reductions (softmax, LayerNorm, RMSNorm),\n  tiled GEMM with triton.autotune, and flash attention. Workflow:\n  design, write, verify (with fast-path for explicit requests).\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Triton Kernel Writing\n\n## Principles\n\n### Correctness First\n\n1. Never benchmark before verification passes.\n2. Always mask loads and stores for non-divisible shapes.\n3. Include `kernel_fn`, `reference_fn`, and `get_inputs()` exports for companion scripts.\n4. Always run `scripts/verify_kernel.py` to validate against the reference.\n\n### FP16/BF16 Precision Rules (LOW FREEDOM -- follow exactly)\n\nTranscendental functions (`tl.exp`, `tl.log`, `tl.math.erf`, `tl.math.tanh`) require fp32 inputs.\n\n```python\n# WRONG -- compilation error or wrong results with fp16/bf16:\nresult = tl.exp(x)\n\n# CORRECT -- cast to fp32, compute, cast back:\nx_fp32 = x.to(tl.float32)\nresult = tl.exp(x_fp32).to(x.dtype)\n```\n\nRule: any math function beyond basic arithmetic (+, -, *, /) requires fp32 cast in, original dtype cast out.\n\nAdditional precision constraints:\n- `tl.sigmoid()` is unavailable in some Triton versions. Use `1.0 / (1.0 + tl.exp(-x_fp32))`.\n- Always cast back to `x.dtype` before `tl.store` -- mismatches cause \"Type mismatch, store Float32 to Float16\".\n- Unlike PyTorch, Triton does NOT auto-promote fp16/bf16 to fp32 for accumulation. Always use `tl.float32` accumulators for `tl.dot`.\n- **TF32 for matmul:** On Ampere+/Hopper, `tl.dot` uses TF32 by default for fp32 inputs (same as `torch.mm`). Do NOT add `input_precision=\"ieee\"` — it is 3-8x slower. TF32 is the correct default. If verification fails due to TF32 precision (~0.01-0.1 abs diff), ensure `reference_fn` also uses TF32 (plain `torch.mm`, no `allow_tf32=False`).\n\n### CPU-GPU Sync Avoidance (LOW FREEDOM)\n\nNever call `.item()` in kernel wrappers. It forces a CPU-GPU sync (~50-100us per call).\n\n| Pitfall | Fix |\n|---------|-----|\n| `tensor.item()` for seed | `x.data_ptr() % (2**31)` |\n| `torch.randint(...).item()` | Use tensor metadata for pseudo-random seed |\n| Allocating output every call | Accept pre-allocated output as parameter |\n| Python loops calling kernel | Batch operations |\n\n### C Integer Division Semantics (CRITICAL)\n\nTriton uses **C semantics** (round toward zero) for `//` and `%`, NOT Python semantics (round toward negative infinity). This only matters when operands can be negative.\n\n| Expression | Python | Triton/C |\n|------------|--------|----------|\n| `-7 // 2` | `-4` | `-3` |\n| `-7 % 2` | `1` | `-1` |\n\n**Safe pattern:** Ensure all index/offset values are non-negative. If negative values are possible, use `(idx % BLOCK + BLOCK) % BLOCK`.\n\nSee [references/concepts-semantics.md](references/concepts-semantics.md) for full rules and scalar-only exception.\n\n### Kernel Design Mental Model\n\n- **Parallelization axis:** Element-wise kernels parallelize over flattened elements. Row-wise kernels (LayerNorm, softmax) parallelize over rows. Matmul kernels tile in 2D (M, N).\n- **Block size:** Power-of-2 only (256, 512, 1024, 2048). Start with 1024 for H100, 512 for V100.\n- **Memory coalescing:** Adjacent threads must access adjacent memory addresses. The compiler handles this automatically from block-level pointer arithmetic.\n- **Grid:** Use `triton.cdiv(n_elements, BLOCK_SIZE)`. With autotune, grid must be a lambda: `lambda meta: (triton.cdiv(n, meta['BLOCK_SIZE']),)`.\n- **Decorator order:** `@triton.autotune` (outermost) -> `@triton.heuristics` -> `@triton.jit` (innermost).\n- **`reset_to_zero`:** Required for autotune on kernels that accumulate (e.g., matmul output). Without it, later configs see leftover values from earlier trials.\n\n## Workflow\n\n**Fast path:** If the user explicitly requests a Triton kernel (e.g., \"Write a Triton kernel for X\", \"Implement softmax in Triton\"), start at **Phase 2**. Only use Phase 0-1 when the request is ambiguous about whether Triton is appropriate.\n\n### Phase 0: Route the Operator (only for ambiguous requests)\n\nSkip this phase if the user explicitly asks for a Triton kernel. Only use when the request is ambiguous (e.g., \"optimize this operation\").\n\nTriton wins when 2+ operations can share registers instead of writing/reading global memory. Quick rules:\n\n| Pattern | Decision |\n|---------|----------|\n| Single element-wise op (`relu`, `sigmoid`) | SKIP — PyTorch already optimal |\n| Standalone matmul | SKIP — cuBLAS is optimized |\n| Standard attention | SKIP — Use FlashAttention |\n| Element-wise chain (2+ ops), reduction, matmul + epilogue | USE TRITON |\n\nIf SKIP, recommend the alternative and STOP. See [references/operator-routing.md](references/operator-routing.md) for edge cases.\n\n### Phase 1: Analyze the Operator (only for ambiguous requests)\n\nFrom the user's request, identify: (1) operation type, (2) parallelization strategy, (3) input shapes and dtypes.\n\n### Phase 2: Design the Kernel\n\nPick the skeleton below that matches your operation. **These skeletons are sufficient for element-wise, reduction, matmul, and fusion kernels — do NOT read reference files for these common patterns.** Only consult `references/` when implementing uncommon patterns (grouped GEMM, TMA, extern functions) or debugging issues.\n\n**Element-wise skeleton** (GELU, dropout, fused ops on flat tensors):\n\n```python\n@triton.jit\ndef kernel(x_ptr, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):\n    pid = tl.program_id(0)\n    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n    mask = offsets < n_elements\n    x = tl.load(x_ptr + offsets, mask=mask)\n    # ... compute ...\n    tl.store(out_ptr + offsets, result, mask=mask)\n```\n\n**Row-wise skeleton** (softmax, LayerNorm, RMSNorm — one program per row):\n\n```python\n@triton.jit\ndef kernel(x_ptr, out_ptr, n_cols, BLOCK_SIZE: tl.constexpr):\n    row_idx = tl.program_id(0)\n    col_offsets = tl.arange(0, BLOCK_SIZE)\n    mask = col_offsets < n_cols\n    x = tl.load(x_ptr + row_idx * n_cols + col_offsets, mask=mask, other=0.0)\n    # ... reduce / normalize ...\n    tl.store(out_ptr + row_idx * n_cols + col_offsets, result, mask=mask)\n```\n\n**Tiled matmul skeleton** (GEMM with 2D tiling, grouped ordering, and autotune):\n\n```python\n@triton.autotune(\n    configs=[\n        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 256, 'BLOCK_K': 64, 'GROUP_SIZE_M': 8}, num_warps=8, num_stages=3),\n        triton.Config({'BLOCK_M': 64, 'BLOCK_N': 256, 'BLOCK_K': 32, 'GROUP_SIZE_M': 8}, num_warps=4, num_stages=4),\n        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 128, 'BLOCK_K': 32, 'GROUP_SIZE_M': 8}, num_warps=4, num_stages=4),\n        triton.Config({'BLOCK_M': 256, 'BLOCK_N': 64, 'BLOCK_K': 32, 'GROUP_SIZE_M': 8}, num_warps=4, num_stages=4),\n    ],\n    key=['M', 'N', 'K'],\n)\n@triton.jit\ndef matmul_kernel(\n    a_ptr, b_ptr, c_ptr, M, N, K,\n    stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn,\n    BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,\n    GROUP_SIZE_M: tl.constexpr,\n):\n    pid = tl.program_id(0)\n    num_m_blocks = tl.cdiv(M, BLOCK_M)\n    num_n_blocks = tl.cdiv(N, BLOCK_N)\n    # Grouped ordering for L2 cache locality\n    num_pid_in_group = GROUP_SIZE_M * num_n_blocks\n    group_id = pid // num_pid_in_group\n    first_pid_m = group_id * GROUP_SIZE_M\n    group_size_m = min(num_m_blocks - first_pid_m, GROUP_SIZE_M)\n    pid_m = first_pid_m + ((pid % num_pid_in_group) % group_size_m)\n    pid_n = (pid % num_pid_in_group) // group_size_m\n\n    offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)\n    offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)\n    offs_k = tl.arange(0, BLOCK_K)\n\n    a_ptrs = a_ptr + offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak\n    b_ptrs = b_ptr + offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn\n    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)\n\n    for k in range(0, tl.cdiv(K, BLOCK_K)):\n        a_mask = (offs_m[:, None] < M) & (offs_k[None, :] < K)\n        b_mask = (offs_k[:, None] < K) & (offs_n[None, :] < N)\n        a = tl.load(a_ptrs, mask=a_mask, other=0.0)\n        b = tl.load(b_ptrs, mask=b_mask, other=0.0)\n        acc += tl.dot(a, b)\n        a_ptrs += BLOCK_K * stride_ak\n        b_ptrs += BLOCK_K * stride_bk\n        offs_k += BLOCK_K\n\n    c_mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)\n    c_ptrs = c_ptr + offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn\n    tl.store(c_ptrs, acc.to(c_ptr.dtype.element_ty), mask=c_mask)\n```\n\n### Phase 3: Write the Kernel\n\nCreate an output directory, then write the kernel file to `{output_dir}/kernel.py`.\n\nThe kernel file MUST include:\n- `@triton.jit` decorated kernel function\n- `@triton.autotune` for production kernels (see [references/api-core.md](references/api-core.md))\n- Python wrapper function (descriptive name for external import)\n- **Fixed contract exports** (companion scripts rely on these exact names):\n  - `kernel_fn` — alias to the wrapper function\n  - `reference_fn(*args)` — PyTorch reference with identical signature\n  - `get_inputs()` — returns `list` of fresh CUDA tensors for testing/benchmarking\n\nConcise example (fused GELU + dropout):\n\n```python\nimport triton\nimport triton.language as tl\nimport torch\n\n@triton.autotune(\n    configs=[\n        triton.Config({'BLOCK_SIZE': 1024}, num_warps=4),\n        triton.Config({'BLOCK_SIZE': 2048}, num_warps=8),\n    ],\n    key=['n_elements'],\n)\n@triton.jit\ndef fused_gelu_dropout_kernel(\n    x_ptr, out_ptr, n_elements, p, seed,\n    BLOCK_SIZE: tl.constexpr,\n):\n    pid = tl.program_id(0)\n    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n    mask = offsets < n_elements\n\n    x = tl.load(x_ptr + offsets, mask=mask)\n    x_fp32 = x.to(tl.float32)\n    x = (0.5 * x_fp32 * (1.0 + tl.math.erf(x_fp32 * 0.7071067811865476))).to(x.dtype)\n\n    random = tl.rand(seed, offsets)\n    x = tl.where(random > p, x / (1.0 - p), 0.0)\n\n    tl.store(out_ptr + offsets, x, mask=mask)\n\n\ndef fused_gelu_dropout_triton(x: torch.Tensor, p: float = 0.1) -> torch.Tensor:\n    n_elements = x.numel()\n    out = torch.empty_like(x)\n    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)\n    seed = (x.data_ptr() % (2**31)) ^ n_elements  # sync-free seed\n    fused_gelu_dropout_kernel[grid](x, out, n_elements, p, seed)\n    return out\n\n\n# --- Fixed contract (companion scripts rely on these names) ---\nkernel_fn = fused_gelu_dropout_triton\n\ndef reference_fn(x, p=0.1):\n    torch.manual_seed((x.data_ptr() % (2**31)) ^ x.numel())\n    return torch.nn.functional.dropout(\n        torch.nn.functional.gelu(x), p, training=True\n    )\n\ndef get_inputs():\n    return [torch.randn(128 * 1024 * 1024, device=\"cuda\")]\n```\n\nFor more patterns (SiLU+mul, RMSNorm, linear+GELU, add+LayerNorm), see [references/patterns-fusion.md](references/patterns-fusion.md). For GEMM patterns, see [references/patterns-gemm.md](references/patterns-gemm.md).\n\n### Phase 4: Verify Correctness\n\nRun the companion verification script:\n\n```bash\npython scripts/verify_kernel.py {output_dir}/kernel.py --rtol 1e-3 --atol 1e-3\n```\n\nOutput:\n```json\n{\"correct\": true, \"max_abs_diff\": 1.2e-7, \"max_rel_diff\": 3.4e-6, \"details\": \"...\"}\n```\n\n**Stop if `correct: false`.** Fix the kernel before benchmarking.\n\n**Tolerance guide:**\n\n| Dtype | rtol | atol | Notes |\n|-------|------|------|-------|\n| float16 | 1e-3 | 1e-3 | |\n| bfloat16 | 1e-2 | 1e-2 | |\n| float32 | 1e-5 | 1e-5 | Element-wise ops |\n| float32 (matmul) | 1e-2 | 1e-1 | TF32 accumulation order differs between Triton tiles and cuBLAS |\n\n### Phase 5: Benchmark Performance (optional)\n\nOnly benchmark if the user explicitly requests performance numbers. Skip this phase for correctness-focused requests.\n\n```bash\npython scripts/benchmark_kernel.py {output_dir}/kernel.py\n```\n\nOutput:\n```json\n{\"kernel_time_ms\": 0.45, \"reference_time_ms\": 1.23, \"speedup\": 2.73, \"warmup_iters\": 10, \"benchmark_iters\": 40}\n```\n\n## References (consult only when stuck)\n\nThe skeletons and principles above cover element-wise, reduction, matmul, and fusion kernels. **Do NOT read reference files for these common patterns.**\n\nOnly consult `references/` when:\n- Implementing **uncommon patterns** (grouped GEMM, TMA, persistent matmul, extern functions)\n- **Debugging** a compile error or incorrect result not covered by the error table below\n- Needing **API details** for an unfamiliar `tl.*` operation\n\n**How to search:** Grep for your keyword across `references/`. Read only the file Grep points to.\n\n| File | When to use |\n|---|---|\n| `references/api-core.md` | Unfamiliar `triton.autotune` / `triton.Config` options |\n| `references/api-language.md` | Unfamiliar `tl.*` operations |\n| `references/patterns-gemm.md` | Grouped GEMM, persistent matmul, TMA, MX formats |\n| `references/patterns-advanced.md` | Flash attention details, backward passes, libdevice |\n| `references/troubleshooting.md` | Debug ops, interpreter mode, env vars |\n\n## Error Handling and Troubleshooting\n\n### Common Errors\n\n| Error / Symptom | Cause | Fix |\n|---------|-------|-----|\n| \"Type mismatch, store Float32 to Float16\" | Missing `.to(x.dtype)` before store | Cast fp32 result back |\n| `BLOCK_SIZE is not a constexpr` | Block size passed as runtime value | Add `: tl.constexpr` annotation |\n| `shape mismatch` in binary op | Tensor shapes don't broadcast | Check with `tl.static_print`; use `[:, None]` / `[None, :]` |\n| Large diffs everywhere | Wrong dtype in `tl.load` | Check load dtype matches input |\n| Matmul 3-8x slower than expected | `input_precision=\"ieee\"` on `tl.dot` | Remove it; use TF32 default. Ensure `reference_fn` also uses TF32 |\n| Matmul ~0.01-0.1 abs diff vs reference | TF32 vs IEEE mismatch | Use same precision in both kernel and reference (TF32 for both) |\n| Diffs at boundaries | Missing mask | Add mask to all load/store ops |\n| Random diffs | Race condition | Check atomics and ordering |\n| NaN/Inf | Division by zero or fp16 overflow | Guard with epsilon; use `tl.float32` accumulator |\n| `grid must be a tuple` | Grid lambda returns int, not tuple | Return `(value,)` with trailing comma |\n| `expected constexpr` in `tl.arange` | Non-constexpr argument | Both args of `tl.arange(start, end)` must be constexpr |\n| `triton.OutOfResources` | Register/shared memory pressure | Reduce BLOCK_SIZE or `num_stages` |\n| Kernel not updating after edit | Stale compilation cache | `rm -rf ~/.triton/cache/` |\n| Mismatched results vs PyTorch | C integer division semantics | Triton uses truncation; see `references/concepts-semantics.md` |\n\nFor extended error table, interpreter mode issues, and environment variables, see [references/troubleshooting.md](references/troubleshooting.md).\n\n### When to Abort\n\nStop and report failure if:\n1. **Not a good fit** -- Pure matmul or complex control flow (Phase 0 should catch this).\n2. **Verification fails after 3 attempts** -- Numerical issues too severe to fix.\n3. **No speedup** -- Reference is already well-optimized (cuBLAS, cuDNN).\n4. **Hardware mismatch** -- Target GPU not available for testing.\n",".claude/skills/kernel-tileir-optimization/SKILL.md":"---\nname: kernel-tileir-optimization\ndescription: >\n  Optimize existing Triton kernels for NVIDIA TileIR backend on Blackwell GPUs (sm_100+).\n  Adds TileIR-specific autotune configs: occupancy, num_ctas, TMA descriptors. Covers\n  kernel classification (dot-related, norm-like, elementwise, reduction), type-specific\n  transformations, and PTX-vs-TileIR benchmarking. Triggered by: \"optimize for TileIR\",\n  \"add TileIR configs\", \"Blackwell optimization\", \"TMA descriptors\", \"2CTA mode\",\n  \"occupancy tuning\". Kernels use standard `import triton`; TileIR activates via\n  ENABLE_TILE=1 when nvtriton is installed.\ncompatibility: >\n  Requires Blackwell GPU (sm_100+) for TileIR execution. Supports development on any GPU.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Triton TileIR Optimization\n\nOptimize EXISTING Triton kernels for NVIDIA's TileIR backend on Blackwell GPUs.\nThis skill does NOT write kernels from scratch -- that is the Triton Specialist's job.\n\n## Principles\n\n### TileIR vs PTX Backend\n\nTileIR is NVIDIA's compiler backend for Triton that generates optimized CUDA code\nusing CGA-level (Cooperative Grid Array) tile representations. Critical differences:\n\n| Parameter | PTX Backend | TileIR Backend |\n|-----------|-------------|----------------|\n| `num_warps` | Strict directive | **Ignored** (compiler decides) |\n| `num_stages` | Strict directive | Cost hint (compiler optimizes) |\n| `occupancy` | Not available | **Critical** tuning param (1-32) |\n| `num_ctas` | Limited | 2CTA mode for Blackwell |\n| Block sizes | Smaller often better | Larger often better |\n| TMA | Not available | Required for dot kernels |\n\n**Key implication**: Do not tune `num_warps` for TileIR -- focus on `occupancy` instead.\n\n### Triton Package Landscape\n\nThree packages share `import triton`:\n\n| Package | Source | Use Case |\n|---------|--------|----------|\n| `pytorch-triton` | PyTorch wheel | `torch.compile`, standard kernels |\n| `triton` | OpenAI PyPI | Official Triton from triton-lang.org |\n| nvtriton | [Triton-to-tile-IR](https://github.com/triton-lang/Triton-to-tile-IR) | TileIR backend for Blackwell |\n\nOnly one triton package should be installed at a time. \"Converting to TileIR\" means\nadding TileIR-specific configs, NOT changing imports. TileIR activates via `ENABLE_TILE=1`.\n\n### When TileIR Applies\n\nTileIR targets Blackwell (sm_100+). Without nvtriton or Blackwell hardware, the\nspecialist still adds TileIR-optimized configs that standard triton safely ignores,\nenabling future deployment.\n\n**Expected speedups** (with nvtriton on Blackwell):\n\n| Kernel Type | Speedup | Key Lever |\n|-------------|---------|-----------|\n| Dot-Related (GEMM, Attention) | 1.2-2.0x | TMA + 2CTA |\n| Norm-Like (LayerNorm, Softmax) | 2.0-5.0x | High occupancy |\n| Element-Wise (ReLU, Add, Exp) | 1.5-3.0x | Occupancy + num_stages |\n| Reduction (Sum, Mean, Max) | 1.8-4.0x | High occupancy |\n\n## Workflow\n\nFive-phase workflow: compatibility, classify, transform, validate, benchmark.\n\n### Phase 1: Compatibility Test (ENABLE_TILE=0)\n\nVerify the kernel works in PTX mode before applying TileIR optimizations.\n\n```bash\npython scripts/tileir_check.py\n```\n\nThen use the **kernel-triton-writing** skill's `verify_kernel.py` to verify with `ENABLE_TILE=0`:\n\n```bash\npython scripts/verify_kernel.py --kernel path/to/kernel.py --reference 'torch reference' --shapes '{\"x\": [32, 512, 4096]}' --dtypes '{\"x\": \"bfloat16\"}'\n```\n\n### Phase 2: Classify Kernel\n\nDetermine kernel type to select the optimization strategy.\n\n```bash\npython scripts/classify_kernel.py --file kernel.py\n```\n\nClassification decision tree:\n\n```\nContains tl.dot()?\n  YES --> dot-related: TMA + 2CTA + occupancy + larger blocks\n  NO  --> Has reduction + normalization?\n            YES --> norm-like: high occupancy (2, 4) + num_warps (4, 8)\n            NO  --> Point-wise only?\n                      YES --> element-wise: occupancy (1-16) + num_stages (2-4)\n                      NO  --> reduction: high occupancy + num_warps\n```\n\n### Phase 3: Apply Transformations\n\nClassify and apply optimizations in one step:\n\n```bash\npython scripts/classify_kernel.py --file kernel.py --apply-optimizations\n```\n\nOutput JSON includes `optimized_code` and `changes_applied` fields.\n\n**Type-specific transformations:**\n\n**Dot-related** (highest priority):\n1. Convert `tl.load`/`tl.store` to TMA descriptors (MANDATORY). See `references/tma-conversion.md`.\n2. Add 2CTA configs (`num_ctas=2`) with SM oversubscription guard in pre-hook.\n3. Add occupancy (1, 2, 4) and extended num_stages (4, 6).\n4. Use larger block sizes (256x256, 256x128).\n\n**Norm-like** (LayerNorm, Softmax, RMSNorm):\n- Add occupancy (2, 4), num_warps (4, 8). No TMA needed.\n\n**Element-wise** (ReLU, GELU, Add, Mul, Exp):\n- Add occupancy (1, 2, 4, 16), num_stages (2, 3, 4). Include extreme configs for small inputs.\n\n**Reduction** (Sum, Mean, Max):\n- Same strategy as norm-like: high occupancy (2, 4), num_warps (4, 8).\n\nGate TileIR-specific configs for sm_100+:\n\n```python\nimport torch\n\ndef get_configs_with_gating(pre_hook=None):\n    configs = get_baseline_configs()\n    if torch.cuda.is_available() and torch.cuda.get_device_capability()[0] >= 10:\n        configs.extend(get_tileir_specific_configs(pre_hook))\n    return configs\n```\n\nSee `references/config-templates.md` for complete config templates per kernel type.\n\n### Phase 4: TileIR Validation (ENABLE_TILE=1)\n\nUse the **kernel-triton-writing** skill's `verify_kernel.py` to verify the optimized kernel with TileIR backend:\n\n```bash\npython scripts/verify_kernel.py --kernel path/to/optimized_kernel.py --reference 'torch reference' --shapes '{\"x\": [32, 512, 4096]}' --dtypes '{\"x\": \"bfloat16\"}'\n```\n\nSet `ENABLE_TILE=1` before running. Check: numerical correctness, no compilation errors,\nTMA/2CTA patterns compile successfully.\n\n### Phase 5: Benchmark\n\nUse `triton.testing.do_bench()` (as documented in the **perf-workload-profiling** skill) to compare PTX (`ENABLE_TILE=0`) vs TileIR (`ENABLE_TILE=1`).\n\nBenchmark across multiple input sizes (128, 1024, 8192) -- performance varies by size.\n\n## Scripts\n\n### tileir_check.py\n\nCheck TileIR availability (nvtriton, ENABLE_TILE, Blackwell GPU):\n\n```bash\npython scripts/tileir_check.py\n```\n\nReturns JSON: `nvtriton_installed`, `tileir_active`, `blackwell_gpu`, `gpu_capability`, `recommendation`.\n\n### classify_kernel.py\n\nClassify kernel type and optionally apply TileIR optimizations:\n\n```bash\n# Classify only\npython scripts/classify_kernel.py --file kernel.py\n\n# Classify + apply optimizations\npython scripts/classify_kernel.py --file kernel.py --apply-optimizations\n\n# From inline code\npython scripts/classify_kernel.py --code '<kernel_code>'\n```\n\nReturns JSON: `classification`, `confidence`, `indicators`, `recommendations`.\nWith `--apply-optimizations`: adds `optimized_code` and `changes_applied`.\n\n## Error Handling\n\n### Common Pitfalls\n\n**TMA descriptor errors** (dot-related kernels):\n- Always pass `pre_hook=tma_set_block_size_hook` to config generation -- without it,\n  TMA descriptors keep dummy block sizes, causing runtime errors or wrong results.\n- For GEMM: pass `b.T.contiguous()` in wrapper and use `tl.dot(a, b.T, accumulator)`\n  in kernel. Transposition mismatch produces incorrect results silently.\n\n**2CTA oversubscription**:\n- Adjust SM count in pre-hook when using `num_ctas=2`:\n  ```python\n  if \"NUM_SMS\" in nargs and \"NUM_CTAS\" in nargs:\n      nargs[\"NUM_SMS\"] = nargs[\"NUM_SMS\"] // nargs[\"NUM_CTAS\"]\n  ```\n\n**Config function signatures**:\n- ALL config helper functions MUST accept `pre_hook=None`, even if unused.\n  Without it: `TypeError: get_autotune_configs() takes 0 positional arguments`.\n\n**Hardware gating**:\n- Gate TileIR configs with `torch.cuda.get_device_capability()[0] >= 10`.\n  TMA/2CTA on pre-Blackwell GPUs causes runtime crashes.\n\n**API availability**:\n- Use `1.0 / (1.0 + tl.exp(-x))` instead of `tl.sigmoid(x)` -- not available in\n  all Triton versions including some nvtriton builds.\n\n**Performance tuning**:\n- Do not over-tune `num_warps` -- TileIR ignores it. Focus on `occupancy`.\n- Use larger block sizes (256x256, 256x128) for TileIR, not PTX-tuned small blocks.\n- Benchmark across small/medium/large inputs; one-size configs underperform.\n- For exp/log heavy kernels, enable approximate math:\n  ```bash\n  export TILEIR_ENABLE_APPROX=1\n  export TILEIR_ENABLE_FTZ=1\n  ```\n\n### When to Abort\n\nStop and report if:\n\n1. **No triton installed** -- cannot proceed.\n2. **Compatibility test fails** -- kernel has syntax/runtime errors before optimization.\n3. **TileIR validation fails** -- optimized kernel produces wrong results.\n4. **No speedup** -- TileIR version is slower than PTX baseline (with nvtriton).\n5. **Not Blackwell GPU** -- still add configs for future deployment, but skip\n   ENABLE_TILE testing and benchmarking.\n\n### Output Format\n\nAfter optimization, return:\n\n```\n## TileIR Optimization: kernel_name\n\n### Classification\n- Kernel type: [dot-related | norm-like | element-wise | reduction]\n- Strategy: [TMA + 2CTA | High occupancy | Occupancy + num_stages]\n\n### Compatibility Check (ENABLE_TILE=0)\n[PASSED | FAILED] — Max difference: X.Xe-Y\n\n### Transformations Applied\n- [List of transformations]\n\n### TileIR Validation (ENABLE_TILE=1)\n[PASSED | FAILED] — Max difference: X.Xe-Y\n\n### Benchmark Comparison\n| Backend | Time (ms) | Speedup |\n|---------|-----------|---------|\n| PTX (ENABLE_TILE=0) | X.XXX | 1.0x |\n| TileIR (ENABLE_TILE=1) | X.XXX | Y.Yx |\n\n### Output\nFile: kernel_name_tileir.py\n```\n",".claude/skills/perf-instability-cookbook/SKILL.md":"# perf-instability-cookbook\n\nPath: `.claude/skills/perf-instability-cookbook/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/SKILL.md)",".claude/skills/perf-nsight-compute-analysis/SKILL.md":"# perf-nsight-compute-analysis\n\nPath: `.claude/skills/perf-nsight-compute-analysis/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/SKILL.md)",".claude/skills/perf-nsight-systems/SKILL.md":"# perf-nsight-systems\n\nPath: `.claude/skills/perf-nsight-systems/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/SKILL.md)",".claude/skills/perf-optimization-casebook/SKILL.md":"# perf-optimization-casebook\n\nPath: `.claude/skills/perf-optimization-casebook/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/SKILL.md)",".claude/skills/perf-optimization/SKILL.md":"# perf-optimization\n\nPath: `.claude/skills/perf-optimization/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization/SKILL.md)",".claude/skills/perf-regression-cookbook/SKILL.md":"# perf-regression-cookbook\n\nPath: `.claude/skills/perf-regression-cookbook/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/SKILL.md)",".claude/skills/perf-torch-cuda-graphs/SKILL.md":"# perf-torch-cuda-graphs\n\nPath: `.claude/skills/perf-torch-cuda-graphs/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/SKILL.md)",".claude/skills/perf-torch-sync-free/SKILL.md":"# perf-torch-sync-free\n\nPath: `.claude/skills/perf-torch-sync-free/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/SKILL.md)",".claude/skills/perf-workload-profiling/SKILL.md":"# perf-workload-profiling\n\nPath: `.claude/skills/perf-workload-profiling/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/SKILL.md)",".claude/skills/trtllm-case-executor/SKILL.md":"# trtllm-case-executor\n\nPath: `.claude/skills/trtllm-case-executor/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-case-executor/SKILL.md)",".claude/skills/trtllm-code-contribution/SKILL.md":"# trtllm-code-contribution\n\nPath: `.claude/skills/trtllm-code-contribution/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-code-contribution/SKILL.md)",".claude/skills/trtllm-codebase-exploration/SKILL.md":"# trtllm-codebase-exploration\n\nPath: `.claude/skills/trtllm-codebase-exploration/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-codebase-exploration/SKILL.md)",".claude/skills/trtllm-flashinfer-upgrade/SKILL.md":"# trtllm-flashinfer-upgrade\n\nPath: `.claude/skills/trtllm-flashinfer-upgrade/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-flashinfer-upgrade/SKILL.md)",".claude/skills/trtllm-model-onboard-multimodal/SKILL.md":"# trtllm-model-onboard-multimodal\n\nPath: `.claude/skills/trtllm-model-onboard-multimodal/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-model-onboard-multimodal/SKILL.md)",".claude/skills/trtllm-moe-develop/SKILL.md":"# trtllm-moe-develop\n\nPath: `.claude/skills/trtllm-moe-develop/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-moe-develop/SKILL.md)",".claude/skills/trtllm-serve-config-guide/SKILL.md":"# trtllm-serve-config-guide\n\nPath: `.claude/skills/trtllm-serve-config-guide/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-serve-config-guide/SKILL.md)",".claude/skills/trtllm-test-script-builder/SKILL.md":"# trtllm-test-script-builder\n\nPath: `.claude/skills/trtllm-test-script-builder/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/SKILL.md)",".claude/skills/trtllm-test-specialist/SKILL.md":"# trtllm-test-specialist\n\nPath: `.claude/skills/trtllm-test-specialist/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/SKILL.md)",".claude/skills/trtllm-visualgen-bench/SKILL.md":"# trtllm-visualgen-bench\n\nPath: `.claude/skills/trtllm-visualgen-bench/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/SKILL.md)",".claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md":"# trtllm-visualgen-lpips-golden-refresh\n\nPath: `.claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md)",".codex/skills/visual-gen-component-test/SKILL.md":"# visual-gen-component-test\n\nPath: `.codex/skills/visual-gen-component-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/skills/visual-gen-component-test/SKILL.md)","agent-flow/.claude/skills/perf-analyze/SKILL.md":"# perf-analyze\n\nPath: `agent-flow/.claude/skills/perf-analyze/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/.claude/skills/perf-analyze/SKILL.md)","agent-flow/.claude/skills/perf-optimize/SKILL.md":"# perf-optimize\n\nPath: `agent-flow/.claude/skills/perf-optimize/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/.claude/skills/perf-optimize/SKILL.md)","AGENTS.md":"# AI Agent Protocol & Instructions\n\nPath: `AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/AGENTS.md)","CLAUDE.md":"# Claude Agent Guidelines & System Prompt\n\nPath: `CLAUDE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/CLAUDE.md)",".codex/AGENTS.md":"# .codex - AI Agent Protocol & Instructions\n\nPath: `.codex/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/AGENTS.md)",".codex/config.toml":"# Codex Directive: config.toml\n\nPath: `.codex/config.toml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/config.toml)","agent-flow/AGENTS.md":"# agent-flow - AI Agent Protocol & Instructions\n\nPath: `agent-flow/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/AGENTS.md)","agent-flow/CLAUDE.md":"# agent-flow - Claude Agent Guidelines & System Prompt\n\nPath: `agent-flow/CLAUDE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/CLAUDE.md)","tests/AGENTS.md":"# tests - AI Agent Protocol & Instructions\n\nPath: `tests/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/AGENTS.md)","tests/CLAUDE.md":"# tests - Claude Agent Guidelines & System Prompt\n\nPath: `tests/CLAUDE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/CLAUDE.md)",".claude/agents/exec-compile-specialist.md":"# Subagent: exec-compile-specialist\n\nPath: `.claude/agents/exec-compile-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-compile-specialist.md)",".claude/agents/exec-local-slurm.md":"# Subagent: exec-local-slurm\n\nPath: `.claude/agents/exec-local-slurm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-local-slurm.md)",".claude/agents/exec-remote-slurm.md":"# Subagent: exec-remote-slurm\n\nPath: `.claude/agents/exec-remote-slurm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-remote-slurm.md)",".claude/agents/kernel-cuda-specialist.md":"# Subagent: kernel-cuda-specialist\n\nPath: `.claude/agents/kernel-cuda-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-cuda-specialist.md)",".claude/agents/kernel-cute-specialist.md":"# Subagent: kernel-cute-specialist\n\nPath: `.claude/agents/kernel-cute-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-cute-specialist.md)",".claude/agents/kernel-tileir-specialist.md":"# Subagent: kernel-tileir-specialist\n\nPath: `.claude/agents/kernel-tileir-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-tileir-specialist.md)",".claude/agents/kernel-triton-specialist.md":"# Subagent: kernel-triton-specialist\n\nPath: `.claude/agents/kernel-triton-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-triton-specialist.md)",".claude/agents/modeling-test-specialist.md":"# Subagent: modeling-test-specialist\n\nPath: `.claude/agents/modeling-test-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/modeling-test-specialist.md)",".claude/agents/perf-profiling-specialist.md":"# Subagent: perf-profiling-specialist\n\nPath: `.claude/agents/perf-profiling-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-profiling-specialist.md)",".claude/agents/perf-test-sync.md":"# Subagent: perf-test-sync\n\nPath: `.claude/agents/perf-test-sync.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-test-sync.md)",".claude/agents/perf-torch-cuda-graph-specialist.md":"# Subagent: perf-torch-cuda-graph-specialist\n\nPath: `.claude/agents/perf-torch-cuda-graph-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-torch-cuda-graph-specialist.md)",".codex/agents/onboard_update_reviewer.toml":"# Codex Directive: onboard_update_reviewer.toml\n\nPath: `.codex/agents/onboard_update_reviewer.toml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/agents/onboard_update_reviewer.toml)","triton_backend/tools/dataset/mini_cnn_eval.json":"# Dataset Skill\n\nPath: `triton_backend/tools/dataset/mini_cnn_eval.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/triton_backend/tools/dataset/mini_cnn_eval.json)","triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json":"# Dataset Skill\n\nPath: `triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json)",".claude/skills/exec-remote-slurm/scripts/build.sh":"# Claude Directive: build.sh\n\nPath: `.claude/skills/exec-remote-slurm/scripts/build.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-remote-slurm/scripts/build.sh)",".claude/skills/exec-remote-slurm/scripts/build.slurm":"# Claude Directive: build.slurm\n\nPath: `.claude/skills/exec-remote-slurm/scripts/build.slurm`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-remote-slurm/scripts/build.slurm)",".claude/skills/exec-slurm-compile/scripts/compile.sh":"# Claude Directive: compile.sh\n\nPath: `.claude/skills/exec-slurm-compile/scripts/compile.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/compile.sh)",".claude/skills/exec-slurm-compile/scripts/compile.slurm":"# Claude Directive: compile.slurm\n\nPath: `.claude/skills/exec-slurm-compile/scripts/compile.slurm`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/compile.slurm)",".claude/skills/exec-slurm-compile/scripts/enroot-import":"# Claude Directive: enroot-import\n\nPath: `.claude/skills/exec-slurm-compile/scripts/enroot-import`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/enroot-import)",".claude/skills/exec-slurm-compile/scripts/submit_compile.sh":"# Claude Directive: submit_compile.sh\n\nPath: `.claude/skills/exec-slurm-compile/scripts/submit_compile.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/submit_compile.sh)",".claude/skills/kernel-cute-writing/references/api-arch.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-arch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-arch.md)",".claude/skills/kernel-cute-writing/references/api-core.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-core.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-core.md)",".claude/skills/kernel-cute-writing/references/api-nvgpu.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-nvgpu.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-nvgpu.md)",".claude/skills/kernel-cute-writing/references/api-runtime-utils.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-runtime-utils.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-runtime-utils.md)",".claude/skills/kernel-cute-writing/references/concepts-architecture.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-architecture.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-architecture.md)",".claude/skills/kernel-cute-writing/references/concepts-layouts.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-layouts.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-layouts.md)",".claude/skills/kernel-cute-writing/references/concepts-mma.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-mma.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-mma.md)",".claude/skills/kernel-cute-writing/references/concepts-tensors.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-tensors.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-tensors.md)",".claude/skills/kernel-cute-writing/references/patterns-compilation.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-compilation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-compilation.md)",".claude/skills/kernel-cute-writing/references/patterns-elementwise.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-elementwise.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-elementwise.md)",".claude/skills/kernel-cute-writing/references/patterns-gemm.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-gemm.md)",".claude/skills/kernel-cute-writing/references/patterns-getting-started.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-getting-started.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-getting-started.md)",".claude/skills/kernel-cute-writing/references/patterns-memory.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-memory.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-memory.md)",".claude/skills/kernel-cute-writing/references/patterns-pipeline.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-pipeline.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-pipeline.md)",".claude/skills/kernel-cute-writing/references/patterns-reduction.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-reduction.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-reduction.md)",".claude/skills/kernel-cute-writing/references/troubleshooting.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/troubleshooting.md)",".claude/skills/kernel-cute-writing/scripts/__init__.py":"# Claude Directive: __init__.py\n\nPath: `.claude/skills/kernel-cute-writing/scripts/__init__.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/__init__.py)",".claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py":"# Claude Directive: benchmark_kernel.py\n\nPath: `.claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py)",".claude/skills/kernel-cute-writing/scripts/verify_kernel.py":"# Claude Directive: verify_kernel.py\n\nPath: `.claude/skills/kernel-cute-writing/scripts/verify_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/verify_kernel.py)",".claude/skills/kernel-tileir-optimization/references/config-templates.md":"# References Skill\n\nPath: `.claude/skills/kernel-tileir-optimization/references/config-templates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/config-templates.md)",".claude/skills/kernel-tileir-optimization/references/tma-conversion.md":"# References Skill\n\nPath: `.claude/skills/kernel-tileir-optimization/references/tma-conversion.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/tma-conversion.md)",".claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py":"# Claude Directive: classify_kernel.py\n\nPath: `.claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py)",".claude/skills/kernel-tileir-optimization/scripts/tileir_check.py":"# Claude Directive: tileir_check.py\n\nPath: `.claude/skills/kernel-tileir-optimization/scripts/tileir_check.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/scripts/tileir_check.py)",".claude/skills/kernel-triton-writing/references/api-core.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/api-core.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/api-core.md)",".claude/skills/kernel-triton-writing/references/api-language.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/api-language.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/api-language.md)",".claude/skills/kernel-triton-writing/references/concepts-semantics.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/concepts-semantics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/concepts-semantics.md)",".claude/skills/kernel-triton-writing/references/operator-routing.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/operator-routing.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/operator-routing.md)",".claude/skills/kernel-triton-writing/references/patterns-advanced.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-advanced.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-advanced.md)",".claude/skills/kernel-triton-writing/references/patterns-basic.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-basic.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-basic.md)",".claude/skills/kernel-triton-writing/references/patterns-fusion.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-fusion.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-fusion.md)",".claude/skills/kernel-triton-writing/references/patterns-gemm.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-gemm.md)",".claude/skills/kernel-triton-writing/references/troubleshooting.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/troubleshooting.md)",".claude/skills/kernel-triton-writing/scripts/__init__.py":"# Claude Directive: __init__.py\n\nPath: `.claude/skills/kernel-triton-writing/scripts/__init__.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/__init__.py)",".claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py":"# Claude Directive: benchmark_kernel.py\n\nPath: `.claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py)",".claude/skills/kernel-triton-writing/scripts/verify_kernel.py":"# Claude Directive: verify_kernel.py\n\nPath: `.claude/skills/kernel-triton-writing/scripts/verify_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/verify_kernel.py)",".claude/skills/perf-host-analysis/references/examples.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/examples.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/examples.md)",".claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md)",".claude/skills/perf-host-analysis/references/kernel-level-analysis.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/kernel-level-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/kernel-level-analysis.md)",".claude/skills/perf-host-analysis/references/metrics.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/metrics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/metrics.md)",".claude/skills/perf-host-analysis/references/nsys-schema.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/nsys-schema.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/nsys-schema.md)",".claude/skills/perf-host-analysis/references/output-format.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/output-format.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/output-format.md)",".claude/skills/perf-host-analysis/references/phase-classification.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/phase-classification.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/phase-classification.md)",".claude/skills/perf-host-analysis/references/thresholds.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/thresholds.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/thresholds.md)",".claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md)",".claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py":"# Claude Directive: analyze_host_overhead.py\n\nPath: `.claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py)",".claude/skills/perf-host-analysis/scripts/detect_host_overhead.py":"# Claude Directive: detect_host_overhead.py\n\nPath: `.claude/skills/perf-host-analysis/scripts/detect_host_overhead.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/scripts/detect_host_overhead.py)",".claude/skills/perf-host-optimization/references/communication-patterns.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/communication-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/communication-patterns.md)",".claude/skills/perf-host-optimization/references/examples.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/examples.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/examples.md)",".claude/skills/perf-host-optimization/references/hot-path-files.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/hot-path-files.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/hot-path-files.md)",".claude/skills/perf-host-optimization/references/hotspot-classification.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/hotspot-classification.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/hotspot-classification.md)",".claude/skills/perf-host-optimization/references/optimization-patterns.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/optimization-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/optimization-patterns.md)",".claude/skills/perf-host-optimization/references/optimization-strategy.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/optimization-strategy.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/optimization-strategy.md)",".claude/skills/perf-instability-cookbook/data/aliases.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/aliases.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/aliases.yaml)",".claude/skills/perf-instability-cookbook/data/modules.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/modules.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/modules.yaml)",".claude/skills/perf-instability-cookbook/data/patterns.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/patterns.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/patterns.yaml)",".claude/skills/perf-instability-cookbook/data/pending.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/pending.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/pending.yaml)",".claude/skills/perf-instability-cookbook/data/removed.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/removed.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/removed.yaml)",".claude/skills/perf-instability-cookbook/data/tags.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/tags.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/tags.yaml)",".claude/skills/perf-instability-cookbook/references/case-template.md":"# References Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/case-template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/case-template.md)",".claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md)",".claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md)",".claude/skills/perf-nsight-compute-analysis/references/cli-reference.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/cli-reference.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/cli-reference.md)",".claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md)",".claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md)",".claude/skills/perf-nsight-compute-analysis/references/python-report-api.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/python-report-api.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/python-report-api.md)",".claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md)",".claude/skills/perf-nsight-compute-analysis/references/sections-guide.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/sections-guide.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/sections-guide.md)",".claude/skills/perf-nsight-systems/references/app-preparation.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/app-preparation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/app-preparation.md)",".claude/skills/perf-nsight-systems/references/cli-post-collection.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/cli-post-collection.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/cli-post-collection.md)",".claude/skills/perf-nsight-systems/references/cli-profiling.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/cli-profiling.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/cli-profiling.md)",".claude/skills/perf-nsight-systems/references/expert-systems.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/expert-systems.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/expert-systems.md)",".claude/skills/perf-nsight-systems/references/nvtx-analysis.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/nvtx-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/nvtx-analysis.md)",".claude/skills/perf-nsight-systems/references/recipes-dl.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/recipes-dl.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/recipes-dl.md)",".claude/skills/perf-nsight-systems/references/stats-reports.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/stats-reports.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/stats-reports.md)",".claude/skills/perf-optimization-casebook/data/aliases.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-optimization-casebook/data/aliases.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/aliases.yaml)",".claude/skills/perf-optimization-casebook/data/patterns.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-optimization-casebook/data/patterns.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/patterns.yaml)",".claude/skills/perf-optimization-casebook/data/tags.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-optimization-casebook/data/tags.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/tags.yaml)",".claude/skills/perf-optimization-casebook/references/case-template.md":"# References Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/case-template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/case-template.md)",".claude/skills/perf-regression-cookbook/data/aliases.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/aliases.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/aliases.yaml)",".claude/skills/perf-regression-cookbook/data/modules.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/modules.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/modules.yaml)",".claude/skills/perf-regression-cookbook/data/patterns.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/patterns.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/patterns.yaml)",".claude/skills/perf-regression-cookbook/data/pending.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/pending.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/pending.yaml)",".claude/skills/perf-regression-cookbook/data/removed.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/removed.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/removed.yaml)",".claude/skills/perf-regression-cookbook/data/tags.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/tags.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/tags.yaml)",".claude/skills/perf-regression-cookbook/references/case-template.md":"# References Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/case-template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/case-template.md)",".claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md)",".claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md)",".claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md)",".claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md)",".claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md)",".claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py":"# Claude Directive: verify_workload.py\n\nPath: `.claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py)",".claude/skills/perf-torch-sync-free/references/sync-patterns.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-sync-free/references/sync-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/references/sync-patterns.md)",".claude/skills/perf-torch-sync-free/scripts/verify_workload.py":"# Claude Directive: verify_workload.py\n\nPath: `.claude/skills/perf-torch-sync-free/scripts/verify_workload.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/scripts/verify_workload.py)",".claude/skills/perf-workload-profiling/references/benchmarking-patterns.md":"# References Skill\n\nPath: `.claude/skills/perf-workload-profiling/references/benchmarking-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/benchmarking-patterns.md)",".claude/skills/perf-workload-profiling/references/nvtx-api.md":"# References Skill\n\nPath: `.claude/skills/perf-workload-profiling/references/nvtx-api.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/nvtx-api.md)",".claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md":"# References Skill\n\nPath: `.claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md)",".claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh":"# Claude Directive: detect_slurm_env.sh\n\nPath: `.claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh)",".claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md":"# References Skill\n\nPath: `.claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md)",".claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md":"# References Skill\n\nPath: `.claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md)",".claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md":"# References Skill\n\nPath: `.claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md)",".claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py":"# Claude Directive: build_slurm_script.py\n\nPath: `.claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py)",".claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh":"# Claude Directive: slurm_run_custom.sh\n\nPath: `.claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh)",".claude/skills/trtllm-test-specialist/references/agg_config_template.yaml":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/agg_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/agg_config_template.yaml)",".claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml)",".claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml)",".claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml":"# Claude Directive: smoke_test_config_template.yml\n\nPath: `.claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml)",".claude/skills/trtllm-test-specialist/references/test_config_template.yaml":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/test_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/test_config_template.yaml)",".claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md)",".claude/skills/trtllm-test-specialist/scripts/build_test_command.py":"# Claude Directive: build_test_command.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/build_test_command.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/build_test_command.py)",".claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py":"# Claude Directive: extract_test_markers.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py)",".claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py":"# Claude Directive: generate_benchmark_config.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py)",".claude/skills/trtllm-test-specialist/scripts/generate_report.py":"# Claude Directive: generate_report.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/generate_report.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/generate_report.py)",".claude/skills/trtllm-test-specialist/scripts/parse_config.py":"# Claude Directive: parse_config.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/parse_config.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/parse_config.py)",".claude/skills/trtllm-visualgen-bench/media/cat_piano.png":"# Claude Directive: cat_piano.png\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/cat_piano.png`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/cat_piano.png)",".claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png":"# Claude Directive: image_with_red_background.png\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png)",".claude/skills/trtllm-visualgen-bench/media/landscape.png":"# Claude Directive: landscape.png\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/landscape.png`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/landscape.png)",".claude/skills/trtllm-visualgen-bench/media/robot_153.jpg":"# Claude Directive: robot_153.jpg\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/robot_153.jpg`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/robot_153.jpg)",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json)",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json)",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json)",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json)",".claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json":"# Reference Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json)",".claude/skills/trtllm-visualgen-bench/scripts/serve_check.py":"# Claude Directive: serve_check.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/scripts/serve_check.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/scripts/serve_check.py)",".claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py":"# Claude Directive: serve_wait.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py)",".claude/skills/trtllm-visualgen-bench/serve/README.md":"# Serve Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/serve/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/serve/README.md)",".claude/skills/trtllm-visualgen-bench/tests/test_discovery.py":"# Claude Directive: test_discovery.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/tests/test_discovery.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_discovery.py)",".claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py":"# Claude Directive: test_serve_wait.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py)",".claude/skills/trtllm-visualgen-bench/tests/test_workloads.py":"# Claude Directive: test_workloads.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/tests/test_workloads.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_workloads.py)",".claude/skills/trtllm-visualgen-bench/workloads/README.md":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/README.md)",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml)","docs/source/commands/trtllm-serve/render-endpoints.md":"# Command: /render-endpoints\n\nPath: `docs/source/commands/trtllm-serve/render-endpoints.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/docs/source/commands/trtllm-serve/render-endpoints.md)","docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md":"# Command: /run-benchmark-with-trtllm-serve\n\nPath: `docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md)","tensorrt_llm/tools/profiler/host_profile_tools/README.md":"# host_profile_tools Documentation\n\nPath: `tensorrt_llm/tools/profiler/host_profile_tools/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/host_profile_tools/README.md)","tensorrt_llm/tools/profiler/nsys_profile_tools/README.md":"# nsys_profile_tools Documentation\n\nPath: `tensorrt_llm/tools/profiler/nsys_profile_tools/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/nsys_profile_tools/README.md)","tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json":"# Nsys_profile_tools Skill\n\nPath: `tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json)","tests/unittest/api_stability/references/batched_logits_processor.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/batched_logits_processor.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/batched_logits_processor.yaml)","tests/unittest/api_stability/references/calib_config.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/calib_config.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/calib_config.yaml)","tests/unittest/api_stability/references/completion_output.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/completion_output.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/completion_output.yaml)","tests/unittest/api_stability/references/conversation_params.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/conversation_params.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/conversation_params.yaml)","tests/unittest/api_stability/references/decoding_base_config.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/decoding_base_config.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/decoding_base_config.yaml)","tests/unittest/api_stability/references/guided_decoding_params.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/guided_decoding_params.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/guided_decoding_params.yaml)","tests/unittest/api_stability/references/llm.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/llm.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/llm.yaml)","tests/unittest/api_stability/references/logits_processor.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/logits_processor.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/logits_processor.yaml)","tests/unittest/api_stability/references/quant_config.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/quant_config.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/quant_config.yaml)","tests/unittest/api_stability/references/request_output.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/request_output.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/request_output.yaml)","tests/unittest/api_stability/references/sampling_params.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/sampling_params.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/sampling_params.yaml)","tests/unittest/api_stability/references/trtllm_serve_api.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/trtllm_serve_api.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/trtllm_serve_api.yaml)","tests/unittest/api_stability/references/trtllm_serve_cli.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/trtllm_serve_cli.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/trtllm_serve_cli.yaml)",".claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md)",".claude/skills/perf-host-optimization/references/patterns/gpu-graph.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/gpu-graph.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/gpu-graph.md)",".claude/skills/perf-host-optimization/references/patterns/loop-iteration.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/loop-iteration.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/loop-iteration.md)",".claude/skills/perf-host-optimization/references/patterns/python-overhead.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/python-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/python-overhead.md)",".claude/skills/perf-host-optimization/references/patterns/sync-alloc.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/sync-alloc.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/sync-alloc.md)",".claude/skills/perf-host-optimization/references/patterns/system.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/system.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/system.md)",".claude/skills/perf-instability-cookbook/references/communication/index.md":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/index.md)",".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md)",".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md)",".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md)",".claude/skills/perf-instability-cookbook/references/moe/index.md":"# Moe Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/moe/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/moe/index.md)",".claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md":"# Moe Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md)",".claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md)",".claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md)",".claude/skills/perf-optimization-casebook/references/communication/deepep.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/deepep.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/deepep.md)",".claude/skills/perf-optimization-casebook/references/communication/index.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/index.md)",".claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md)",".claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md)",".claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md)",".claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/index.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/index.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/index.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/index.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md)",".claude/skills/perf-regression-cookbook/references/autotuner/index.md":"# Autotuner Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/autotuner/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/autotuner/index.md)",".claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md":"# Autotuner Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md)",".claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md)",".claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md)",".claude/skills/perf-regression-cookbook/references/communication/index.md":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/index.md)",".claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md)",".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md":"# Disagg-orchestrator Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md)",".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md":"# Disagg-orchestrator Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md)",".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md":"# Disagg-orchestrator Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md)",".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md":"# Gemm-and-quantization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md)",".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md":"# Gemm-and-quantization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md)",".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md":"# Gemm-and-quantization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md)",".claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md":"# Kernel-fusion Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md)",".claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md":"# Kernel-fusion Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md)",".claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md":"# Kernel-fusion Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md":"# Kv-cache-transceiver Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md":"# Kv-cache-transceiver Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md":"# Kv-cache-transceiver Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md)",".claude/skills/perf-regression-cookbook/references/model-definition/index.md":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/index.md)",".claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md)",".claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md)",".claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md)",".claude/skills/perf-regression-cookbook/references/moe/index.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/index.md)",".claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md)",".claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md)",".claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md)",".claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md)",".claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md)",".claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md)",".claude/skills/perf-regression-cookbook/references/sampler/index.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/index.md)",".claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md)",".claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md)",".claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md)",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md)",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md)",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md)",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md)",".claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md":"# Spec-decode Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md)",".claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md":"# Spec-decode Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md)",".claude/skills/perf-regression-cookbook/references/spec-decode/index.md":"# Spec-decode Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/spec-decode/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/index.md)",".claude/skills/perf-regression-cookbook/references/tokenization/index.md":"# Tokenization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/tokenization/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/tokenization/index.md)",".claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md":"# Tokenization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md)","jenkins/scripts/cbts/cbts/rules/README.md":"# Rules Skill\n\nPath: `jenkins/scripts/cbts/cbts/rules/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/jenkins/scripts/cbts/cbts/rules/README.md)","tests/integration/defs/accuracy/references/acceptance_length.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/acceptance_length.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/acceptance_length.yaml)","tests/integration/defs/accuracy/references/cnn_dailymail.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/cnn_dailymail.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/cnn_dailymail.yaml)","tests/integration/defs/accuracy/references/gpqa_diamond.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/gpqa_diamond.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gpqa_diamond.yaml)","tests/integration/defs/accuracy/references/gsm8k.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/gsm8k.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gsm8k.yaml)","tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml)","tests/integration/defs/accuracy/references/humaneval.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/humaneval.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/humaneval.yaml)","tests/integration/defs/accuracy/references/json_mode_eval.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/json_mode_eval.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/json_mode_eval.yaml)","tests/integration/defs/accuracy/references/longbench_v1.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/longbench_v1.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/longbench_v1.yaml)","tests/integration/defs/accuracy/references/longbench_v2.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/longbench_v2.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/longbench_v2.yaml)","tests/integration/defs/accuracy/references/mmlu.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/mmlu.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/mmlu.yaml)","tests/integration/defs/accuracy/references/mmmu.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/mmmu.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/mmmu.yaml)","tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml)","tests/integration/defs/accuracy/references/videomme.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/videomme.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/videomme.yaml)","tests/integration/defs/accuracy/references/voxpopuli.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/voxpopuli.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/voxpopuli.yaml)","tests/integration/defs/accuracy/references/zero_scrolls.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/zero_scrolls.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/zero_scrolls.yaml)"},"files":{".claude/skills/exec-local-slurm/SKILL.md":"---\nname: exec-local-slurm\ndescription: >-\n  Submit and monitor a Slurm job on a local cluster. Supports two modes:\n  (1) Persistent allocation (default) — allocates nodes once via nohup salloc,\n  imports the container once, installs once, and reuses across runs by setting\n  SLURM env vars and running the sbatch script via bash. (2) One-shot sbatch —\n  submits a fully-generated Slurm script via sbatch, polls job status, reads\n  logs on completion, and reports results. Workflow-agnostic — handles pytest,\n  eval, benchmark, and custom scripts identically.\ntags: [slurm, execution, infrastructure]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Local Slurm Executor\n\nSubmit a Slurm job locally, monitor it, and report results. Uses persistent\nallocation by default to eliminate queue wait, container import, and install\noverhead on repeated runs.\n\n## Input (from orchestrator prompt)\n\nThe orchestrator passes these fields in the skill prompt:\n\n| Field | Description |\n|-------|-------------|\n| `script_path` | Full local path to the generated `.slurm` script |\n| `work_dir` | Local work directory for logs and artifacts |\n| `model_name` | Short model name for reporting |\n| `workflow_type` | `pytest`, `eval`, `custom`, or `benchmark` — for output parsing hints |\n| `success_patterns` | Comma-separated patterns indicating success |\n| `failure_patterns` | Comma-separated patterns indicating failure |\n| `log_file_pattern` | Log filename pattern with `%j` placeholder (e.g., `llama_auto_test_%j.out`) |\n| `persistent_mode` | Default: `true` for all local slurm workflows. Set to `false` to force one-shot sbatch (opt-out). |\n| `release_allocation` | `true` to release the current allocation and stop. Only set when the user explicitly says no more jobs are needed. Default: `false`. Never auto-release. |\n| `alloc_time_limit` | Walltime for the persistent allocation. Default: `04:00:00`. |\n| `docker_image` | Container image for persistent container import. From `job_spec.json`. |\n| `container_name` | Container name used in the `.slurm` script (e.g., `llama_auto_test`). Must match exactly. From `job_spec.json`. |\n| `container_mounts` | Comma-separated mount mappings. From `job_spec.json`. |\n| `repo_root` | Repo root path (for locating state file and project path). |\n| `slurm_params` | Slurm parameters object: `partition`, `account`, `nodes`, `ntasks`, `ntasks_per_node`, `gpus_per_node`. From `job_spec.json`. |\n\n## Procedure\n\n### Pre-step: Resolve Image and Build (when `build_project=true`)\n\nThis executor owns image selection and the build for the local SLURM cluster. Skip the entire pre-step when `build_project=false`.\n\n1. **Read the cluster GPU type** from `device_type` in `job_spec.json` (resolved upstream by env-check / internal-env-info). If absent, probe from any compute node via `srun -p <partition> --ntasks-per-node=1 nvidia-smi --query-gpu=name --format=csv,noheader | head -1`. If `skills/internal-env-info/` is not installed, skip the upstream lookup silently and rely on the srun probe — do not report the missing skill as an error.\n2. **Determine the partition's CPU arch.** Look up `job_spec.slurm_env.partitions[]` for the entry where `name == partition` and read its `arch` (`x86_64` / `aarch64`). Case-executor's Step 2.5 already detected this — do **not** run `scontrol show node` here. If `slurm_env` is absent or the entry's `arch` is null (case-executor ran without SLURM tools), fall back to inferring from `device_type`: Grace-based parts (`GB*` / `GH*`) → `aarch64`; everything else → `x86_64`.\n3. **Resolve the container image.** If `docker_image` is already set in `job_spec.json`, use it directly. Otherwise read `<repo_root>/jenkins/current_image_tags.properties` and pick the tag matching the partition's CPU arch.\n4. **Map GPU → build arch (`-a` flag):** `H100`/`H200` → `90-real`; `B200`/`GB200`/`B300`/`GB300` → `100-real`; `A100` → `80-real`; `L40S` → `89-real`. Default `100-real`.\n5. **Compile.** Invoke the `exec-slurm-compile` skill with `repo_dir=<repo_root>`, `partition`, `account`, `container_image=<resolved tag>`, `user_root_dir`, `arch`. Wait for completion.\n6. **On failure**, do not allocate or run the workload. Report `BUILD_FAILED` with the last 100 lines of the build log.\n\n`build_project`, `device_type`, `repo_root`, `partition`, `account`, `user_root_dir`, and (optionally) `docker_image` come from `job_spec.json`.\n\n### Step 0: Allocation Management\n\nThis step runs before any execution. It determines whether to reuse an existing\npersistent allocation, create a new one, or fall through to one-shot sbatch.\n\n#### Step 0A — Release mode\n\nIf `release_allocation=true`:\n\n1. Read `<repo_root>/work_dirs/.slurm_alloc.json`\n2. If file exists and `job_id` is present:\n   ```bash\n   scancel <job_id>\n   ```\n3. Delete the state file and `.salloc.log`\n4. Report: \"Allocation <job_id> released.\" Stop.\n\n#### Step 0B — One-shot mode\n\nIf `persistent_mode=false`, skip entirely to Step 1 (One-Shot Path).\n\n#### Step 0C — Validate existing allocation\n\nRead `<repo_root>/work_dirs/.slurm_alloc.json`.\n\n- If state file **does not exist** → go to Step 0D.\n- If state file exists, validate the allocation:\n  ```bash\n  squeue -j <job_id> -h -o \"%T %L\"\n  ```\n  - **RUNNING + remaining > 5 min + params compatible** → **reuse** (skip to Step 0F)\n  - **RUNNING + remaining <= 5 min** → warn user (\"Allocation expiring soon\"), `scancel <job_id>`, delete state file, go to Step 0D\n  - **RUNNING + params incompatible** → `scancel <job_id>`, delete state file, go to Step 0D\n  - **Empty output / error** (job gone) → stale state file, delete it, go to Step 0D\n\n**Params compatibility check:**\n- `partition` must match\n- `nodes` in state must be **>=** requested nodes (a 2-node allocation can serve 1-node jobs)\n- `docker_image` must match\n\n#### Step 0D — Justify and prepare allocation\n\nBefore allocating, ensure no orphaned allocations exist and log the reason:\n\n1. Check for existing persistent jobs:\n   ```bash\n   squeue -u $(whoami) -h -o \"%i %T %j\" --name=<account>-trtllm.persistent\n   ```\n2. If found → orphan (state file was missing/corrupt). Cancel it:\n   ```bash\n   scancel <orphan_job_id>\n   ```\n   Log: \"Released orphaned allocation <orphan_job_id> — state file was missing.\"\n3. Log justification: \"No reusable allocation found. Allocating <nodes> node(s) on partition <partition> for <alloc_time_limit>.\"\n\n#### Step 0E — Allocate new\n\nConvert `alloc_time_limit` from `HH:MM:SS` to seconds (e.g., `04:00:00` → `14400`).\n\n```bash\nmkdir -p <repo_root>/work_dirs\nnohup salloc --partition=<partition> --account=<account> \\\n  --nodes=<nodes> --time=<alloc_time_limit> \\\n  --job-name=<account>-trtllm.persistent \\\n  sleep <alloc_time_limit_seconds> \\\n  > <repo_root>/work_dirs/.salloc.log 2>&1 &\n```\n\nRetrieve job ID:\n```bash\nsqueue -u $(whoami) -h -o \"%i\" --name=<account>-trtllm.persistent\n```\n\nIf no job appears after 10 seconds, read `<repo_root>/work_dirs/.salloc.log`\nfor errors (bad partition, invalid account, etc.) and report to user.\n\nOtherwise, poll until RUNNING (every 10s, max 60 polls). Once RUNNING, get\nthe nodelist:\n```bash\nsqueue -j <JOB_ID> -h -o \"%N\"\n```\n\n**Import container** on all nodes (using the **same container name** as the\n`.slurm` script — critical for pyxis reuse):\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-image=<docker_image> \\\n  --container-name=<container_name> true\n```\n\n**Warm up filesystem mounts** on all nodes — `ls` each mounted path so\nLustre/NFS metadata is cached for later use:\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-name=<container_name> \\\n  --container-mounts=<container_mounts> \\\n  bash -c 'for p in <mount_target_1> <mount_target_2> ...; do ls \"$p\" > /dev/null 2>&1; done'\n```\nParse mount targets from `container_mounts` — the right-hand side of each\n`host:container` pair.\n\n**Check GPU status** on all nodes (no container needed — `nvidia-smi` is on\nthe host):\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  bash -c 'echo \"=== $(hostname) ===\" && nvidia-smi --query-compute-apps=pid,name,used_memory --format=csv,noheader'\n```\n- If output shows no processes → GPUs are clean, proceed\n- If unexpected processes found → warn user: \"GPU processes found on\n  <node>: <process_list>. These may interfere with the job.\"\n\n**Run install** on all nodes with `--container-writable` (so packages persist\nin the named container across srun calls). Install all common requirements\nupfront so that subsequent jobs of any workflow type can skip install:\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-name=<container_name> --container-writable \\\n  --container-mounts=<container_mounts> \\\n  bash -c 'cd <repo_root> && pip install -e . && pip install -r requirements-dev.txt && \\\n    if [ -f examples/trtllm-eval/requirements.txt ]; then pip install -r examples/trtllm-eval/requirements.txt; fi'\n```\nFor custom workflow with `skip_install=true`, skip this step entirely.\n\n**Write state file** `<repo_root>/work_dirs/.slurm_alloc.json`:\n```json\n{\n  \"job_id\": \"<JOB_ID>\",\n  \"container_name\": \"<container_name>\",\n  \"nodelist\": \"<NODELIST>\",\n  \"partition\": \"<partition>\",\n  \"account\": \"<account>\",\n  \"nodes\": <nodes>,\n  \"gpus_per_node\": <gpus_per_node>,\n  \"docker_image\": \"<docker_image>\",\n  \"container_mounts\": \"<container_mounts>\",\n  \"allocated_at\": \"<ISO timestamp>\",\n  \"time_limit\": \"<alloc_time_limit>\",\n  \"installed\": true\n}\n```\n\nProceed to Step 1 (Persistent Path).\n\n#### Step 0F — Reuse path validation\n\nAllocation is valid. Check if container name matches current request:\n\n- If `container_name` in state file **matches** `container_name` from\n  `job_spec.json` → proceed to Step 1 (Persistent Path). Container and install\n  are already set up.\n- If **different** (model changed) → import the new container, warm up mounts,\n  and install:\n  ```bash\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-image=<docker_image> \\\n    --container-name=<new_container_name> true\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-name=<new_container_name> \\\n    --container-mounts=<container_mounts> \\\n    bash -c 'for p in <mount_targets>; do ls \"$p\" > /dev/null 2>&1; done'\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-name=<new_container_name> --container-writable \\\n    --container-mounts=<container_mounts> \\\n    bash -c 'cd <repo_root> && pip install -e . && pip install -r requirements-dev.txt'\n  ```\n  Update `container_name` in state file. Proceed to Step 1 (Persistent Path).\n\n---\n\n### Step 1: Execute\n\nTwo execution paths depending on `persistent_mode`.\n\n#### Persistent Path (persistent_mode=true)\n\n**Pre-flight time check** — verify the allocation has enough remaining time:\n```bash\nsqueue -j <job_id> -h -o \"%L\"\n```\nIf remaining time < job's `time_limit` from `slurm_params`, warn user:\n\"Allocation has <remaining> left but job expects <time_limit>. The job may\nbe killed early.\"\n\n**Pre-flight GPU check** — verify GPUs are not occupied by leftover processes:\n```bash\nsrun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n  bash -c 'procs=$(nvidia-smi --query-compute-apps=pid,name,used_memory --format=csv,noheader 2>/dev/null); [ -n \"$procs\" ] && echo \"WARNING: GPU processes on $(hostname): $procs\"'\n```\nIf unexpected processes found, warn user before proceeding.\n\n**Run the sbatch script** with SLURM env vars — the `#SBATCH` directives are\ncomments when run via `bash`; the inner `srun` inherits the env vars and uses\nthe persistent allocation:\n```bash\nexport SLURM_JOB_ID=<job_id>\nexport SLURM_JOB_NUM_NODES=<nodes>\nexport SLURM_NNODES=<nodes>\nexport SLURM_NTASKS=<ntasks>\nexport SLURM_NTASKS_PER_NODE=<ntasks_per_node>\nexport SLURM_NODELIST=<nodelist>\nexport SLURM_JOB_NODELIST=<nodelist>\nbash <script_path> 2>&1 | tee <work_dir>/<model_name>_<workflow_type>.log\n```\n\nRun with `run_in_background`. Container reuse is automatic (pyxis skips\nre-import when `--container-name` already exists on the node). The install\nstep (Step 1 srun) always runs to keep the container up to date.\n\n**Step cancellation:** If a job hangs or the user wants to abort, cancel just\nthe srun step without killing the allocation:\n```bash\n# List active steps:\nsqueue -s -j <job_id>\n# Cancel a specific step (e.g., step 0):\nscancel <job_id>.0\n```\nThe allocation stays RUNNING — new jobs can run immediately. The job name\n`<account>-trtllm.persistent` uniquely identifies allocations created by this\nskill via `squeue --name=<account>-trtllm.persistent`.\n\n#### One-Shot Path (persistent_mode=false)\n\n```bash\nsbatch <script_path>\n```\n\nParse the job ID from output: `Submitted batch job <JOB_ID>`.\n\nIf sbatch fails, report the error immediately and stop.\n\n### Step 2: Report Submission\n\n**One-shot mode:**\n```\nJob ID: <JOB_ID>\nScript: <script_path>\nWork directory: <work_dir>\nLog files: <work_dir>/<log_file_pattern with %j replaced>\nMonitor: squeue -j <JOB_ID>\n```\n\n**Persistent mode:**\n```\nAllocation: <job_id> (persistent, remaining: <time>)\nContainer: <container_name>\nScript: <script_path>\nWork directory: <work_dir>\nLog file: <work_dir>/<model_name>_<workflow_type>.log\n```\n\n### Step 3: Poll / Wait for Completion (with Hang Monitoring)\n\nBoth polling modes must also monitor the log file for hang indicators. Every\n60 seconds, check the log file for the pattern `hang detected`\n(case-insensitive). If found, terminate the job and skip to Step 5 with\n`HANG_DETECTED` status.\n\n**One-shot mode:** Poll every 30 seconds, max 60 polls (30 minutes):\n\n```bash\nsqueue -j <JOB_ID> -h -o \"%T %M %R\"\n```\n\n- Report state transitions: `PENDING` → `RUNNING` → `COMPLETING` → done\n- If `PENDING` for > 5 minutes, check the reason:\n  ```bash\n  squeue -j <JOB_ID> -o \"%i %T %r %S\"\n  ```\n- The job is done when `squeue -j <JOB_ID>` returns empty (no rows)\n- If max polls exceeded, report the job is still running and provide manual\n  check commands\n- **Hang check** (every other poll, i.e., every 60s): Once the job is RUNNING,\n  resolve the log file path (`<work_dir>/<log_file_pattern>` with `%j` replaced\n  by `<JOB_ID>`) and check:\n  ```bash\n  grep -qi \"hang detected\" <log_file> 2>/dev/null && echo \"HANG_DETECTED\"\n  ```\n  If `HANG_DETECTED` is printed, immediately cancel the job:\n  ```bash\n  scancel <JOB_ID>\n  ```\n  Then read the last 200 lines of the log file and skip to Step 5.\n\n**Persistent mode:** While waiting for `run_in_background` completion\nnotification, launch a hang monitoring loop using `run_in_background`:\n\n```bash\nwhile true; do\n  sleep 60\n  LOG_FILE=\"<work_dir>/<model_name>_<workflow_type>.log\"\n  if [ -f \"$LOG_FILE\" ] && grep -qi \"hang detected\" \"$LOG_FILE\"; then\n    echo \"HANG_DETECTED: Found 'hang detected' in log file\"\n    squeue -s -j <job_id> -h -o \"%i\" | head -1 | xargs -r scancel 2>/dev/null\n    exit 1\n  fi\ndone\n```\n\nWhen the main process completes normally, kill the monitoring loop. If the\nmonitor fires first, it cancels the active Slurm step (not the allocation\nitself — the allocation stays for future runs).\n\n### Step 4: Collect Results\n\nOn completion:\n\n1. **Get exit code:**\n   - **One-shot:** `sacct -j <JOB_ID> --format=JobID,State,ExitCode,Elapsed,MaxRSS -n`\n   - **Persistent:** From the `run_in_background` result exit code\n\n2. **Resolve log file paths:**\n   - **One-shot:** Replace `%j` in `log_file_pattern` with `<JOB_ID>`:\n     - Output: `<work_dir>/<MODEL_NAME>_..._<JOB_ID>.out`\n     - Error: `<work_dir>/<MODEL_NAME>_..._<JOB_ID>.err`\n   - **Persistent:** Single log at `<work_dir>/<model_name>_<workflow_type>.log`\n\n3. **Read log files:**\n   - Read the last 100 lines of the output log\n   - If job failed, also read the first 50 lines to catch early errors\n   - **One-shot only:** If job failed, read the last 50 lines of the `.err` file\n\n4. **Search for patterns:**\n   - Grep the log file for each `success_patterns` entry\n   - Grep the log file for each `failure_patterns` entry\n\n### Step 5: Report\n\nReturn a structured result:\n\n```\nJob ID: <JOB_ID> (or allocation <job_id> for persistent mode)\nStatus: COMPLETED | FAILED | TIMEOUT | CANCELLED | OUT_OF_MEMORY | HANG_DETECTED\nExit code: <exit code>\nElapsed: <time>\nLog files:\n  stdout: <path>\n  stderr: <path> (one-shot only)\nWork directory: <work_dir>\nSummary: <relevant output — pytest summary, accuracy score, throughput, or last lines>\nErrors: <if failed, relevant error output>\n```\n\n### Output Parsing by Workflow Type\n\n- **pytest**: Look for `X passed, Y failed in Zs` summary line. Status maps:\n  `COMPLETED` + 0 exit → `PASSED`; non-zero → `FAILED`\n- **eval**: Look for `accuracy:` lines. If `--check_accuracy` was used,\n  assertion error means threshold not met\n- **benchmark**: Look for throughput/latency numbers. Check for\n  `8_done_*.txt` completion marker\n- **custom**: No specific patterns — report last 10 lines of output\n\n### Failure Diagnosis\n\n| Symptom | Likely Cause | Action |\n|---------|-------------|--------|\n| `TIMEOUT` | Job exceeded time limit | Report and suggest increasing time |\n| `OUT_OF_MEMORY` | OOM on node | Report and suggest reducing batch size |\n| Exit `1:0` + `FAILED` | Non-zero exit from srun | Read `.err` file for details |\n| `CANCELLED` | Preempted or user cancelled | Check `sacct --format=JobID,State,Reason` |\n| Empty output log | Container failed to start | Check `.err` / `.salloc.log` for mount or image errors |\n| `HANG_DETECTED` in log | Process/GPU hang (deadlock, NCCL timeout) | Job auto-cancelled; check last 200 log lines for root cause |\n| `srun: error: Unable to create step` | Allocation expired | Delete state file, re-run (will re-allocate) |\n| `srun: error: ... does not exist` | Named container missing | Delete state file, re-run (will re-import) |\n\n## Rules\n\n- Never use `squeue` with the `-uall` flag\n- Always use `sacct` (not `squeue`) for final job status in one-shot mode\n- Poll with `run_in_background` + `sleep 30 &&` to avoid blocking (one-shot)\n- Never read full log files — use Read with offset/limit\n- If the Slurm script path doesn't exist, report the error immediately\n- For cluster hardware info (GPU types, partitions, constraints, storage\n  paths), use the `trtllm-agent-toolkit:internal-env-info` skill\n  when it is installed. If the skill is absent, skip the lookup silently and\n  use the caller-supplied fields or the srun probe in Step 1 — do not report\n  the missing skill as an error.\n- Never auto-release an allocation after job completion — it should persist\n  for potential future runs until walltime expires or user explicitly releases\n- Must pass `-N <nodes>` to every direct `srun --jobid` call — without it,\n  srun defaults to 1 node even in a multi-node allocation\n- The container name used in setup steps must **exactly match** the\n  `--container-name` in the `.slurm` script, otherwise pyxis will re-import\n  the image and lose `--container-writable` state\n",".claude/skills/exec-local-compile/SKILL.md":"---\nname: exec-local-compile\ndescription: Compile TensorRT-LLM on a compute node inside a Docker container. Use this when already on a compute node with GPUs visible.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Compile TensorRT-LLM (Local / Compute Node)\n\nCompile TensorRT-LLM from source on a compute node inside a Docker container.\n\n## When to Use\n\n| Scenario | Use This Skill? |\n|----------|----------------|\n| On a compute node with GPUs visible (`nvidia-smi` works) | Yes |\n| On a SLURM login node (no GPUs) | No — use `exec-slurm-compile` instead |\n\n## Prerequisites\n\n- You are inside a Docker/enroot container on a compute node\n- `nvidia-smi` succeeds (GPUs visible)\n- `/usr/local/tensorrt` exists (TensorRT installation in the container)\n\n## Instructions\n\n### Step 1: Verify Environment\n\nRun `nvidia-smi` to confirm you are on a compute node with GPU access.\n\n### Step 2: Locate the Codebase\n\n`cd` to the TensorRT-LLM repository. If the path is not provided by the user, ask for it.\n\n### Step 3: (Optional) Checkout Branch\n\nIf the user specifies a branch (e.g., \"compile ToT\"), checkout and pull:\n```bash\ngit checkout main && git pull\n```\n\n### Step 4: Build\n\nRun the build command (**incremental by default** — omit `-c`/`--clean` unless explicitly requested or the incremental build fails):\n\n```bash\n./scripts/build_wheel.py --use_ccache -a \"<arch>\" -f --nvtx\n```\n\nReplace `<arch>` with the target GPU architecture (see Architecture Reference below). If not specified by the user, auto-detect from `nvidia-smi`.\n\n### Step 5: Install\n\n```bash\npip install -e .[devel]\n```\n\n### Step 6: Verify\n\n```bash\npython3 -c \"import tensorrt_llm; print(tensorrt_llm.__version__)\"\n```\n\n## Build Flags\n\n| Flag | Description |\n|------|-------------|\n| `-a \"<arch>\"` | Target GPU architecture(s) |\n| `--nvtx` | Enable NVTX markers for profiling |\n| `--use_ccache` | Use ccache for faster recompilation |\n| `-f` / `--fast_build` | Skip some kernels for faster dev compilation. **Always use for dev builds.** |\n| `-c` / `--clean` | Clean build directory before building. Only when needed (see below). |\n| `--skip_building_wheel` | Build in-place without creating a wheel file |\n| `--no-venv` | Skip virtual environment creation |\n\n## Architecture Reference\n\n| Value | GPU Family |\n|-------|-----------|\n| `\"100-real\"` | Blackwell (B200, GB200) |\n| `\"90-real\"` | Hopper (H100, H200) |\n| `\"89-real\"` | Ada Lovelace (L40S) |\n| `\"80-real\"` | Ampere (A100) |\n| `\"90;100-real\"` | Multiple architectures |\n\n## Incremental vs. Clean Builds\n\n**Default to incremental builds** — CMake only recompiles changed files, saving significant time.\n\nUse a **clean build** (`-c`) only when:\n- The user explicitly requests a clean/fresh build\n- An incremental build fails with linker errors, stale object files, or CMake cache issues\n- Major branch changes (e.g., rebasing across many commits) that may invalidate the build cache\n- Build system files changed (`CMakeLists.txt`, `*.cmake`)\n",".claude/skills/exec-env-check/SKILL.md":"---\nname: exec-env-check\ndescription: >-\n  Check the local execution environment for GPU availability, Docker support,\n  and Slurm access. Returns the execution scenario (`satisfied, local, docker`,\n  `satisfied, local, direct`, `satisfied, slurm, local`, or `not_satisfied`),\n  the number of available GPUs, and the GPU type. On Slurm login nodes\n  without local GPUs, the cluster is identified by delegating the hostname\n  to internal-env-info (hostname-based mode), which owns the\n  hostname → cluster_name patterns; GPU type and gpus_per_node then come\n  from that skill's reference files. If internal-env-info is not\n  installed, the scenario falls back to `not_satisfied` without probing\n  compute nodes via srun.\ntags: [infrastructure, slurm, environment]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# TensorRT-LLM Environment Check\n\nDetect whether the current machine can run a GPU workload locally (Docker) or via Slurm, and report hardware details.\n\n## Input\n\n| Field | Description | Required |\n|-------|-------------|----------|\n| `required_devices` | Minimum number of GPUs needed | Yes |\n| `account` | Slurm account (unused for GPU probing, kept for compatibility) | No |\n\n## Procedure\n\n### 1. Check local GPUs\n\n```bash\ntimeout 5 nvidia-smi --query-gpu=name,memory.total --format=csv,noheader 2>/dev/null\n```\n\nIf this succeeds:\n- Count the number of GPU lines → `available_gpus`\n- Extract the GPU name from the first line → `device_type` (e.g., `NVIDIA B200`, `NVIDIA H100 80GB HBM3`)\n- Normalize `device_type`: strip `NVIDIA ` prefix and trailing memory info to get the short name (e.g., `B200`, `H100`, `A100`, `L40S`, `RTX 6000`)\n\nIf `nvidia-smi` fails or returns no GPUs → `available_gpus = 0`, `device_type = null`.\n\n### 2. Check if local GPUs are sufficient\n\nIf `available_gpus >= required_devices`, continue to step 2a to determine whether Docker is available on this host.\n\n### 2a. Check Docker availability\n\n```bash\ncommand -v docker >/dev/null 2>&1 && timeout 3 docker info >/dev/null 2>&1\n```\n\n- If the command succeeds (Docker CLI exists **and** the daemon responds) → **Result**: `satisfied, local, docker`\n- Otherwise (Docker CLI missing, daemon not running, or permission denied) → **Result**: `satisfied, local, direct`\n\nInclude `available_gpus` in the result. Do NOT include `device_type` — local execution does not need it.\n\n### 3. Check Slurm availability\n\nIf local GPUs are insufficient (or as additional detection), check for Slurm:\n\n```bash\nwhich squeue 2>/dev/null && squeue --version 2>/dev/null\n```\n\nIf Slurm is NOT available → go to step 5.\n\n### 4. Resolve GPU type\n\n**Optional dependency.** Before doing anything else in this step, check whether `skills/internal-env-info/` exists in the toolkit. If it does **not**, skip this step entirely and return `scenario: not_satisfied` with `available_gpus = 0`, `device_type = null`, `gpus_per_node = null`, `cluster_name = null`, `default_models_repo = null`, and `default_user_root_dir = null`. Do **not** report this as an error — the skill is an internal-only dependency.\n\nWhen Slurm is available but local `nvidia-smi` returned no GPUs or `device_type` is null (login nodes typically have no GPUs), capture the hostname and delegate cluster identification to `internal-env-info`:\n\n```bash\nhostname -f 2>/dev/null || hostname\n```\n\n- Pass the captured hostname to `internal-env-info` in **hostname-based mode**. That skill owns the NVIDIA-internal login-host patterns and the hostname → `<cluster>` mapping; do **not** parse the hostname or hard-code any cluster identifier here.\n- It returns the standard output template (`cluster_name`, `device_type`, `gpus_per_node`, `default_models_repo`, `default_user_root_dir`) plus supplementary field (`mfa_style`).\n- Set `available_gpus` = `gpus_per_node` (if resolved).\n- Set `cluster_name` = the value returned (a placeholder `<cluster>` token in this skill). The orchestrator uses this to fetch per-cluster info (`mfa_style`, `default_models_repo`, `default_user_root_dir`, `gpus_per_node`) from `internal-env-info`; connection fields (`mounts`, `ssh_host`, `partition`, etc.) come from caller-supplied inputs in `job_spec.json` (with `internal-env-info` default values / ask-the-user fallbacks).\n- Set `default_user_root_dir` = the user root directory returned (with `<user_name>` substituted with the actual SLURM username). Set to `null` if not found.\n- If `internal-env-info` returns `null` for `cluster_name` (no pattern matched), set `device_type = null`, `cluster_name = null`, `default_user_root_dir = null` and fall through to Step 5 (`not_satisfied`).\n\n**Result**: `satisfied, slurm, local` with `device_type`, `gpus_per_node`, `cluster_name`, and `default_user_root_dir`.\n\n### 5. Not satisfied\n\nIf neither local GPUs nor Slurm is available:\n- **Result**: `not_satisfied`\n\n## Output\n\nReturn a single structured result:\n\n```\nscenario: <satisfied, local, docker | satisfied, local, direct | satisfied, slurm, local | not_satisfied>\navailable_gpus: <N>\ndevice_type: <short GPU/device name or null>\ngpus_per_node: <N or null>\ncluster_name: <cluster name or null>\ndefault_models_repo: <host path to llm-models directory, or null>\ndefault_user_root_dir: <per-user root directory on cluster storage, or null>\n```\n\n**Examples:**\n\n```\nscenario: satisfied, local, docker\navailable_gpus: 4\n```\n\n```\nscenario: satisfied, local, direct\navailable_gpus: 4\n```\n\n```\nscenario: satisfied, slurm, local\navailable_gpus: 4\ndevice_type: B200\ngpus_per_node: 4\ncluster_name: <cluster>\ndefault_user_root_dir: <default_user_root_dir>/<user_name>\n```\n\n```\nscenario: not_satisfied\navailable_gpus: 0\ndevice_type: null\ncluster_name: null\n```\n\n## Rules\n\n- Never install drivers or modify the system\n- If `nvidia-smi` hangs, use a 5-second timeout: `timeout 5 nvidia-smi ...`\n- GPU type is derived from the hostname by matching the cluster name via the `internal-env-info` skill when it is installed — no `srun` allocation needed. If that skill is absent, the scenario falls back to `not_satisfied` (see Step 4); do not error.\n- Report the GPU type exactly as found in the mapping (do not guess or fabricate)\n",".claude/skills/exec-slurm-compile/SKILL.md":"---\nname: exec-slurm-compile\ndescription: Compile TensorRT-LLM on a SLURM cluster. Covers submitting a batch job with a container image, monitoring the job, and verifying the build. Use when the user wants to compile TRT-LLM remotely via SLURM rather than on a local compute node.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Compile TensorRT-LLM on SLURM Cluster\n\nSubmit, monitor, and verify a TensorRT-LLM compilation job on a SLURM cluster using enroot containers.\n\n## When to Use\n\n| Scenario | Use This Skill? |\n|----------|----------------|\n| User wants to compile TRT-LLM on a SLURM cluster | Yes |\n| User is already on a compute node and wants to compile | No — use `exec-local-compile` skill instead |\n\n## Finding the Docker Image\n\nThe official Docker image tag for a given TensorRT-LLM version is recorded in the repo itself:\n\n```\n<repo_dir>/jenkins/current_image_tags.properties\n```\n\nRead this file to find the current image URL. The format is `<registry-host>/<namespace>/tensorrt-llm:<tag>`, for example `<your-registry>/<your-namespace>/tensorrt-llm:pytorch-25.12-py3-aarch64-ubuntu24.04-trt10.14.1.48-skip-tritondevel-202602011118-10901`. Substitute `<your-registry>` and `<your-namespace>` with your own container registry coordinates.\n\n\n## Pre-dumping the Container Image (enroot import)\n\nSLURM clusters using enroot/pyxis require a `.sqsh` container image. To avoid download overhead at compile time, **pre-dump the image in advance** using the `enroot-import` companion script:\n\n```bash\n# Basic usage — submits a SLURM job on a CPU partition to import the image\nenroot-import --partition cpu_datamover --debug <docker_image_url>\n```\n\nThe script submits an `sbatch` job that runs `enroot import docker://<image_url>` and produces a `.sqsh` file in the current directory. The output on stdout is the SLURM job ID.\n\n### enroot-import flags\n\n| Flag | Description |\n|------|-------------|\n| `-p, --partition` | SLURM partition for the import job (use a CPU partition like `cpu_datamover`) |\n| `-d, --debug` | Enable debug output and preserve the SLURM log (recommended) |\n| `-o, --output` | Custom output path for the `.sqsh` file |\n| `-A, --account` | SLURM account (defaults to user's first account) |\n| `-t, --time` | Time limit for the import job (default: 1 hour) |\n| `-n, --just-print` | Print the sbatch command without executing |\n| `-J, --job-name` | Custom job name |\n\n### enroot-import workflow\n\n1. Read the image tag from `jenkins/current_image_tags.properties` in the TRT-LLM repo.\n2. Run `enroot-import` to submit the import job:\n   ```bash\n   cd <directory_where_sqsh_should_be_stored>\n   <path_to>/enroot-import --partition cpu_datamover --debug <image_url>\n   ```\n   **IMPORTANT:** Replace the first `/` after the registry host with `#` to avoid credential issues. For example, `<your-registry>/<your-namespace>/tensorrt-llm:xxx` becomes `<your-registry>#<your-namespace>/tensorrt-llm:xxx`.\n3. Wait for the import job to complete (`squeue -j <job_id>`).\n4. The resulting `.sqsh` file is the `container_image` used in the compile step.\n\n\n## Prerequisites\n\nThe user must provide (or you must ask for) these values:\n\n| Parameter | Description | Example |\n|-----------|-------------|---------|\n| `container_image` | Path to `.sqsh` container image (see enroot import above) | `/path/to/pytorch.sqsh` |\n| `repo_dir` | Path to the TensorRT-LLM repository | `/path/to/TensorRT-LLM` |\n| `user_root_dir` | Top-level directory to bind-mount into the container | `/shared/users` |\n| `partition` | SLURM partition | `batch` |\n| `account` | SLURM account | `my_account` |\n\nOptional parameters:\n\n| Parameter | Description | Default |\n|-----------|-------------|---------|\n| `jobname` | SLURM job name | `trtllm-compile.<username>` |\n| `gpu_count` | Number of GPUs to request | `4` |\n| `time_limit` | Job time limit | `02:00:00` |\n| `arch` | GPU architecture(s) for `-a` flag | `100-real` |\n| `extra_build_args` | Extra flags for `build_wheel.py` | (none) |\n\n## Companion Scripts\n\nThis skill includes three companion scripts in `scripts/`:\n\n| Script | Purpose |\n|--------|---------|\n| `enroot-import` | Pre-dump a Docker image to `.sqsh` via a SLURM batch job |\n| `submit_compile.sh` | Template for submitting the SLURM job — copy and customize |\n| `compile.slurm` | SLURM batch script — launches the container and calls `compile.sh` |\n| `compile.sh` | Runs inside the container — executes `build_wheel.py` |\n\nScripts directory: `skills/exec-slurm-compile/scripts/`\n\n## Instructions\n\nFollow these steps in order:\n\n### Step 0: Resolve the Container Image (if needed)\n\nIf the user does not already have a `.sqsh` container image:\n\n1. Read the Docker image tag from `<repo_dir>/jenkins/current_image_tags.properties`.\n2. Use `enroot-import` to pre-dump it:\n   ```bash\n   cd <directory_for_sqsh_files>\n   <scripts_dir>/enroot-import --partition cpu_datamover --debug <image_url>\n   ```\n3. Monitor the import job with `squeue -j <job_id>`.\n4. Once complete, the `.sqsh` file path becomes the `container_image` parameter.\n\nIf the user already has a `.sqsh` file, skip this step.\n\n### Step 1: Gather Information\n\nAsk the user for any missing prerequisite values listed above. At minimum you need:\n- `container_image` (or the Docker image URL — then run Step 0 first)\n- `repo_dir`\n- `user_root_dir`\n- `partition` and `account`\n\nIf the user has used this workflow before, check if previous values are stored in memory files.\n\n### Step 2: Prepare the Scripts Directory\n\nThe compile scripts must be accessible from inside the container (i.e., under `user_root_dir`). Either:\n\n**Option A** — Copy companion scripts to a location under `user_root_dir`:\n```bash\nscripts_dir=<user_root_dir>/<username>/workspace/tensorrt_llm_scripts\nmkdir -p ${scripts_dir}/log\ncp skills/exec-slurm-compile/scripts/compile.sh ${scripts_dir}/\ncp skills/exec-slurm-compile/scripts/compile.slurm ${scripts_dir}/\nchmod +x ${scripts_dir}/compile.sh ${scripts_dir}/compile.slurm\n```\n\n**Option B** — If the user already has scripts at a known location, use those directly.\n\n### Step 3: Submit the Job\n\nRun `sbatch` from the login node (or a node with SLURM client access):\n\n```bash\nsbatch \\\n    --nodes=1 --ntasks=1 --ntasks-per-node=1 \\\n    --gres=gpu:<gpu_count> \\\n    --partition=<partition> \\\n    --account=<account> \\\n    --job-name=<jobname> \\\n    --time=<time_limit> \\\n    <scripts_dir>/compile.slurm \\\n    <container_image> <user_root_dir> <scripts_dir> <repo_dir>\n```\n\nCapture and report the job ID from the `sbatch` output.\n\n### Step 4: Monitor the Job (Proactive — Do NOT Wait for User)\n\n**You MUST actively poll the job until it completes.** Do not submit and walk away.\n\n```bash\n# Check job status (repeat every 30-60 seconds)\nsqueue -j <job_id> -o \"%.18i %.9P %.30j %.8u %.2t %.10M %.6D %R\"\n\n# Once running, periodically tail the log (do NOT use tail -f, use tail -30 instead)\ntail -30 <scripts_dir>/log/compile_<job_id>.srun.log\n```\n\n**Monitoring loop:**\n1. Poll `squeue -j <job_id>` to check state\n2. If `PD` (pending) — report the reason, keep polling every 30-60s\n3. If `R` (running) — tail the build log every 30-60s; look for `[XX%] Building`, errors, or completion\n4. If the job disappears from `squeue`, it has finished — proceed to Step 5\n5. If `F` (failed) — immediately read the full log and report the error\n\n**Progress indicators to look for in the log:**\n- `[XX%] Building CXX object...` — compilation progress\n- `Linking CXX...` — link phase\n- `FAILED:`, `error:`, `fatal error:` — build failure\n- `Successfully built` — success\n\n### Step 5: Verify the Build\n\nOnce the job completes, check for success:\n\n```bash\n# Check SLURM exit code\nsacct -j <job_id> --format=JobID,State,ExitCode,Elapsed\n\n# Check the build log for errors\ntail -50 <scripts_dir>/log/compile_<job_id>.srun.log\n```\n\nA successful build ends with a message like `Successfully built tensorrt_llm` or completes without error.\n\n## Common Build Flags Reference\n\n| Flag | Description |\n|------|-------------|\n| `-a \"100-real\"` | Target architecture — `100` for Blackwell, `90` for Hopper, etc. |\n| `--nvtx` | Enable NVTX markers for profiling |\n| `--no-venv` | Skip virtual environment creation |\n| `--use_ccache` | Use ccache to speed up recompilation |\n| `--skip_building_wheel` | Build in-place without creating a wheel file |\n| `-f` | Fast build — skip some kernels for faster dev compilation |\n| `-c` | Clean build — wipe build directory before building |\n\nCommon architecture values:\n- `\"100-real\"` — Blackwell (B200, GB200)\n- `\"90-real\"` — Hopper (H100, H200)\n- `\"89-real\"` — Ada Lovelace (L40S)\n- `\"80-real\"` — Ampere (A100)\n- `\"90;100-real\"` — Multiple architectures\n\n## Troubleshooting\n\n| Issue | Solution |\n|-------|----------|\n| `sbatch: error: invalid partition` | Verify partition name with `sinfo -s` |\n| `sbatch: error: invalid account` | Check available accounts with `sacctmgr show assoc user=$USER` |\n| Container image not found | Verify the `.sqsh` path exists and is readable |\n| Build OOM (out of memory) | Reduce parallelism with `-j <N>` flag to `build_wheel.py` |\n| `srun: error: Unable to create step` | The node may lack enroot/pyxis — check with cluster admin |\n| Job stuck in `PD` state | Check `squeue -j <id> -o %R` for the reason (e.g., resource limits, priority) |\n| `enroot import` fails with auth error | Check `~/.config/enroot/.credentials` has the correct registry credentials |\n| `enroot import` produces empty/corrupt `.sqsh` | Re-run with `--debug` and check the SLURM log; verify the image URL has no `https://` prefix |\n| Weird compile issues | Retry with a clean build (`-c` flag) |\n| `QOSGrpNodeLimit` shown in `NODELIST(REASON)` | Not a blocker, just wait for the job to get scheduled |\n\n## Example Interaction\n\n**User**: \"Compile TRT-LLM on the OCI cluster\"\n\n**Agent actions**:\n1. Ask for container image path, repo path, mount dir (if not known)\n2. Confirm partition/account for OCI cluster\n3. Copy scripts to accessible location under user_root_dir\n4. Submit with `sbatch`\n5. Report job ID\n6. Monitor with `squeue` until complete\n7. Check logs and report success/failure\n",".claude/skills/exec-local-docker/SKILL.md":"---\nname: exec-local-docker\ndescription: >-\n  Execute a TensorRT-LLM workload locally in Docker. Runs a fully-resolved\n  Docker command in background, monitors completion, reads logs, and reports\n  results. Workflow-agnostic — does not need to know if the workload is pytest,\n  eval, benchmark, or a custom script.\ntags: [docker, execution, infrastructure]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Local Docker Executor\n\nRun a Docker command locally, monitor it, and report results.\n\n## Input (from orchestrator prompt)\n\nThe orchestrator passes these fields in the skill prompt:\n\n| Field | Description |\n|-------|-------------|\n| `docker_cmd` | Complete `docker run` command string, ready to execute |\n| `work_dir` | Local work directory for logs and artifacts |\n| `log_file` | Full path to the log file (output redirected here) |\n| `model_name` | Short model name for reporting |\n| `workflow_type` | `pytest`, `eval`, `custom`, or `benchmark` — for output parsing hints |\n| `success_patterns` | Comma-separated patterns indicating success (e.g., `passed,accuracy:`) |\n| `failure_patterns` | Comma-separated patterns indicating failure (e.g., `FAILED,Error,AssertionError`) |\n\n## Procedure\n\n### Step 0: Resolve Image and Build (when `build_project=true`)\n\nThis executor owns image selection and the build for the local Docker target. Skip the entire step when `build_project=false`.\n\n1. **Detect the target GPU type.** Use `gpu_type` from `job_spec.json` if upstream env-check resolved it; otherwise probe locally with `nvidia-smi --query-gpu=name --format=csv,noheader | head -1`.\n2. **Detect the host CPU arch** with `uname -m` (`x86_64` or `aarch64`).\n3. **Resolve the container image.** Read `<repo_root>/jenkins/current_image_tags.properties` and pick the tag whose CPU-arch flavor matches the host. If the orchestrator already passed a `container_image` field in the job spec, use that and skip the lookup.\n4. **Map GPU → build arch (`-a` flag):** `H100`/`H200` → `90-real`; `B200`/`GB200`/`B300`/`GB300` → `100-real`; `A100` → `80-real`; `L40S` → `89-real`. Default `100-real` when the GPU is unknown.\n5. **Compile.** Invoke the `exec-local-compile` skill with `repo_dir=<repo_root>`, `image=<resolved tag>`, `arch=<arch>`. Wait for completion.\n6. **On failure**, do not launch the workload. Report `BUILD_FAILED` with the last 100 lines of the compile log.\n\n`build_project`, `gpu_type`, `repo_root`, and (optionally) `container_image` come from `job_spec.json`.\n\n### Step 1: Launch\n\nRun the Docker command in background using `run_in_background`:\n\n```bash\n<docker_cmd> 2>&1 | tee <log_file>\n```\n\nReport to the orchestrator: \"Launched locally, log at `<log_file>`\"\n\n### Step 2: Monitor for Hangs\n\nWhile waiting for the background process to complete, actively monitor the log\nfile for hang indicators. Launch a monitoring loop using `run_in_background`:\n\n```bash\nwhile true; do\n  sleep 60\n  if [ -f \"<log_file>\" ] && grep -qi \"hang detected\" \"<log_file>\"; then\n    echo \"HANG_DETECTED: Found 'hang detected' in log file\"\n    docker ps --filter \"ancestor=<image>\" -q | xargs -r docker kill 2>/dev/null\n    exit 1\n  fi\ndone\n```\n\n- If the monitor detects a hang, it kills the Docker container and exits with\n  code 1. The main background process will also terminate.\n- When the main process completes normally (background notification received),\n  kill the monitoring loop (it is no longer needed).\n- If a hang is detected, skip to Step 4 and report `HANG_DETECTED` status\n  instead of proceeding to normal result collection.\n\n### Step 3: Wait for Completion\n\nThe Bash tool's `run_in_background` will notify when the process finishes\n(either normally or because the container was killed by the hang monitor).\n\n### Step 4: Read Results\n\nOn completion:\n\n1. **Read exit code** from the background command result.\n2. **Read the last 100 lines** of `<log_file>` using the Read tool.\n3. **If exit code != 0**, also read the first 50 lines to catch early errors (import failures, setup crashes).\n4. **Search for patterns**:\n   - Grep `<log_file>` for each `success_patterns` entry\n   - Grep `<log_file>` for each `failure_patterns` entry\n\n### Step 5: Report\n\nReturn a structured result:\n\n```\nStatus: PASSED | FAILED | ERROR | HANG_DETECTED\nExit code: <N>\nLog file: <log_file>\nWork directory: <work_dir>\nSummary: <relevant output lines — pytest summary, accuracy score, or last few lines>\nErrors: <if failed, the relevant error output>\n```\n\n### Output Parsing by Workflow Type\n\n- **pytest**: Look for `X passed, Y failed in Zs` summary line\n- **eval**: Look for `accuracy:` or `score:` lines; check for `Expected accuracy >= X, but got Y` assertion\n- **custom**: No specific patterns — report last 10 lines of output\n- **benchmark**: Look for throughput/latency numbers\n\n## Rules\n\n- Never run the Docker command in foreground — always use `run_in_background`\n- Never `cat` the full log file — use Read with offset/limit or tail\n- If the Docker command fails immediately (exit code within seconds), check if the image exists locally\n- Report results even if the log file is empty (container may have failed to start)\n",".claude/skills/exec-remote-slurm/SKILL.md":"---\nname: exec-remote-slurm\ndescription: >-\n  Remote SLURM cluster development via SSH. Use when running jobs,\n  profiling, or developing on a remote SLURM cluster with pyxis/enroot\n  containers. Covers SSH connection management, srun/sbatch/salloc job\n  patterns, tmux-based allocation persistence, file transfer, and safe\n  remote file access. Works with any SLURM cluster accessible via SSH.\ntags: [infrastructure, remote, slurm]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Remote SLURM Development\n\nDevelop and execute on remote SLURM clusters via SSH commands. No filesystem\nsync needed — all remote operations go through SSH.\n\n## Step 1: Resolve Cluster\n\n**Optional dependency — `internal-env-info`.** Before doing\nanything else in Step 1, check whether `skills/internal-env-info/`\nexists in the toolkit. If it does **not**, skip every sub-step that invokes\nthe skill or reads its `references/` files and proceed using only the\ncluster fields the orchestrator wrote into `job_spec.json` (`ssh_host`,\n`slurm_user`, `remote_cwd`, `partition`, `account`, `container_image`,\n`mounts`, `gpus_per_node`, `mfa_style`). If a required field is also\nmissing, stop and ask the user to supply it — do **not** emit an error\nabout the missing skill.\n\nThe orchestrator supplies one of two input shapes; pick the right path:\n\n- **Explicit cluster** — `slurm_cluster` name is provided. Skip\n  auto-selection and go straight to **Load the per-cluster info**.\n- **Auto-select cluster** — only hardware constraints are provided\n  (`device_type`, `total_required_devices`, `required_devices_per_node`).\n  Resolve a cluster from the constraints, then load its per-cluster info.\n\n### 1a. Auto-select (only when no explicit cluster was provided)\n\nInvoke the `internal-env-info` skill in **constraint-based mode**, passing\n`device_type`, `required_devices_per_node`, and (optionally)\n`total_required_devices`. That skill owns the filter rules and either\nreturns a single resolved cluster entry or stops and asks the user. Do not\nimplement the filter here.\n\nUse `node_count` from `job_spec.json` directly as `--nodes` and\n`required_devices_per_node` as `--ntasks-per-node`. **Do not recompute\n`node_count`** — `trtllm-case-executor` is the single source of truth and\nwrites the resolved value into `job_spec.json` for every executor to read.\n\nThe returned entry's `cluster_name`, `mfa_style`, `default_user_root_dir`,\n`default_models_repo`, and `gpus_per_node` feed Step 1b.\n\n### 1b. Load the per-cluster info\n\nConnection fields are read **only** from `job_spec.json` (written by\n`trtllm-case-executor`): `ssh_host`, `slurm_user`, `remote_cwd`, `account`,\n`partition`, `mounts`, `container_image`, `gpus_per_node`, and `packages`.\nThere is no per-cluster connection-config file to parse — if a required\nfield is absent from `job_spec.json`, stop and ask the user to supply it.\nDo not silently default.\n\nInvoke the `internal-env-info` skill in single-cluster mode with\nthe resolved `cluster_name` to retrieve per-cluster info: partition limits,\ntopology-aware scheduling guidance, `default_models_repo`,\n`default_user_root_dir`, `gpus_per_node`, `mfa_style` (used by Step 2 to\nchoose the SSH path — see below), and shared constraint and tuning\npatterns. The runtime `container_image` always comes from `job_spec.json`,\nnever from this skill.\n\n## Step 2: SSH Preflight\n\nThe cluster's `mfa_style` (returned by `internal-env-info` in Step 1b) is\nthe **single source of truth for which SSH path to use**. Switch on the\nvalue — do **not** probe-and-fall-back on the SSH error string.\n\n| `mfa_style` | Path | Why |\n|-------------|------|-----|\n| `false` | Direct (Step 2a) | Cluster accepts key-based / passwordless SSH; no special handling needed. |\n| `true`  | MFA (Step 2b)    | Cluster requires interactive MFA (Duo / device-code / one-time PIN). |\n| `null`  | Unknown (Step 2c)| `mfa_style` not declared in the cluster reference (or `internal-env-info` is not installed). Probe direct first, fall back to MFA. When the reference exists, also ask the user to update it so future runs skip the probe; otherwise just continue silently. |\n\n### 2.0. Build `<ssh_host>` from `slurm_user` + `cluster_name`\n\nBefore any SSH attempt, derive the SSH destination. Priority order:\n\n1. **Caller-supplied `ssh_host`** (passed in through case-executor) → use\n   verbatim, skip everything below. This is the canonical escape hatch.\n2. **`internal-env-info` skill installed** → delegate for the\n   `ssh_host_pattern` key (defined in the skill's `Default Environment Values`\n   section). It returns a single-line snippet that concatenates `slurm_user`\n   and `cluster_name` into `SSH_HOST`. Do **not** re-derive or re-suffix\n   `slurm_user` here — `internal-env-info` already applied whatever its\n   `slurm_user` construction rule requires for that cluster.\n\n   Inputs to substitute:\n   - `slurm_user` — pre-resolved from `internal-env-info` (already in its\n     final, SSH-ready form for that cluster). Read it from\n     `job_spec.json`.\n   - `cluster_name` — from `internal-env-info` Step 1b.\n3. **Skill not installed** → stop and ask the user to supply `ssh_host`\n   explicitly. Do **not** hard-code a login-host pattern in this skill —\n   the NVIDIA-internal convention lives only in `internal-env-info`.\n\nThe resolved value is the canonical `<ssh_host>` placeholder used everywhere\nin the rest of this skill (Steps 2a/2b/2c, Step 3, recipes, file-transfer,\nhang detection).\n\n### 2a. Direct path — `mfa_style: false`\n\nVerify connectivity using the user's existing SSH config:\n\n```bash\nssh -o ConnectTimeout=5 <ssh_host> \"echo ok\"\n```\n\nIf this succeeds, **skip to Step 3**. Do not create a trtllm-managed\nControlMaster — reuse whatever the user already has.\n\nIf it fails, do **not** silently switch to the MFA path. Stop and ask the\nuser to verify their `ssh_host` value or their key-based SSH setup.\nA direct-path cluster failing here is a configuration issue, not an MFA\nprompt to interpret.\n\n### 2b. MFA path — `mfa_style: true`\n\nFirst check whether a ControlMaster is already live (the user may have set\none up earlier in this session):\n\n```bash\nssh -O check <ssh_host> 2>&1 || ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -O check <ssh_host> 2>&1\n```\n\nIf either succeeds, **skip to Step 3** — reuse it.\n\nOtherwise the user must complete an interactive MFA handshake. MFA uses\nbrowser-based device-code verification that needs an interactive SSH\nsession — the skill cannot complete this on the user's behalf because the\nsession must stay alive while the user authenticates in their browser and\npresses ENTER.\n\nTell the user to run the ControlMaster command:\n\n> Run this in your terminal:\n>\n>     ssh -fNM -o ControlMaster=auto -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -o ControlPersist=24h <ssh_host>\n>\n> If the cluster uses a device-code flow you will see a PIN and a URL like:\n>\n>     Authenticate with PIN XXXXXX at <your-identity-provider-device-url>\n>\n> Open that URL, enter the PIN, approve, then press ENTER. The SSH session\n> will background itself as a persistent ControlMaster. Let me know when done.\n\nAfter the user confirms, verify with\n`ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -O check <ssh_host>` and\nproceed to Step 3.\n\n### 2c. Unknown — `mfa_style: null`\n\nThe cluster reference (`references/<cluster>.md`) does not declare an\n`mfa_style` value. Try the direct path (Step 2a) first. If it fails with\n`Permission denied (keyboard-interactive)` — and only that — switch to\nthe MFA path (Step 2b). Any other failure stops here and is reported to\nthe user.\n\nAfter connecting successfully, **ask the user to update the cluster\nreference file** with the discovered value (`mfa_style: true` or\n`false`, in the `## Authentication` table) so that future runs skip\nthis probe and use the deterministic switch in Step 2a / 2b.\n\n## SSH Command Shorthand\n\nEvery SSH and SCP invocation in the rest of this skill uses the placeholders below. Both pin `ControlPath` to the trtllm-managed socket so they reuse the ControlMaster set up in Step 2 — without this, OpenSSH starts a fresh handshake on every call, which on MFA clusters re-prompts for keyboard-interactive auth (and defeats the whole point of multiplexing):\n\n| Placeholder | Expansion |\n|-------------|-----------|\n| `<ssh_cmd>` | `ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p` |\n| `<scp_cmd>` | `scp -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p` |\n\nWhen the user has their own ControlMaster configured in `~/.ssh/config` (no trtllm-managed socket needed), the explicit `-o ControlPath` is harmless: if the path has no live master, OpenSSH falls back to a fresh connection. The connectivity probe in Step 2 Check 2 (`ssh -o ConnectTimeout=5 <ssh_host> \"echo ok\"`) deliberately uses bare `ssh` to test un-multiplexed access; everywhere else use `<ssh_cmd>` / `<scp_cmd>`.\n\n## Step 3: Remote Environment Check\n\nVerify the remote environment is ready:\n\n```bash\n<ssh_cmd> <ssh_host> \"test -d <remote_cwd> && echo ok\"\n<ssh_cmd> <ssh_host> \"which tmux\"\n<ssh_cmd> <ssh_host> \"squeue --version\"\n```\n\nIf `remote_cwd` does not exist, create it or ask the user. If tmux or SLURM\ncommands are not available, stop and inform the user.\n\n## Step 4: Prepare Remote Environment\n\nThis step sets up the user's working directory and clones the TensorRT-LLM\nrepository on the remote cluster. It runs once before any job submission.\n\n> **MANDATORY: two-mechanism sync — both must run.**\n> - **Tracked changes** → temp branch (**4c–4f**): commit on a session-scoped\n>   branch, push to `origin`, fetch + checkout on remote, delete the ref both\n>   sides.\n> - **Untracked files** (`git ls-files --others --exclude-standard`) →\n>   `tar | ssh` pipe (**4e.1**), right after the remote checkout. Never touches\n>   `origin`.\n>\n> **Do not** shortcut with `git fetch origin <repo_branch> && git checkout <repo_branch>` —\n> `repo_branch` may not exist on `origin`, may be stale, or may lack uncommitted\n> edits. Always run 4c–4f **and** 4e.1.\n\n**Inputs received from `trtllm-case-executor`:**\n\n| Field | Description |\n|-------|-------------|\n| `default_user_root_dir` | Per-user root directory on the remote cluster (from `exec-env-check` / `internal-env-info`) |\n| `repo_url` | Git remote URL of the local TensorRT-LLM repo (from `git remote get-url origin`). **Used to clone the bare repository on the remote (Step 4b) and to push the temp branch to (Step 4d).** |\n| `repo_branch` | **Informational only** — the local current branch name (from `git rev-parse --abbrev-ref HEAD`). Used in Step 4c to remember which branch to restore locally after the temp branch is deleted. **Never used as the ref the remote checks out** — the remote always checks out the temp branch from Step 4c. |\n| `device_type` | Required short GPU name (e.g., `B200`, `B300`). Used by Step 1a auto-selection. Ignored when an explicit cluster was provided. |\n| `total_required_devices` | Total GPU processes across all nodes. Used by Step 1a auto-selection only. SLURM job sizing reads `node_count` from `job_spec.json` (precomputed by `trtllm-case-executor`) — do not recompute it here. |\n| `required_devices_per_node` | Minimum GPUs per node. Used by Step 1a auto-selection (cluster must have `gpus_per_node >= required_devices_per_node`) and as `--ntasks-per-node`. |\n\n**Procedure:**\n\n### 4a. Create user root directory\n\n```bash\n<ssh_cmd> <ssh_host> \"mkdir -p <default_user_root_dir>\"\n```\n\nIf `default_user_root_dir` is null or empty, skip this step and use `remote_cwd` directly.\n\n### 4b. Clone or update the repository\n\nDetermine the repo directory name from `repo_url` (strip `.git` suffix if present):\n\n```bash\nREPO_DIR_NAME=$(basename <repo_url> .git)\nREMOTE_REPO_PATH=\"<default_user_root_dir>/${REPO_DIR_NAME}\"\n```\n\nCheck if it already exists:\n\n```bash\n<ssh_cmd> <ssh_host> \"test -d <remote_repo_path>/.git && echo exists\"\n```\n\nIf the repository does **not** exist, clone it:\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <default_user_root_dir> && git clone <repo_url>\"\n```\n\nIf the repository **already exists**, refresh remote refs only — do **not**\ncheckout or reset to any branch here:\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <remote_repo_path> && git fetch origin --prune\"\n```\n\nThe actual ref the remote will check out is the temp branch created in 4c\nand pushed in 4d; the remote checkout happens in 4e. **Do not** add a\n`git checkout <repo_branch>` or `git reset --hard origin/<repo_branch>`\nhere as a shortcut — that path is forbidden (see the MANDATORY callout\nat the top of Step 4).\n\n### 4c. Create a temporary local branch capturing all local changes\n\nGenerate a unique temp branch name tied to this session. The branch name\n**must start with `<user_name>_<model_name>`** so it is easy to identify\nthe owner and the workload it was created for. A timestamp suffix keeps\nit unique across sessions:\n\n```bash\n# Resolve the user name (prefer the slurm_user input over whoami,\n# so the temp branch is attributed to the cluster user rather than the\n# host that happens to dispatch the job).\nUSER_NAME=\"${slurm_user:-$(whoami)}\"\n\n# Resolve the workload name. job_spec.json's `model_name` is authoritative;\n# fall back to the repo basename if it is missing.\nMODEL_NAME=\"${model_name:-$(basename \"<REPO_ROOT>\")}\"\n\n# Sanitize to a git-ref-safe identifier (alphanum + `_`/`-` only).\n# Any other character is replaced with `_`.\nsanitize() { printf '%s' \"$1\" | tr -c 'A-Za-z0-9_-' '_'; }\nUSER_NAME_SAFE=$(sanitize \"$USER_NAME\")\nMODEL_NAME_SAFE=$(sanitize \"$MODEL_NAME\")\n\nTEMP_BRANCH=\"${USER_NAME_SAFE}_${MODEL_NAME_SAFE}_$(date +%Y%m%d%H%M%S)\"\n```\n\nExample resolved value: `<user_name>_dwdp_deepseek_v3_lite_20260430145200`.\n\nStash the current branch name so it can be restored later:\n\n```bash\nORIGINAL_BRANCH=$(git -C \"<REPO_ROOT>\" rev-parse --abbrev-ref HEAD)\n```\n\nCreate the temp branch at HEAD and stage tracked changes only\n(`-u`, not `-A` — untracked files ship in 4e.1):\n\n```bash\ngit -C \"<REPO_ROOT>\" checkout -b \"$TEMP_BRANCH\"\ngit -C \"<REPO_ROOT>\" add -u\ngit -C \"<REPO_ROOT>\" commit --allow-empty -m \"sync: temp snapshot for remote execution\"\n```\n\n`--allow-empty` covers the no-tracked-changes case; the branch still exists\nat HEAD and 4e.1 still runs.\n\n### 4d. Push temp branch to the remote origin\n\nPush the temp branch so the remote cluster can fetch it:\n\n```bash\ngit -C \"<REPO_ROOT>\" push origin \"$TEMP_BRANCH\"\n```\n\n### 4e. Fetch and check out temp branch on the remote cluster\n\nOn the remote: stash prior-session leftovers, fetch, force-checkout the temp\nbranch, and re-materialize `cpp/` from HEAD (stale `cpp/` produces\nmissing-symbol `ImportError`s in Step 5):\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <remote_repo_path> && \\\n    git stash push -u -m 'auto-stash before remote sync' || true && \\\n    git fetch origin --prune && \\\n    git checkout -B $TEMP_BRANCH origin/$TEMP_BRANCH && \\\n    rm -rf cpp/ && \\\n    git checkout HEAD -- cpp/\"\n```\n\n- `stash -u` sweeps both tracked edits and untracked files left from a prior\n  session; `|| true` covers the no-op.\n- `checkout -B` force-resets even if a same-named branch already exists.\n- `rm -rf cpp/ && git checkout HEAD -- cpp/` always runs, even when checkout\n  reports no diffs.\n\n`remote_cwd` is updated to `<remote_repo_path>`. Untracked files for this\nsession arrive next in 4e.1.\n\n### 4e.1. Transfer untracked files via tar pipe over SSH\n\nShip untracked, non-gitignored files on top of the freshly checked-out tree:\n\n```bash\nUNTRACKED_COUNT=$(git -C \"<REPO_ROOT>\" ls-files --others --exclude-standard | wc -l)\nif [ \"$UNTRACKED_COUNT\" -gt 0 ]; then\n    git -C \"<REPO_ROOT>\" ls-files --others --exclude-standard -z | \\\n        tar -C \"<REPO_ROOT>\" --null -T - -cf - | \\\n        <ssh_cmd> <ssh_host> \"tar -C <remote_repo_path> -xf -\"\nfi\n```\n\n`-z` / `--null` keeps filenames with spaces or newlines intact; `tar -x`\nauto-creates parent dirs; the guard skips the round-trip when nothing to\nship. Do **not** substitute `rsync` or `sshfs` (not reliably available on HPC).\n\n### 4f. Delete the temp branch (local and remote)\n\nAfter the remote checkout succeeds, immediately delete the temp branch from\nthe local repo and the remote origin, then restore the original local branch:\n\n```bash\ngit -C \"<REPO_ROOT>\" checkout \"$ORIGINAL_BRANCH\"\ngit -C \"<REPO_ROOT>\" branch -D \"$TEMP_BRANCH\"\ngit -C \"<REPO_ROOT>\" push origin --delete \"$TEMP_BRANCH\"\n```\n\nThe remote cluster keeps the checked-out working tree from the deleted branch\n— only the ref is removed, not the files.\n\n### 4f.1 Create remote work_dir under the cloned repo\n\nThe remote work_dir must mirror the local layout (`<REPO_ROOT>/work_dirs/<name>`)\nso that all per-run artefacts — slurm script, stdout/stderr logs, intermediate\nfiles — live inside the cloned repo and are co-located with the code that\nproduced them.\n\n`remote_work_dir` is precomputed by `trtllm-case-executor` (Step 1) and read\nverbatim from `job_spec.json` — it is `<remote_repo_path>/work_dirs/<basename(local_work_dir)>`.\n**Do not** fall back to `<remote_cwd>/work_dirs/...` or any other root; that\nolder layout is deprecated.\n\n```bash\n<ssh_cmd> <ssh_host> \"mkdir -p <remote_work_dir>\"\n```\n\nAll subsequent steps (file uploads, sbatch submission, log paths) reference\n`<remote_work_dir>` from `job_spec.json`.\n\n### 4g. Copy build scripts to remote node\n\nThe skill ships two build scripts in `scripts/` alongside this SKILL.md.\nCopy them to a dedicated directory on the remote so they are accessible\nfrom inside the SLURM container:\n\n```bash\nREMOTE_SCRIPTS_DIR=\"<default_user_root_dir>/trtllm-build-scripts\"\nSKILL_SCRIPTS_DIR=\"<LOCAL_SKILL_DIR>/scripts\"   # local path to this skill's scripts/\n\n<ssh_cmd> <ssh_host> \"mkdir -p ${REMOTE_SCRIPTS_DIR}/log\"\n\n<ssh_cmd> <ssh_host> \"cat > ${REMOTE_SCRIPTS_DIR}/build.slurm\" < \"${SKILL_SCRIPTS_DIR}/build.slurm\"\n<ssh_cmd> <ssh_host> \"cat > ${REMOTE_SCRIPTS_DIR}/build.sh\"    < \"${SKILL_SCRIPTS_DIR}/build.sh\"\n<ssh_cmd> <ssh_host> \"chmod +x ${REMOTE_SCRIPTS_DIR}/build.slurm ${REMOTE_SCRIPTS_DIR}/build.sh\"\n```\n\n`<LOCAL_SKILL_DIR>` is the directory containing this SKILL.md file\n(`.claude/skills/exec-remote-slurm/`).\n\n---\n\n## Step 5: Build on Remote SLURM Node\n\n> **Mandatory unless `build_project: false`.** Step 4e wipes `cpp/`, so\n> the remote has no compiled C++ extensions until Step 5 succeeds (exit\n> code `0:0`). Submitting any test/bench job before then fails on the\n> first `import tensorrt_llm` with missing-symbol errors.\n>\n> Do **not** replace this with `pip install -e .` inside the user's\n> `.slurm` script — it does not reliably recompile `cpp/`, and when it\n> does it blocks the GPU allocation behind a 30+ minute compile. Step 5\n> runs the build as a 1-node side job and installs into the persistent\n> container volume.\n>\n> If the caller sets `build_project: false`, skip Step 5 (the caller is\n> asserting the install is current). Do not silently downgrade `true` →\n> `false` because Step 5 looks expensive.\n\nSubmit a SLURM batch job that compiles TensorRT-LLM on the remote cluster\ninside the official container. Wait for it to complete before proceeding to\njob submission.\n\n### 5a. Read the container image\n\nThe container image is resolved by `trtllm-case-executor` (Step 3) and\nforwarded in `job_spec.json` as `container_image`. Read it from there —\n**do not grep `current_image_tags.properties` and do not re-derive from\nthe cluster config**. case-executor is the single source of truth for the\nimage URI; this skill only consumes it.\n\nApply the enroot URI rewrite at use-time so the registry path is not\nmistaken for a Docker-style credential lookup. The rewrite recipe lives\nin the `internal-env-info` skill's `Default Environment Values` section\n(`enroot_uri_rewrite` key):\n\n- **Skill installed** → delegate to it, passing the canonical\n  `<container_image>` from `job_spec.json`. The skill returns the bash\n  snippet (an `sed` over the registry-host prefix); run it here to\n  produce `<IMAGE_PATH>`.\n- **Skill not installed** → use `<container_image>` verbatim as `<IMAGE_PATH>`.\n\nUse `<IMAGE_PATH>` as the value passed to `--container-image=` in the\nbuild sbatch below and in every `srun`/`sbatch` recipe later in this\nskill. The canonical (un-rewritten) URI in `job_spec.json` is unchanged;\nthe rewrite is purely a transport adaptation for enroot.\n\n### 5b. Derive build architecture\n\nMap `device_type` (from the env-check output) to the `build_wheel.py` `-a` flag:\n\n| `device_type` | `-a` value |\n|------------|------------|\n| B200, GB200, GB300, B300 | `100-real` |\n| H100, H200 | `90-real` |\n| L40S | `89-real` |\n| A100 | `80-real` |\n\nDefault to `100-real` if `device_type` is unknown.\n\n### 5c. Submit the build job\n\n```bash\n<ssh_cmd> <ssh_host> \"sbatch \\\n    --nodes=1 --ntasks=1 --ntasks-per-node=1 \\\n    --partition=<partition> \\\n    --account=<account> \\\n    --job-name=<account>-trtllm.build \\\n    --time=02:00:00 \\\n    ${REMOTE_SCRIPTS_DIR}/build.slurm \\\n    <IMAGE_PATH> <default_user_root_dir> ${REMOTE_SCRIPTS_DIR} <remote_repo_path> \\\n    -a \\\"<arch>\\\" --nvtx\"\n```\n\nCapture the job ID from `sbatch` stdout (format: `Submitted batch job <id>`).\n\nIf `gpus_per_node` is set in the cluster config, add\n`--gpus-per-node=<gpus_per_node>` to the `sbatch` flags.\n\n### 5d. Monitor the build job\n\nPoll every 30–60 seconds until the job leaves the queue:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <job_id> -o '%.18i %.9P %.30j %.8u %.2t %.10M %R'\"\n```\n\nOnce running, tail the log to track compile progress:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -30 ${REMOTE_SCRIPTS_DIR}/log/build_<job_id>.srun.log\"\n```\n\nProgress indicators in the log:\n- `[XX%] Building CXX object ...` — compilation in progress\n- `Linking CXX ...` — link phase\n- `FAILED:` / `error:` / `fatal error:` — build failure\n- `Successfully built` — success\n\nIf the job disappears from `squeue`, proceed to Step 5e.\n\n### 5e. Verify the build\n\n```bash\n<ssh_cmd> <ssh_host> \"sacct -j <job_id> --format=JobID,State,ExitCode,Elapsed -n\"\n<ssh_cmd> <ssh_host> \"tail -50 ${REMOTE_SCRIPTS_DIR}/log/build_<job_id>.srun.log\"\n```\n\nA successful build shows exit code `0:0` and ends with\n`Successfully built tensorrt_llm` or no error lines.\n\nIf the build fails, report the last 100 lines of the log to the user and\nstop — do not proceed to job submission.\n\n## Build Scripts Reference\n\n| Script | Location | Purpose |\n|--------|----------|---------|\n| `build.slurm` | `scripts/build.slurm` | SLURM batch script — launches container, calls `build.sh` |\n| `build.sh` | `scripts/build.sh` | Runs inside container — calls `build_wheel.py` with defaults |\n\nBoth scripts accept the same positional arguments:\n`<container_image> <user_root_dir> <scripts_dir> <repo_dir> [build_wheel_args...]`\n\n---\n\n## Recipes\n\nPick the recipe that matches your workflow. Each recipe is self-contained —\ncopy it, substitute the `<placeholders>` from config, and run.\n\n**Placeholders** (all from `job_spec.json` unless noted):\n\n| Placeholder | Source |\n|-------------|--------|\n| `<ssh_host>`, `<account>`, `<partition>`, `<mounts>` | `job_spec.json` |\n| `<image>` | `container_image` from `job_spec.json` (resolved by `trtllm-case-executor` Step 3). Apply enroot rewrite (`/` → `#`) at use-time per Step 5a — substitute the rewritten `<IMAGE_PATH>` here. Do not re-resolve from the cluster config or `current_image_tags.properties`. |\n| `<job_name>` | Pre-resolved by `trtllm-case-executor` and forwarded as `job_name` in `job_spec.json`. Substitute directly into `-J <job_name>`; do not reconstruct from `account` / `subproject` / `detail`. |\n| `<nodes>` | `node_count` from `job_spec.json` (computed by `trtllm-case-executor`). Do not recompute. |\n| `<gpus>` | `required_devices_per_node` from `job_spec.json` — probe only if absent (see bottom of this section) |\n| `<remote_repo_path>` | `remote_repo_path` from `job_spec.json` — `<default_user_root_dir>/<repo_dir_name>` precomputed by `trtllm-case-executor` Step 1. The cloned repo and `--container-workdir` for srun. Do not recompute. |\n| `<remote_work_dir>` | `remote_work_dir` from `job_spec.json` — `<remote_repo_path>/work_dirs/<basename(local_work_dir)>` precomputed by `trtllm-case-executor` Step 1. Used for `--output`/`--error` slurm log paths and any per-run artefact uploads. Do not recompute or fall back to `<remote_cwd>/work_dirs/...`. |\n| `<time>` | Walltime in HH:MM:SS |\n\n**Job naming:** Every recipe includes `-J <job_name>` using the\npre-resolved `job_name` value from `job_spec.json` (built by\n`trtllm-case-executor`). This is required — without `-J`, SLURM emits\nwarnings and jobs are hard to identify in `squeue`. Do not reconstruct\nthe name from `account` / `subproject` / `detail` here; that convention\nlives in `trtllm-case-executor`.\n\n**GPU allocation:** Cluster behavior varies:\n\n- **Exclusive-access clusters** (most NVIDIA internal): Nodes are allocated\n  exclusively with all GPUs available. No `--gpus-per-node` needed.\n- **Shared/QOS clusters**: Require `--gpus-per-node=<gpus_per_node>`. Without\n  it, you may get a CPU-only node (no GPUs at all), a \"Cannot find GPU\n  specification\" error, or a \"QOSMinGRES\" error (below minimum).\n\nIf `gpus_per_node` is set in the cluster config, add `--gpus-per-node=<gpus_per_node>`\nto every `srun`, `salloc`, and `sbatch` command. If not set, omit it.\n\nIf a job fails with GPU-related errors and `gpus_per_node` is not configured,\nprobe the node's GPU count, set `gpus_per_node` in the cluster config, and retry.\n\nDo not set `CUDA_VISIBLE_DEVICES` for MPI workloads — the MPI runtime handles\nGPU assignment.\n\n---\n\n### Recipe 1: Single-process command (srun)\n\nOne-off commands on a single node: profiling, inspecting, lightweight scripts.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N 1 --ntasks-per-node=1 \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <command>\"\n```\n\n---\n\n### Recipe 2: MPI workload (srun)\n\nMPI applications across nodes: NCCL tests, HPCX benchmarks. The binary must\nbe MPI-aware (e.g., `all_reduce_perf_mpi`, not `all_reduce_perf`).\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=<gpus> --mpi=pmix \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <command>\"\n```\n\n---\n\n### Recipe 3: Distributed training (srun)\n\nFramework-launched distributed training (PyTorch DDP, DeepSpeed). SLURM starts\n`<nodes> × <gpus>` processes. Pyxis/enroot sets `RANK`, `LOCAL_RANK`, and\n`WORLD_SIZE`. Each process runs the training script directly.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=<gpus> \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  python train.py\"\n```\n\n---\n\n### Recipe 4: Launcher-based training (srun)\n\nOne process per node, an in-container launcher forks locally (torchrun,\ndeepspeed launcher).\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=1 \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  torchrun --nproc-per-node=<gpus> train.py\"\n```\n\n---\n\n### Recipe 5: Batch job (sbatch)\n\nFire-and-forget background job:\n\n```bash\n<ssh_cmd> <ssh_host> \"sbatch -A <account> -p <partition> -N <nodes> \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <script-path>\"\n```\n\nCheck status:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <jobid>\"\n<ssh_cmd> <ssh_host> \"sacct -j <jobid> --format=JobID,State,ExitCode,Elapsed -n\"\n```\n\n---\n\n### Recipe 6: Multi-command session (salloc + srun)\n\nInteractive development with persistent container state. Use this when you need\nmultiple commands on the same allocation — installing packages, compiling,\nrunning experiments iteratively.\n\n> **Relationship to `persistent_mode`:** The `persistent_mode` field in\n> `job_spec.json` is consumed only by `exec-local-slurm`. This\n> skill does **not** read `persistent_mode`, `container_name`, or\n> `container_mounts` from `job_spec.json` for allocation reuse — remote\n> persistence is driven by this recipe (tmux + named `salloc` + named\n> container) and controlled by inputs to this skill directly (e.g.,\n> `release_allocation`, `alloc_time_limit`). If a caller forwards a\n> `persistent_mode=true` job spec into a `remote_slurm` run, treat the field\n> as informational only and use this recipe's flow.\n\n**Step 1: Allocate via tmux**\n\ntmux keeps the allocation alive on the frontend — invisible plumbing.\n\n```bash\n<ssh_cmd> <ssh_host> \"tmux new-session -d -s trtllm-<id> \\\n  'salloc -A <account> -p <partition> -N <nodes> \\\n  -J <job_name> -t <time>'\"\n```\n\n**Step 2: Check allocation status**\n\nWait ~30 seconds for the scheduler, then get job ID and status:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -u <user> -h -o '%i %T %S' --state=RUNNING,PENDING\"\n```\n\nIf RUNNING → proceed to Step 3.\n\nIf PENDING → follow the **walltime negotiation flow** (see Walltime strategy\nsection). Only proceed to Step 3 once the job is RUNNING.\n\n**Step 3: Create a persistent container**\n\nImport the image once. This creates a named container on each node that\npersists across `srun` calls (container persistence):\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-image=<image> --container-name=trtllm-<jobid> true\"\n```\n\n**Container packages:** If `packages` is defined in config, install them\nusing the exact names listed — do not guess or substitute package names:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> --container-writable \\\n  pip install <packages>\"\n```\n\nOnly install when the workload actually needs them — not for simple\ncommands like `nvidia-smi` or `hostname`.\n\n**Step 4: Run commands**\n\nUse `--container-name` (without `--container-image`) to reuse the persistent\ncontainer. All in-container state is preserved between commands.\n\nSingle-process command:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nPer-node operation (mkdir, pip install) — use `--ntasks-per-node=1`:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nMPI workload within the allocation:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --ntasks-per-node=<gpus> --mpi=pmix \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nDistributed training within the allocation:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --ntasks-per-node=<gpus> \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  python train.py\"\n```\n\nAdd `--container-writable` if the workload writes to container paths outside\nof mounted directories (e.g., `pip install` into site-packages).\n\n**Step 5: Release the allocation**\n\nClean up the named container, then kill tmux. If the allocation has already\nexpired (walltime exceeded), the `srun` will fail — that is fine, the container\nis already gone. Just kill the tmux session.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  bash -c 'enroot remove -f trtllm-<jobid> 2>/dev/null; true'\" 2>/dev/null\n<ssh_cmd> <ssh_host> \"tmux kill-session -t trtllm-<id>\"\n```\n\n---\n\n### Probing GPUs per node\n\nIf you do not know `<gpus>`, probe once and record in cluster notes:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> bash -c 'nvidia-smi -L | wc -l'\"\n```\n\n## Cluster-Specific Flags\n\nAfter substituting recipe placeholders, apply the cluster-specific flags\nreturned by the `internal-env-info` skill in Step 1:\n\n- **SHARP**: Add `--network=sharp` if the cluster supports it\n- **Topology constraints**: Add `--constraint=<feature>` for topology-aware\n  placement (e.g., `--constraint=nvlblk01` for a single rack)\n- **GPU clocks**: Run `sudo nvidia-smi -ac <mem>,<graphics>` as a separate\n  `srun --ntasks-per-node=1` pre-step in sbatch scripts\n- **Comments**: Add `--comment=metrics` for Prometheus/Grafana collection\n\n## File Transfer\n\nUse SSH pipes for file transfer — no extra tools needed, works on every cluster.\n\n### Local → Remote (writing files)\n\nPipe file content directly via SSH heredoc. Never write locally then scp:\n\n```bash\n<ssh_cmd> <ssh_host> \"cat > <remote_cwd>/script.py\" << 'PYEOF'\nimport torch\n# ... file content ...\nPYEOF\n```\n\nFor multiple files, repeat the pattern — each file is one SSH command.\n\n### Remote → Local (reading results)\n\nRead remote files via SSH stdout:\n\n```bash\n<ssh_cmd> <ssh_host> \"cat <remote_cwd>/output.json\"\n```\n\nFor large files, read only what you need:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/slurm-<jobid>.out\"\n<ssh_cmd> <ssh_host> \"head -n 50 <remote_cwd>/results.csv\"\n```\n\n### When to use scp\n\nFall back to `scp` only for binary files or large directories that cannot be\npiped through SSH:\n\n```bash\n<scp_cmd> <ssh_host>:<remote_cwd>/profile.nsys-rep .    # binary profile\n<scp_cmd> -r local_dir/ <ssh_host>:<remote_cwd>/        # whole directory\n```\n\n**Decision rule:** text files → SSH pipe. Binary files or directories → scp.\nNever use rsync or sshfs — not reliably available on HPC clusters.\n\n## Rules\n\n### Never read full remote files\n\nRemote output can be gigabytes. Always use `tail`, `head`, or `grep`:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/output.log\"\n<ssh_cmd> <ssh_host> \"head -n 50 <remote_cwd>/script.py\"\n<ssh_cmd> <ssh_host> \"grep -n 'error\\|Error\\|ERROR' <remote_cwd>/output.log\"\n```\n\nNever use `cat` on remote files unless you are certain the file is small.\n\n### Never run heavy computation on the frontend\n\nThe frontend node is shared. Only use it for SLURM commands, lightweight file\noperations, and tmux management.\n\n### Check before creating\n\nBefore creating resources, check if they already exist:\n- SSH ControlMaster: `ssh -O check <ssh_host> 2>&1 || <ssh_cmd> -O check <ssh_host>` (probe both the user's default socket and the trtllm-managed socket from Step 2)\n- tmux sessions: `<ssh_cmd> <ssh_host> \"tmux has-session -t trtllm-<id> 2>/dev/null && echo exists\"`\n- Existing allocations: `<ssh_cmd> <ssh_host> \"squeue -u <user> -h -o '%i %j' --state=RUNNING\"`\n\n### Clean up when done\n\n- Remove named containers: `enroot remove -f` (see Recipe 6, Step 5)\n- Kill tmux sessions: `<ssh_cmd> <ssh_host> \"tmux kill-session -t trtllm-<id>\"`\n- Cancel unneeded jobs: `<ssh_cmd> <ssh_host> \"scancel <jobid>\"`\n\n### Walltime strategy\n\n- **Production (sbatch, srun):** Tight walltime with ~20% margin — gets\n  scheduled sooner.\n- **Dev/debug (salloc):** Use **walltime negotiation** — submit with max\n  walltime, then reduce if the queue wait is too long:\n\n**Walltime negotiation flow:**\n\n1. Submit `salloc` with **max walltime** for the partition.\n2. Wait ~30 seconds for the scheduler to process, then check:\n   ```bash\n   <ssh_cmd> <ssh_host> \"squeue -j <jobid> -h -o '%i %T %S'\"\n   ```\n   - `%S` = estimated start time. Parse it to calculate wait.\n3. **If RUNNING** → proceed to next step.\n4. **If PENDING and wait ≤ 10 minutes** → wait for it, proceed.\n5. **If PENDING and wait > 10 minutes** (or no start time shown):\n   a. Infer how long you actually need from context (profiling → ~30 min,\n      iterative dev → ~2h, training run → match expected duration).\n      Round up to the next hour boundary.\n   b. Reduce walltime in-place:\n      ```bash\n      <ssh_cmd> <ssh_host> \"scontrol update JobId=<jobid> TimeLimit=<new_time>\"\n      ```\n   c. Tell the user: \"Reduced walltime from Xh to Yh for faster scheduling —\n      you'll have less time on the allocation.\"\n   d. Wait ~15 seconds, re-check estimated start time.\n   e. **If still > 10 minutes:** inform user with the ETA and suggest\n      alternatives — different partition (e.g., backfill), fewer nodes.\n\n## Hang Detection Monitoring\n\nWhen monitoring a running job (polling via `squeue`), also check the remote log\nfile every 60 seconds for the pattern `hang detected` (case-insensitive). This\ncatches GPU hangs, NCCL deadlocks, and other stalls that keep the job running\nindefinitely without making progress.\n\n**Check for hang (run alongside each `squeue` poll once job is RUNNING):**\n\n```bash\n<ssh_cmd> <ssh_host> \"grep -qi 'hang detected' <remote_work_dir>/slurm-<jobid>.out 2>/dev/null && echo HANG_DETECTED\"\n```\n\nAlso check the `.err` file if the output file has no match:\n\n```bash\n<ssh_cmd> <ssh_host> \"grep -qi 'hang detected' <remote_work_dir>/slurm-<jobid>.err 2>/dev/null && echo HANG_DETECTED\"\n```\n\n**If `HANG_DETECTED` is printed:**\n\n1. Immediately cancel the job:\n   ```bash\n   <ssh_cmd> <ssh_host> \"scancel <jobid>\"\n   ```\n2. Read the last 200 lines of the log file for context:\n   ```bash\n   <ssh_cmd> <ssh_host> \"tail -n 200 <remote_work_dir>/slurm-<jobid>.out\"\n   ```\n3. Report status as `HANG_DETECTED` with the relevant log lines. Do not\n   continue polling — the job is terminated.\n\n## Recovering from Failures\n\n**SSH connection drops:** Re-run SSH preflight (Step 2). tmux-held allocations\nsurvive disconnects.\n\n**Job stuck in PENDING:**\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <jobid> -o '%i %T %r %S'\"\n```\n\nCheck the reason (`%r`) and estimated start (`%S`). If wait > 10 minutes,\nfollow the **walltime negotiation flow** (see Walltime strategy). If already\nnegotiated and still stuck, suggest alternatives: backfill partition, fewer\nnodes, or a different cluster.\n\n**Job fails:** Read last lines, never dump full output:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/slurm-<jobid>.out\"\n<ssh_cmd> <ssh_host> \"tail -n 50 <remote_cwd>/slurm-<jobid>.err\"\n```\n\n**Hang detected:** If the log file contains `hang detected`, the job is\nauto-cancelled (see Hang Detection Monitoring section above). Report\n`HANG_DETECTED` with the last 200 lines of log context.\n",".claude/skills/kernel-cute-writing/SKILL.md":"---\nname: kernel-cute-writing\ndescription: >\n  Write and implement GPU kernels using NVIDIA CuTe DSL (CUTLASS 4.x Python\n  API) — NOT for Triton, CUDA C++, or conceptual explanations.\n  Trigger only when the user wants to write or implement a kernel, not when\n  asking questions about CuTe DSL concepts or layouts.\n  CuTe DSL uses cute.jit/cute.kernel decorators and cutlass.cute imports.\n  Covers element-wise kernels, GEMM patterns, reductions, memory hierarchy\n  (global/shared/register/TMA), MMA tensor core operations, software\n  pipelining, and framework integration.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# CuTe DSL\n\nCuTe DSL is a Python-based domain-specific language for GPU kernel development,\npart of CUTLASS 4.x. It provides Python abstractions over CUTLASS C++ templates\nwith JIT compilation to optimized CUDA kernels via MLIR and ptxas.\n\n## When to Use\n\n**Triggers:**\n- Writing CUDA kernels in Python (element-wise, GEMM, custom ops)\n- Optimizing GPU memory access patterns (vectorized loads, TMA, shared memory)\n- Building tensor core (MMA) kernels for Ampere/Hopper/Blackwell\n- Integrating custom GPU kernels with PyTorch or JAX\n- Prototyping high-performance kernels without C++ metaprogramming\n\n**Symptoms (wrong tool otherwise):**\n- Need shared memory coordination or tensor core MMA → use CuTe DSL (not Triton for complex patterns)\n- Need simple element-wise ops with no shared memory → CuTe DSL or Triton both work\n- Need to call existing CUTLASS C++ kernels → use CUTLASS C++ APIs instead\n- Need reductions, scans, or non-GEMM collective ops → consider CUB/Thrust\n\n**Keywords:** cute, cutlass, cute.jit, cute.kernel, from_dlpack, zipped_divide,\nTiledMMA, TiledCopy, TMA, WGMMA, tcgen05, pipeline, mbarrier\n\n## Requirements\n\n| Requirement | Detail |\n|-------------|--------|\n| Platform | Linux x86_64 only |\n| Python | 3.10–3.13 |\n| GPU | NVIDIA Ampere+ (SM80, SM90, SM100) |\n| CUDA Driver | ≥ 575.51.03 (Toolkit 12.9 compat) |\n| Install | `pip install nvidia-cutlass-dsl` |\n| Optional | `apache-tvm-ffi`, `torch-c-dlpack-ext` |\n\n## Workflows\n\n### Workflow 0: Starting from Examples (Recommended)\n\nFor any non-trivial kernel (GEMM, attention, pipelined, fused ops), start by\nfinding the most similar existing example to use as a **starting point** — study\nits structure, then rework it for your use case. Do not copy examples verbatim;\nthey target specific dtypes, architectures, and problem shapes that likely differ.\n\n1. **Pick the closest example** from the index below.\n   **Prefer examples matching the target GPU architecture** (check with\n   `torch.cuda.get_device_capability()`) when the operation is similar.\n\n   Fetch via `web_fetch` with base URL\n   `https://raw.githubusercontent.com/NVIDIA/cutlass/main/examples/python/CuTeDSL`\n\n   | Operation | Arch | Example path (append to base URL) |\n   |-----------|------|-----------------------------------|\n   | Element-wise add | SM80 | `ampere/elementwise_add.py` |\n   | Element-wise + autotune | SM80 | `ampere/elementwise_add_autotune.py` |\n   | Element-wise apply | SM80 | `ampere/elementwise_apply.py` |\n   | SGEMM (scalar) | SM80 | `ampere/sgemm.py` |\n   | Tensor-core GEMM | SM80 | `ampere/tensorop_gemm.py` |\n   | Flash Attention v2 | SM80 | `ampere/flash_attention_v2.py` |\n   | HSTU Attention | SM80 | `ampere/hstu_attention.py` |\n   | Shared memory allocator | SM80 | `ampere/smem_allocator.py` |\n   | CTA norm (LayerNorm) | SM90 | `hopper/cta_norm.py` |\n   | Dense GEMM | SM90 | `hopper/dense_gemm.py` |\n   | Dense GEMM persistent | SM90 | `hopper/dense_gemm_persistent.py` |\n   | Flash MHA | SM90 | `hopper/fmha.py` |\n   | Dense GEMM | SM100 | `blackwell/dense_gemm.py` |\n   | Dense GEMM persistent | SM100 | `blackwell/dense_gemm_persistent.py` |\n   | Dense GEMM + alpha/beta | SM100 | `blackwell/dense_gemm_alpha_beta_persistent.py` |\n   | RMSNorm | SM100 | `blackwell/rmsnorm.py` |\n   | Reduce | SM100 | `blackwell/reduce.py` |\n   | Flash MHA | SM100 | `blackwell/fmha.py` |\n   | Grouped GEMM | SM100 | `blackwell/grouped_gemm.py` |\n   | Mamba2 SSD | SM100 | `blackwell/mamba2_ssd/` |\n   | GEMM tutorial (notebook) | SM100 | `notebooks/tour_to_sol_gemm.ipynb` |\n\n   **Example:** To fetch the Hopper dense GEMM:\n   ```bash\n   web_fetch https://raw.githubusercontent.com/NVIDIA/cutlass/main/examples/python/CuTeDSL/hopper/dense_gemm.py\n   ```\n\n2. **Read reference materials first** — before diving into example code, read\n   the relevant `references/` docs to understand the patterns and APIs:\n   - For GEMM: `references/patterns-gemm.md` (3-level tiling, epilogue fusion,\n     `cute.compile` with `mark_layout_dynamic`, shared memory layouts)\n   - For reductions: `references/patterns-reduction.md` (warp reductions,\n     `cute.compile` cache pattern)\n   - For element-wise: `references/patterns-elementwise.md` (variations A–E)\n   - Always: `references/api-arch.md` (available APIs, arch-specific caveats)\n\n   This gives you the conceptual foundation so you can rework the example\n   intelligently rather than trying to copy-paste complex pipelines.\n\n3. **Fetch and study the example source** — read for structure, not to copy:\n   - Identify: decorators, tiling strategy, shared memory usage, mainloop flow\n   - Note which dtype/arch it targets (many examples are fp16/bf16-specific)\n   - Check if it uses APIs tied to a specific arch (TMA → SM90+, tcgen05 → SM100)\n\n4. **Rework for the user's workload** (do not copy-paste):\n   - Change shapes, data types, tile sizes to match requirements\n   - Replace compute logic (epilogue, activation fusion) as needed\n   - If dtype differs (e.g., example is fp16, need fp32), expect vectorization\n     and layout changes — the scalar-loop patterns in `references/` may be a\n     better starting point than adapting a vectorized example\n   - **Runtime wrapper must be lightweight**: `kernel_fn()` should only call\n     `from_dlpack()` + the compiled kernel. Never allocate intermediate tensors,\n     copy data, or re-compile per call — these belong in one-time setup\n   - Apply optimizations from this skill's reference docs\n\n   **⛔ Blackwell/Hopper GEMM + extra tensors — STOP:**\n   If the target GPU is SM90+ (Hopper/Blackwell) **and** the GEMM requires\n   extra tensors beyond A, B, C in the epilogue (e.g., bias vector, activation\n   inputs), **do not attempt it**. These examples use TMA descriptors for all\n   data movement — adding tensors requires modifying TMA descriptor setup,\n   which is prohibitively complex. Instead, tell the user this limitation and\n   suggest a **two-kernel approach**: run the GEMM kernel as-is, then apply\n   bias + activation in a separate element-wise kernel (Workflow 1).\n   Plain GEMM (just A×B→C with scalar alpha/beta) on Hopper/Blackwell is fine.\n\n5. **Validate and benchmark** using companion scripts:\n   ```bash\n   python scripts/verify_kernel.py kernel.py --rtol 1e-3 --atol 1e-3\n   python scripts/benchmark_kernel.py kernel.py\n   ```\n   The kernel file must export `kernel_fn`, `reference_fn`, and `get_inputs()`.\n\n**When to skip examples:** Pure element-wise operations (Workflow 1) have\ncomplete patterns in `references/patterns-elementwise.md` — no need to fetch\nexternal examples.\n\n**Reduction kernels** (softmax, layernorm, RMSNorm): Use\n`references/patterns-reduction.md` which provides complete, proven patterns\nfor float32 reductions using scalar loops + butterfly shuffle + shared memory.\n\n### Workflow 1: Element-wise Kernel\n\nFor unary/binary/in-place operations that map inputs to outputs 1:1.\n\n1. **Determine kernel structure**: inputs/outputs count, tensor rank, target arch\n2. **Select pattern** from `references/patterns-elementwise.md` (Variations A–E)\n3. **Write kernel** applying all four invariant principles:\n   - P1: `from_dlpack(tensor, assumed_align=16)` for vector loads\n   - P2: Derive `vec_size` from `element_type.width`\n   - P3: `cute.zipped_divide(mA, tiler)` for coalesced access\n   - P4: `cutlass.dynamic_expr(thread_idx < total)` for bounds\n4. **Critical rules**: No early return, no `a * 2` (use `a + a`), no `cute.math.sigmoid`\n5. **Pre-compile with `cute.compile()`**: Always pre-compile the kernel once\n   using `cute.compile()` so that `kernel_fn` calls the compiled object, not\n   `@cute.jit` directly.  Without pre-compilation, every call recompiles\n   (~20-50ms overhead).  Use `.mark_layout_dynamic()` so a single compiled\n   kernel handles arbitrary input shapes without recompilation:\n   ```python\n   # Compile once with dynamic layouts — works for any shape\n   fake_x = from_dlpack(torch.empty(1, 1, dtype=torch.float16, device=\"cuda\"),\n                         assumed_align=16).mark_layout_dynamic()\n   fake_out = from_dlpack(torch.empty(1, 1, dtype=torch.float16, device=\"cuda\"),\n                           assumed_align=16).mark_layout_dynamic()\n   compiled_kernel = cute.compile(host_fn, fake_x, fake_out)\n\n   def kernel_fn(x):\n       out = torch.empty_like(x)\n       compiled_kernel(from_dlpack(x, assumed_align=16).mark_layout_dynamic(),\n                       from_dlpack(out, assumed_align=16).mark_layout_dynamic())\n       return out\n   ```\n6. **Verify correctness** using companion script:\n   ```bash\n   python scripts/verify_kernel.py kernel.py --rtol 1e-3 --atol 1e-3\n   ```\n   The kernel file must export `kernel_fn`, `reference_fn`, and `get_inputs()`.\n7. **Benchmark** using companion script:\n   ```bash\n   python scripts/benchmark_kernel.py kernel.py\n   ```\n\n### Workflow 2: GEMM Kernel\n\nFor matrix multiplication with tiling, shared memory, and tensor cores.\n\n1. **Define problem**: shapes (M, N, K), data types, target architecture\n2. **Choose tiling**: CTA tile (bM, bN, bK), pipeline stages, cluster shape\n3. **Three-level partitioning** (see `references/patterns-gemm.md`):\n   - Level 1: CTA tiling with `local_tile()`\n   - Level 2: Copy partitioning (global → shared) with `TiledCopy`\n   - Level 3: Compute partitioning (shared → register) with `TiledMMA`\n4. **Shared memory**: Use swizzled layouts (`make_smem_layout_atom`) to avoid bank conflicts\n5. **Mainloop**: K-tile loop with copy → sync → MMA → sync\n6. **Pipeline**: Use `PipelineTmaAsync` (Hopper) or `PipelineTmaUmma` (Blackwell).\n   ⚠️ TMA-based pipelines manage data movement via TMA descriptors — adding\n   extra tensors (bias, activation inputs) to the epilogue requires modifying\n   descriptor setup, which is prohibitively complex. See the stop condition in\n   Workflow 0 step 4.\n7. **Epilogue**: Predicated store with alpha/beta scaling\n8. **Pre-compile with `cute.compile()`**: Always pre-compile the GEMM kernel\n   so `kernel_fn` calls the compiled object, not `@cute.jit` directly.\n   Without pre-compilation, every call recompiles (~20-50ms overhead).\n9. **Autotune**: Search over tile sizes, cluster shapes, pipeline depths\n\n### Workflow 3: Framework Integration\n\nFor wrapping CuTe DSL kernels as PyTorch/JAX custom operators.\n\n1. **Write kernel** using Workflow 1 or 2\n2. **Create wrapper**: Accept `torch.Tensor`, convert via `from_dlpack`, call host fn\n3. **For production**: Compile with TVM FFI for zero-overhead tensor passing:\n   ```python\n   compiled = cute.compile(host_fn, *fake_tensors, options=\"--enable-tvm-ffi\")\n   compiled(torch_a, torch_b)  # Direct torch.Tensor, no from_dlpack\n   ```\n4. **For deployment**: Use AOT compilation → export to `.o` → load at runtime\n\n### Workflow 4: Debugging & Profiling\n\n1. **Set environment**: `CUTE_DSL_PRINT_IR=1`, `CUTE_DSL_KEEP_PTX=1`\n2. **Use `cute.printf()`** for runtime values (not Python `print`)\n3. **Inspect generated code**: `compiled.__ptx__`, `compiled.__mlir__`\n4. **Profile**: Enable `CUTE_DSL_LINEINFO=1`, use Nsight Compute/Systems\n5. **Debug memory**: Run with `compute-sanitizer python script.py`\n\n## Output Formats\n\nA typical CuTe DSL kernel project:\n\n```\nkernel_dir/\n  kernel.py          # @cute.kernel + @cute.jit functions\n  test_kernel.py     # Correctness test vs PyTorch reference\n  bench_kernel.py    # Benchmark with cute.compile() setup\n```\n\n**Success indicators:**\n- Correctness test passes (`torch.testing.assert_close`)\n- Nsight shows vector loads (LDG.128/LDG.256), not scalar loads\n- For GEMM: tensor core utilization > 80% in Nsight Compute\n\n## Companion Script Contract\n\nKernel files used with `scripts/verify_kernel.py` and `scripts/benchmark_kernel.py`\nmust export three names:\n\n- `kernel_fn(*inputs)` — the CuTe DSL kernel wrapper (calls `cute.compile` + runs kernel)\n- `reference_fn(*inputs)` — PyTorch reference implementation (same signature)\n- `get_inputs()` — returns a list of CUDA tensors for testing\n\n```python\n# Example kernel.py contract\nimport torch\nimport cutlass.cute as cute\nfrom cutlass.cute.runtime import from_dlpack\n\ndef kernel_fn(x):\n    out = torch.empty_like(x)\n    # ... call compiled cute kernel ...\n    return out\n\ndef reference_fn(x):\n    return torch.nn.functional.gelu(x)\n\ndef get_inputs():\n    return [torch.randn(1024, 512, dtype=torch.float16, device=\"cuda\")]\n```\n\n## Examples\n\n### Example: 2D Unary Element-wise (ReLU)\n\n```python\nimport torch, cutlass, cutlass.cute as cute\nfrom cutlass.cute.runtime import from_dlpack\n\n@cute.kernel\ndef relu_kernel(gA: cute.Tensor, gC: cute.Tensor):\n    tidx, _, _ = cute.arch.thread_idx()\n    bidx, _, _ = cute.arch.block_idx()\n    bdim, _, _ = cute.arch.block_dim()\n    idx = bidx * bdim + tidx\n    m, n = gA.shape[1]\n    total = m * n\n    if cutlass.dynamic_expr(idx < total):\n        a = gA[(None, (idx // n, idx % n))].load()\n        gC[(None, (idx // n, idx % n))] = cute.where(a > 0, a, 0)\n\n@cute.jit\ndef relu_host(mA: cute.Tensor, mC: cute.Tensor):\n    vec = 16 // (mA.element_type.width // 8)\n    gA = cute.zipped_divide(mA, (1, vec))\n    gC = cute.zipped_divide(mC, (1, vec))\n    T = 256\n    N = cute.size(gA.shape[1])\n    relu_kernel(gA, gC).launch(grid=((N+T-1)//T,1,1), block=(T,1,1))\n\nx = torch.randn(1024, 512, dtype=torch.float16, device=\"cuda\")\nout = torch.empty_like(x)\nrelu_host(from_dlpack(x, assumed_align=16), from_dlpack(out, assumed_align=16))\n```\n\n## Error Handling\n\n| Error | Cause | Fix |\n|-------|-------|-----|\n| `MLIR function requires a Context` | Called @kernel from Python | Launch via @cute.jit host function |\n| `DSLAstPreprocessorError` on return | Early return in @kernel | Use `if cutlass.dynamic_expr(cond):` |\n| Type mismatch on store | `a * 2` promotes FP16→FP32 | Use `a + a` or `.to(cutlass.Float16)` |\n| `could not get source code` | Kernel in `exec()` context | Write to file and import |\n| Scalar loads in Nsight | Missing alignment hint | Add `assumed_align=16` to `from_dlpack` |\n| `Missing required argument` | Not all @jit params passed | Pass ALL declared parameters |\n| `AttributeError: sigmoid` | No `cute.math.sigmoid` | Use `1.0/(1.0+cute.math.exp(-x))` |\n\nSee `references/troubleshooting.md` for the full error table and limitations.\n\n**Debugging rule:** Never delete kernel.py during debugging. Use `backup_file`\nto save a checkpoint, then `edit_file` to iterate. If stuck, `revert_file` to\nrestore the backup. A partially-working kernel is always better than no kernel.\n\n## Finding More Information\n\n### Tier 1: This File (SKILL.md)\n\nWorkflows above cover element-wise kernels, GEMM, framework integration, and\ndebugging. Search this file first for procedural questions.\n\n### Tier 2: references/ Directory\n\nGrep for keywords across `references/`. Headers are grep-friendly.\n\n| File | Content |\n|------|---------|\n| `concepts-architecture.md` | Core abstractions, terminology, compilation pipeline |\n| `concepts-layouts.md` | Layout algebra: composition, complement, divide, swizzle |\n| `concepts-tensors.md` | Tensor types, partitioning, tiling, predication |\n| `concepts-mma.md` | MMA atoms, TiledMMA, per-architecture tensor core ops |\n| `patterns-getting-started.md` | Installation, decorators, first kernel walkthrough |\n| `patterns-elementwise.md` | Invariant principles, pattern variations, reference impl |\n| `patterns-gemm.md` | 3-level tiling, shared memory, pipelining, autotuning |\n| `patterns-memory.md` | from_dlpack, TMA, cp.async, TMEM, copy atoms |\n| `patterns-compilation.md` | Control flow, JIT caching, TVM FFI, AOT compilation |\n| `patterns-pipeline.md` | Producer-consumer, pipeline classes, barriers, warp specialization |\n| `api-core.md` | cute module: layouts, tensors, math, copy, gemm, printing |\n| `api-arch.md` | cute.arch: thread indexing, sync, atomics, memory ops |\n| `api-nvgpu.md` | cute.nvgpu: warp/warpgroup/cpasync/tcgen05 MMA and copy |\n| `api-runtime-utils.md` | Runtime: from_dlpack, fake tensors, utils, schedulers |\n| `troubleshooting.md` | Debugging, env vars, common errors, limitations, FAQ |\n\n**How to search:** Grep for your keyword across `references/`. Read only the\nfile and section that Grep points to.\n\n### Tier 3: Original Documentation\n\nIf Tiers 1–2 don't answer, consult the source:\n- **Web**: https://docs.nvidia.com/cutlass/latest/\n- **GitHub**: https://github.com/NVIDIA/cutlass\n- Fetch specific doc pages or search for \"CUTLASS CuTe DSL <topic>\"\n- Consider distilling the answer back into references/\n",".claude/skills/perf-host-analysis/SKILL.md":"---\nname: perf-host-analysis\ndescription: >\n  Analyze host/CPU overhead in TensorRT-LLM inference from nsys traces.\n  Detect whether host overhead is the bottleneck using GPU idle ratio, host prep\n  exposed ratio, and per-phase evidence. For regressions, isolate forward steps\n  via allreduce/NVTX patterns, compare host operation breakdowns across versions,\n  and identify scheduling or request-management overhead. Supports optional\n  inter-kernel gap, eager-vs-graph, pattern mapping, and multi-rank straggler\n  drill-down. Use standalone or within perf-analysis. Triggers: host overhead,\n  inter-step gap, scheduling overhead, forward step isolation, nsys iteration\n  analysis, NVTX breakdown, request management overhead, GPU idle, host bottleneck,\n  host prep exposed, inter-kernel gap, bubble analysis, graph coverage, eager kernel,\n  rank imbalance, straggler detection.\nlicense: Apache-2.0\ntags:\n  - analysis\n  - profiling\n  - host-overhead\n  - nsys\n  - inference\ndependencies:\n  - workload-instrumentation\n  - trace-interpretation\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Host Performance Analysis\n\nAnalyze host/CPU overhead in TensorRT-LLM inference workloads from nsys traces. This skill operates in two phases:\n\n| Phase | Question | Input | Output |\n|-------|----------|-------|--------|\n| **Detection** | Is host overhead the bottleneck? | Single nsys trace | YES/NO verdict with metric evidence |\n| **Root Cause** | What specifically regressed? | One or two nsys traces | NVTX per-step breakdown, regression sources, optional kernel-level drill-down |\n\n## When to Use\n\n- Before starting host optimization work -- confirms the bottleneck is real (Detection)\n- As a sub-step of `perf-analysis` for bottleneck classification (Detection)\n- When GPU utilization is suspiciously low and you need to know why (Detection)\n- When throughput regressed but GPU kernel execution times are unchanged (Root Cause)\n- When the gap between forward step iterations has increased (Root Cause)\n- To compare inter-iteration overhead between two versions of the inference engine (Root Cause)\n- When you need sub-operation granularity on inter-kernel gaps or graph coverage (Root Cause, kernel-level drill-down)\n- When piecewise CUDA graph coverage is unexpectedly low (Root Cause, kernel-level drill-down)\n- When multi-rank inference shows unexplained performance asymmetry (Root Cause, kernel-level drill-down)\n\nDo NOT use when:\n- The regression is in individual kernel performance (use `perf-nsight-compute-analysis`)\n- You need to profile a workload from scratch (use `workload-instrumentation` first)\n- The issue is NCCL communication (use distributed analysis)\n\n## Prerequisites\n\n- An nsys trace file (`.sqlite` or `.nsys-rep`) from a TRT-LLM benchmark run\n- For Root Cause comparison: two traces (baseline and target)\n- Python 3 with sqlite3 support\n\n---\n\n## Key Concepts\n\n### Host Overhead in LLM Inference\n\nIn an LLM inference loop, each iteration consists of:\n```\n[inter-step gap] -> [_forward_step] -> [inter-step gap] -> [_forward_step] -> ...\n```\n\nThe **forward step** includes GPU kernel execution (GEMM, attention, normalization, allreduce) plus host-side preparation. The **inter-step gap** includes host-side work between forward steps (scheduling, request fetching, broadcasting, sampling, response handling).\n\nSee [references/trtllm-nvtx-ranges.md](references/trtllm-nvtx-ranges.md) for the full per-operation breakdown and timing ranges.\n\n### Hidden vs Exposed Host Overhead\n\nHost overhead only hurts performance when it is **exposed** -- the GPU is idle waiting for work. When host prep overlaps with GPU execution, it is **hidden** and free. See [references/metrics.md](references/metrics.md) (M3 section) for diagrams and the exposed/hidden computation.\n\n### Forward Step Isolation\n\nIn TP configurations, forward steps are isolated via allreduce kernel grouping (deterministic count per transformer layer). For TP=1, NVTX `_forward_step` ranges are used directly. See [references/iteration-isolation-techniques.md](references/iteration-isolation-techniques.md) for the full algorithm.\n\n### Phase Classification (Context vs Generation)\n\nIterations are classified by NVTX marker text into **context** (eager, no CUDA graphs) and **generation** (CUDA graph replay). Per-phase analysis is critical because aggregate metrics can mask phase-specific bottlenecks. See [references/phase-classification.md](references/phase-classification.md).\n\n---\n\n## Phase 1: Detection (YES/NO Verdict)\n\nDetermine whether host overhead is the primary bottleneck.\n\n### Detection Metrics\n\nSix metrics in four categories. See [references/metrics.md](references/metrics.md) for full definitions, formulas, and SQL queries.\n\n| # | Metric | Threshold | What it answers |\n|---|--------|-----------|-----------------|\n| M1 | GPU idle ratio | > 0.30 | Is the GPU starved for work? |\n| M2 | Launch overhead ratio | > 0.10 | Is kernel launch itself expensive? |\n| M3a | Host prep exposed ratio | > 0.50 | How well is host prep pipelined? |\n| M3b | Host prep perf impact | > 0.05 | How much throughput does exposed prep cost? |\n| M3c | Host prep idle attribution | > 0.50 | Is host prep the main cause of GPU idle? |\n| M4 | GPU utilization | < 0.60 | Is GPU utilization too low? |\n| M5 | NCCL ratio (caveat) | > 0.20 | Is communication a confounding factor? |\n\n**Host prep confirmation rule**: Host prep is a confirmed bottleneck only when **both** M3b AND M3c cross their thresholds.\n\nThresholds are configurable with per-phase variants. See [references/thresholds.md](references/thresholds.md).\n\n### Detection Workflow\n\n#### Step 1: Input Validation\n\n```bash\n# Accept .sqlite or .nsys-rep\nls -la <trace_file>\n\n# If .nsys-rep, export to SQLite first\nnsys export -t sqlite -o <output.sqlite> <input.nsys-rep>\n```\n\n#### Step 2: Run Detection Script\n\n```bash\npython scripts/detect_host_overhead.py \\\n  --trace /path/to/trace.sqlite \\\n  --output /path/to/verdict.json\n```\n\nThe script computes M1, M2, M4, M5 from SQL, optionally M3 via range intersection, applies the verdict logic, and outputs structured JSON. See [references/output-format.md](references/output-format.md) for the output schema.\n\nFor manual metric extraction via SQL, see [references/nsys-schema.md](references/nsys-schema.md).\n\n#### Step 3: Interpret Verdict\n\n**Overall Verdict:**\n```\nif aggregate_verdict == YES or context_verdict == YES or generation_verdict == YES:\n    overall_verdict = YES\n```\n\nPer-phase analysis can **elevate** the verdict but never **demote** it.\n\nFormat using the template in [references/output-format.md](references/output-format.md).\n\n**Next Steps:**\n- **If YES** -> Proceed to Phase 2 (Root Cause) below, then use `perf-host-optimization` skill\n- **If NO** -> Use `perf-nsight-compute-analysis` for kernel SOL% or `perf-analysis` for full classification\n\n---\n\n## Phase 2: Root Cause Analysis\n\nIdentify which specific host operations regressed and by how much. Works with a single trace (breakdown) or two traces (comparison).\n\n### Principles\n\n1. **Isolate forward steps, not the full trace.** nsys traces contain warmup, JIT, model loading, and teardown.\n2. **Use structural kernel patterns for iteration detection.** Allreduce grouping is more robust than kernel density.\n3. **Compare steady-state iterations.** Filter to identical workload (same batch size, same ctx/gen mix).\n4. **Per-step metrics, not totals.** Always compare per-step averages.\n\n### Root Cause Workflow\n\n#### Step 1: Collect nsys Traces\n\nProfile both versions (if comparing) with identical settings:\n\n```bash\nnsys profile -o /path/to/trace \\\n  -t cuda,nvtx,osrt \\\n  --force-overwrite=true \\\n  --cuda-memory-usage=true \\\n  -w true \\\n  <benchmark_command> --num_requests 500\n```\n\n#### Step 2: Export to SQLite\n\n```bash\nnsys export --type=sqlite --force-overwrite=true -o trace.sqlite trace.nsys-rep\n```\n\n#### Step 3: Run Host Overhead Analysis\n\n```bash\n# Two-trace comparison\npython scripts/analyze_host_overhead.py \\\n  --baseline /path/to/baseline/trace.sqlite \\\n  --target /path/to/target/trace.sqlite \\\n  --baseline-label \"v1.1\" \\\n  --target-label \"main\" \\\n  --output /path/to/output/analysis.txt\n\n# Single-trace breakdown\npython scripts/analyze_host_overhead.py \\\n  --baseline /path/to/trace.sqlite \\\n  --baseline-label \"current\"\n```\n\n#### Step 4: Interpret Results\n\nThe script produces:\n1. **Allreduce-based iteration detection** -- confirms forward step boundaries\n2. **Per-step wall time comparison** -- quantifies the regression\n3. **NVTX per-step breakdown** -- identifies which host operations regressed\n4. **GPU kernel comparison** -- confirms GPU execution is unchanged\n5. **CUDA API comparison** -- detects kernel launch overhead changes\n\n### Reading the Output\n\n**Per-Step Wall Time:**\n```\nAvg wall time per step: 3,317 us (baseline) vs 3,978 us (target)  +19.9%\n```\nThis is the primary regression metric.\n\n**NVTX Breakdown:**\n```\nOperation           | baseline (us/step) | target (us/step) | Delta    | Status\n_fetch_new_requests |               36   |             270  | +234     | REGRESSION\nbroadcast_requests  |                -   |             250  | +250     | NEW\n_update_requests    |              413   |             723  | +310     | REGRESSION\n```\nFocus on operations with large absolute deltas.\n\n**GPU Kernel Comparison:**\n```\nKernels per step (launched): 6.2 (baseline) vs 21.9 (target)  +253%\n```\nMore individual launches = more host-side launch overhead.\n\n### Step 5: Kernel-Level Drill-Down (Optional)\n\nWhen the NVTX breakdown identifies a regressing operation but does not reveal *why* (the overhead is inside the GPU dispatch, not between NVTX ranges), drill below NVTX operations into individual GPU kernel launches.\n\nSee [references/kernel-level-analysis.md](references/kernel-level-analysis.md) for full technique details, SQL queries, and examples.\n\n**When to drill down:**\n- An operation has high wall time but the overhead is inside GPU dispatch, not between NVTX ranges\n- You need to understand how much of the forward pass is graph-captured vs eager\n- Per-layer overhead is significant and you need to map kernels to functional groups\n- Multi-rank inference shows unexplained performance asymmetry\n\n#### Kernel-Level Techniques\n\n| Technique | Question | Key Output |\n|-----------|----------|------------|\n| **Inter-Kernel Gap Analysis** | Where is the GPU idle between kernels? | Gap bucket distribution, top-N largest gaps with source mapping |\n| **Eager vs Graph Classification** | What fraction of kernels are graph-captured? | Graph coverage ratio, list of eager kernels with source attribution |\n| **Repeating-Pattern Mapping** | Which functional group within a layer has the most overhead? | Per-group gap totals, priority ranking |\n| **Straggler Detection** | Is one rank consistently slower? | Straggler rank ID, root cause (extra host work, queue depth feedback loop) |\n\n#### Workflow\n\n1. **Start with Inter-Kernel Gap Analysis** — bucket the gap distribution to understand the dominant overhead type (graph dispatch, Python interpreter, host-device sync)\n2. **If piecewise graph is in use**, run Eager vs Graph Classification to measure graph coverage and identify unnecessary eager kernels\n3. **For per-layer overhead**, use Repeating-Pattern Mapping to isolate the highest-overhead functional group within a single layer\n4. **For multi-rank setups**, run Straggler Detection if per-step wall time varies across ranks\n\n#### Kernel-Level Findings to Optimization Patterns\n\n| Finding | Optimization Pattern |\n|---------|---------------------|\n| Large gaps from Python tensor view chains | CUSTOM_OP — replace with C++ custom op |\n| Graph-capturable kernels running eagerly | GRAPH_EXPAND — fix partition poisoning |\n| Monolithic custom op blocking graph capture | GRAPH_SPLIT — split into capturable + eager parts |\n| Host-device sync (`.item()`) in per-layer code | SYNC (Pattern 1: pre-compute on CPU) + HOIST (Variant B: pass from step level) |\n| Per-layer buffer allocation | ALLOC — pre-allocate at init |\n| Straggler rank with extra host work | Apply targeted optimization to coordinator-only code paths |\n\n---\n\n## Common Patterns and Root Causes\n\n### Pattern 1: Request Management Refactor\n**Symptom**: `_fetch_new_requests` regressed 5-10x, new `broadcast_requests` operation.\n**Cause**: Request fetching refactored for multi-rank broadcasting in TP.\n**Mitigation**: Optimize broadcast path; batch request state updates.\n\n### Pattern 2: Increased Kernel Launch Count\n**Symptom**: 3-5x more `cudaLaunchKernel` calls per step, similar GPU time.\n**Cause**: Operations that were fused or graph-captured are now individual launches.\n**Mitigation**: Re-fuse kernels; extend CUDA graph capture scope.\n\n### Pattern 3: New Bookkeeping Operations\n**Symptom**: New NVTX ranges like `_write_finish_reasons`, `handle_additional_outputs`.\n**Cause**: New features added to the inference loop without overhead budgeting.\n**Mitigation**: Defer non-critical bookkeeping to async paths; batch updates.\n\n### Pattern 4: Flashinfer JIT Warmup Masquerading as Inference\n**Symptom**: Massive elementwise/reduce kernel counts in \"steady state\" analysis.\n**Cause**: Analysis window includes flashinfer JIT compilation phase.\n**Fix**: Use allreduce-based iteration isolation, not kernel density or time windows.\n\n### Pattern 5: Context-Only Bottleneck (Masked by Aggregate)\n**Symptom**: Aggregate metrics below threshold, but context iterations have 50% GPU idle.\n**Cause**: Generation iterations dilute the context-phase bottleneck.\n**Fix**: Per-phase analysis in Detection phase catches this.\n\n---\n\n## Pitfalls\n\n### 1. shortName is an Integer ID\nIn `CUPTI_ACTIVITY_KIND_KERNEL`, `shortName` is an integer referencing `StringIds.id`. Always join. See [references/nsys-schema.md](references/nsys-schema.md).\n\n### 2. NVTX textId vs text\nMost NVTX events have `textId` (integer) but NULL `text`. Join with StringIds. See [references/nsys-schema.md](references/nsys-schema.md).\n\n### 3. Duplicate NVTX Ranges from TP Ranks\nIn TP configurations, each rank reports NVTX ranges independently. De-duplicate by grouping entries within 100us of each other.\n\n### 4. Negative Inter-Step Gaps\nWhen TP ranks report overlapping NVTX ranges, `gap = next_start - prev_end` can be negative. Use the maximum end time when de-duplicating.\n\n### 5. Benchmark Window Selection\nThe allreduce-based window captures context+generation phases; steady-state NVTX filtering captures generation-only. Both are valid; use the appropriate one for your comparison goal.\n\n---\n\n## Handoff to Optimization\n\nWhen analysis is complete and the verdict is **YES**, hand off to the `perf-host-optimization` skill with:\n\n1. **Detection verdict and evidence**: Which metrics crossed thresholds (M1-M5), whether host prep was confirmed (M3b+M3c), and per-phase breakdown.\n\n2. **NVTX-based triage** (from Root Cause): Top regressing operations by absolute delta (us/step). Map NVTX range names to source functions -- see [references/trtllm-nvtx-ranges.md](references/trtllm-nvtx-ranges.md).\n\n3. **Handoff data block**: Include structured data from [references/output-format.md](references/output-format.md) (see \"Handoff to Optimization\" section).\n\n4. **Kernel-level findings** (from drill-down, if performed): Inter-kernel gap distribution, graph coverage ratio, per-group overhead map, and straggler rank identification. Map findings to optimization patterns using the table in the Root Cause kernel-level drill-down section above.\n\n---\n\n## Reference\n\n| File | Contents |\n|------|----------|\n| [references/metrics.md](references/metrics.md) | Full metric definitions, formulas, SQL queries, M3 sub-metric analysis |\n| [references/thresholds.md](references/thresholds.md) | Aggregate and per-phase threshold tables |\n| [references/phase-classification.md](references/phase-classification.md) | NVTX marker parsing, iteration classification, per-phase aggregation |\n| [references/output-format.md](references/output-format.md) | Report template and integration JSON schema |\n| [references/examples.md](references/examples.md) | Worked scenarios (aggregate, phase-specific, and case study) |\n| [references/iteration-isolation-techniques.md](references/iteration-isolation-techniques.md) | Allreduce, NVTX, and kernel-density iteration isolation techniques |\n| [references/trtllm-nvtx-ranges.md](references/trtllm-nvtx-ranges.md) | TRT-LLM NVTX range reference with per-operation timings |\n| [references/kernel-level-analysis.md](references/kernel-level-analysis.md) | Kernel-level drill-down techniques: gap analysis, graph classification, pattern mapping, straggler detection |\n| [references/nsys-schema.md](references/nsys-schema.md) | nsys SQLite schema reference and useful queries |\n| [scripts/analyze_host_overhead.py](scripts/analyze_host_overhead.py) | Python script for Phase 2 root cause analysis |\n| [scripts/detect_host_overhead.py](scripts/detect_host_overhead.py) | Python script for Phase 1 detection verdict |\n",".claude/skills/perf-analysis/SKILL.md":"---\nname: perf-analysis\ndescription: >\n  Performance analysis coordination workflow. Guides profiling delegation,\n  bottleneck classification (compute/memory/launch/communication/sync),\n  and structured report generation. Use when the user asks to analyze\n  performance, profile a workload, check MFU/SOL, or diagnose bottlenecks.\ntags:\n  - analysis\n  - profiling\n  - bottleneck\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Performance Analysis\n\n## Principles\n\n1. **Delegate profiling, own analysis.** You coordinate the analysis workflow\n   but do not run profiling tools directly. Delegate all profiling and\n   measurement tasks to **perf-profiling-specialist** or other domain specialists.\n2. **Metrics from tools, never invented.** All performance numbers must come\n   from profiling tool output. Never fabricate metrics.\n3. **Classify before recommending.** Identify the bottleneck type before\n   suggesting optimizations. The wrong classification leads to wasted effort.\n4. **Structured reports.** Every analysis produces a report with Summary,\n   Metrics, Findings, and Recommendations.\n\n## Key Performance Metrics\n\n- **Throughput**: samples/sec, tokens/sec, iterations/sec\n- **Latency**: end-to-end time, kernel time, communication time\n- **MFU (Model FLOPs Utilization)**: actual FLOPs / theoretical peak FLOPs\n- **% of SOL (Speed of Light)**: current perf / hardware peak perf\n- **GPU Utilization**: SM occupancy, tensor core usage\n- **Memory Bandwidth**: DRAM bandwidth utilization vs peak\n\n## Analysis Workflow\n\n1. **Understand**: Clarify what metrics the user needs (MFU, SOL, latency, etc.)\n2. **Plan**: Plan your profiling and analysis steps before starting\n3. **Profile**: Delegate to **perf-profiling-specialist** for actual measurements\n4. **Measure**: Extract requested metrics from profiling results\n5. **Classify**: If diagnosing issues, determine primary bottleneck type\n6. **Report**: Generate performance analysis report with findings\n\n## Bottleneck Classification\n\nWhen diagnosing performance issues, classify the primary bottleneck:\n\n| Type | Indicator | Description |\n|------|-----------|-------------|\n| **Compute-bound** | High GPU utilization, low memory bandwidth usage | Limited by compute capacity (FLOPs) |\n| **Memory-bound** | High memory bandwidth, low compute utilization | Limited by DRAM throughput |\n| **Launch-overhead** | Many small kernels, high CPU time | CPU becoming bottleneck from kernel launch overhead |\n| **Communication-bound** | Significant time in collective operations | Limited by inter-GPU or inter-node communication |\n| **Sync-bound** | Excessive CPU-GPU synchronization points | Stalls from unnecessary synchronization |\n\nAfter classifying, consult the **perf-optimization-casebook** skill for\ncandidate optimizations (decision precedents) matching the bottleneck class,\nmodel, and hardware. Use it to populate the Recommendations section with\nproven, adapted approaches rather than inventing them — but report only\nmeasured numbers, never a case's expected effect as an achieved result.\n\n## Delegation Guidelines\n\nWhen delegating to specialists, describe the **desired outcome** -- not the\ntool methodology.\n\n**DO include:**\n- The workload: file path, code snippet, or command to profile\n- Problem context: dimensions, dtypes, FLOPs calculations, batch sizes\n- Desired metrics: SOL%, MFU, throughput, occupancy, bottleneck classification\n- Any constraints: specific kernel to target, profiling region markers\n\n**DO NOT include:**\n- Specific tool flags or command patterns (e.g., `--set=full`, `--section SpeedOfLight`)\n- Step-by-step tool usage instructions\n- Fallback strategies for tool failures\n- Example commands\n- Output file paths or artifact locations (specialists create their own workspace artifacts)\n\nSpecialists have their own skills that encode best practices for tool usage\nand their own workspace artifacts for output. Prescribing commands in the\ndelegation overrides their skills and may lead to suboptimal profiling\nstrategies (e.g., collecting 8000+ metrics with `--set=full` when a targeted\nsection analysis would be faster and more surgical).\n\n### Good Example\n\n```\nProfile the batched GEMM kernel in bmm_workload.py with NCU.\nThe workload uses cudaProfilerStart/Stop markers to isolate the region of interest.\nCollect kernel-level metrics: SOL%, compute/memory throughput, DRAM bandwidth,\ntensor core utilization, occupancy, warp stall reasons, and roofline classification.\nThe batched GEMM performs 68.72 GFLOP per call (B=32, M=512, N=1024, K=2048, FP16).\nCalculate MFU against the GPU's peak FP16 tensor core TFLOP/s.\n```\n\n### Bad Example\n\n```\nRun NCU with --set=full --profile-from-start off --target-processes all.\nIf --set=full fails, try --set=detailed. Parse the CSV output for\nsm__throughput.avg.pct_of_peak_sustained_elapsed.\nSave raw NCU output to ncu_output.txt.\n```\n\n### Remote Profiling\n\nWhen profiling on a remote SLURM cluster, include the\n**Remote Execution Context** block in the delegation prompt with the SSH+srun\nwrapper for the target cluster. The perf-profiling-specialist will prefix its\ncommands (nsys, ncu, nvidia-smi) with this wrapper.\n\nThe perf-profiling-specialist does not need the `exec-remote-slurm` skill — the\ncontext block provides everything it needs to execute remotely.\n\n## Available Specialists\n\nDelegate profiling and domain-specific analysis to these specialists:\n\n- **perf-profiling-specialist**: Runs nvidia-smi, nsys, ncu, torch.profiler. Use for ALL profiling tasks.\n- **perf-torch-cuda-graph-specialist**: Analyzes CUDA Graph compatibility and applies capture workflows\n\n## Report Format\n\nStructure every analysis report with these four sections:\n\n1. **Summary**: High-level performance status or bottleneck classification\n2. **Metrics**: Key performance numbers from profiling\n3. **Findings**: Detailed observations with evidence\n4. **Recommendations**: Prioritized list of optimizations (if applicable)\n\n### Example Report\n\n```\n## Summary\nTraining at 42% MFU, memory-bound due to large attention tensors.\n\n## Metrics\n- Throughput: 1,247 samples/sec\n- MFU: 42% (vs 65% theoretical for this model)\n- % of SOL: 58% (room for 1.7x improvement)\n- GPU Utilization: 45%\n- Memory Bandwidth: 850 GB/s (89% of peak)\n- Kernel Count: 1,247 per iteration\n\n## Findings\n1. Self-attention consumes 60% of memory bandwidth\n2. Optimizer step has 3 unnecessary synchronizations\n3. Batch size could be increased by 2x\n\n## Recommendations\n1. Enable FlashAttention (expected: +15% MFU)\n2. Remove synchronizations in optimizer (expected: +5% throughput)\n3. Increase batch size to improve GPU utilization\n```\n",".claude/skills/perf-host-optimization/SKILL.md":"---\nname: perf-host-optimization\ndescription: Profiles and optimizes TensorRT-LLM host/CPU overhead using line_profiler (with nsys support planned). Runs iterative profile-analyze-optimize-validate rounds. Use when GPU utilization is low or optimizing PyExecutor throughput.\nlicense: Apache-2.0\ntags:\n  - optimization\n  - profiling\n  - host-overhead\n  - line-profiler\n  - inference\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Host Performance Optimization Skill\n\nAutomates detection and optimization of host-side (CPU) overhead in TensorRT-LLM's PyTorch backend.\n\n## When to Use\n\n- GPU utilization is low during inference (CPU bottleneck suspected)\n- User asks to reduce host overhead or CPU latency\n- Optimizing PyExecutor throughput (requests/sec)\n- Need line-by-line profiling of specific Python functions\n\n### Confirming the Bottleneck\n\nline_profiler measures *where* CPU time is spent but not *whether* CPU is the bottleneck.\nIf you need to confirm CPU is the limiting factor, run the `perf-host-analysis` skill first -- it provides a YES/NO verdict with metric evidence.\n\nAs a rough heuristic without nsys: if doubling the batch size does not proportionally increase GPU utilization or throughput, CPU overhead is likely the bottleneck.\n\n### Using Analysis Skill Results\n\nIf the `perf-host-analysis` skill has already been run, use its output to skip the confirmation step and prioritize targets:\n\n1. **Detection verdict**: If YES with host_prep_confirmed, start with `_prepare_tp_inputs`.\n2. **NVTX triage** (from Root Cause): The `top_regressing_ops` in the handoff data block maps NVTX range names to source functions. Profile the function with the largest absolute delta first.\n3. **Cross-function triage**: When the top NVTX regression is NOT in `_prepare_tp_inputs` (e.g., `_fetch_new_requests`, `broadcast_requests`, `_update_requests`), target that function's source file directly instead of defaulting to `_prepare_tp_inputs`. See [references/trtllm-nvtx-ranges.md](../perf-host-analysis/references/trtllm-nvtx-ranges.md) for the NVTX-to-source mapping.\n\n---\n\n## Profiling Setup\n\n### line_profiler (Primary Method)\n\n**Environment Variables:**\n- `TLLM_LINE_PROFILER_ENABLED=True` — Enable the profiler\n- `TLLM_LINE_PROFILER_PATH` — Output file path\n- `TLLM_LINE_PROFILER_FUNCTIONS` — Additional functions to profile (comma-separated)\n\n**Function specification format:**\n```bash\n# Class methods: module.path.ClassName.method_name\nTLLM_LINE_PROFILER_FUNCTIONS=\"tensorrt_llm._torch.pyexecutor.model_engine.PyTorchModelEngine._prepare_tp_inputs\"\n\n# Standalone functions: module.path::function_name\nTLLM_LINE_PROFILER_FUNCTIONS=\"tensorrt_llm._torch.pyexecutor.sampler::_group_requests_by_strategy_key\"\n\n# Multiple functions (comma-separated)\nTLLM_LINE_PROFILER_FUNCTIONS=\"module.Class.method1,module.Class.method2\"\n```\n\n### CPU Affinity (Environment Factor)\n\nCPU core affinity can significantly affect host overhead measurements,\nespecially on multi-socket systems (e.g., B300). Pinning processes to cores\nnear the GPU's NUMA node reduces cross-socket memory access latency.\n\n- Check current affinity: `taskset -p <pid>` or `numactl --show`\n- Pin to local NUMA node: `numactl --cpunodebind=<node> --membind=<node>`\n- **Impact**: Up to 2x difference in host overhead on B300 systems\n\nWhen comparing profiling results across runs, ensure CPU affinity is consistent.\nDo not externally modify the affinity, unless user requires to do this to examine the affects of this part.\nDocument the affinity setting in each round's report if it varies.\n\n### Workspace & Suffix Management\n\nEach profiling run should have a unique suffix to track progress across rounds:\n```bash\nEXTRA_SUFFIX=round0_baseline bash profile.sh\nEXTRA_SUFFIX=round1_eliminate_redundant_iter bash profile.sh\n```\n\n---\n\n## Autonomous Optimization Loop\n\nBefore starting the loop, review [references/optimization-strategy.md](references/optimization-strategy.md) for strategic guidance on ordering (zero-risk-first), measurement traps, and overhead scoping.\n\n**Key insight**: Optimizations are NOT independent. Fixing a 50ms bottleneck may reveal a 30ms bottleneck that was previously masked (hidden behind the larger one). Always re-profile after each significant change — the bottleneck landscape shifts.\n\n**Ordering principle**: Within each round, prefer zero-risk optimizations (caching, pre-allocation, hoisting invariants) over medium/high-risk ones (graph partition changes, algorithm fusion). Zero-risk changes provide free gains and make subsequent profiling cleaner.\n\nRun N rounds (default 3) of the following cycle:\n\n```\nFOR round = 1 to MAX_ROUNDS:\n\n  1. PROFILE (with Drill-Down)\n  2. ANALYZE (Multi-Option)\n  3. OPTIMIZE (Apply Change — prefer zero-risk first)\n  4. TEST (Unit Test Validation)\n  5. VALIDATE (Re-Profile — expect bottleneck landscape to shift)\n  6. REPORT\n\nEND FOR → FINAL SUMMARY\n```\n\n### Phase 1: PROFILE (with Drill-Down)\n\n- Run workload with profiler enabled\n- Parse output: identify functions with highest Total time and lines with highest % Time\n- **CRITICAL: Drill down into sub-functions that are not yet profiled** (see below)\n\n#### Drill-Down Profiling\n\nThe default profiler covers top-level executor functions but **not all sub-functions**. When a profiled function shows most time in a single sub-call, you must drill down.\n\n**When:** A single line consumes >80% of a function's time calling an unprofiled sub-function:\n```\nLine #      Hits         Time    Per Hit   % Time  Line Contents\n==============================================================\n  2848      4100  59200000000.0  14439024.4   98.7      output = self.model_engine.forward(...)\n```\n\n**How:**\n1. Identify the sub-function's full qualified path (e.g., `tensorrt_llm._torch.pyexecutor.model_engine.PyTorchModelEngine._prepare_tp_inputs`)\n2. Add it to `TLLM_LINE_PROFILER_FUNCTIONS`\n3. Re-profile to get line-level data inside it\n4. Now analyze the **inner** hotspots\n\nFor common drill-down targets, see [references/hot-path-files.md](references/hot-path-files.md).\n\n### Phase 2: ANALYZE (Multi-Option)\n\nFor the chosen hotspot:\n\n1. **Identify** the top hotspots by **absolute time** (not just %) within the target function\n2. **Classify** each hotspot by type. Summary table:\n\n| Type | Indicators | Severity |\n|------|------------|----------|\n| **HOST_SYNC** | `.item()`, `.cpu()` in per-layer forward path | Critical |\n| **SYNC** | `.item()`, `.cpu()`, `synchronize()` in step-level code | Critical |\n| **CUSTOM_OP** | Chain of Python tensor ops (view/slice/cast) before kernel launch | Critical |\n| **GRAPH_BREAK** | Op that prevents CUDA graph capture of surrounding code (fix via GRAPH_EXPAND / GRAPH_SPLIT) | High |\n| **ALLOC** | `torch.zeros/empty/tensor()` in loops, `.clone()` | High |\n| **HOIST** | Per-layer recomputation of step-invariant values | High |\n| **PYLOOP** | `for x in collection:` with many iterations | High |\n| **REDUNDANT_ITER** | Multiple passes over the same collection | High |\n| **DEAD_WORK** | Object construction whose results are always discarded | High |\n| **CONTAINER** | Dict/set lookups in hot loops | Medium |\n| **FUNCALL** | Repeated method/property calls | Medium |\n| **COMM** | `dist.all_reduce`, `dist.barrier`, NCCL overhead in TP/PP paths | Medium |\n| **GIL** | Lock/queue contention | Medium |\n| **SERIALIZE** | `pickle.dumps/loads`, `json.dumps/loads` in request processing | Medium |\n| **GC** | Periodic latency spikes, non-deterministic pauses (tail latency) | Low |\n| **COMPUTE** | Actual computation (may not be optimizable) | Low |\n\nFor detailed classification with code examples, see [references/hotspot-classification.md](references/hotspot-classification.md).\n\n3. **Propose 2-4 optimization options** in a table:\n\n| Option | Description | Estimated Savings | Risk | Complexity |\n|--------|-------------|-------------------|------|------------|\n| A | ... | ... | Low/Med/High | ... |\n| B | ... | ... | ... | ... |\n\n4. **Select the best option** and explain reasoning (prefer high-savings + low-risk; follow zero-risk-first ordering from [references/optimization-strategy.md](references/optimization-strategy.md))\n\nFor optimization patterns by type, see [references/optimization-patterns.md](references/optimization-patterns.md) (index) — it links to the relevant sub-file for each hotspot type. For GPU-specific patterns (CUSTOM_OP, GRAPH_SPLIT, GRAPH_EXPAND), see [references/patterns/gpu-graph.md](references/patterns/gpu-graph.md).\n\n### Phase 3: OPTIMIZE (Apply Change)\n\n- Apply the selected code change with Edit tool\n- **One optimization per round** — keep changes minimal and targeted\n- Record the exact change (file, line range, before/after) for potential rollback\n\n### Phase 4: TEST (Unit Test Validation)\n\n**Mandatory** after each optimization. Find and run related UTs to verify correctness.\n\n**Finding related tests:**\n```bash\n# Search by modified file name\ngrep -rl \"model_engine\\|PyTorchModelEngine\" tests/unittest/_torch/executor/\n\n# Search by modified function name\ngrep -rl \"_prepare_tp_inputs\\|prepare_inputs\" tests/\n```\n\n**Running tests:**\n```bash\n# Run specific test file with stop-on-first-failure\npytest tests/unittest/_torch/executor/test_pytorch_model_engine.py -v -x --timeout=120\n\n# Run specific test method\npytest tests/unittest/_torch/executor/test_pytorch_model_engine.py::PyTorchModelEngineTestCase::test_position_id_preparation -v -x\n```\n\nFor the full UT-to-file mapping, see [references/hot-path-files.md](references/hot-path-files.md).\n\n**If tests fail:**\n1. Read the failure message\n2. Rollback immediately (`git checkout -- <file>`)\n3. Analyze why the optimization broke correctness\n4. Try the next-best option from Phase 2\n\n### Phase 5: VALIDATE (Re-Profile)\n\n- Re-run profiler with identical workload, using suffix `round<N>_<description>`\n- Compare three things:\n  1. Did the **target hotspot** time decrease?\n  2. Did the **overall function** Total time decrease?\n  3. Did **benchmark metrics** (TPOT, throughput) improve?\n\n**If regression detected** (function time increased or metrics worsened):\n- The \"optimization\" may have triggered a CPython pitfall — see [references/patterns/compound-pitfalls.md](references/patterns/compound-pitfalls.md) (CPython Pitfalls section)\n- Rollback and try the next-best option from Phase 2\n\n### Phase 6: REPORT\n\nLog for this round:\n- Round number\n- Hotspot location (file:line) and classification\n- Optimization applied (with before/after code summary)\n- Time delta: function Total time before → after\n- Benchmark delta: TPOT, throughput before → after\n\n---\n\n## Reading Profile Output\n\n```\nTimer unit: 1e-06 s\nTotal time: 1.234 s\nFile: /path/to/file.py\nFunction: my_function at line 100\n\nLine #      Hits         Time  Per Hit   % Time  Line Contents\n==============================================================\n   100                                           def my_function(self):\n   101       500    890000.0   1780.0     72.1       result = tensor.item()\n   102       500    234567.0    469.1     19.0       return result\n```\n\n**How to read effectively:**\n1. Start with **Total time** for each function — this is the overall budget\n2. Sort lines mentally by **absolute Time**, not just % Time (3% of a 60s function = 1.8s)\n3. Check **Hits count** to understand iteration patterns:\n   - Hits = 2 × expected count → `for x in range(1):` loop overhead (2 hits = enter + exit check)\n   - Hits ≫ expected → the line is inside a nested loop\n4. Look for **repeated patterns**: if 10 lines each take 3% in a loop body, the loop itself costs 30%\n\n---\n\n## Stopping Criteria\n\nStop the optimization loop when:\n1. **Iteration limit reached**: Completed N rounds (default 3)\n2. **No actionable hotspots**: Top hotspots are pure GPU compute (COMPUTE type)\n3. **Diminishing returns**: < 5% improvement in last 2 rounds\n4. **Risk threshold**: Further optimizations require architectural changes (e.g., Cython, struct-of-arrays)\n5. **Test failures**: Cannot find an optimization that passes UTs\n\n**Primary success metric**: Benchmark throughput (requests/sec or tokens/sec) as measured by the profiling script. line_profiler time reductions are leading indicators, but throughput is the ground truth — a function-level speedup that doesn't improve throughput is not a real win.\n\n---\n\n## Final Summary Output\n\nThe final report should include:\n- **Rounds executed**: Number of profile-optimize cycles completed\n- **Cumulative improvement**: Total host time reduction (percentage and absolute)\n- **Benchmark metrics**: Before/after comparison table (TPOT, throughput, ITL, E2EL)\n- **Optimizations applied**: List of changes with file:line locations and classification\n- **Failed attempts**: Any optimizations tried and reverted (with why)\n- **Remaining hotspots**: Top bottlenecks that couldn't be optimized (with classification)\n- **Recommendations**: Suggested follow-up for architectural changes if needed\n\nFor a concrete multi-round example, see [references/examples.md](references/examples.md).\n\n---\n\n## Reference Files\n\n| File | Contents |\n|------|----------|\n| [references/optimization-patterns.md](references/optimization-patterns.md) | Pattern index — links to 6 sub-files: [sync-alloc](references/patterns/sync-alloc.md), [loop-iteration](references/patterns/loop-iteration.md), [python-overhead](references/patterns/python-overhead.md), [gpu-graph](references/patterns/gpu-graph.md), [system](references/patterns/system.md), [compound-pitfalls](references/patterns/compound-pitfalls.md) |\n| [references/optimization-strategy.md](references/optimization-strategy.md) | Zero-risk-first ordering, metric traps, three scopes of host overhead, pattern selection guide |\n| [references/hotspot-classification.md](references/hotspot-classification.md) | Extended per-type indicators and code examples (including CUSTOM_OP, GRAPH_BREAK, HOST_SYNC) |\n| [references/communication-patterns.md](references/communication-patterns.md) | Communication overhead patterns (NCCL batching, barrier removal, async overlap, reduce_scatter) |\n| [references/hot-path-files.md](references/hot-path-files.md) | Key file tables, drill-down targets, UT mapping |\n| [references/examples.md](references/examples.md) | Usage examples and multi-round walkthrough |\n| [trtllm-nvtx-ranges.md](../perf-host-analysis/references/trtllm-nvtx-ranges.md) | TRT-LLM NVTX range reference (from analysis skill) — maps range names to source functions |\n",".claude/skills/kernel-triton-writing/SKILL.md":"---\nname: kernel-triton-writing\ntags: [triton]\ndescription: >\n  ONLY for OpenAI Triton (@triton.jit) kernel development. NEVER use for\n  CUDA C++ kernels, CuTe DSL kernels, TileIR, or profiling tools (ncu, nsys).\n  The user's request must involve Triton explicitly. Covers Triton-specific\n  patterns: fused elementwise, reductions (softmax, LayerNorm, RMSNorm),\n  tiled GEMM with triton.autotune, and flash attention. Workflow:\n  design, write, verify (with fast-path for explicit requests).\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Triton Kernel Writing\n\n## Principles\n\n### Correctness First\n\n1. Never benchmark before verification passes.\n2. Always mask loads and stores for non-divisible shapes.\n3. Include `kernel_fn`, `reference_fn`, and `get_inputs()` exports for companion scripts.\n4. Always run `scripts/verify_kernel.py` to validate against the reference.\n\n### FP16/BF16 Precision Rules (LOW FREEDOM -- follow exactly)\n\nTranscendental functions (`tl.exp`, `tl.log`, `tl.math.erf`, `tl.math.tanh`) require fp32 inputs.\n\n```python\n# WRONG -- compilation error or wrong results with fp16/bf16:\nresult = tl.exp(x)\n\n# CORRECT -- cast to fp32, compute, cast back:\nx_fp32 = x.to(tl.float32)\nresult = tl.exp(x_fp32).to(x.dtype)\n```\n\nRule: any math function beyond basic arithmetic (+, -, *, /) requires fp32 cast in, original dtype cast out.\n\nAdditional precision constraints:\n- `tl.sigmoid()` is unavailable in some Triton versions. Use `1.0 / (1.0 + tl.exp(-x_fp32))`.\n- Always cast back to `x.dtype` before `tl.store` -- mismatches cause \"Type mismatch, store Float32 to Float16\".\n- Unlike PyTorch, Triton does NOT auto-promote fp16/bf16 to fp32 for accumulation. Always use `tl.float32` accumulators for `tl.dot`.\n- **TF32 for matmul:** On Ampere+/Hopper, `tl.dot` uses TF32 by default for fp32 inputs (same as `torch.mm`). Do NOT add `input_precision=\"ieee\"` — it is 3-8x slower. TF32 is the correct default. If verification fails due to TF32 precision (~0.01-0.1 abs diff), ensure `reference_fn` also uses TF32 (plain `torch.mm`, no `allow_tf32=False`).\n\n### CPU-GPU Sync Avoidance (LOW FREEDOM)\n\nNever call `.item()` in kernel wrappers. It forces a CPU-GPU sync (~50-100us per call).\n\n| Pitfall | Fix |\n|---------|-----|\n| `tensor.item()` for seed | `x.data_ptr() % (2**31)` |\n| `torch.randint(...).item()` | Use tensor metadata for pseudo-random seed |\n| Allocating output every call | Accept pre-allocated output as parameter |\n| Python loops calling kernel | Batch operations |\n\n### C Integer Division Semantics (CRITICAL)\n\nTriton uses **C semantics** (round toward zero) for `//` and `%`, NOT Python semantics (round toward negative infinity). This only matters when operands can be negative.\n\n| Expression | Python | Triton/C |\n|------------|--------|----------|\n| `-7 // 2` | `-4` | `-3` |\n| `-7 % 2` | `1` | `-1` |\n\n**Safe pattern:** Ensure all index/offset values are non-negative. If negative values are possible, use `(idx % BLOCK + BLOCK) % BLOCK`.\n\nSee [references/concepts-semantics.md](references/concepts-semantics.md) for full rules and scalar-only exception.\n\n### Kernel Design Mental Model\n\n- **Parallelization axis:** Element-wise kernels parallelize over flattened elements. Row-wise kernels (LayerNorm, softmax) parallelize over rows. Matmul kernels tile in 2D (M, N).\n- **Block size:** Power-of-2 only (256, 512, 1024, 2048). Start with 1024 for H100, 512 for V100.\n- **Memory coalescing:** Adjacent threads must access adjacent memory addresses. The compiler handles this automatically from block-level pointer arithmetic.\n- **Grid:** Use `triton.cdiv(n_elements, BLOCK_SIZE)`. With autotune, grid must be a lambda: `lambda meta: (triton.cdiv(n, meta['BLOCK_SIZE']),)`.\n- **Decorator order:** `@triton.autotune` (outermost) -> `@triton.heuristics` -> `@triton.jit` (innermost).\n- **`reset_to_zero`:** Required for autotune on kernels that accumulate (e.g., matmul output). Without it, later configs see leftover values from earlier trials.\n\n## Workflow\n\n**Fast path:** If the user explicitly requests a Triton kernel (e.g., \"Write a Triton kernel for X\", \"Implement softmax in Triton\"), start at **Phase 2**. Only use Phase 0-1 when the request is ambiguous about whether Triton is appropriate.\n\n### Phase 0: Route the Operator (only for ambiguous requests)\n\nSkip this phase if the user explicitly asks for a Triton kernel. Only use when the request is ambiguous (e.g., \"optimize this operation\").\n\nTriton wins when 2+ operations can share registers instead of writing/reading global memory. Quick rules:\n\n| Pattern | Decision |\n|---------|----------|\n| Single element-wise op (`relu`, `sigmoid`) | SKIP — PyTorch already optimal |\n| Standalone matmul | SKIP — cuBLAS is optimized |\n| Standard attention | SKIP — Use FlashAttention |\n| Element-wise chain (2+ ops), reduction, matmul + epilogue | USE TRITON |\n\nIf SKIP, recommend the alternative and STOP. See [references/operator-routing.md](references/operator-routing.md) for edge cases.\n\n### Phase 1: Analyze the Operator (only for ambiguous requests)\n\nFrom the user's request, identify: (1) operation type, (2) parallelization strategy, (3) input shapes and dtypes.\n\n### Phase 2: Design the Kernel\n\nPick the skeleton below that matches your operation. **These skeletons are sufficient for element-wise, reduction, matmul, and fusion kernels — do NOT read reference files for these common patterns.** Only consult `references/` when implementing uncommon patterns (grouped GEMM, TMA, extern functions) or debugging issues.\n\n**Element-wise skeleton** (GELU, dropout, fused ops on flat tensors):\n\n```python\n@triton.jit\ndef kernel(x_ptr, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):\n    pid = tl.program_id(0)\n    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n    mask = offsets < n_elements\n    x = tl.load(x_ptr + offsets, mask=mask)\n    # ... compute ...\n    tl.store(out_ptr + offsets, result, mask=mask)\n```\n\n**Row-wise skeleton** (softmax, LayerNorm, RMSNorm — one program per row):\n\n```python\n@triton.jit\ndef kernel(x_ptr, out_ptr, n_cols, BLOCK_SIZE: tl.constexpr):\n    row_idx = tl.program_id(0)\n    col_offsets = tl.arange(0, BLOCK_SIZE)\n    mask = col_offsets < n_cols\n    x = tl.load(x_ptr + row_idx * n_cols + col_offsets, mask=mask, other=0.0)\n    # ... reduce / normalize ...\n    tl.store(out_ptr + row_idx * n_cols + col_offsets, result, mask=mask)\n```\n\n**Tiled matmul skeleton** (GEMM with 2D tiling, grouped ordering, and autotune):\n\n```python\n@triton.autotune(\n    configs=[\n        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 256, 'BLOCK_K': 64, 'GROUP_SIZE_M': 8}, num_warps=8, num_stages=3),\n        triton.Config({'BLOCK_M': 64, 'BLOCK_N': 256, 'BLOCK_K': 32, 'GROUP_SIZE_M': 8}, num_warps=4, num_stages=4),\n        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 128, 'BLOCK_K': 32, 'GROUP_SIZE_M': 8}, num_warps=4, num_stages=4),\n        triton.Config({'BLOCK_M': 256, 'BLOCK_N': 64, 'BLOCK_K': 32, 'GROUP_SIZE_M': 8}, num_warps=4, num_stages=4),\n    ],\n    key=['M', 'N', 'K'],\n)\n@triton.jit\ndef matmul_kernel(\n    a_ptr, b_ptr, c_ptr, M, N, K,\n    stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn,\n    BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,\n    GROUP_SIZE_M: tl.constexpr,\n):\n    pid = tl.program_id(0)\n    num_m_blocks = tl.cdiv(M, BLOCK_M)\n    num_n_blocks = tl.cdiv(N, BLOCK_N)\n    # Grouped ordering for L2 cache locality\n    num_pid_in_group = GROUP_SIZE_M * num_n_blocks\n    group_id = pid // num_pid_in_group\n    first_pid_m = group_id * GROUP_SIZE_M\n    group_size_m = min(num_m_blocks - first_pid_m, GROUP_SIZE_M)\n    pid_m = first_pid_m + ((pid % num_pid_in_group) % group_size_m)\n    pid_n = (pid % num_pid_in_group) // group_size_m\n\n    offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)\n    offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)\n    offs_k = tl.arange(0, BLOCK_K)\n\n    a_ptrs = a_ptr + offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak\n    b_ptrs = b_ptr + offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn\n    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)\n\n    for k in range(0, tl.cdiv(K, BLOCK_K)):\n        a_mask = (offs_m[:, None] < M) & (offs_k[None, :] < K)\n        b_mask = (offs_k[:, None] < K) & (offs_n[None, :] < N)\n        a = tl.load(a_ptrs, mask=a_mask, other=0.0)\n        b = tl.load(b_ptrs, mask=b_mask, other=0.0)\n        acc += tl.dot(a, b)\n        a_ptrs += BLOCK_K * stride_ak\n        b_ptrs += BLOCK_K * stride_bk\n        offs_k += BLOCK_K\n\n    c_mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)\n    c_ptrs = c_ptr + offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn\n    tl.store(c_ptrs, acc.to(c_ptr.dtype.element_ty), mask=c_mask)\n```\n\n### Phase 3: Write the Kernel\n\nCreate an output directory, then write the kernel file to `{output_dir}/kernel.py`.\n\nThe kernel file MUST include:\n- `@triton.jit` decorated kernel function\n- `@triton.autotune` for production kernels (see [references/api-core.md](references/api-core.md))\n- Python wrapper function (descriptive name for external import)\n- **Fixed contract exports** (companion scripts rely on these exact names):\n  - `kernel_fn` — alias to the wrapper function\n  - `reference_fn(*args)` — PyTorch reference with identical signature\n  - `get_inputs()` — returns `list` of fresh CUDA tensors for testing/benchmarking\n\nConcise example (fused GELU + dropout):\n\n```python\nimport triton\nimport triton.language as tl\nimport torch\n\n@triton.autotune(\n    configs=[\n        triton.Config({'BLOCK_SIZE': 1024}, num_warps=4),\n        triton.Config({'BLOCK_SIZE': 2048}, num_warps=8),\n    ],\n    key=['n_elements'],\n)\n@triton.jit\ndef fused_gelu_dropout_kernel(\n    x_ptr, out_ptr, n_elements, p, seed,\n    BLOCK_SIZE: tl.constexpr,\n):\n    pid = tl.program_id(0)\n    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n    mask = offsets < n_elements\n\n    x = tl.load(x_ptr + offsets, mask=mask)\n    x_fp32 = x.to(tl.float32)\n    x = (0.5 * x_fp32 * (1.0 + tl.math.erf(x_fp32 * 0.7071067811865476))).to(x.dtype)\n\n    random = tl.rand(seed, offsets)\n    x = tl.where(random > p, x / (1.0 - p), 0.0)\n\n    tl.store(out_ptr + offsets, x, mask=mask)\n\n\ndef fused_gelu_dropout_triton(x: torch.Tensor, p: float = 0.1) -> torch.Tensor:\n    n_elements = x.numel()\n    out = torch.empty_like(x)\n    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)\n    seed = (x.data_ptr() % (2**31)) ^ n_elements  # sync-free seed\n    fused_gelu_dropout_kernel[grid](x, out, n_elements, p, seed)\n    return out\n\n\n# --- Fixed contract (companion scripts rely on these names) ---\nkernel_fn = fused_gelu_dropout_triton\n\ndef reference_fn(x, p=0.1):\n    torch.manual_seed((x.data_ptr() % (2**31)) ^ x.numel())\n    return torch.nn.functional.dropout(\n        torch.nn.functional.gelu(x), p, training=True\n    )\n\ndef get_inputs():\n    return [torch.randn(128 * 1024 * 1024, device=\"cuda\")]\n```\n\nFor more patterns (SiLU+mul, RMSNorm, linear+GELU, add+LayerNorm), see [references/patterns-fusion.md](references/patterns-fusion.md). For GEMM patterns, see [references/patterns-gemm.md](references/patterns-gemm.md).\n\n### Phase 4: Verify Correctness\n\nRun the companion verification script:\n\n```bash\npython scripts/verify_kernel.py {output_dir}/kernel.py --rtol 1e-3 --atol 1e-3\n```\n\nOutput:\n```json\n{\"correct\": true, \"max_abs_diff\": 1.2e-7, \"max_rel_diff\": 3.4e-6, \"details\": \"...\"}\n```\n\n**Stop if `correct: false`.** Fix the kernel before benchmarking.\n\n**Tolerance guide:**\n\n| Dtype | rtol | atol | Notes |\n|-------|------|------|-------|\n| float16 | 1e-3 | 1e-3 | |\n| bfloat16 | 1e-2 | 1e-2 | |\n| float32 | 1e-5 | 1e-5 | Element-wise ops |\n| float32 (matmul) | 1e-2 | 1e-1 | TF32 accumulation order differs between Triton tiles and cuBLAS |\n\n### Phase 5: Benchmark Performance (optional)\n\nOnly benchmark if the user explicitly requests performance numbers. Skip this phase for correctness-focused requests.\n\n```bash\npython scripts/benchmark_kernel.py {output_dir}/kernel.py\n```\n\nOutput:\n```json\n{\"kernel_time_ms\": 0.45, \"reference_time_ms\": 1.23, \"speedup\": 2.73, \"warmup_iters\": 10, \"benchmark_iters\": 40}\n```\n\n## References (consult only when stuck)\n\nThe skeletons and principles above cover element-wise, reduction, matmul, and fusion kernels. **Do NOT read reference files for these common patterns.**\n\nOnly consult `references/` when:\n- Implementing **uncommon patterns** (grouped GEMM, TMA, persistent matmul, extern functions)\n- **Debugging** a compile error or incorrect result not covered by the error table below\n- Needing **API details** for an unfamiliar `tl.*` operation\n\n**How to search:** Grep for your keyword across `references/`. Read only the file Grep points to.\n\n| File | When to use |\n|---|---|\n| `references/api-core.md` | Unfamiliar `triton.autotune` / `triton.Config` options |\n| `references/api-language.md` | Unfamiliar `tl.*` operations |\n| `references/patterns-gemm.md` | Grouped GEMM, persistent matmul, TMA, MX formats |\n| `references/patterns-advanced.md` | Flash attention details, backward passes, libdevice |\n| `references/troubleshooting.md` | Debug ops, interpreter mode, env vars |\n\n## Error Handling and Troubleshooting\n\n### Common Errors\n\n| Error / Symptom | Cause | Fix |\n|---------|-------|-----|\n| \"Type mismatch, store Float32 to Float16\" | Missing `.to(x.dtype)` before store | Cast fp32 result back |\n| `BLOCK_SIZE is not a constexpr` | Block size passed as runtime value | Add `: tl.constexpr` annotation |\n| `shape mismatch` in binary op | Tensor shapes don't broadcast | Check with `tl.static_print`; use `[:, None]` / `[None, :]` |\n| Large diffs everywhere | Wrong dtype in `tl.load` | Check load dtype matches input |\n| Matmul 3-8x slower than expected | `input_precision=\"ieee\"` on `tl.dot` | Remove it; use TF32 default. Ensure `reference_fn` also uses TF32 |\n| Matmul ~0.01-0.1 abs diff vs reference | TF32 vs IEEE mismatch | Use same precision in both kernel and reference (TF32 for both) |\n| Diffs at boundaries | Missing mask | Add mask to all load/store ops |\n| Random diffs | Race condition | Check atomics and ordering |\n| NaN/Inf | Division by zero or fp16 overflow | Guard with epsilon; use `tl.float32` accumulator |\n| `grid must be a tuple` | Grid lambda returns int, not tuple | Return `(value,)` with trailing comma |\n| `expected constexpr` in `tl.arange` | Non-constexpr argument | Both args of `tl.arange(start, end)` must be constexpr |\n| `triton.OutOfResources` | Register/shared memory pressure | Reduce BLOCK_SIZE or `num_stages` |\n| Kernel not updating after edit | Stale compilation cache | `rm -rf ~/.triton/cache/` |\n| Mismatched results vs PyTorch | C integer division semantics | Triton uses truncation; see `references/concepts-semantics.md` |\n\nFor extended error table, interpreter mode issues, and environment variables, see [references/troubleshooting.md](references/troubleshooting.md).\n\n### When to Abort\n\nStop and report failure if:\n1. **Not a good fit** -- Pure matmul or complex control flow (Phase 0 should catch this).\n2. **Verification fails after 3 attempts** -- Numerical issues too severe to fix.\n3. **No speedup** -- Reference is already well-optimized (cuBLAS, cuDNN).\n4. **Hardware mismatch** -- Target GPU not available for testing.\n",".claude/skills/kernel-tileir-optimization/SKILL.md":"---\nname: kernel-tileir-optimization\ndescription: >\n  Optimize existing Triton kernels for NVIDIA TileIR backend on Blackwell GPUs (sm_100+).\n  Adds TileIR-specific autotune configs: occupancy, num_ctas, TMA descriptors. Covers\n  kernel classification (dot-related, norm-like, elementwise, reduction), type-specific\n  transformations, and PTX-vs-TileIR benchmarking. Triggered by: \"optimize for TileIR\",\n  \"add TileIR configs\", \"Blackwell optimization\", \"TMA descriptors\", \"2CTA mode\",\n  \"occupancy tuning\". Kernels use standard `import triton`; TileIR activates via\n  ENABLE_TILE=1 when nvtriton is installed.\ncompatibility: >\n  Requires Blackwell GPU (sm_100+) for TileIR execution. Supports development on any GPU.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Triton TileIR Optimization\n\nOptimize EXISTING Triton kernels for NVIDIA's TileIR backend on Blackwell GPUs.\nThis skill does NOT write kernels from scratch -- that is the Triton Specialist's job.\n\n## Principles\n\n### TileIR vs PTX Backend\n\nTileIR is NVIDIA's compiler backend for Triton that generates optimized CUDA code\nusing CGA-level (Cooperative Grid Array) tile representations. Critical differences:\n\n| Parameter | PTX Backend | TileIR Backend |\n|-----------|-------------|----------------|\n| `num_warps` | Strict directive | **Ignored** (compiler decides) |\n| `num_stages` | Strict directive | Cost hint (compiler optimizes) |\n| `occupancy` | Not available | **Critical** tuning param (1-32) |\n| `num_ctas` | Limited | 2CTA mode for Blackwell |\n| Block sizes | Smaller often better | Larger often better |\n| TMA | Not available | Required for dot kernels |\n\n**Key implication**: Do not tune `num_warps` for TileIR -- focus on `occupancy` instead.\n\n### Triton Package Landscape\n\nThree packages share `import triton`:\n\n| Package | Source | Use Case |\n|---------|--------|----------|\n| `pytorch-triton` | PyTorch wheel | `torch.compile`, standard kernels |\n| `triton` | OpenAI PyPI | Official Triton from triton-lang.org |\n| nvtriton | [Triton-to-tile-IR](https://github.com/triton-lang/Triton-to-tile-IR) | TileIR backend for Blackwell |\n\nOnly one triton package should be installed at a time. \"Converting to TileIR\" means\nadding TileIR-specific configs, NOT changing imports. TileIR activates via `ENABLE_TILE=1`.\n\n### When TileIR Applies\n\nTileIR targets Blackwell (sm_100+). Without nvtriton or Blackwell hardware, the\nspecialist still adds TileIR-optimized configs that standard triton safely ignores,\nenabling future deployment.\n\n**Expected speedups** (with nvtriton on Blackwell):\n\n| Kernel Type | Speedup | Key Lever |\n|-------------|---------|-----------|\n| Dot-Related (GEMM, Attention) | 1.2-2.0x | TMA + 2CTA |\n| Norm-Like (LayerNorm, Softmax) | 2.0-5.0x | High occupancy |\n| Element-Wise (ReLU, Add, Exp) | 1.5-3.0x | Occupancy + num_stages |\n| Reduction (Sum, Mean, Max) | 1.8-4.0x | High occupancy |\n\n## Workflow\n\nFive-phase workflow: compatibility, classify, transform, validate, benchmark.\n\n### Phase 1: Compatibility Test (ENABLE_TILE=0)\n\nVerify the kernel works in PTX mode before applying TileIR optimizations.\n\n```bash\npython scripts/tileir_check.py\n```\n\nThen use the **kernel-triton-writing** skill's `verify_kernel.py` to verify with `ENABLE_TILE=0`:\n\n```bash\npython scripts/verify_kernel.py --kernel path/to/kernel.py --reference 'torch reference' --shapes '{\"x\": [32, 512, 4096]}' --dtypes '{\"x\": \"bfloat16\"}'\n```\n\n### Phase 2: Classify Kernel\n\nDetermine kernel type to select the optimization strategy.\n\n```bash\npython scripts/classify_kernel.py --file kernel.py\n```\n\nClassification decision tree:\n\n```\nContains tl.dot()?\n  YES --> dot-related: TMA + 2CTA + occupancy + larger blocks\n  NO  --> Has reduction + normalization?\n            YES --> norm-like: high occupancy (2, 4) + num_warps (4, 8)\n            NO  --> Point-wise only?\n                      YES --> element-wise: occupancy (1-16) + num_stages (2-4)\n                      NO  --> reduction: high occupancy + num_warps\n```\n\n### Phase 3: Apply Transformations\n\nClassify and apply optimizations in one step:\n\n```bash\npython scripts/classify_kernel.py --file kernel.py --apply-optimizations\n```\n\nOutput JSON includes `optimized_code` and `changes_applied` fields.\n\n**Type-specific transformations:**\n\n**Dot-related** (highest priority):\n1. Convert `tl.load`/`tl.store` to TMA descriptors (MANDATORY). See `references/tma-conversion.md`.\n2. Add 2CTA configs (`num_ctas=2`) with SM oversubscription guard in pre-hook.\n3. Add occupancy (1, 2, 4) and extended num_stages (4, 6).\n4. Use larger block sizes (256x256, 256x128).\n\n**Norm-like** (LayerNorm, Softmax, RMSNorm):\n- Add occupancy (2, 4), num_warps (4, 8). No TMA needed.\n\n**Element-wise** (ReLU, GELU, Add, Mul, Exp):\n- Add occupancy (1, 2, 4, 16), num_stages (2, 3, 4). Include extreme configs for small inputs.\n\n**Reduction** (Sum, Mean, Max):\n- Same strategy as norm-like: high occupancy (2, 4), num_warps (4, 8).\n\nGate TileIR-specific configs for sm_100+:\n\n```python\nimport torch\n\ndef get_configs_with_gating(pre_hook=None):\n    configs = get_baseline_configs()\n    if torch.cuda.is_available() and torch.cuda.get_device_capability()[0] >= 10:\n        configs.extend(get_tileir_specific_configs(pre_hook))\n    return configs\n```\n\nSee `references/config-templates.md` for complete config templates per kernel type.\n\n### Phase 4: TileIR Validation (ENABLE_TILE=1)\n\nUse the **kernel-triton-writing** skill's `verify_kernel.py` to verify the optimized kernel with TileIR backend:\n\n```bash\npython scripts/verify_kernel.py --kernel path/to/optimized_kernel.py --reference 'torch reference' --shapes '{\"x\": [32, 512, 4096]}' --dtypes '{\"x\": \"bfloat16\"}'\n```\n\nSet `ENABLE_TILE=1` before running. Check: numerical correctness, no compilation errors,\nTMA/2CTA patterns compile successfully.\n\n### Phase 5: Benchmark\n\nUse `triton.testing.do_bench()` (as documented in the **perf-workload-profiling** skill) to compare PTX (`ENABLE_TILE=0`) vs TileIR (`ENABLE_TILE=1`).\n\nBenchmark across multiple input sizes (128, 1024, 8192) -- performance varies by size.\n\n## Scripts\n\n### tileir_check.py\n\nCheck TileIR availability (nvtriton, ENABLE_TILE, Blackwell GPU):\n\n```bash\npython scripts/tileir_check.py\n```\n\nReturns JSON: `nvtriton_installed`, `tileir_active`, `blackwell_gpu`, `gpu_capability`, `recommendation`.\n\n### classify_kernel.py\n\nClassify kernel type and optionally apply TileIR optimizations:\n\n```bash\n# Classify only\npython scripts/classify_kernel.py --file kernel.py\n\n# Classify + apply optimizations\npython scripts/classify_kernel.py --file kernel.py --apply-optimizations\n\n# From inline code\npython scripts/classify_kernel.py --code '<kernel_code>'\n```\n\nReturns JSON: `classification`, `confidence`, `indicators`, `recommendations`.\nWith `--apply-optimizations`: adds `optimized_code` and `changes_applied`.\n\n## Error Handling\n\n### Common Pitfalls\n\n**TMA descriptor errors** (dot-related kernels):\n- Always pass `pre_hook=tma_set_block_size_hook` to config generation -- without it,\n  TMA descriptors keep dummy block sizes, causing runtime errors or wrong results.\n- For GEMM: pass `b.T.contiguous()` in wrapper and use `tl.dot(a, b.T, accumulator)`\n  in kernel. Transposition mismatch produces incorrect results silently.\n\n**2CTA oversubscription**:\n- Adjust SM count in pre-hook when using `num_ctas=2`:\n  ```python\n  if \"NUM_SMS\" in nargs and \"NUM_CTAS\" in nargs:\n      nargs[\"NUM_SMS\"] = nargs[\"NUM_SMS\"] // nargs[\"NUM_CTAS\"]\n  ```\n\n**Config function signatures**:\n- ALL config helper functions MUST accept `pre_hook=None`, even if unused.\n  Without it: `TypeError: get_autotune_configs() takes 0 positional arguments`.\n\n**Hardware gating**:\n- Gate TileIR configs with `torch.cuda.get_device_capability()[0] >= 10`.\n  TMA/2CTA on pre-Blackwell GPUs causes runtime crashes.\n\n**API availability**:\n- Use `1.0 / (1.0 + tl.exp(-x))` instead of `tl.sigmoid(x)` -- not available in\n  all Triton versions including some nvtriton builds.\n\n**Performance tuning**:\n- Do not over-tune `num_warps` -- TileIR ignores it. Focus on `occupancy`.\n- Use larger block sizes (256x256, 256x128) for TileIR, not PTX-tuned small blocks.\n- Benchmark across small/medium/large inputs; one-size configs underperform.\n- For exp/log heavy kernels, enable approximate math:\n  ```bash\n  export TILEIR_ENABLE_APPROX=1\n  export TILEIR_ENABLE_FTZ=1\n  ```\n\n### When to Abort\n\nStop and report if:\n\n1. **No triton installed** -- cannot proceed.\n2. **Compatibility test fails** -- kernel has syntax/runtime errors before optimization.\n3. **TileIR validation fails** -- optimized kernel produces wrong results.\n4. **No speedup** -- TileIR version is slower than PTX baseline (with nvtriton).\n5. **Not Blackwell GPU** -- still add configs for future deployment, but skip\n   ENABLE_TILE testing and benchmarking.\n\n### Output Format\n\nAfter optimization, return:\n\n```\n## TileIR Optimization: kernel_name\n\n### Classification\n- Kernel type: [dot-related | norm-like | element-wise | reduction]\n- Strategy: [TMA + 2CTA | High occupancy | Occupancy + num_stages]\n\n### Compatibility Check (ENABLE_TILE=0)\n[PASSED | FAILED] — Max difference: X.Xe-Y\n\n### Transformations Applied\n- [List of transformations]\n\n### TileIR Validation (ENABLE_TILE=1)\n[PASSED | FAILED] — Max difference: X.Xe-Y\n\n### Benchmark Comparison\n| Backend | Time (ms) | Speedup |\n|---------|-----------|---------|\n| PTX (ENABLE_TILE=0) | X.XXX | 1.0x |\n| TileIR (ENABLE_TILE=1) | X.XXX | Y.Yx |\n\n### Output\nFile: kernel_name_tileir.py\n```\n",".claude/skills/perf-instability-cookbook/SKILL.md":"# perf-instability-cookbook\n\nPath: `.claude/skills/perf-instability-cookbook/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/SKILL.md)",".claude/skills/perf-nsight-compute-analysis/SKILL.md":"# perf-nsight-compute-analysis\n\nPath: `.claude/skills/perf-nsight-compute-analysis/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/SKILL.md)",".claude/skills/perf-nsight-systems/SKILL.md":"# perf-nsight-systems\n\nPath: `.claude/skills/perf-nsight-systems/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/SKILL.md)",".claude/skills/perf-optimization-casebook/SKILL.md":"# perf-optimization-casebook\n\nPath: `.claude/skills/perf-optimization-casebook/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/SKILL.md)",".claude/skills/perf-optimization/SKILL.md":"# perf-optimization\n\nPath: `.claude/skills/perf-optimization/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization/SKILL.md)",".claude/skills/perf-regression-cookbook/SKILL.md":"# perf-regression-cookbook\n\nPath: `.claude/skills/perf-regression-cookbook/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/SKILL.md)",".claude/skills/perf-torch-cuda-graphs/SKILL.md":"# perf-torch-cuda-graphs\n\nPath: `.claude/skills/perf-torch-cuda-graphs/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/SKILL.md)",".claude/skills/perf-torch-sync-free/SKILL.md":"# perf-torch-sync-free\n\nPath: `.claude/skills/perf-torch-sync-free/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/SKILL.md)",".claude/skills/perf-workload-profiling/SKILL.md":"# perf-workload-profiling\n\nPath: `.claude/skills/perf-workload-profiling/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/SKILL.md)",".claude/skills/trtllm-case-executor/SKILL.md":"# trtllm-case-executor\n\nPath: `.claude/skills/trtllm-case-executor/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-case-executor/SKILL.md)",".claude/skills/trtllm-code-contribution/SKILL.md":"# trtllm-code-contribution\n\nPath: `.claude/skills/trtllm-code-contribution/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-code-contribution/SKILL.md)",".claude/skills/trtllm-codebase-exploration/SKILL.md":"# trtllm-codebase-exploration\n\nPath: `.claude/skills/trtllm-codebase-exploration/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-codebase-exploration/SKILL.md)",".claude/skills/trtllm-flashinfer-upgrade/SKILL.md":"# trtllm-flashinfer-upgrade\n\nPath: `.claude/skills/trtllm-flashinfer-upgrade/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-flashinfer-upgrade/SKILL.md)",".claude/skills/trtllm-model-onboard-multimodal/SKILL.md":"# trtllm-model-onboard-multimodal\n\nPath: `.claude/skills/trtllm-model-onboard-multimodal/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-model-onboard-multimodal/SKILL.md)",".claude/skills/trtllm-moe-develop/SKILL.md":"# trtllm-moe-develop\n\nPath: `.claude/skills/trtllm-moe-develop/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-moe-develop/SKILL.md)",".claude/skills/trtllm-serve-config-guide/SKILL.md":"# trtllm-serve-config-guide\n\nPath: `.claude/skills/trtllm-serve-config-guide/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-serve-config-guide/SKILL.md)",".claude/skills/trtllm-test-script-builder/SKILL.md":"# trtllm-test-script-builder\n\nPath: `.claude/skills/trtllm-test-script-builder/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/SKILL.md)",".claude/skills/trtllm-test-specialist/SKILL.md":"# trtllm-test-specialist\n\nPath: `.claude/skills/trtllm-test-specialist/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/SKILL.md)",".claude/skills/trtllm-visualgen-bench/SKILL.md":"# trtllm-visualgen-bench\n\nPath: `.claude/skills/trtllm-visualgen-bench/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/SKILL.md)",".claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md":"# trtllm-visualgen-lpips-golden-refresh\n\nPath: `.claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md)",".codex/skills/visual-gen-component-test/SKILL.md":"# visual-gen-component-test\n\nPath: `.codex/skills/visual-gen-component-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/skills/visual-gen-component-test/SKILL.md)","agent-flow/.claude/skills/perf-analyze/SKILL.md":"# perf-analyze\n\nPath: `agent-flow/.claude/skills/perf-analyze/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/.claude/skills/perf-analyze/SKILL.md)","agent-flow/.claude/skills/perf-optimize/SKILL.md":"# perf-optimize\n\nPath: `agent-flow/.claude/skills/perf-optimize/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/.claude/skills/perf-optimize/SKILL.md)","AGENTS.md":"# AI Agent Protocol & Instructions\n\nPath: `AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/AGENTS.md)","CLAUDE.md":"# Claude Agent Guidelines & System Prompt\n\nPath: `CLAUDE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/CLAUDE.md)",".codex/AGENTS.md":"# .codex - AI Agent Protocol & Instructions\n\nPath: `.codex/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/AGENTS.md)",".codex/config.toml":"# Codex Directive: config.toml\n\nPath: `.codex/config.toml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/config.toml)","agent-flow/AGENTS.md":"# agent-flow - AI Agent Protocol & Instructions\n\nPath: `agent-flow/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/AGENTS.md)","agent-flow/CLAUDE.md":"# agent-flow - Claude Agent Guidelines & System Prompt\n\nPath: `agent-flow/CLAUDE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/CLAUDE.md)","tests/AGENTS.md":"# tests - AI Agent Protocol & Instructions\n\nPath: `tests/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/AGENTS.md)","tests/CLAUDE.md":"# tests - Claude Agent Guidelines & System Prompt\n\nPath: `tests/CLAUDE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/CLAUDE.md)",".claude/agents/exec-compile-specialist.md":"# Subagent: exec-compile-specialist\n\nPath: `.claude/agents/exec-compile-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-compile-specialist.md)",".claude/agents/exec-local-slurm.md":"# Subagent: exec-local-slurm\n\nPath: `.claude/agents/exec-local-slurm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-local-slurm.md)",".claude/agents/exec-remote-slurm.md":"# Subagent: exec-remote-slurm\n\nPath: `.claude/agents/exec-remote-slurm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-remote-slurm.md)",".claude/agents/kernel-cuda-specialist.md":"# Subagent: kernel-cuda-specialist\n\nPath: `.claude/agents/kernel-cuda-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-cuda-specialist.md)",".claude/agents/kernel-cute-specialist.md":"# Subagent: kernel-cute-specialist\n\nPath: `.claude/agents/kernel-cute-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-cute-specialist.md)",".claude/agents/kernel-tileir-specialist.md":"# Subagent: kernel-tileir-specialist\n\nPath: `.claude/agents/kernel-tileir-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-tileir-specialist.md)",".claude/agents/kernel-triton-specialist.md":"# Subagent: kernel-triton-specialist\n\nPath: `.claude/agents/kernel-triton-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-triton-specialist.md)",".claude/agents/modeling-test-specialist.md":"# Subagent: modeling-test-specialist\n\nPath: `.claude/agents/modeling-test-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/modeling-test-specialist.md)",".claude/agents/perf-profiling-specialist.md":"# Subagent: perf-profiling-specialist\n\nPath: `.claude/agents/perf-profiling-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-profiling-specialist.md)",".claude/agents/perf-test-sync.md":"# Subagent: perf-test-sync\n\nPath: `.claude/agents/perf-test-sync.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-test-sync.md)",".claude/agents/perf-torch-cuda-graph-specialist.md":"# Subagent: perf-torch-cuda-graph-specialist\n\nPath: `.claude/agents/perf-torch-cuda-graph-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-torch-cuda-graph-specialist.md)",".codex/agents/onboard_update_reviewer.toml":"# Codex Directive: onboard_update_reviewer.toml\n\nPath: `.codex/agents/onboard_update_reviewer.toml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/agents/onboard_update_reviewer.toml)","triton_backend/tools/dataset/mini_cnn_eval.json":"# Dataset Skill\n\nPath: `triton_backend/tools/dataset/mini_cnn_eval.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/triton_backend/tools/dataset/mini_cnn_eval.json)","triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json":"# Dataset Skill\n\nPath: `triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json)",".claude/skills/exec-remote-slurm/scripts/build.sh":"# Claude Directive: build.sh\n\nPath: `.claude/skills/exec-remote-slurm/scripts/build.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-remote-slurm/scripts/build.sh)",".claude/skills/exec-remote-slurm/scripts/build.slurm":"# Claude Directive: build.slurm\n\nPath: `.claude/skills/exec-remote-slurm/scripts/build.slurm`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-remote-slurm/scripts/build.slurm)",".claude/skills/exec-slurm-compile/scripts/compile.sh":"# Claude Directive: compile.sh\n\nPath: `.claude/skills/exec-slurm-compile/scripts/compile.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/compile.sh)",".claude/skills/exec-slurm-compile/scripts/compile.slurm":"# Claude Directive: compile.slurm\n\nPath: `.claude/skills/exec-slurm-compile/scripts/compile.slurm`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/compile.slurm)",".claude/skills/exec-slurm-compile/scripts/enroot-import":"# Claude Directive: enroot-import\n\nPath: `.claude/skills/exec-slurm-compile/scripts/enroot-import`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/enroot-import)",".claude/skills/exec-slurm-compile/scripts/submit_compile.sh":"# Claude Directive: submit_compile.sh\n\nPath: `.claude/skills/exec-slurm-compile/scripts/submit_compile.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/submit_compile.sh)",".claude/skills/kernel-cute-writing/references/api-arch.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-arch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-arch.md)",".claude/skills/kernel-cute-writing/references/api-core.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-core.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-core.md)",".claude/skills/kernel-cute-writing/references/api-nvgpu.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-nvgpu.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-nvgpu.md)",".claude/skills/kernel-cute-writing/references/api-runtime-utils.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-runtime-utils.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-runtime-utils.md)",".claude/skills/kernel-cute-writing/references/concepts-architecture.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-architecture.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-architecture.md)",".claude/skills/kernel-cute-writing/references/concepts-layouts.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-layouts.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-layouts.md)",".claude/skills/kernel-cute-writing/references/concepts-mma.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-mma.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-mma.md)",".claude/skills/kernel-cute-writing/references/concepts-tensors.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-tensors.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-tensors.md)",".claude/skills/kernel-cute-writing/references/patterns-compilation.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-compilation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-compilation.md)",".claude/skills/kernel-cute-writing/references/patterns-elementwise.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-elementwise.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-elementwise.md)",".claude/skills/kernel-cute-writing/references/patterns-gemm.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-gemm.md)",".claude/skills/kernel-cute-writing/references/patterns-getting-started.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-getting-started.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-getting-started.md)",".claude/skills/kernel-cute-writing/references/patterns-memory.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-memory.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-memory.md)",".claude/skills/kernel-cute-writing/references/patterns-pipeline.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-pipeline.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-pipeline.md)",".claude/skills/kernel-cute-writing/references/patterns-reduction.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-reduction.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-reduction.md)",".claude/skills/kernel-cute-writing/references/troubleshooting.md":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/troubleshooting.md)",".claude/skills/kernel-cute-writing/scripts/__init__.py":"# Claude Directive: __init__.py\n\nPath: `.claude/skills/kernel-cute-writing/scripts/__init__.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/__init__.py)",".claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py":"# Claude Directive: benchmark_kernel.py\n\nPath: `.claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py)",".claude/skills/kernel-cute-writing/scripts/verify_kernel.py":"# Claude Directive: verify_kernel.py\n\nPath: `.claude/skills/kernel-cute-writing/scripts/verify_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/verify_kernel.py)",".claude/skills/kernel-tileir-optimization/references/config-templates.md":"# References Skill\n\nPath: `.claude/skills/kernel-tileir-optimization/references/config-templates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/config-templates.md)",".claude/skills/kernel-tileir-optimization/references/tma-conversion.md":"# References Skill\n\nPath: `.claude/skills/kernel-tileir-optimization/references/tma-conversion.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/tma-conversion.md)",".claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py":"# Claude Directive: classify_kernel.py\n\nPath: `.claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py)",".claude/skills/kernel-tileir-optimization/scripts/tileir_check.py":"# Claude Directive: tileir_check.py\n\nPath: `.claude/skills/kernel-tileir-optimization/scripts/tileir_check.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/scripts/tileir_check.py)",".claude/skills/kernel-triton-writing/references/api-core.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/api-core.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/api-core.md)",".claude/skills/kernel-triton-writing/references/api-language.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/api-language.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/api-language.md)",".claude/skills/kernel-triton-writing/references/concepts-semantics.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/concepts-semantics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/concepts-semantics.md)",".claude/skills/kernel-triton-writing/references/operator-routing.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/operator-routing.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/operator-routing.md)",".claude/skills/kernel-triton-writing/references/patterns-advanced.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-advanced.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-advanced.md)",".claude/skills/kernel-triton-writing/references/patterns-basic.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-basic.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-basic.md)",".claude/skills/kernel-triton-writing/references/patterns-fusion.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-fusion.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-fusion.md)",".claude/skills/kernel-triton-writing/references/patterns-gemm.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-gemm.md)",".claude/skills/kernel-triton-writing/references/troubleshooting.md":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/troubleshooting.md)",".claude/skills/kernel-triton-writing/scripts/__init__.py":"# Claude Directive: __init__.py\n\nPath: `.claude/skills/kernel-triton-writing/scripts/__init__.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/__init__.py)",".claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py":"# Claude Directive: benchmark_kernel.py\n\nPath: `.claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py)",".claude/skills/kernel-triton-writing/scripts/verify_kernel.py":"# Claude Directive: verify_kernel.py\n\nPath: `.claude/skills/kernel-triton-writing/scripts/verify_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/verify_kernel.py)",".claude/skills/perf-host-analysis/references/examples.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/examples.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/examples.md)",".claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md)",".claude/skills/perf-host-analysis/references/kernel-level-analysis.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/kernel-level-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/kernel-level-analysis.md)",".claude/skills/perf-host-analysis/references/metrics.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/metrics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/metrics.md)",".claude/skills/perf-host-analysis/references/nsys-schema.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/nsys-schema.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/nsys-schema.md)",".claude/skills/perf-host-analysis/references/output-format.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/output-format.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/output-format.md)",".claude/skills/perf-host-analysis/references/phase-classification.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/phase-classification.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/phase-classification.md)",".claude/skills/perf-host-analysis/references/thresholds.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/thresholds.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/thresholds.md)",".claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md)",".claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py":"# Claude Directive: analyze_host_overhead.py\n\nPath: `.claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py)",".claude/skills/perf-host-analysis/scripts/detect_host_overhead.py":"# Claude Directive: detect_host_overhead.py\n\nPath: `.claude/skills/perf-host-analysis/scripts/detect_host_overhead.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/scripts/detect_host_overhead.py)",".claude/skills/perf-host-optimization/references/communication-patterns.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/communication-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/communication-patterns.md)",".claude/skills/perf-host-optimization/references/examples.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/examples.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/examples.md)",".claude/skills/perf-host-optimization/references/hot-path-files.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/hot-path-files.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/hot-path-files.md)",".claude/skills/perf-host-optimization/references/hotspot-classification.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/hotspot-classification.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/hotspot-classification.md)",".claude/skills/perf-host-optimization/references/optimization-patterns.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/optimization-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/optimization-patterns.md)",".claude/skills/perf-host-optimization/references/optimization-strategy.md":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/optimization-strategy.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/optimization-strategy.md)",".claude/skills/perf-instability-cookbook/data/aliases.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/aliases.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/aliases.yaml)",".claude/skills/perf-instability-cookbook/data/modules.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/modules.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/modules.yaml)",".claude/skills/perf-instability-cookbook/data/patterns.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/patterns.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/patterns.yaml)",".claude/skills/perf-instability-cookbook/data/pending.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/pending.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/pending.yaml)",".claude/skills/perf-instability-cookbook/data/removed.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/removed.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/removed.yaml)",".claude/skills/perf-instability-cookbook/data/tags.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/tags.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/tags.yaml)",".claude/skills/perf-instability-cookbook/references/case-template.md":"# References Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/case-template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/case-template.md)",".claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md)",".claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md)",".claude/skills/perf-nsight-compute-analysis/references/cli-reference.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/cli-reference.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/cli-reference.md)",".claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md)",".claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md)",".claude/skills/perf-nsight-compute-analysis/references/python-report-api.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/python-report-api.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/python-report-api.md)",".claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md)",".claude/skills/perf-nsight-compute-analysis/references/sections-guide.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/sections-guide.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/sections-guide.md)",".claude/skills/perf-nsight-systems/references/app-preparation.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/app-preparation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/app-preparation.md)",".claude/skills/perf-nsight-systems/references/cli-post-collection.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/cli-post-collection.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/cli-post-collection.md)",".claude/skills/perf-nsight-systems/references/cli-profiling.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/cli-profiling.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/cli-profiling.md)",".claude/skills/perf-nsight-systems/references/expert-systems.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/expert-systems.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/expert-systems.md)",".claude/skills/perf-nsight-systems/references/nvtx-analysis.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/nvtx-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/nvtx-analysis.md)",".claude/skills/perf-nsight-systems/references/recipes-dl.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/recipes-dl.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/recipes-dl.md)",".claude/skills/perf-nsight-systems/references/stats-reports.md":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/stats-reports.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/stats-reports.md)",".claude/skills/perf-optimization-casebook/data/aliases.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-optimization-casebook/data/aliases.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/aliases.yaml)",".claude/skills/perf-optimization-casebook/data/patterns.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-optimization-casebook/data/patterns.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/patterns.yaml)",".claude/skills/perf-optimization-casebook/data/tags.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-optimization-casebook/data/tags.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/tags.yaml)",".claude/skills/perf-optimization-casebook/references/case-template.md":"# References Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/case-template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/case-template.md)",".claude/skills/perf-regression-cookbook/data/aliases.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/aliases.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/aliases.yaml)",".claude/skills/perf-regression-cookbook/data/modules.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/modules.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/modules.yaml)",".claude/skills/perf-regression-cookbook/data/patterns.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/patterns.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/patterns.yaml)",".claude/skills/perf-regression-cookbook/data/pending.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/pending.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/pending.yaml)",".claude/skills/perf-regression-cookbook/data/removed.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/removed.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/removed.yaml)",".claude/skills/perf-regression-cookbook/data/tags.yaml":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/tags.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/tags.yaml)",".claude/skills/perf-regression-cookbook/references/case-template.md":"# References Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/case-template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/case-template.md)",".claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md)",".claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md)",".claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md)",".claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md)",".claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md)",".claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py":"# Claude Directive: verify_workload.py\n\nPath: `.claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py)",".claude/skills/perf-torch-sync-free/references/sync-patterns.md":"# References Skill\n\nPath: `.claude/skills/perf-torch-sync-free/references/sync-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/references/sync-patterns.md)",".claude/skills/perf-torch-sync-free/scripts/verify_workload.py":"# Claude Directive: verify_workload.py\n\nPath: `.claude/skills/perf-torch-sync-free/scripts/verify_workload.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/scripts/verify_workload.py)",".claude/skills/perf-workload-profiling/references/benchmarking-patterns.md":"# References Skill\n\nPath: `.claude/skills/perf-workload-profiling/references/benchmarking-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/benchmarking-patterns.md)",".claude/skills/perf-workload-profiling/references/nvtx-api.md":"# References Skill\n\nPath: `.claude/skills/perf-workload-profiling/references/nvtx-api.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/nvtx-api.md)",".claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md":"# References Skill\n\nPath: `.claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md)",".claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh":"# Claude Directive: detect_slurm_env.sh\n\nPath: `.claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh)",".claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md":"# References Skill\n\nPath: `.claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md)",".claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md":"# References Skill\n\nPath: `.claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md)",".claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md":"# References Skill\n\nPath: `.claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md)",".claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py":"# Claude Directive: build_slurm_script.py\n\nPath: `.claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py)",".claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh":"# Claude Directive: slurm_run_custom.sh\n\nPath: `.claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh)",".claude/skills/trtllm-test-specialist/references/agg_config_template.yaml":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/agg_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/agg_config_template.yaml)",".claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml)",".claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml)",".claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml":"# Claude Directive: smoke_test_config_template.yml\n\nPath: `.claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml)",".claude/skills/trtllm-test-specialist/references/test_config_template.yaml":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/test_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/test_config_template.yaml)",".claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md)",".claude/skills/trtllm-test-specialist/scripts/build_test_command.py":"# Claude Directive: build_test_command.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/build_test_command.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/build_test_command.py)",".claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py":"# Claude Directive: extract_test_markers.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py)",".claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py":"# Claude Directive: generate_benchmark_config.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py)",".claude/skills/trtllm-test-specialist/scripts/generate_report.py":"# Claude Directive: generate_report.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/generate_report.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/generate_report.py)",".claude/skills/trtllm-test-specialist/scripts/parse_config.py":"# Claude Directive: parse_config.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/parse_config.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/parse_config.py)",".claude/skills/trtllm-visualgen-bench/media/cat_piano.png":"# Claude Directive: cat_piano.png\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/cat_piano.png`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/cat_piano.png)",".claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png":"# Claude Directive: image_with_red_background.png\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png)",".claude/skills/trtllm-visualgen-bench/media/landscape.png":"# Claude Directive: landscape.png\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/landscape.png`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/landscape.png)",".claude/skills/trtllm-visualgen-bench/media/robot_153.jpg":"# Claude Directive: robot_153.jpg\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/robot_153.jpg`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/robot_153.jpg)",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json)",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json)",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json)",".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json)",".claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json":"# Reference Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json)",".claude/skills/trtllm-visualgen-bench/scripts/serve_check.py":"# Claude Directive: serve_check.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/scripts/serve_check.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/scripts/serve_check.py)",".claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py":"# Claude Directive: serve_wait.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py)",".claude/skills/trtllm-visualgen-bench/serve/README.md":"# Serve Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/serve/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/serve/README.md)",".claude/skills/trtllm-visualgen-bench/tests/test_discovery.py":"# Claude Directive: test_discovery.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/tests/test_discovery.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_discovery.py)",".claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py":"# Claude Directive: test_serve_wait.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py)",".claude/skills/trtllm-visualgen-bench/tests/test_workloads.py":"# Claude Directive: test_workloads.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/tests/test_workloads.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_workloads.py)",".claude/skills/trtllm-visualgen-bench/workloads/README.md":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/README.md)",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml)",".claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml)","docs/source/commands/trtllm-serve/render-endpoints.md":"# Command: /render-endpoints\n\nPath: `docs/source/commands/trtllm-serve/render-endpoints.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/docs/source/commands/trtllm-serve/render-endpoints.md)","docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md":"# Command: /run-benchmark-with-trtllm-serve\n\nPath: `docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md)","tensorrt_llm/tools/profiler/host_profile_tools/README.md":"# host_profile_tools Documentation\n\nPath: `tensorrt_llm/tools/profiler/host_profile_tools/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/host_profile_tools/README.md)","tensorrt_llm/tools/profiler/nsys_profile_tools/README.md":"# nsys_profile_tools Documentation\n\nPath: `tensorrt_llm/tools/profiler/nsys_profile_tools/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/nsys_profile_tools/README.md)","tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json":"# Nsys_profile_tools Skill\n\nPath: `tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json)","tests/unittest/api_stability/references/batched_logits_processor.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/batched_logits_processor.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/batched_logits_processor.yaml)","tests/unittest/api_stability/references/calib_config.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/calib_config.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/calib_config.yaml)","tests/unittest/api_stability/references/completion_output.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/completion_output.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/completion_output.yaml)","tests/unittest/api_stability/references/conversation_params.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/conversation_params.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/conversation_params.yaml)","tests/unittest/api_stability/references/decoding_base_config.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/decoding_base_config.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/decoding_base_config.yaml)","tests/unittest/api_stability/references/guided_decoding_params.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/guided_decoding_params.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/guided_decoding_params.yaml)","tests/unittest/api_stability/references/llm.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/llm.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/llm.yaml)","tests/unittest/api_stability/references/logits_processor.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/logits_processor.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/logits_processor.yaml)","tests/unittest/api_stability/references/quant_config.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/quant_config.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/quant_config.yaml)","tests/unittest/api_stability/references/request_output.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/request_output.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/request_output.yaml)","tests/unittest/api_stability/references/sampling_params.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/sampling_params.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/sampling_params.yaml)","tests/unittest/api_stability/references/trtllm_serve_api.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/trtllm_serve_api.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/trtllm_serve_api.yaml)","tests/unittest/api_stability/references/trtllm_serve_cli.yaml":"# References Skill\n\nPath: `tests/unittest/api_stability/references/trtllm_serve_cli.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/trtllm_serve_cli.yaml)",".claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md)",".claude/skills/perf-host-optimization/references/patterns/gpu-graph.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/gpu-graph.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/gpu-graph.md)",".claude/skills/perf-host-optimization/references/patterns/loop-iteration.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/loop-iteration.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/loop-iteration.md)",".claude/skills/perf-host-optimization/references/patterns/python-overhead.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/python-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/python-overhead.md)",".claude/skills/perf-host-optimization/references/patterns/sync-alloc.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/sync-alloc.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/sync-alloc.md)",".claude/skills/perf-host-optimization/references/patterns/system.md":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/system.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/system.md)",".claude/skills/perf-instability-cookbook/references/communication/index.md":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/index.md)",".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md)",".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md)",".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md)",".claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md)",".claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md)",".claude/skills/perf-instability-cookbook/references/moe/index.md":"# Moe Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/moe/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/moe/index.md)",".claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md":"# Moe Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md)",".claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md)",".claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md)",".claude/skills/perf-optimization-casebook/references/communication/deepep.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/deepep.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/deepep.md)",".claude/skills/perf-optimization-casebook/references/communication/index.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/index.md)",".claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md)",".claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md)",".claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md)",".claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md)",".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/index.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/index.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md)",".claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/index.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/index.md)",".claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md)",".claude/skills/perf-regression-cookbook/references/autotuner/index.md":"# Autotuner Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/autotuner/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/autotuner/index.md)",".claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md":"# Autotuner Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md)",".claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md)",".claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md)",".claude/skills/perf-regression-cookbook/references/communication/index.md":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/index.md)",".claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md)",".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md)",".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md":"# Disagg-orchestrator Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md)",".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md":"# Disagg-orchestrator Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md)",".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md":"# Disagg-orchestrator Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md)",".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md":"# Gemm-and-quantization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md)",".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md":"# Gemm-and-quantization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md)",".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md":"# Gemm-and-quantization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md)",".claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md)",".claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md":"# Kernel-fusion Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md)",".claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md":"# Kernel-fusion Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md)",".claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md":"# Kernel-fusion Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md":"# Kv-cache-transceiver Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md":"# Kv-cache-transceiver Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md)",".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md":"# Kv-cache-transceiver Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md)",".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md)",".claude/skills/perf-regression-cookbook/references/model-definition/index.md":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/index.md)",".claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md)",".claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md)",".claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md)",".claude/skills/perf-regression-cookbook/references/moe/index.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/index.md)",".claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md)",".claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md)",".claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md)",".claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md)",".claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md)",".claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md)",".claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md)",".claude/skills/perf-regression-cookbook/references/sampler/index.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/index.md)",".claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md)",".claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md)",".claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md)",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md)",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md)",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md)",".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md)",".claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md":"# Spec-decode Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md)",".claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md":"# Spec-decode Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md)",".claude/skills/perf-regression-cookbook/references/spec-decode/index.md":"# Spec-decode Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/spec-decode/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/index.md)",".claude/skills/perf-regression-cookbook/references/tokenization/index.md":"# Tokenization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/tokenization/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/tokenization/index.md)",".claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md":"# Tokenization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md)","jenkins/scripts/cbts/cbts/rules/README.md":"# Rules Skill\n\nPath: `jenkins/scripts/cbts/cbts/rules/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/jenkins/scripts/cbts/cbts/rules/README.md)","tests/integration/defs/accuracy/references/acceptance_length.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/acceptance_length.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/acceptance_length.yaml)","tests/integration/defs/accuracy/references/cnn_dailymail.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/cnn_dailymail.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/cnn_dailymail.yaml)","tests/integration/defs/accuracy/references/gpqa_diamond.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/gpqa_diamond.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gpqa_diamond.yaml)","tests/integration/defs/accuracy/references/gsm8k.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/gsm8k.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gsm8k.yaml)","tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml)","tests/integration/defs/accuracy/references/humaneval.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/humaneval.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/humaneval.yaml)","tests/integration/defs/accuracy/references/json_mode_eval.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/json_mode_eval.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/json_mode_eval.yaml)","tests/integration/defs/accuracy/references/longbench_v1.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/longbench_v1.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/longbench_v1.yaml)","tests/integration/defs/accuracy/references/longbench_v2.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/longbench_v2.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/longbench_v2.yaml)","tests/integration/defs/accuracy/references/mmlu.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/mmlu.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/mmlu.yaml)","tests/integration/defs/accuracy/references/mmmu.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/mmmu.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/mmmu.yaml)","tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml)","tests/integration/defs/accuracy/references/videomme.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/videomme.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/videomme.yaml)","tests/integration/defs/accuracy/references/voxpopuli.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/voxpopuli.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/voxpopuli.yaml)","tests/integration/defs/accuracy/references/zero_scrolls.yaml":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/zero_scrolls.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/zero_scrolls.yaml)"},"items":[{"name":"AGENTS.md","path":".codex/AGENTS.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/AGENTS.md","title":".codex - AI Agent Protocol & Instructions","category":"root-instruction","format":"markdown","content":"# .codex - AI Agent Protocol & Instructions\n\nPath: `.codex/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/AGENTS.md)","tokens":300,"sizeBytes":176},{"name":"AGENTS.md","path":"agent-flow/AGENTS.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/AGENTS.md","title":"agent-flow - AI Agent Protocol & Instructions","category":"root-instruction","format":"markdown","content":"# agent-flow - AI Agent Protocol & Instructions\n\nPath: `agent-flow/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/AGENTS.md)","tokens":300,"sizeBytes":188},{"name":"AGENTS.md","path":"AGENTS.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/AGENTS.md","title":"AI Agent Protocol & Instructions","category":"root-instruction","format":"markdown","content":"# AI Agent Protocol & Instructions\n\nPath: `AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/AGENTS.md)","tokens":300,"sizeBytes":153},{"name":"AGENTS.md","path":"tests/AGENTS.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/AGENTS.md","title":"tests - AI Agent Protocol & Instructions","category":"root-instruction","format":"markdown","content":"# tests - AI Agent Protocol & Instructions\n\nPath: `tests/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/AGENTS.md)","tokens":300,"sizeBytes":173},{"name":"SKILL.md","path":".claude/skills/exec-env-check/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-env-check/SKILL.md","title":"exec-env-check","category":"anthropic-skill","format":"markdown","content":"---\nname: exec-env-check\ndescription: >-\n  Check the local execution environment for GPU availability, Docker support,\n  and Slurm access. Returns the execution scenario (`satisfied, local, docker`,\n  `satisfied, local, direct`, `satisfied, slurm, local`, or `not_satisfied`),\n  the number of available GPUs, and the GPU type. On Slurm login nodes\n  without local GPUs, the cluster is identified by delegating the hostname\n  to internal-env-info (hostname-based mode), which owns the\n  hostname → cluster_name patterns; GPU type and gpus_per_node then come\n  from that skill's reference files. If internal-env-info is not\n  installed, the scenario falls back to `not_satisfied` without probing\n  compute nodes via srun.\ntags: [infrastructure, slurm, environment]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# TensorRT-LLM Environment Check\n\nDetect whether the current machine can run a GPU workload locally (Docker) or via Slurm, and report hardware details.\n\n## Input\n\n| Field | Description | Required |\n|-------|-------------|----------|\n| `required_devices` | Minimum number of GPUs needed | Yes |\n| `account` | Slurm account (unused for GPU probing, kept for compatibility) | No |\n\n## Procedure\n\n### 1. Check local GPUs\n\n```bash\ntimeout 5 nvidia-smi --query-gpu=name,memory.total --format=csv,noheader 2>/dev/null\n```\n\nIf this succeeds:\n- Count the number of GPU lines → `available_gpus`\n- Extract the GPU name from the first line → `device_type` (e.g., `NVIDIA B200`, `NVIDIA H100 80GB HBM3`)\n- Normalize `device_type`: strip `NVIDIA ` prefix and trailing memory info to get the short name (e.g., `B200`, `H100`, `A100`, `L40S`, `RTX 6000`)\n\nIf `nvidia-smi` fails or returns no GPUs → `available_gpus = 0`, `device_type = null`.\n\n### 2. Check if local GPUs are sufficient\n\nIf `available_gpus >= required_devices`, continue to step 2a to determine whether Docker is available on this host.\n\n### 2a. Check Docker availability\n\n```bash\ncommand -v docker >/dev/null 2>&1 && timeout 3 docker info >/dev/null 2>&1\n```\n\n- If the command succeeds (Docker CLI exists **and** the daemon responds) → **Result**: `satisfied, local, docker`\n- Otherwise (Docker CLI missing, daemon not running, or permission denied) → **Result**: `satisfied, local, direct`\n\nInclude `available_gpus` in the result. Do NOT include `device_type` — local execution does not need it.\n\n### 3. Check Slurm availability\n\nIf local GPUs are insufficient (or as additional detection), check for Slurm:\n\n```bash\nwhich squeue 2>/dev/null && squeue --version 2>/dev/null\n```\n\nIf Slurm is NOT available → go to step 5.\n\n### 4. Resolve GPU type\n\n**Optional dependency.** Before doing anything else in this step, check whether `skills/internal-env-info/` exists in the toolkit. If it does **not**, skip this step entirely and return `scenario: not_satisfied` with `available_gpus = 0`, `device_type = null`, `gpus_per_node = null`, `cluster_name = null`, `default_models_repo = null`, and `default_user_root_dir = null`. Do **not** report this as an error — the skill is an internal-only dependency.\n\nWhen Slurm is available but local `nvidia-smi` returned no GPUs or `device_type` is null (login nodes typically have no GPUs), capture the hostname and delegate cluster identification to `internal-env-info`:\n\n```bash\nhostname -f 2>/dev/null || hostname\n```\n\n- Pass the captured hostname to `internal-env-info` in **hostname-based mode**. That skill owns the NVIDIA-internal login-host patterns and the hostname → `<cluster>` mapping; do **not** parse the hostname or hard-code any cluster identifier here.\n- It returns the standard output template (`cluster_name`, `device_type`, `gpus_per_node`, `default_models_repo`, `default_user_root_dir`) plus supplementary field (`mfa_style`).\n- Set `available_gpus` = `gpus_per_node` (if resolved).\n- Set `cluster_name` = the value returned (a placeholder `<cluster>` token in this skill). The orchestrator uses this to fetch per-cluster info (`mfa_style`, `default_models_repo`, `default_user_root_dir`, `gpus_per_node`) from `internal-env-info`; connection fields (`mounts`, `ssh_host`, `partition`, etc.) come from caller-supplied inputs in `job_spec.json` (with `internal-env-info` default values / ask-the-user fallbacks).\n- Set `default_user_root_dir` = the user root directory returned (with `<user_name>` substituted with the actual SLURM username). Set to `null` if not found.\n- If `internal-env-info` returns `null` for `cluster_name` (no pattern matched), set `device_type = null`, `cluster_name = null`, `default_user_root_dir = null` and fall through to Step 5 (`not_satisfied`).\n\n**Result**: `satisfied, slurm, local` with `device_type`, `gpus_per_node`, `cluster_name`, and `default_user_root_dir`.\n\n### 5. Not satisfied\n\nIf neither local GPUs nor Slurm is available:\n- **Result**: `not_satisfied`\n\n## Output\n\nReturn a single structured result:\n\n```\nscenario: <satisfied, local, docker | satisfied, local, direct | satisfied, slurm, local | not_satisfied>\navailable_gpus: <N>\ndevice_type: <short GPU/device name or null>\ngpus_per_node: <N or null>\ncluster_name: <cluster name or null>\ndefault_models_repo: <host path to llm-models directory, or null>\ndefault_user_root_dir: <per-user root directory on cluster storage, or null>\n```\n\n**Examples:**\n\n```\nscenario: satisfied, local, docker\navailable_gpus: 4\n```\n\n```\nscenario: satisfied, local, direct\navailable_gpus: 4\n```\n\n```\nscenario: satisfied, slurm, local\navailable_gpus: 4\ndevice_type: B200\ngpus_per_node: 4\ncluster_name: <cluster>\ndefault_user_root_dir: <default_user_root_dir>/<user_name>\n```\n\n```\nscenario: not_satisfied\navailable_gpus: 0\ndevice_type: null\ncluster_name: null\n```\n\n## Rules\n\n- Never install drivers or modify the system\n- If `nvidia-smi` hangs, use a 5-second timeout: `timeout 5 nvidia-smi ...`\n- GPU type is derived from the hostname by matching the cluster name via the `internal-env-info` skill when it is installed — no `srun` allocation needed. If that skill is absent, the scenario falls back to `not_satisfied` (see Step 4); do not error.\n- Report the GPU type exactly as found in the mapping (do not guess or fabricate)\n","frontmatter":{"name":"exec-env-check","description":"Check the local execution environment for GPU availability, Docker support, and Slurm access. Returns the execution scenario (`satisfied, local, docker`, `satisfied, local, direct`, `satisfied, slurm, local`, or `not_satisfied`), the number of available GPUs, and the GPU type. On Slurm login nodes without local GPUs, the cluster is identified by delegating the hostname to internal-env-info (hostname-based mode), which owns the hostname → cluster_name patterns; GPU type and gpus_per_node then come from that skill's reference files. If internal-env-info is not installed, the scenario falls back to `not_satisfied` without probing compute nodes via srun.","tags":["infrastructure","slurm","environment"],"license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":1570,"sizeBytes":6162},{"name":"SKILL.md","path":".claude/skills/exec-local-compile/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-local-compile/SKILL.md","title":"exec-local-compile","category":"anthropic-skill","format":"markdown","content":"---\nname: exec-local-compile\ndescription: Compile TensorRT-LLM on a compute node inside a Docker container. Use this when already on a compute node with GPUs visible.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Compile TensorRT-LLM (Local / Compute Node)\n\nCompile TensorRT-LLM from source on a compute node inside a Docker container.\n\n## When to Use\n\n| Scenario | Use This Skill? |\n|----------|----------------|\n| On a compute node with GPUs visible (`nvidia-smi` works) | Yes |\n| On a SLURM login node (no GPUs) | No — use `exec-slurm-compile` instead |\n\n## Prerequisites\n\n- You are inside a Docker/enroot container on a compute node\n- `nvidia-smi` succeeds (GPUs visible)\n- `/usr/local/tensorrt` exists (TensorRT installation in the container)\n\n## Instructions\n\n### Step 1: Verify Environment\n\nRun `nvidia-smi` to confirm you are on a compute node with GPU access.\n\n### Step 2: Locate the Codebase\n\n`cd` to the TensorRT-LLM repository. If the path is not provided by the user, ask for it.\n\n### Step 3: (Optional) Checkout Branch\n\nIf the user specifies a branch (e.g., \"compile ToT\"), checkout and pull:\n```bash\ngit checkout main && git pull\n```\n\n### Step 4: Build\n\nRun the build command (**incremental by default** — omit `-c`/`--clean` unless explicitly requested or the incremental build fails):\n\n```bash\n./scripts/build_wheel.py --use_ccache -a \"<arch>\" -f --nvtx\n```\n\nReplace `<arch>` with the target GPU architecture (see Architecture Reference below). If not specified by the user, auto-detect from `nvidia-smi`.\n\n### Step 5: Install\n\n```bash\npip install -e .[devel]\n```\n\n### Step 6: Verify\n\n```bash\npython3 -c \"import tensorrt_llm; print(tensorrt_llm.__version__)\"\n```\n\n## Build Flags\n\n| Flag | Description |\n|------|-------------|\n| `-a \"<arch>\"` | Target GPU architecture(s) |\n| `--nvtx` | Enable NVTX markers for profiling |\n| `--use_ccache` | Use ccache for faster recompilation |\n| `-f` / `--fast_build` | Skip some kernels for faster dev compilation. **Always use for dev builds.** |\n| `-c` / `--clean` | Clean build directory before building. Only when needed (see below). |\n| `--skip_building_wheel` | Build in-place without creating a wheel file |\n| `--no-venv` | Skip virtual environment creation |\n\n## Architecture Reference\n\n| Value | GPU Family |\n|-------|-----------|\n| `\"100-real\"` | Blackwell (B200, GB200) |\n| `\"90-real\"` | Hopper (H100, H200) |\n| `\"89-real\"` | Ada Lovelace (L40S) |\n| `\"80-real\"` | Ampere (A100) |\n| `\"90;100-real\"` | Multiple architectures |\n\n## Incremental vs. Clean Builds\n\n**Default to incremental builds** — CMake only recompiles changed files, saving significant time.\n\nUse a **clean build** (`-c`) only when:\n- The user explicitly requests a clean/fresh build\n- An incremental build fails with linker errors, stale object files, or CMake cache issues\n- Major branch changes (e.g., rebasing across many commits) that may invalidate the build cache\n- Build system files changed (`CMakeLists.txt`, `*.cmake`)\n","frontmatter":{"name":"exec-local-compile","description":"Compile TensorRT-LLM on a compute node inside a Docker container. Use this when already on a compute node with GPUs visible.","license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":786,"sizeBytes":2985},{"name":"SKILL.md","path":".claude/skills/exec-local-docker/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-local-docker/SKILL.md","title":"exec-local-docker","category":"anthropic-skill","format":"markdown","content":"---\nname: exec-local-docker\ndescription: >-\n  Execute a TensorRT-LLM workload locally in Docker. Runs a fully-resolved\n  Docker command in background, monitors completion, reads logs, and reports\n  results. Workflow-agnostic — does not need to know if the workload is pytest,\n  eval, benchmark, or a custom script.\ntags: [docker, execution, infrastructure]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Local Docker Executor\n\nRun a Docker command locally, monitor it, and report results.\n\n## Input (from orchestrator prompt)\n\nThe orchestrator passes these fields in the skill prompt:\n\n| Field | Description |\n|-------|-------------|\n| `docker_cmd` | Complete `docker run` command string, ready to execute |\n| `work_dir` | Local work directory for logs and artifacts |\n| `log_file` | Full path to the log file (output redirected here) |\n| `model_name` | Short model name for reporting |\n| `workflow_type` | `pytest`, `eval`, `custom`, or `benchmark` — for output parsing hints |\n| `success_patterns` | Comma-separated patterns indicating success (e.g., `passed,accuracy:`) |\n| `failure_patterns` | Comma-separated patterns indicating failure (e.g., `FAILED,Error,AssertionError`) |\n\n## Procedure\n\n### Step 0: Resolve Image and Build (when `build_project=true`)\n\nThis executor owns image selection and the build for the local Docker target. Skip the entire step when `build_project=false`.\n\n1. **Detect the target GPU type.** Use `gpu_type` from `job_spec.json` if upstream env-check resolved it; otherwise probe locally with `nvidia-smi --query-gpu=name --format=csv,noheader | head -1`.\n2. **Detect the host CPU arch** with `uname -m` (`x86_64` or `aarch64`).\n3. **Resolve the container image.** Read `<repo_root>/jenkins/current_image_tags.properties` and pick the tag whose CPU-arch flavor matches the host. If the orchestrator already passed a `container_image` field in the job spec, use that and skip the lookup.\n4. **Map GPU → build arch (`-a` flag):** `H100`/`H200` → `90-real`; `B200`/`GB200`/`B300`/`GB300` → `100-real`; `A100` → `80-real`; `L40S` → `89-real`. Default `100-real` when the GPU is unknown.\n5. **Compile.** Invoke the `exec-local-compile` skill with `repo_dir=<repo_root>`, `image=<resolved tag>`, `arch=<arch>`. Wait for completion.\n6. **On failure**, do not launch the workload. Report `BUILD_FAILED` with the last 100 lines of the compile log.\n\n`build_project`, `gpu_type`, `repo_root`, and (optionally) `container_image` come from `job_spec.json`.\n\n### Step 1: Launch\n\nRun the Docker command in background using `run_in_background`:\n\n```bash\n<docker_cmd> 2>&1 | tee <log_file>\n```\n\nReport to the orchestrator: \"Launched locally, log at `<log_file>`\"\n\n### Step 2: Monitor for Hangs\n\nWhile waiting for the background process to complete, actively monitor the log\nfile for hang indicators. Launch a monitoring loop using `run_in_background`:\n\n```bash\nwhile true; do\n  sleep 60\n  if [ -f \"<log_file>\" ] && grep -qi \"hang detected\" \"<log_file>\"; then\n    echo \"HANG_DETECTED: Found 'hang detected' in log file\"\n    docker ps --filter \"ancestor=<image>\" -q | xargs -r docker kill 2>/dev/null\n    exit 1\n  fi\ndone\n```\n\n- If the monitor detects a hang, it kills the Docker container and exits with\n  code 1. The main background process will also terminate.\n- When the main process completes normally (background notification received),\n  kill the monitoring loop (it is no longer needed).\n- If a hang is detected, skip to Step 4 and report `HANG_DETECTED` status\n  instead of proceeding to normal result collection.\n\n### Step 3: Wait for Completion\n\nThe Bash tool's `run_in_background` will notify when the process finishes\n(either normally or because the container was killed by the hang monitor).\n\n### Step 4: Read Results\n\nOn completion:\n\n1. **Read exit code** from the background command result.\n2. **Read the last 100 lines** of `<log_file>` using the Read tool.\n3. **If exit code != 0**, also read the first 50 lines to catch early errors (import failures, setup crashes).\n4. **Search for patterns**:\n   - Grep `<log_file>` for each `success_patterns` entry\n   - Grep `<log_file>` for each `failure_patterns` entry\n\n### Step 5: Report\n\nReturn a structured result:\n\n```\nStatus: PASSED | FAILED | ERROR | HANG_DETECTED\nExit code: <N>\nLog file: <log_file>\nWork directory: <work_dir>\nSummary: <relevant output lines — pytest summary, accuracy score, or last few lines>\nErrors: <if failed, the relevant error output>\n```\n\n### Output Parsing by Workflow Type\n\n- **pytest**: Look for `X passed, Y failed in Zs` summary line\n- **eval**: Look for `accuracy:` or `score:` lines; check for `Expected accuracy >= X, but got Y` assertion\n- **custom**: No specific patterns — report last 10 lines of output\n- **benchmark**: Look for throughput/latency numbers\n\n## Rules\n\n- Never run the Docker command in foreground — always use `run_in_background`\n- Never `cat` the full log file — use Read with offset/limit or tail\n- If the Docker command fails immediately (exit code within seconds), check if the image exists locally\n- Report results even if the log file is empty (container may have failed to start)\n","frontmatter":{"name":"exec-local-docker","description":"Execute a TensorRT-LLM workload locally in Docker. Runs a fully-resolved Docker command in background, monitors completion, reads logs, and reports results. Workflow-agnostic — does not need to know if the workload is pytest, eval, benchmark, or a custom script.","tags":["docker","execution","infrastructure"],"license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":1307,"sizeBytes":5159},{"name":"SKILL.md","path":".claude/skills/exec-local-slurm/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-local-slurm/SKILL.md","title":"exec-local-slurm","category":"anthropic-skill","format":"markdown","content":"---\nname: exec-local-slurm\ndescription: >-\n  Submit and monitor a Slurm job on a local cluster. Supports two modes:\n  (1) Persistent allocation (default) — allocates nodes once via nohup salloc,\n  imports the container once, installs once, and reuses across runs by setting\n  SLURM env vars and running the sbatch script via bash. (2) One-shot sbatch —\n  submits a fully-generated Slurm script via sbatch, polls job status, reads\n  logs on completion, and reports results. Workflow-agnostic — handles pytest,\n  eval, benchmark, and custom scripts identically.\ntags: [slurm, execution, infrastructure]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Local Slurm Executor\n\nSubmit a Slurm job locally, monitor it, and report results. Uses persistent\nallocation by default to eliminate queue wait, container import, and install\noverhead on repeated runs.\n\n## Input (from orchestrator prompt)\n\nThe orchestrator passes these fields in the skill prompt:\n\n| Field | Description |\n|-------|-------------|\n| `script_path` | Full local path to the generated `.slurm` script |\n| `work_dir` | Local work directory for logs and artifacts |\n| `model_name` | Short model name for reporting |\n| `workflow_type` | `pytest`, `eval`, `custom`, or `benchmark` — for output parsing hints |\n| `success_patterns` | Comma-separated patterns indicating success |\n| `failure_patterns` | Comma-separated patterns indicating failure |\n| `log_file_pattern` | Log filename pattern with `%j` placeholder (e.g., `llama_auto_test_%j.out`) |\n| `persistent_mode` | Default: `true` for all local slurm workflows. Set to `false` to force one-shot sbatch (opt-out). |\n| `release_allocation` | `true` to release the current allocation and stop. Only set when the user explicitly says no more jobs are needed. Default: `false`. Never auto-release. |\n| `alloc_time_limit` | Walltime for the persistent allocation. Default: `04:00:00`. |\n| `docker_image` | Container image for persistent container import. From `job_spec.json`. |\n| `container_name` | Container name used in the `.slurm` script (e.g., `llama_auto_test`). Must match exactly. From `job_spec.json`. |\n| `container_mounts` | Comma-separated mount mappings. From `job_spec.json`. |\n| `repo_root` | Repo root path (for locating state file and project path). |\n| `slurm_params` | Slurm parameters object: `partition`, `account`, `nodes`, `ntasks`, `ntasks_per_node`, `gpus_per_node`. From `job_spec.json`. |\n\n## Procedure\n\n### Pre-step: Resolve Image and Build (when `build_project=true`)\n\nThis executor owns image selection and the build for the local SLURM cluster. Skip the entire pre-step when `build_project=false`.\n\n1. **Read the cluster GPU type** from `device_type` in `job_spec.json` (resolved upstream by env-check / internal-env-info). If absent, probe from any compute node via `srun -p <partition> --ntasks-per-node=1 nvidia-smi --query-gpu=name --format=csv,noheader | head -1`. If `skills/internal-env-info/` is not installed, skip the upstream lookup silently and rely on the srun probe — do not report the missing skill as an error.\n2. **Determine the partition's CPU arch.** Look up `job_spec.slurm_env.partitions[]` for the entry where `name == partition` and read its `arch` (`x86_64` / `aarch64`). Case-executor's Step 2.5 already detected this — do **not** run `scontrol show node` here. If `slurm_env` is absent or the entry's `arch` is null (case-executor ran without SLURM tools), fall back to inferring from `device_type`: Grace-based parts (`GB*` / `GH*`) → `aarch64`; everything else → `x86_64`.\n3. **Resolve the container image.** If `docker_image` is already set in `job_spec.json`, use it directly. Otherwise read `<repo_root>/jenkins/current_image_tags.properties` and pick the tag matching the partition's CPU arch.\n4. **Map GPU → build arch (`-a` flag):** `H100`/`H200` → `90-real`; `B200`/`GB200`/`B300`/`GB300` → `100-real`; `A100` → `80-real`; `L40S` → `89-real`. Default `100-real`.\n5. **Compile.** Invoke the `exec-slurm-compile` skill with `repo_dir=<repo_root>`, `partition`, `account`, `container_image=<resolved tag>`, `user_root_dir`, `arch`. Wait for completion.\n6. **On failure**, do not allocate or run the workload. Report `BUILD_FAILED` with the last 100 lines of the build log.\n\n`build_project`, `device_type`, `repo_root`, `partition`, `account`, `user_root_dir`, and (optionally) `docker_image` come from `job_spec.json`.\n\n### Step 0: Allocation Management\n\nThis step runs before any execution. It determines whether to reuse an existing\npersistent allocation, create a new one, or fall through to one-shot sbatch.\n\n#### Step 0A — Release mode\n\nIf `release_allocation=true`:\n\n1. Read `<repo_root>/work_dirs/.slurm_alloc.json`\n2. If file exists and `job_id` is present:\n   ```bash\n   scancel <job_id>\n   ```\n3. Delete the state file and `.salloc.log`\n4. Report: \"Allocation <job_id> released.\" Stop.\n\n#### Step 0B — One-shot mode\n\nIf `persistent_mode=false`, skip entirely to Step 1 (One-Shot Path).\n\n#### Step 0C — Validate existing allocation\n\nRead `<repo_root>/work_dirs/.slurm_alloc.json`.\n\n- If state file **does not exist** → go to Step 0D.\n- If state file exists, validate the allocation:\n  ```bash\n  squeue -j <job_id> -h -o \"%T %L\"\n  ```\n  - **RUNNING + remaining > 5 min + params compatible** → **reuse** (skip to Step 0F)\n  - **RUNNING + remaining <= 5 min** → warn user (\"Allocation expiring soon\"), `scancel <job_id>`, delete state file, go to Step 0D\n  - **RUNNING + params incompatible** → `scancel <job_id>`, delete state file, go to Step 0D\n  - **Empty output / error** (job gone) → stale state file, delete it, go to Step 0D\n\n**Params compatibility check:**\n- `partition` must match\n- `nodes` in state must be **>=** requested nodes (a 2-node allocation can serve 1-node jobs)\n- `docker_image` must match\n\n#### Step 0D — Justify and prepare allocation\n\nBefore allocating, ensure no orphaned allocations exist and log the reason:\n\n1. Check for existing persistent jobs:\n   ```bash\n   squeue -u $(whoami) -h -o \"%i %T %j\" --name=<account>-trtllm.persistent\n   ```\n2. If found → orphan (state file was missing/corrupt). Cancel it:\n   ```bash\n   scancel <orphan_job_id>\n   ```\n   Log: \"Released orphaned allocation <orphan_job_id> — state file was missing.\"\n3. Log justification: \"No reusable allocation found. Allocating <nodes> node(s) on partition <partition> for <alloc_time_limit>.\"\n\n#### Step 0E — Allocate new\n\nConvert `alloc_time_limit` from `HH:MM:SS` to seconds (e.g., `04:00:00` → `14400`).\n\n```bash\nmkdir -p <repo_root>/work_dirs\nnohup salloc --partition=<partition> --account=<account> \\\n  --nodes=<nodes> --time=<alloc_time_limit> \\\n  --job-name=<account>-trtllm.persistent \\\n  sleep <alloc_time_limit_seconds> \\\n  > <repo_root>/work_dirs/.salloc.log 2>&1 &\n```\n\nRetrieve job ID:\n```bash\nsqueue -u $(whoami) -h -o \"%i\" --name=<account>-trtllm.persistent\n```\n\nIf no job appears after 10 seconds, read `<repo_root>/work_dirs/.salloc.log`\nfor errors (bad partition, invalid account, etc.) and report to user.\n\nOtherwise, poll until RUNNING (every 10s, max 60 polls). Once RUNNING, get\nthe nodelist:\n```bash\nsqueue -j <JOB_ID> -h -o \"%N\"\n```\n\n**Import container** on all nodes (using the **same container name** as the\n`.slurm` script — critical for pyxis reuse):\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-image=<docker_image> \\\n  --container-name=<container_name> true\n```\n\n**Warm up filesystem mounts** on all nodes — `ls` each mounted path so\nLustre/NFS metadata is cached for later use:\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-name=<container_name> \\\n  --container-mounts=<container_mounts> \\\n  bash -c 'for p in <mount_target_1> <mount_target_2> ...; do ls \"$p\" > /dev/null 2>&1; done'\n```\nParse mount targets from `container_mounts` — the right-hand side of each\n`host:container` pair.\n\n**Check GPU status** on all nodes (no container needed — `nvidia-smi` is on\nthe host):\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  bash -c 'echo \"=== $(hostname) ===\" && nvidia-smi --query-compute-apps=pid,name,used_memory --format=csv,noheader'\n```\n- If output shows no processes → GPUs are clean, proceed\n- If unexpected processes found → warn user: \"GPU processes found on\n  <node>: <process_list>. These may interfere with the job.\"\n\n**Run install** on all nodes with `--container-writable` (so packages persist\nin the named container across srun calls). Install all common requirements\nupfront so that subsequent jobs of any workflow type can skip install:\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-name=<container_name> --container-writable \\\n  --container-mounts=<container_mounts> \\\n  bash -c 'cd <repo_root> && pip install -e . && pip install -r requirements-dev.txt && \\\n    if [ -f examples/trtllm-eval/requirements.txt ]; then pip install -r examples/trtllm-eval/requirements.txt; fi'\n```\nFor custom workflow with `skip_install=true`, skip this step entirely.\n\n**Write state file** `<repo_root>/work_dirs/.slurm_alloc.json`:\n```json\n{\n  \"job_id\": \"<JOB_ID>\",\n  \"container_name\": \"<container_name>\",\n  \"nodelist\": \"<NODELIST>\",\n  \"partition\": \"<partition>\",\n  \"account\": \"<account>\",\n  \"nodes\": <nodes>,\n  \"gpus_per_node\": <gpus_per_node>,\n  \"docker_image\": \"<docker_image>\",\n  \"container_mounts\": \"<container_mounts>\",\n  \"allocated_at\": \"<ISO timestamp>\",\n  \"time_limit\": \"<alloc_time_limit>\",\n  \"installed\": true\n}\n```\n\nProceed to Step 1 (Persistent Path).\n\n#### Step 0F — Reuse path validation\n\nAllocation is valid. Check if container name matches current request:\n\n- If `container_name` in state file **matches** `container_name` from\n  `job_spec.json` → proceed to Step 1 (Persistent Path). Container and install\n  are already set up.\n- If **different** (model changed) → import the new container, warm up mounts,\n  and install:\n  ```bash\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-image=<docker_image> \\\n    --container-name=<new_container_name> true\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-name=<new_container_name> \\\n    --container-mounts=<container_mounts> \\\n    bash -c 'for p in <mount_targets>; do ls \"$p\" > /dev/null 2>&1; done'\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-name=<new_container_name> --container-writable \\\n    --container-mounts=<container_mounts> \\\n    bash -c 'cd <repo_root> && pip install -e . && pip install -r requirements-dev.txt'\n  ```\n  Update `container_name` in state file. Proceed to Step 1 (Persistent Path).\n\n---\n\n### Step 1: Execute\n\nTwo execution paths depending on `persistent_mode`.\n\n#### Persistent Path (persistent_mode=true)\n\n**Pre-flight time check** — verify the allocation has enough remaining time:\n```bash\nsqueue -j <job_id> -h -o \"%L\"\n```\nIf remaining time < job's `time_limit` from `slurm_params`, warn user:\n\"Allocation has <remaining> left but job expects <time_limit>. The job may\nbe killed early.\"\n\n**Pre-flight GPU check** — verify GPUs are not occupied by leftover processes:\n```bash\nsrun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n  bash -c 'procs=$(nvidia-smi --query-compute-apps=pid,name,used_memory --format=csv,noheader 2>/dev/null); [ -n \"$procs\" ] && echo \"WARNING: GPU processes on $(hostname): $procs\"'\n```\nIf unexpected processes found, warn user before proceeding.\n\n**Run the sbatch script** with SLURM env vars — the `#SBATCH` directives are\ncomments when run via `bash`; the inner `srun` inherits the env vars and uses\nthe persistent allocation:\n```bash\nexport SLURM_JOB_ID=<job_id>\nexport SLURM_JOB_NUM_NODES=<nodes>\nexport SLURM_NNODES=<nodes>\nexport SLURM_NTASKS=<ntasks>\nexport SLURM_NTASKS_PER_NODE=<ntasks_per_node>\nexport SLURM_NODELIST=<nodelist>\nexport SLURM_JOB_NODELIST=<nodelist>\nbash <script_path> 2>&1 | tee <work_dir>/<model_name>_<workflow_type>.log\n```\n\nRun with `run_in_background`. Container reuse is automatic (pyxis skips\nre-import when `--container-name` already exists on the node). The install\nstep (Step 1 srun) always runs to keep the container up to date.\n\n**Step cancellation:** If a job hangs or the user wants to abort, cancel just\nthe srun step without killing the allocation:\n```bash\n# List active steps:\nsqueue -s -j <job_id>\n# Cancel a specific step (e.g., step 0):\nscancel <job_id>.0\n```\nThe allocation stays RUNNING — new jobs can run immediately. The job name\n`<account>-trtllm.persistent` uniquely identifies allocations created by this\nskill via `squeue --name=<account>-trtllm.persistent`.\n\n#### One-Shot Path (persistent_mode=false)\n\n```bash\nsbatch <script_path>\n```\n\nParse the job ID from output: `Submitted batch job <JOB_ID>`.\n\nIf sbatch fails, report the error immediately and stop.\n\n### Step 2: Report Submission\n\n**One-shot mode:**\n```\nJob ID: <JOB_ID>\nScript: <script_path>\nWork directory: <work_dir>\nLog files: <work_dir>/<log_file_pattern with %j replaced>\nMonitor: squeue -j <JOB_ID>\n```\n\n**Persistent mode:**\n```\nAllocation: <job_id> (persistent, remaining: <time>)\nContainer: <container_name>\nScript: <script_path>\nWork directory: <work_dir>\nLog file: <work_dir>/<model_name>_<workflow_type>.log\n```\n\n### Step 3: Poll / Wait for Completion (with Hang Monitoring)\n\nBoth polling modes must also monitor the log file for hang indicators. Every\n60 seconds, check the log file for the pattern `hang detected`\n(case-insensitive). If found, terminate the job and skip to Step 5 with\n`HANG_DETECTED` status.\n\n**One-shot mode:** Poll every 30 seconds, max 60 polls (30 minutes):\n\n```bash\nsqueue -j <JOB_ID> -h -o \"%T %M %R\"\n```\n\n- Report state transitions: `PENDING` → `RUNNING` → `COMPLETING` → done\n- If `PENDING` for > 5 minutes, check the reason:\n  ```bash\n  squeue -j <JOB_ID> -o \"%i %T %r %S\"\n  ```\n- The job is done when `squeue -j <JOB_ID>` returns empty (no rows)\n- If max polls exceeded, report the job is still running and provide manual\n  check commands\n- **Hang check** (every other poll, i.e., every 60s): Once the job is RUNNING,\n  resolve the log file path (`<work_dir>/<log_file_pattern>` with `%j` replaced\n  by `<JOB_ID>`) and check:\n  ```bash\n  grep -qi \"hang detected\" <log_file> 2>/dev/null && echo \"HANG_DETECTED\"\n  ```\n  If `HANG_DETECTED` is printed, immediately cancel the job:\n  ```bash\n  scancel <JOB_ID>\n  ```\n  Then read the last 200 lines of the log file and skip to Step 5.\n\n**Persistent mode:** While waiting for `run_in_background` completion\nnotification, launch a hang monitoring loop using `run_in_background`:\n\n```bash\nwhile true; do\n  sleep 60\n  LOG_FILE=\"<work_dir>/<model_name>_<workflow_type>.log\"\n  if [ -f \"$LOG_FILE\" ] && grep -qi \"hang detected\" \"$LOG_FILE\"; then\n    echo \"HANG_DETECTED: Found 'hang detected' in log file\"\n    squeue -s -j <job_id> -h -o \"%i\" | head -1 | xargs -r scancel 2>/dev/null\n    exit 1\n  fi\ndone\n```\n\nWhen the main process completes normally, kill the monitoring loop. If the\nmonitor fires first, it cancels the active Slurm step (not the allocation\nitself — the allocation stays for future runs).\n\n### Step 4: Collect Results\n\nOn completion:\n\n1. **Get exit code:**\n   - **One-shot:** `sacct -j <JOB_ID> --format=JobID,State,ExitCode,Elapsed,MaxRSS -n`\n   - **Persistent:** From the `run_in_background` result exit code\n\n2. **Resolve log file paths:**\n   - **One-shot:** Replace `%j` in `log_file_pattern` with `<JOB_ID>`:\n     - Output: `<work_dir>/<MODEL_NAME>_..._<JOB_ID>.out`\n     - Error: `<work_dir>/<MODEL_NAME>_..._<JOB_ID>.err`\n   - **Persistent:** Single log at `<work_dir>/<model_name>_<workflow_type>.log`\n\n3. **Read log files:**\n   - Read the last 100 lines of the output log\n   - If job failed, also read the first 50 lines to catch early errors\n   - **One-shot only:** If job failed, read the last 50 lines of the `.err` file\n\n4. **Search for patterns:**\n   - Grep the log file for each `success_patterns` entry\n   - Grep the log file for each `failure_patterns` entry\n\n### Step 5: Report\n\nReturn a structured result:\n\n```\nJob ID: <JOB_ID> (or allocation <job_id> for persistent mode)\nStatus: COMPLETED | FAILED | TIMEOUT | CANCELLED | OUT_OF_MEMORY | HANG_DETECTED\nExit code: <exit code>\nElapsed: <time>\nLog files:\n  stdout: <path>\n  stderr: <path> (one-shot only)\nWork directory: <work_dir>\nSummary: <relevant output — pytest summary, accuracy score, throughput, or last lines>\nErrors: <if failed, relevant error output>\n```\n\n### Output Parsing by Workflow Type\n\n- **pytest**: Look for `X passed, Y failed in Zs` summary line. Status maps:\n  `COMPLETED` + 0 exit → `PASSED`; non-zero → `FAILED`\n- **eval**: Look for `accuracy:` lines. If `--check_accuracy` was used,\n  assertion error means threshold not met\n- **benchmark**: Look for throughput/latency numbers. Check for\n  `8_done_*.txt` completion marker\n- **custom**: No specific patterns — report last 10 lines of output\n\n### Failure Diagnosis\n\n| Symptom | Likely Cause | Action |\n|---------|-------------|--------|\n| `TIMEOUT` | Job exceeded time limit | Report and suggest increasing time |\n| `OUT_OF_MEMORY` | OOM on node | Report and suggest reducing batch size |\n| Exit `1:0` + `FAILED` | Non-zero exit from srun | Read `.err` file for details |\n| `CANCELLED` | Preempted or user cancelled | Check `sacct --format=JobID,State,Reason` |\n| Empty output log | Container failed to start | Check `.err` / `.salloc.log` for mount or image errors |\n| `HANG_DETECTED` in log | Process/GPU hang (deadlock, NCCL timeout) | Job auto-cancelled; check last 200 log lines for root cause |\n| `srun: error: Unable to create step` | Allocation expired | Delete state file, re-run (will re-allocate) |\n| `srun: error: ... does not exist` | Named container missing | Delete state file, re-run (will re-import) |\n\n## Rules\n\n- Never use `squeue` with the `-uall` flag\n- Always use `sacct` (not `squeue`) for final job status in one-shot mode\n- Poll with `run_in_background` + `sleep 30 &&` to avoid blocking (one-shot)\n- Never read full log files — use Read with offset/limit\n- If the Slurm script path doesn't exist, report the error immediately\n- For cluster hardware info (GPU types, partitions, constraints, storage\n  paths), use the `trtllm-agent-toolkit:internal-env-info` skill\n  when it is installed. If the skill is absent, skip the lookup silently and\n  use the caller-supplied fields or the srun probe in Step 1 — do not report\n  the missing skill as an error.\n- Never auto-release an allocation after job completion — it should persist\n  for potential future runs until walltime expires or user explicitly releases\n- Must pass `-N <nodes>` to every direct `srun --jobid` call — without it,\n  srun defaults to 1 node even in a multi-node allocation\n- The container name used in setup steps must **exactly match** the\n  `--container-name` in the `.slurm` script, otherwise pyxis will re-import\n  the image and lose `--container-writable` state\n","frontmatter":{"name":"exec-local-slurm","description":"Submit and monitor a Slurm job on a local cluster. Supports two modes: (1) Persistent allocation (default) — allocates nodes once via nohup salloc, imports the container once, installs once, and reuses across runs by setting SLURM env vars and running the sbatch script via bash. (2) One-shot sbatch — submits a fully-generated Slurm script via sbatch, polls job status, reads logs on completion, and reports results. Workflow-agnostic — handles pytest, eval, benchmark, and custom scripts identically.","tags":["slurm","execution","infrastructure"],"license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":5160,"sizeBytes":18949},{"name":"SKILL.md","path":".claude/skills/exec-remote-slurm/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-remote-slurm/SKILL.md","title":"exec-remote-slurm","category":"anthropic-skill","format":"markdown","content":"---\nname: exec-remote-slurm\ndescription: >-\n  Remote SLURM cluster development via SSH. Use when running jobs,\n  profiling, or developing on a remote SLURM cluster with pyxis/enroot\n  containers. Covers SSH connection management, srun/sbatch/salloc job\n  patterns, tmux-based allocation persistence, file transfer, and safe\n  remote file access. Works with any SLURM cluster accessible via SSH.\ntags: [infrastructure, remote, slurm]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Remote SLURM Development\n\nDevelop and execute on remote SLURM clusters via SSH commands. No filesystem\nsync needed — all remote operations go through SSH.\n\n## Step 1: Resolve Cluster\n\n**Optional dependency — `internal-env-info`.** Before doing\nanything else in Step 1, check whether `skills/internal-env-info/`\nexists in the toolkit. If it does **not**, skip every sub-step that invokes\nthe skill or reads its `references/` files and proceed using only the\ncluster fields the orchestrator wrote into `job_spec.json` (`ssh_host`,\n`slurm_user`, `remote_cwd`, `partition`, `account`, `container_image`,\n`mounts`, `gpus_per_node`, `mfa_style`). If a required field is also\nmissing, stop and ask the user to supply it — do **not** emit an error\nabout the missing skill.\n\nThe orchestrator supplies one of two input shapes; pick the right path:\n\n- **Explicit cluster** — `slurm_cluster` name is provided. Skip\n  auto-selection and go straight to **Load the per-cluster info**.\n- **Auto-select cluster** — only hardware constraints are provided\n  (`device_type`, `total_required_devices`, `required_devices_per_node`).\n  Resolve a cluster from the constraints, then load its per-cluster info.\n\n### 1a. Auto-select (only when no explicit cluster was provided)\n\nInvoke the `internal-env-info` skill in **constraint-based mode**, passing\n`device_type`, `required_devices_per_node`, and (optionally)\n`total_required_devices`. That skill owns the filter rules and either\nreturns a single resolved cluster entry or stops and asks the user. Do not\nimplement the filter here.\n\nUse `node_count` from `job_spec.json` directly as `--nodes` and\n`required_devices_per_node` as `--ntasks-per-node`. **Do not recompute\n`node_count`** — `trtllm-case-executor` is the single source of truth and\nwrites the resolved value into `job_spec.json` for every executor to read.\n\nThe returned entry's `cluster_name`, `mfa_style`, `default_user_root_dir`,\n`default_models_repo`, and `gpus_per_node` feed Step 1b.\n\n### 1b. Load the per-cluster info\n\nConnection fields are read **only** from `job_spec.json` (written by\n`trtllm-case-executor`): `ssh_host`, `slurm_user`, `remote_cwd`, `account`,\n`partition`, `mounts`, `container_image`, `gpus_per_node`, and `packages`.\nThere is no per-cluster connection-config file to parse — if a required\nfield is absent from `job_spec.json`, stop and ask the user to supply it.\nDo not silently default.\n\nInvoke the `internal-env-info` skill in single-cluster mode with\nthe resolved `cluster_name` to retrieve per-cluster info: partition limits,\ntopology-aware scheduling guidance, `default_models_repo`,\n`default_user_root_dir`, `gpus_per_node`, `mfa_style` (used by Step 2 to\nchoose the SSH path — see below), and shared constraint and tuning\npatterns. The runtime `container_image` always comes from `job_spec.json`,\nnever from this skill.\n\n## Step 2: SSH Preflight\n\nThe cluster's `mfa_style` (returned by `internal-env-info` in Step 1b) is\nthe **single source of truth for which SSH path to use**. Switch on the\nvalue — do **not** probe-and-fall-back on the SSH error string.\n\n| `mfa_style` | Path | Why |\n|-------------|------|-----|\n| `false` | Direct (Step 2a) | Cluster accepts key-based / passwordless SSH; no special handling needed. |\n| `true`  | MFA (Step 2b)    | Cluster requires interactive MFA (Duo / device-code / one-time PIN). |\n| `null`  | Unknown (Step 2c)| `mfa_style` not declared in the cluster reference (or `internal-env-info` is not installed). Probe direct first, fall back to MFA. When the reference exists, also ask the user to update it so future runs skip the probe; otherwise just continue silently. |\n\n### 2.0. Build `<ssh_host>` from `slurm_user` + `cluster_name`\n\nBefore any SSH attempt, derive the SSH destination. Priority order:\n\n1. **Caller-supplied `ssh_host`** (passed in through case-executor) → use\n   verbatim, skip everything below. This is the canonical escape hatch.\n2. **`internal-env-info` skill installed** → delegate for the\n   `ssh_host_pattern` key (defined in the skill's `Default Environment Values`\n   section). It returns a single-line snippet that concatenates `slurm_user`\n   and `cluster_name` into `SSH_HOST`. Do **not** re-derive or re-suffix\n   `slurm_user` here — `internal-env-info` already applied whatever its\n   `slurm_user` construction rule requires for that cluster.\n\n   Inputs to substitute:\n   - `slurm_user` — pre-resolved from `internal-env-info` (already in its\n     final, SSH-ready form for that cluster). Read it from\n     `job_spec.json`.\n   - `cluster_name` — from `internal-env-info` Step 1b.\n3. **Skill not installed** → stop and ask the user to supply `ssh_host`\n   explicitly. Do **not** hard-code a login-host pattern in this skill —\n   the NVIDIA-internal convention lives only in `internal-env-info`.\n\nThe resolved value is the canonical `<ssh_host>` placeholder used everywhere\nin the rest of this skill (Steps 2a/2b/2c, Step 3, recipes, file-transfer,\nhang detection).\n\n### 2a. Direct path — `mfa_style: false`\n\nVerify connectivity using the user's existing SSH config:\n\n```bash\nssh -o ConnectTimeout=5 <ssh_host> \"echo ok\"\n```\n\nIf this succeeds, **skip to Step 3**. Do not create a trtllm-managed\nControlMaster — reuse whatever the user already has.\n\nIf it fails, do **not** silently switch to the MFA path. Stop and ask the\nuser to verify their `ssh_host` value or their key-based SSH setup.\nA direct-path cluster failing here is a configuration issue, not an MFA\nprompt to interpret.\n\n### 2b. MFA path — `mfa_style: true`\n\nFirst check whether a ControlMaster is already live (the user may have set\none up earlier in this session):\n\n```bash\nssh -O check <ssh_host> 2>&1 || ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -O check <ssh_host> 2>&1\n```\n\nIf either succeeds, **skip to Step 3** — reuse it.\n\nOtherwise the user must complete an interactive MFA handshake. MFA uses\nbrowser-based device-code verification that needs an interactive SSH\nsession — the skill cannot complete this on the user's behalf because the\nsession must stay alive while the user authenticates in their browser and\npresses ENTER.\n\nTell the user to run the ControlMaster command:\n\n> Run this in your terminal:\n>\n>     ssh -fNM -o ControlMaster=auto -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -o ControlPersist=24h <ssh_host>\n>\n> If the cluster uses a device-code flow you will see a PIN and a URL like:\n>\n>     Authenticate with PIN XXXXXX at <your-identity-provider-device-url>\n>\n> Open that URL, enter the PIN, approve, then press ENTER. The SSH session\n> will background itself as a persistent ControlMaster. Let me know when done.\n\nAfter the user confirms, verify with\n`ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -O check <ssh_host>` and\nproceed to Step 3.\n\n### 2c. Unknown — `mfa_style: null`\n\nThe cluster reference (`references/<cluster>.md`) does not declare an\n`mfa_style` value. Try the direct path (Step 2a) first. If it fails with\n`Permission denied (keyboard-interactive)` — and only that — switch to\nthe MFA path (Step 2b). Any other failure stops here and is reported to\nthe user.\n\nAfter connecting successfully, **ask the user to update the cluster\nreference file** with the discovered value (`mfa_style: true` or\n`false`, in the `## Authentication` table) so that future runs skip\nthis probe and use the deterministic switch in Step 2a / 2b.\n\n## SSH Command Shorthand\n\nEvery SSH and SCP invocation in the rest of this skill uses the placeholders below. Both pin `ControlPath` to the trtllm-managed socket so they reuse the ControlMaster set up in Step 2 — without this, OpenSSH starts a fresh handshake on every call, which on MFA clusters re-prompts for keyboard-interactive auth (and defeats the whole point of multiplexing):\n\n| Placeholder | Expansion |\n|-------------|-----------|\n| `<ssh_cmd>` | `ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p` |\n| `<scp_cmd>` | `scp -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p` |\n\nWhen the user has their own ControlMaster configured in `~/.ssh/config` (no trtllm-managed socket needed), the explicit `-o ControlPath` is harmless: if the path has no live master, OpenSSH falls back to a fresh connection. The connectivity probe in Step 2 Check 2 (`ssh -o ConnectTimeout=5 <ssh_host> \"echo ok\"`) deliberately uses bare `ssh` to test un-multiplexed access; everywhere else use `<ssh_cmd>` / `<scp_cmd>`.\n\n## Step 3: Remote Environment Check\n\nVerify the remote environment is ready:\n\n```bash\n<ssh_cmd> <ssh_host> \"test -d <remote_cwd> && echo ok\"\n<ssh_cmd> <ssh_host> \"which tmux\"\n<ssh_cmd> <ssh_host> \"squeue --version\"\n```\n\nIf `remote_cwd` does not exist, create it or ask the user. If tmux or SLURM\ncommands are not available, stop and inform the user.\n\n## Step 4: Prepare Remote Environment\n\nThis step sets up the user's working directory and clones the TensorRT-LLM\nrepository on the remote cluster. It runs once before any job submission.\n\n> **MANDATORY: two-mechanism sync — both must run.**\n> - **Tracked changes** → temp branch (**4c–4f**): commit on a session-scoped\n>   branch, push to `origin`, fetch + checkout on remote, delete the ref both\n>   sides.\n> - **Untracked files** (`git ls-files --others --exclude-standard`) →\n>   `tar | ssh` pipe (**4e.1**), right after the remote checkout. Never touches\n>   `origin`.\n>\n> **Do not** shortcut with `git fetch origin <repo_branch> && git checkout <repo_branch>` —\n> `repo_branch` may not exist on `origin`, may be stale, or may lack uncommitted\n> edits. Always run 4c–4f **and** 4e.1.\n\n**Inputs received from `trtllm-case-executor`:**\n\n| Field | Description |\n|-------|-------------|\n| `default_user_root_dir` | Per-user root directory on the remote cluster (from `exec-env-check` / `internal-env-info`) |\n| `repo_url` | Git remote URL of the local TensorRT-LLM repo (from `git remote get-url origin`). **Used to clone the bare repository on the remote (Step 4b) and to push the temp branch to (Step 4d).** |\n| `repo_branch` | **Informational only** — the local current branch name (from `git rev-parse --abbrev-ref HEAD`). Used in Step 4c to remember which branch to restore locally after the temp branch is deleted. **Never used as the ref the remote checks out** — the remote always checks out the temp branch from Step 4c. |\n| `device_type` | Required short GPU name (e.g., `B200`, `B300`). Used by Step 1a auto-selection. Ignored when an explicit cluster was provided. |\n| `total_required_devices` | Total GPU processes across all nodes. Used by Step 1a auto-selection only. SLURM job sizing reads `node_count` from `job_spec.json` (precomputed by `trtllm-case-executor`) — do not recompute it here. |\n| `required_devices_per_node` | Minimum GPUs per node. Used by Step 1a auto-selection (cluster must have `gpus_per_node >= required_devices_per_node`) and as `--ntasks-per-node`. |\n\n**Procedure:**\n\n### 4a. Create user root directory\n\n```bash\n<ssh_cmd> <ssh_host> \"mkdir -p <default_user_root_dir>\"\n```\n\nIf `default_user_root_dir` is null or empty, skip this step and use `remote_cwd` directly.\n\n### 4b. Clone or update the repository\n\nDetermine the repo directory name from `repo_url` (strip `.git` suffix if present):\n\n```bash\nREPO_DIR_NAME=$(basename <repo_url> .git)\nREMOTE_REPO_PATH=\"<default_user_root_dir>/${REPO_DIR_NAME}\"\n```\n\nCheck if it already exists:\n\n```bash\n<ssh_cmd> <ssh_host> \"test -d <remote_repo_path>/.git && echo exists\"\n```\n\nIf the repository does **not** exist, clone it:\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <default_user_root_dir> && git clone <repo_url>\"\n```\n\nIf the repository **already exists**, refresh remote refs only — do **not**\ncheckout or reset to any branch here:\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <remote_repo_path> && git fetch origin --prune\"\n```\n\nThe actual ref the remote will check out is the temp branch created in 4c\nand pushed in 4d; the remote checkout happens in 4e. **Do not** add a\n`git checkout <repo_branch>` or `git reset --hard origin/<repo_branch>`\nhere as a shortcut — that path is forbidden (see the MANDATORY callout\nat the top of Step 4).\n\n### 4c. Create a temporary local branch capturing all local changes\n\nGenerate a unique temp branch name tied to this session. The branch name\n**must start with `<user_name>_<model_name>`** so it is easy to identify\nthe owner and the workload it was created for. A timestamp suffix keeps\nit unique across sessions:\n\n```bash\n# Resolve the user name (prefer the slurm_user input over whoami,\n# so the temp branch is attributed to the cluster user rather than the\n# host that happens to dispatch the job).\nUSER_NAME=\"${slurm_user:-$(whoami)}\"\n\n# Resolve the workload name. job_spec.json's `model_name` is authoritative;\n# fall back to the repo basename if it is missing.\nMODEL_NAME=\"${model_name:-$(basename \"<REPO_ROOT>\")}\"\n\n# Sanitize to a git-ref-safe identifier (alphanum + `_`/`-` only).\n# Any other character is replaced with `_`.\nsanitize() { printf '%s' \"$1\" | tr -c 'A-Za-z0-9_-' '_'; }\nUSER_NAME_SAFE=$(sanitize \"$USER_NAME\")\nMODEL_NAME_SAFE=$(sanitize \"$MODEL_NAME\")\n\nTEMP_BRANCH=\"${USER_NAME_SAFE}_${MODEL_NAME_SAFE}_$(date +%Y%m%d%H%M%S)\"\n```\n\nExample resolved value: `<user_name>_dwdp_deepseek_v3_lite_20260430145200`.\n\nStash the current branch name so it can be restored later:\n\n```bash\nORIGINAL_BRANCH=$(git -C \"<REPO_ROOT>\" rev-parse --abbrev-ref HEAD)\n```\n\nCreate the temp branch at HEAD and stage tracked changes only\n(`-u`, not `-A` — untracked files ship in 4e.1):\n\n```bash\ngit -C \"<REPO_ROOT>\" checkout -b \"$TEMP_BRANCH\"\ngit -C \"<REPO_ROOT>\" add -u\ngit -C \"<REPO_ROOT>\" commit --allow-empty -m \"sync: temp snapshot for remote execution\"\n```\n\n`--allow-empty` covers the no-tracked-changes case; the branch still exists\nat HEAD and 4e.1 still runs.\n\n### 4d. Push temp branch to the remote origin\n\nPush the temp branch so the remote cluster can fetch it:\n\n```bash\ngit -C \"<REPO_ROOT>\" push origin \"$TEMP_BRANCH\"\n```\n\n### 4e. Fetch and check out temp branch on the remote cluster\n\nOn the remote: stash prior-session leftovers, fetch, force-checkout the temp\nbranch, and re-materialize `cpp/` from HEAD (stale `cpp/` produces\nmissing-symbol `ImportError`s in Step 5):\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <remote_repo_path> && \\\n    git stash push -u -m 'auto-stash before remote sync' || true && \\\n    git fetch origin --prune && \\\n    git checkout -B $TEMP_BRANCH origin/$TEMP_BRANCH && \\\n    rm -rf cpp/ && \\\n    git checkout HEAD -- cpp/\"\n```\n\n- `stash -u` sweeps both tracked edits and untracked files left from a prior\n  session; `|| true` covers the no-op.\n- `checkout -B` force-resets even if a same-named branch already exists.\n- `rm -rf cpp/ && git checkout HEAD -- cpp/` always runs, even when checkout\n  reports no diffs.\n\n`remote_cwd` is updated to `<remote_repo_path>`. Untracked files for this\nsession arrive next in 4e.1.\n\n### 4e.1. Transfer untracked files via tar pipe over SSH\n\nShip untracked, non-gitignored files on top of the freshly checked-out tree:\n\n```bash\nUNTRACKED_COUNT=$(git -C \"<REPO_ROOT>\" ls-files --others --exclude-standard | wc -l)\nif [ \"$UNTRACKED_COUNT\" -gt 0 ]; then\n    git -C \"<REPO_ROOT>\" ls-files --others --exclude-standard -z | \\\n        tar -C \"<REPO_ROOT>\" --null -T - -cf - | \\\n        <ssh_cmd> <ssh_host> \"tar -C <remote_repo_path> -xf -\"\nfi\n```\n\n`-z` / `--null` keeps filenames with spaces or newlines intact; `tar -x`\nauto-creates parent dirs; the guard skips the round-trip when nothing to\nship. Do **not** substitute `rsync` or `sshfs` (not reliably available on HPC).\n\n### 4f. Delete the temp branch (local and remote)\n\nAfter the remote checkout succeeds, immediately delete the temp branch from\nthe local repo and the remote origin, then restore the original local branch:\n\n```bash\ngit -C \"<REPO_ROOT>\" checkout \"$ORIGINAL_BRANCH\"\ngit -C \"<REPO_ROOT>\" branch -D \"$TEMP_BRANCH\"\ngit -C \"<REPO_ROOT>\" push origin --delete \"$TEMP_BRANCH\"\n```\n\nThe remote cluster keeps the checked-out working tree from the deleted branch\n— only the ref is removed, not the files.\n\n### 4f.1 Create remote work_dir under the cloned repo\n\nThe remote work_dir must mirror the local layout (`<REPO_ROOT>/work_dirs/<name>`)\nso that all per-run artefacts — slurm script, stdout/stderr logs, intermediate\nfiles — live inside the cloned repo and are co-located with the code that\nproduced them.\n\n`remote_work_dir` is precomputed by `trtllm-case-executor` (Step 1) and read\nverbatim from `job_spec.json` — it is `<remote_repo_path>/work_dirs/<basename(local_work_dir)>`.\n**Do not** fall back to `<remote_cwd>/work_dirs/...` or any other root; that\nolder layout is deprecated.\n\n```bash\n<ssh_cmd> <ssh_host> \"mkdir -p <remote_work_dir>\"\n```\n\nAll subsequent steps (file uploads, sbatch submission, log paths) reference\n`<remote_work_dir>` from `job_spec.json`.\n\n### 4g. Copy build scripts to remote node\n\nThe skill ships two build scripts in `scripts/` alongside this SKILL.md.\nCopy them to a dedicated directory on the remote so they are accessible\nfrom inside the SLURM container:\n\n```bash\nREMOTE_SCRIPTS_DIR=\"<default_user_root_dir>/trtllm-build-scripts\"\nSKILL_SCRIPTS_DIR=\"<LOCAL_SKILL_DIR>/scripts\"   # local path to this skill's scripts/\n\n<ssh_cmd> <ssh_host> \"mkdir -p ${REMOTE_SCRIPTS_DIR}/log\"\n\n<ssh_cmd> <ssh_host> \"cat > ${REMOTE_SCRIPTS_DIR}/build.slurm\" < \"${SKILL_SCRIPTS_DIR}/build.slurm\"\n<ssh_cmd> <ssh_host> \"cat > ${REMOTE_SCRIPTS_DIR}/build.sh\"    < \"${SKILL_SCRIPTS_DIR}/build.sh\"\n<ssh_cmd> <ssh_host> \"chmod +x ${REMOTE_SCRIPTS_DIR}/build.slurm ${REMOTE_SCRIPTS_DIR}/build.sh\"\n```\n\n`<LOCAL_SKILL_DIR>` is the directory containing this SKILL.md file\n(`.claude/skills/exec-remote-slurm/`).\n\n---\n\n## Step 5: Build on Remote SLURM Node\n\n> **Mandatory unless `build_project: false`.** Step 4e wipes `cpp/`, so\n> the remote has no compiled C++ extensions until Step 5 succeeds (exit\n> code `0:0`). Submitting any test/bench job before then fails on the\n> first `import tensorrt_llm` with missing-symbol errors.\n>\n> Do **not** replace this with `pip install -e .` inside the user's\n> `.slurm` script — it does not reliably recompile `cpp/`, and when it\n> does it blocks the GPU allocation behind a 30+ minute compile. Step 5\n> runs the build as a 1-node side job and installs into the persistent\n> container volume.\n>\n> If the caller sets `build_project: false`, skip Step 5 (the caller is\n> asserting the install is current). Do not silently downgrade `true` →\n> `false` because Step 5 looks expensive.\n\nSubmit a SLURM batch job that compiles TensorRT-LLM on the remote cluster\ninside the official container. Wait for it to complete before proceeding to\njob submission.\n\n### 5a. Read the container image\n\nThe container image is resolved by `trtllm-case-executor` (Step 3) and\nforwarded in `job_spec.json` as `container_image`. Read it from there —\n**do not grep `current_image_tags.properties` and do not re-derive from\nthe cluster config**. case-executor is the single source of truth for the\nimage URI; this skill only consumes it.\n\nApply the enroot URI rewrite at use-time so the registry path is not\nmistaken for a Docker-style credential lookup. The rewrite recipe lives\nin the `internal-env-info` skill's `Default Environment Values` section\n(`enroot_uri_rewrite` key):\n\n- **Skill installed** → delegate to it, passing the canonical\n  `<container_image>` from `job_spec.json`. The skill returns the bash\n  snippet (an `sed` over the registry-host prefix); run it here to\n  produce `<IMAGE_PATH>`.\n- **Skill not installed** → use `<container_image>` verbatim as `<IMAGE_PATH>`.\n\nUse `<IMAGE_PATH>` as the value passed to `--container-image=` in the\nbuild sbatch below and in every `srun`/`sbatch` recipe later in this\nskill. The canonical (un-rewritten) URI in `job_spec.json` is unchanged;\nthe rewrite is purely a transport adaptation for enroot.\n\n### 5b. Derive build architecture\n\nMap `device_type` (from the env-check output) to the `build_wheel.py` `-a` flag:\n\n| `device_type` | `-a` value |\n|------------|------------|\n| B200, GB200, GB300, B300 | `100-real` |\n| H100, H200 | `90-real` |\n| L40S | `89-real` |\n| A100 | `80-real` |\n\nDefault to `100-real` if `device_type` is unknown.\n\n### 5c. Submit the build job\n\n```bash\n<ssh_cmd> <ssh_host> \"sbatch \\\n    --nodes=1 --ntasks=1 --ntasks-per-node=1 \\\n    --partition=<partition> \\\n    --account=<account> \\\n    --job-name=<account>-trtllm.build \\\n    --time=02:00:00 \\\n    ${REMOTE_SCRIPTS_DIR}/build.slurm \\\n    <IMAGE_PATH> <default_user_root_dir> ${REMOTE_SCRIPTS_DIR} <remote_repo_path> \\\n    -a \\\"<arch>\\\" --nvtx\"\n```\n\nCapture the job ID from `sbatch` stdout (format: `Submitted batch job <id>`).\n\nIf `gpus_per_node` is set in the cluster config, add\n`--gpus-per-node=<gpus_per_node>` to the `sbatch` flags.\n\n### 5d. Monitor the build job\n\nPoll every 30–60 seconds until the job leaves the queue:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <job_id> -o '%.18i %.9P %.30j %.8u %.2t %.10M %R'\"\n```\n\nOnce running, tail the log to track compile progress:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -30 ${REMOTE_SCRIPTS_DIR}/log/build_<job_id>.srun.log\"\n```\n\nProgress indicators in the log:\n- `[XX%] Building CXX object ...` — compilation in progress\n- `Linking CXX ...` — link phase\n- `FAILED:` / `error:` / `fatal error:` — build failure\n- `Successfully built` — success\n\nIf the job disappears from `squeue`, proceed to Step 5e.\n\n### 5e. Verify the build\n\n```bash\n<ssh_cmd> <ssh_host> \"sacct -j <job_id> --format=JobID,State,ExitCode,Elapsed -n\"\n<ssh_cmd> <ssh_host> \"tail -50 ${REMOTE_SCRIPTS_DIR}/log/build_<job_id>.srun.log\"\n```\n\nA successful build shows exit code `0:0` and ends with\n`Successfully built tensorrt_llm` or no error lines.\n\nIf the build fails, report the last 100 lines of the log to the user and\nstop — do not proceed to job submission.\n\n## Build Scripts Reference\n\n| Script | Location | Purpose |\n|--------|----------|---------|\n| `build.slurm` | `scripts/build.slurm` | SLURM batch script — launches container, calls `build.sh` |\n| `build.sh` | `scripts/build.sh` | Runs inside container — calls `build_wheel.py` with defaults |\n\nBoth scripts accept the same positional arguments:\n`<container_image> <user_root_dir> <scripts_dir> <repo_dir> [build_wheel_args...]`\n\n---\n\n## Recipes\n\nPick the recipe that matches your workflow. Each recipe is self-contained —\ncopy it, substitute the `<placeholders>` from config, and run.\n\n**Placeholders** (all from `job_spec.json` unless noted):\n\n| Placeholder | Source |\n|-------------|--------|\n| `<ssh_host>`, `<account>`, `<partition>`, `<mounts>` | `job_spec.json` |\n| `<image>` | `container_image` from `job_spec.json` (resolved by `trtllm-case-executor` Step 3). Apply enroot rewrite (`/` → `#`) at use-time per Step 5a — substitute the rewritten `<IMAGE_PATH>` here. Do not re-resolve from the cluster config or `current_image_tags.properties`. |\n| `<job_name>` | Pre-resolved by `trtllm-case-executor` and forwarded as `job_name` in `job_spec.json`. Substitute directly into `-J <job_name>`; do not reconstruct from `account` / `subproject` / `detail`. |\n| `<nodes>` | `node_count` from `job_spec.json` (computed by `trtllm-case-executor`). Do not recompute. |\n| `<gpus>` | `required_devices_per_node` from `job_spec.json` — probe only if absent (see bottom of this section) |\n| `<remote_repo_path>` | `remote_repo_path` from `job_spec.json` — `<default_user_root_dir>/<repo_dir_name>` precomputed by `trtllm-case-executor` Step 1. The cloned repo and `--container-workdir` for srun. Do not recompute. |\n| `<remote_work_dir>` | `remote_work_dir` from `job_spec.json` — `<remote_repo_path>/work_dirs/<basename(local_work_dir)>` precomputed by `trtllm-case-executor` Step 1. Used for `--output`/`--error` slurm log paths and any per-run artefact uploads. Do not recompute or fall back to `<remote_cwd>/work_dirs/...`. |\n| `<time>` | Walltime in HH:MM:SS |\n\n**Job naming:** Every recipe includes `-J <job_name>` using the\npre-resolved `job_name` value from `job_spec.json` (built by\n`trtllm-case-executor`). This is required — without `-J`, SLURM emits\nwarnings and jobs are hard to identify in `squeue`. Do not reconstruct\nthe name from `account` / `subproject` / `detail` here; that convention\nlives in `trtllm-case-executor`.\n\n**GPU allocation:** Cluster behavior varies:\n\n- **Exclusive-access clusters** (most NVIDIA internal): Nodes are allocated\n  exclusively with all GPUs available. No `--gpus-per-node` needed.\n- **Shared/QOS clusters**: Require `--gpus-per-node=<gpus_per_node>`. Without\n  it, you may get a CPU-only node (no GPUs at all), a \"Cannot find GPU\n  specification\" error, or a \"QOSMinGRES\" error (below minimum).\n\nIf `gpus_per_node` is set in the cluster config, add `--gpus-per-node=<gpus_per_node>`\nto every `srun`, `salloc`, and `sbatch` command. If not set, omit it.\n\nIf a job fails with GPU-related errors and `gpus_per_node` is not configured,\nprobe the node's GPU count, set `gpus_per_node` in the cluster config, and retry.\n\nDo not set `CUDA_VISIBLE_DEVICES` for MPI workloads — the MPI runtime handles\nGPU assignment.\n\n---\n\n### Recipe 1: Single-process command (srun)\n\nOne-off commands on a single node: profiling, inspecting, lightweight scripts.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N 1 --ntasks-per-node=1 \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <command>\"\n```\n\n---\n\n### Recipe 2: MPI workload (srun)\n\nMPI applications across nodes: NCCL tests, HPCX benchmarks. The binary must\nbe MPI-aware (e.g., `all_reduce_perf_mpi`, not `all_reduce_perf`).\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=<gpus> --mpi=pmix \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <command>\"\n```\n\n---\n\n### Recipe 3: Distributed training (srun)\n\nFramework-launched distributed training (PyTorch DDP, DeepSpeed). SLURM starts\n`<nodes> × <gpus>` processes. Pyxis/enroot sets `RANK`, `LOCAL_RANK`, and\n`WORLD_SIZE`. Each process runs the training script directly.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=<gpus> \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  python train.py\"\n```\n\n---\n\n### Recipe 4: Launcher-based training (srun)\n\nOne process per node, an in-container launcher forks locally (torchrun,\ndeepspeed launcher).\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=1 \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  torchrun --nproc-per-node=<gpus> train.py\"\n```\n\n---\n\n### Recipe 5: Batch job (sbatch)\n\nFire-and-forget background job:\n\n```bash\n<ssh_cmd> <ssh_host> \"sbatch -A <account> -p <partition> -N <nodes> \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <script-path>\"\n```\n\nCheck status:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <jobid>\"\n<ssh_cmd> <ssh_host> \"sacct -j <jobid> --format=JobID,State,ExitCode,Elapsed -n\"\n```\n\n---\n\n### Recipe 6: Multi-command session (salloc + srun)\n\nInteractive development with persistent container state. Use this when you need\nmultiple commands on the same allocation — installing packages, compiling,\nrunning experiments iteratively.\n\n> **Relationship to `persistent_mode`:** The `persistent_mode` field in\n> `job_spec.json` is consumed only by `exec-local-slurm`. This\n> skill does **not** read `persistent_mode`, `container_name`, or\n> `container_mounts` from `job_spec.json` for allocation reuse — remote\n> persistence is driven by this recipe (tmux + named `salloc` + named\n> container) and controlled by inputs to this skill directly (e.g.,\n> `release_allocation`, `alloc_time_limit`). If a caller forwards a\n> `persistent_mode=true` job spec into a `remote_slurm` run, treat the field\n> as informational only and use this recipe's flow.\n\n**Step 1: Allocate via tmux**\n\ntmux keeps the allocation alive on the frontend — invisible plumbing.\n\n```bash\n<ssh_cmd> <ssh_host> \"tmux new-session -d -s trtllm-<id> \\\n  'salloc -A <account> -p <partition> -N <nodes> \\\n  -J <job_name> -t <time>'\"\n```\n\n**Step 2: Check allocation status**\n\nWait ~30 seconds for the scheduler, then get job ID and status:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -u <user> -h -o '%i %T %S' --state=RUNNING,PENDING\"\n```\n\nIf RUNNING → proceed to Step 3.\n\nIf PENDING → follow the **walltime negotiation flow** (see Walltime strategy\nsection). Only proceed to Step 3 once the job is RUNNING.\n\n**Step 3: Create a persistent container**\n\nImport the image once. This creates a named container on each node that\npersists across `srun` calls (container persistence):\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-image=<image> --container-name=trtllm-<jobid> true\"\n```\n\n**Container packages:** If `packages` is defined in config, install them\nusing the exact names listed — do not guess or substitute package names:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> --container-writable \\\n  pip install <packages>\"\n```\n\nOnly install when the workload actually needs them — not for simple\ncommands like `nvidia-smi` or `hostname`.\n\n**Step 4: Run commands**\n\nUse `--container-name` (without `--container-image`) to reuse the persistent\ncontainer. All in-container state is preserved between commands.\n\nSingle-process command:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nPer-node operation (mkdir, pip install) — use `--ntasks-per-node=1`:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nMPI workload within the allocation:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --ntasks-per-node=<gpus> --mpi=pmix \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nDistributed training within the allocation:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --ntasks-per-node=<gpus> \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  python train.py\"\n```\n\nAdd `--container-writable` if the workload writes to container paths outside\nof mounted directories (e.g., `pip install` into site-packages).\n\n**Step 5: Release the allocation**\n\nClean up the named container, then kill tmux. If the allocation has already\nexpired (walltime exceeded), the `srun` will fail — that is fine, the container\nis already gone. Just kill the tmux session.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  bash -c 'enroot remove -f trtllm-<jobid> 2>/dev/null; true'\" 2>/dev/null\n<ssh_cmd> <ssh_host> \"tmux kill-session -t trtllm-<id>\"\n```\n\n---\n\n### Probing GPUs per node\n\nIf you do not know `<gpus>`, probe once and record in cluster notes:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> bash -c 'nvidia-smi -L | wc -l'\"\n```\n\n## Cluster-Specific Flags\n\nAfter substituting recipe placeholders, apply the cluster-specific flags\nreturned by the `internal-env-info` skill in Step 1:\n\n- **SHARP**: Add `--network=sharp` if the cluster supports it\n- **Topology constraints**: Add `--constraint=<feature>` for topology-aware\n  placement (e.g., `--constraint=nvlblk01` for a single rack)\n- **GPU clocks**: Run `sudo nvidia-smi -ac <mem>,<graphics>` as a separate\n  `srun --ntasks-per-node=1` pre-step in sbatch scripts\n- **Comments**: Add `--comment=metrics` for Prometheus/Grafana collection\n\n## File Transfer\n\nUse SSH pipes for file transfer — no extra tools needed, works on every cluster.\n\n### Local → Remote (writing files)\n\nPipe file content directly via SSH heredoc. Never write locally then scp:\n\n```bash\n<ssh_cmd> <ssh_host> \"cat > <remote_cwd>/script.py\" << 'PYEOF'\nimport torch\n# ... file content ...\nPYEOF\n```\n\nFor multiple files, repeat the pattern — each file is one SSH command.\n\n### Remote → Local (reading results)\n\nRead remote files via SSH stdout:\n\n```bash\n<ssh_cmd> <ssh_host> \"cat <remote_cwd>/output.json\"\n```\n\nFor large files, read only what you need:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/slurm-<jobid>.out\"\n<ssh_cmd> <ssh_host> \"head -n 50 <remote_cwd>/results.csv\"\n```\n\n### When to use scp\n\nFall back to `scp` only for binary files or large directories that cannot be\npiped through SSH:\n\n```bash\n<scp_cmd> <ssh_host>:<remote_cwd>/profile.nsys-rep .    # binary profile\n<scp_cmd> -r local_dir/ <ssh_host>:<remote_cwd>/        # whole directory\n```\n\n**Decision rule:** text files → SSH pipe. Binary files or directories → scp.\nNever use rsync or sshfs — not reliably available on HPC clusters.\n\n## Rules\n\n### Never read full remote files\n\nRemote output can be gigabytes. Always use `tail`, `head`, or `grep`:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/output.log\"\n<ssh_cmd> <ssh_host> \"head -n 50 <remote_cwd>/script.py\"\n<ssh_cmd> <ssh_host> \"grep -n 'error\\|Error\\|ERROR' <remote_cwd>/output.log\"\n```\n\nNever use `cat` on remote files unless you are certain the file is small.\n\n### Never run heavy computation on the frontend\n\nThe frontend node is shared. Only use it for SLURM commands, lightweight file\noperations, and tmux management.\n\n### Check before creating\n\nBefore creating resources, check if they already exist:\n- SSH ControlMaster: `ssh -O check <ssh_host> 2>&1 || <ssh_cmd> -O check <ssh_host>` (probe both the user's default socket and the trtllm-managed socket from Step 2)\n- tmux sessions: `<ssh_cmd> <ssh_host> \"tmux has-session -t trtllm-<id> 2>/dev/null && echo exists\"`\n- Existing allocations: `<ssh_cmd> <ssh_host> \"squeue -u <user> -h -o '%i %j' --state=RUNNING\"`\n\n### Clean up when done\n\n- Remove named containers: `enroot remove -f` (see Recipe 6, Step 5)\n- Kill tmux sessions: `<ssh_cmd> <ssh_host> \"tmux kill-session -t trtllm-<id>\"`\n- Cancel unneeded jobs: `<ssh_cmd> <ssh_host> \"scancel <jobid>\"`\n\n### Walltime strategy\n\n- **Production (sbatch, srun):** Tight walltime with ~20% margin — gets\n  scheduled sooner.\n- **Dev/debug (salloc):** Use **walltime negotiation** — submit with max\n  walltime, then reduce if the queue wait is too long:\n\n**Walltime negotiation flow:**\n\n1. Submit `salloc` with **max walltime** for the partition.\n2. Wait ~30 seconds for the scheduler to process, then check:\n   ```bash\n   <ssh_cmd> <ssh_host> \"squeue -j <jobid> -h -o '%i %T %S'\"\n   ```\n   - `%S` = estimated start time. Parse it to calculate wait.\n3. **If RUNNING** → proceed to next step.\n4. **If PENDING and wait ≤ 10 minutes** → wait for it, proceed.\n5. **If PENDING and wait > 10 minutes** (or no start time shown):\n   a. Infer how long you actually need from context (profiling → ~30 min,\n      iterative dev → ~2h, training run → match expected duration).\n      Round up to the next hour boundary.\n   b. Reduce walltime in-place:\n      ```bash\n      <ssh_cmd> <ssh_host> \"scontrol update JobId=<jobid> TimeLimit=<new_time>\"\n      ```\n   c. Tell the user: \"Reduced walltime from Xh to Yh for faster scheduling —\n      you'll have less time on the allocation.\"\n   d. Wait ~15 seconds, re-check estimated start time.\n   e. **If still > 10 minutes:** inform user with the ETA and suggest\n      alternatives — different partition (e.g., backfill), fewer nodes.\n\n## Hang Detection Monitoring\n\nWhen monitoring a running job (polling via `squeue`), also check the remote log\nfile every 60 seconds for the pattern `hang detected` (case-insensitive). This\ncatches GPU hangs, NCCL deadlocks, and other stalls that keep the job running\nindefinitely without making progress.\n\n**Check for hang (run alongside each `squeue` poll once job is RUNNING):**\n\n```bash\n<ssh_cmd> <ssh_host> \"grep -qi 'hang detected' <remote_work_dir>/slurm-<jobid>.out 2>/dev/null && echo HANG_DETECTED\"\n```\n\nAlso check the `.err` file if the output file has no match:\n\n```bash\n<ssh_cmd> <ssh_host> \"grep -qi 'hang detected' <remote_work_dir>/slurm-<jobid>.err 2>/dev/null && echo HANG_DETECTED\"\n```\n\n**If `HANG_DETECTED` is printed:**\n\n1. Immediately cancel the job:\n   ```bash\n   <ssh_cmd> <ssh_host> \"scancel <jobid>\"\n   ```\n2. Read the last 200 lines of the log file for context:\n   ```bash\n   <ssh_cmd> <ssh_host> \"tail -n 200 <remote_work_dir>/slurm-<jobid>.out\"\n   ```\n3. Report status as `HANG_DETECTED` with the relevant log lines. Do not\n   continue polling — the job is terminated.\n\n## Recovering from Failures\n\n**SSH connection drops:** Re-run SSH preflight (Step 2). tmux-held allocations\nsurvive disconnects.\n\n**Job stuck in PENDING:**\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <jobid> -o '%i %T %r %S'\"\n```\n\nCheck the reason (`%r`) and estimated start (`%S`). If wait > 10 minutes,\nfollow the **walltime negotiation flow** (see Walltime strategy). If already\nnegotiated and still stuck, suggest alternatives: backfill partition, fewer\nnodes, or a different cluster.\n\n**Job fails:** Read last lines, never dump full output:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/slurm-<jobid>.out\"\n<ssh_cmd> <ssh_host> \"tail -n 50 <remote_cwd>/slurm-<jobid>.err\"\n```\n\n**Hang detected:** If the log file contains `hang detected`, the job is\nauto-cancelled (see Hang Detection Monitoring section above). Report\n`HANG_DETECTED` with the last 200 lines of log context.\n","frontmatter":{"name":"exec-remote-slurm","description":"Remote SLURM cluster development via SSH. Use when running jobs, profiling, or developing on a remote SLURM cluster with pyxis/enroot containers. Covers SSH connection management, srun/sbatch/salloc job patterns, tmux-based allocation persistence, file transfer, and safe remote file access. Works with any SLURM cluster accessible via SSH.","tags":["infrastructure","remote","slurm"],"license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":10397,"sizeBytes":37863},{"name":"SKILL.md","path":".claude/skills/exec-slurm-compile/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/SKILL.md","title":"exec-slurm-compile","category":"anthropic-skill","format":"markdown","content":"---\nname: exec-slurm-compile\ndescription: Compile TensorRT-LLM on a SLURM cluster. Covers submitting a batch job with a container image, monitoring the job, and verifying the build. Use when the user wants to compile TRT-LLM remotely via SLURM rather than on a local compute node.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Compile TensorRT-LLM on SLURM Cluster\n\nSubmit, monitor, and verify a TensorRT-LLM compilation job on a SLURM cluster using enroot containers.\n\n## When to Use\n\n| Scenario | Use This Skill? |\n|----------|----------------|\n| User wants to compile TRT-LLM on a SLURM cluster | Yes |\n| User is already on a compute node and wants to compile | No — use `exec-local-compile` skill instead |\n\n## Finding the Docker Image\n\nThe official Docker image tag for a given TensorRT-LLM version is recorded in the repo itself:\n\n```\n<repo_dir>/jenkins/current_image_tags.properties\n```\n\nRead this file to find the current image URL. The format is `<registry-host>/<namespace>/tensorrt-llm:<tag>`, for example `<your-registry>/<your-namespace>/tensorrt-llm:pytorch-25.12-py3-aarch64-ubuntu24.04-trt10.14.1.48-skip-tritondevel-202602011118-10901`. Substitute `<your-registry>` and `<your-namespace>` with your own container registry coordinates.\n\n\n## Pre-dumping the Container Image (enroot import)\n\nSLURM clusters using enroot/pyxis require a `.sqsh` container image. To avoid download overhead at compile time, **pre-dump the image in advance** using the `enroot-import` companion script:\n\n```bash\n# Basic usage — submits a SLURM job on a CPU partition to import the image\nenroot-import --partition cpu_datamover --debug <docker_image_url>\n```\n\nThe script submits an `sbatch` job that runs `enroot import docker://<image_url>` and produces a `.sqsh` file in the current directory. The output on stdout is the SLURM job ID.\n\n### enroot-import flags\n\n| Flag | Description |\n|------|-------------|\n| `-p, --partition` | SLURM partition for the import job (use a CPU partition like `cpu_datamover`) |\n| `-d, --debug` | Enable debug output and preserve the SLURM log (recommended) |\n| `-o, --output` | Custom output path for the `.sqsh` file |\n| `-A, --account` | SLURM account (defaults to user's first account) |\n| `-t, --time` | Time limit for the import job (default: 1 hour) |\n| `-n, --just-print` | Print the sbatch command without executing |\n| `-J, --job-name` | Custom job name |\n\n### enroot-import workflow\n\n1. Read the image tag from `jenkins/current_image_tags.properties` in the TRT-LLM repo.\n2. Run `enroot-import` to submit the import job:\n   ```bash\n   cd <directory_where_sqsh_should_be_stored>\n   <path_to>/enroot-import --partition cpu_datamover --debug <image_url>\n   ```\n   **IMPORTANT:** Replace the first `/` after the registry host with `#` to avoid credential issues. For example, `<your-registry>/<your-namespace>/tensorrt-llm:xxx` becomes `<your-registry>#<your-namespace>/tensorrt-llm:xxx`.\n3. Wait for the import job to complete (`squeue -j <job_id>`).\n4. The resulting `.sqsh` file is the `container_image` used in the compile step.\n\n\n## Prerequisites\n\nThe user must provide (or you must ask for) these values:\n\n| Parameter | Description | Example |\n|-----------|-------------|---------|\n| `container_image` | Path to `.sqsh` container image (see enroot import above) | `/path/to/pytorch.sqsh` |\n| `repo_dir` | Path to the TensorRT-LLM repository | `/path/to/TensorRT-LLM` |\n| `user_root_dir` | Top-level directory to bind-mount into the container | `/shared/users` |\n| `partition` | SLURM partition | `batch` |\n| `account` | SLURM account | `my_account` |\n\nOptional parameters:\n\n| Parameter | Description | Default |\n|-----------|-------------|---------|\n| `jobname` | SLURM job name | `trtllm-compile.<username>` |\n| `gpu_count` | Number of GPUs to request | `4` |\n| `time_limit` | Job time limit | `02:00:00` |\n| `arch` | GPU architecture(s) for `-a` flag | `100-real` |\n| `extra_build_args` | Extra flags for `build_wheel.py` | (none) |\n\n## Companion Scripts\n\nThis skill includes three companion scripts in `scripts/`:\n\n| Script | Purpose |\n|--------|---------|\n| `enroot-import` | Pre-dump a Docker image to `.sqsh` via a SLURM batch job |\n| `submit_compile.sh` | Template for submitting the SLURM job — copy and customize |\n| `compile.slurm` | SLURM batch script — launches the container and calls `compile.sh` |\n| `compile.sh` | Runs inside the container — executes `build_wheel.py` |\n\nScripts directory: `skills/exec-slurm-compile/scripts/`\n\n## Instructions\n\nFollow these steps in order:\n\n### Step 0: Resolve the Container Image (if needed)\n\nIf the user does not already have a `.sqsh` container image:\n\n1. Read the Docker image tag from `<repo_dir>/jenkins/current_image_tags.properties`.\n2. Use `enroot-import` to pre-dump it:\n   ```bash\n   cd <directory_for_sqsh_files>\n   <scripts_dir>/enroot-import --partition cpu_datamover --debug <image_url>\n   ```\n3. Monitor the import job with `squeue -j <job_id>`.\n4. Once complete, the `.sqsh` file path becomes the `container_image` parameter.\n\nIf the user already has a `.sqsh` file, skip this step.\n\n### Step 1: Gather Information\n\nAsk the user for any missing prerequisite values listed above. At minimum you need:\n- `container_image` (or the Docker image URL — then run Step 0 first)\n- `repo_dir`\n- `user_root_dir`\n- `partition` and `account`\n\nIf the user has used this workflow before, check if previous values are stored in memory files.\n\n### Step 2: Prepare the Scripts Directory\n\nThe compile scripts must be accessible from inside the container (i.e., under `user_root_dir`). Either:\n\n**Option A** — Copy companion scripts to a location under `user_root_dir`:\n```bash\nscripts_dir=<user_root_dir>/<username>/workspace/tensorrt_llm_scripts\nmkdir -p ${scripts_dir}/log\ncp skills/exec-slurm-compile/scripts/compile.sh ${scripts_dir}/\ncp skills/exec-slurm-compile/scripts/compile.slurm ${scripts_dir}/\nchmod +x ${scripts_dir}/compile.sh ${scripts_dir}/compile.slurm\n```\n\n**Option B** — If the user already has scripts at a known location, use those directly.\n\n### Step 3: Submit the Job\n\nRun `sbatch` from the login node (or a node with SLURM client access):\n\n```bash\nsbatch \\\n    --nodes=1 --ntasks=1 --ntasks-per-node=1 \\\n    --gres=gpu:<gpu_count> \\\n    --partition=<partition> \\\n    --account=<account> \\\n    --job-name=<jobname> \\\n    --time=<time_limit> \\\n    <scripts_dir>/compile.slurm \\\n    <container_image> <user_root_dir> <scripts_dir> <repo_dir>\n```\n\nCapture and report the job ID from the `sbatch` output.\n\n### Step 4: Monitor the Job (Proactive — Do NOT Wait for User)\n\n**You MUST actively poll the job until it completes.** Do not submit and walk away.\n\n```bash\n# Check job status (repeat every 30-60 seconds)\nsqueue -j <job_id> -o \"%.18i %.9P %.30j %.8u %.2t %.10M %.6D %R\"\n\n# Once running, periodically tail the log (do NOT use tail -f, use tail -30 instead)\ntail -30 <scripts_dir>/log/compile_<job_id>.srun.log\n```\n\n**Monitoring loop:**\n1. Poll `squeue -j <job_id>` to check state\n2. If `PD` (pending) — report the reason, keep polling every 30-60s\n3. If `R` (running) — tail the build log every 30-60s; look for `[XX%] Building`, errors, or completion\n4. If the job disappears from `squeue`, it has finished — proceed to Step 5\n5. If `F` (failed) — immediately read the full log and report the error\n\n**Progress indicators to look for in the log:**\n- `[XX%] Building CXX object...` — compilation progress\n- `Linking CXX...` — link phase\n- `FAILED:`, `error:`, `fatal error:` — build failure\n- `Successfully built` — success\n\n### Step 5: Verify the Build\n\nOnce the job completes, check for success:\n\n```bash\n# Check SLURM exit code\nsacct -j <job_id> --format=JobID,State,ExitCode,Elapsed\n\n# Check the build log for errors\ntail -50 <scripts_dir>/log/compile_<job_id>.srun.log\n```\n\nA successful build ends with a message like `Successfully built tensorrt_llm` or completes without error.\n\n## Common Build Flags Reference\n\n| Flag | Description |\n|------|-------------|\n| `-a \"100-real\"` | Target architecture — `100` for Blackwell, `90` for Hopper, etc. |\n| `--nvtx` | Enable NVTX markers for profiling |\n| `--no-venv` | Skip virtual environment creation |\n| `--use_ccache` | Use ccache to speed up recompilation |\n| `--skip_building_wheel` | Build in-place without creating a wheel file |\n| `-f` | Fast build — skip some kernels for faster dev compilation |\n| `-c` | Clean build — wipe build directory before building |\n\nCommon architecture values:\n- `\"100-real\"` — Blackwell (B200, GB200)\n- `\"90-real\"` — Hopper (H100, H200)\n- `\"89-real\"` — Ada Lovelace (L40S)\n- `\"80-real\"` — Ampere (A100)\n- `\"90;100-real\"` — Multiple architectures\n\n## Troubleshooting\n\n| Issue | Solution |\n|-------|----------|\n| `sbatch: error: invalid partition` | Verify partition name with `sinfo -s` |\n| `sbatch: error: invalid account` | Check available accounts with `sacctmgr show assoc user=$USER` |\n| Container image not found | Verify the `.sqsh` path exists and is readable |\n| Build OOM (out of memory) | Reduce parallelism with `-j <N>` flag to `build_wheel.py` |\n| `srun: error: Unable to create step` | The node may lack enroot/pyxis — check with cluster admin |\n| Job stuck in `PD` state | Check `squeue -j <id> -o %R` for the reason (e.g., resource limits, priority) |\n| `enroot import` fails with auth error | Check `~/.config/enroot/.credentials` has the correct registry credentials |\n| `enroot import` produces empty/corrupt `.sqsh` | Re-run with `--debug` and check the SLURM log; verify the image URL has no `https://` prefix |\n| Weird compile issues | Retry with a clean build (`-c` flag) |\n| `QOSGrpNodeLimit` shown in `NODELIST(REASON)` | Not a blocker, just wait for the job to get scheduled |\n\n## Example Interaction\n\n**User**: \"Compile TRT-LLM on the OCI cluster\"\n\n**Agent actions**:\n1. Ask for container image path, repo path, mount dir (if not known)\n2. Confirm partition/account for OCI cluster\n3. Copy scripts to accessible location under user_root_dir\n4. Submit with `sbatch`\n5. Report job ID\n6. Monitor with `squeue` until complete\n7. Check logs and report success/failure\n","frontmatter":{"name":"exec-slurm-compile","description":"Compile TensorRT-LLM on a SLURM cluster. Covers submitting a batch job with a container image, monitoring the job, and verifying the build. Use when the user wants to compile TRT-LLM remotely via SLURM rather than on a local compute node.","license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":2733,"sizeBytes":10159},{"name":"api-arch.md","path":".claude/skills/kernel-cute-writing/references/api-arch.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-arch.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-arch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-arch.md)","tokens":300,"sizeBytes":233},{"name":"api-core.md","path":".claude/skills/kernel-cute-writing/references/api-core.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-core.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-core.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-core.md)","tokens":300,"sizeBytes":233},{"name":"api-nvgpu.md","path":".claude/skills/kernel-cute-writing/references/api-nvgpu.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-nvgpu.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-nvgpu.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-nvgpu.md)","tokens":300,"sizeBytes":235},{"name":"api-runtime-utils.md","path":".claude/skills/kernel-cute-writing/references/api-runtime-utils.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-runtime-utils.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-runtime-utils.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-runtime-utils.md)","tokens":300,"sizeBytes":251},{"name":"concepts-architecture.md","path":".claude/skills/kernel-cute-writing/references/concepts-architecture.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-architecture.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-architecture.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-architecture.md)","tokens":300,"sizeBytes":259},{"name":"concepts-layouts.md","path":".claude/skills/kernel-cute-writing/references/concepts-layouts.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-layouts.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-layouts.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-layouts.md)","tokens":300,"sizeBytes":249},{"name":"concepts-mma.md","path":".claude/skills/kernel-cute-writing/references/concepts-mma.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-mma.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-mma.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-mma.md)","tokens":300,"sizeBytes":241},{"name":"concepts-tensors.md","path":".claude/skills/kernel-cute-writing/references/concepts-tensors.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-tensors.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-tensors.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-tensors.md)","tokens":300,"sizeBytes":249},{"name":"patterns-compilation.md","path":".claude/skills/kernel-cute-writing/references/patterns-compilation.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-compilation.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-compilation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-compilation.md)","tokens":300,"sizeBytes":257},{"name":"patterns-elementwise.md","path":".claude/skills/kernel-cute-writing/references/patterns-elementwise.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-elementwise.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-elementwise.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-elementwise.md)","tokens":300,"sizeBytes":257},{"name":"patterns-gemm.md","path":".claude/skills/kernel-cute-writing/references/patterns-gemm.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-gemm.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-gemm.md)","tokens":300,"sizeBytes":243},{"name":"patterns-getting-started.md","path":".claude/skills/kernel-cute-writing/references/patterns-getting-started.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-getting-started.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-getting-started.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-getting-started.md)","tokens":300,"sizeBytes":265},{"name":"patterns-memory.md","path":".claude/skills/kernel-cute-writing/references/patterns-memory.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-memory.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-memory.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-memory.md)","tokens":300,"sizeBytes":247},{"name":"patterns-pipeline.md","path":".claude/skills/kernel-cute-writing/references/patterns-pipeline.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-pipeline.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-pipeline.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-pipeline.md)","tokens":300,"sizeBytes":251},{"name":"patterns-reduction.md","path":".claude/skills/kernel-cute-writing/references/patterns-reduction.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-reduction.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-reduction.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-reduction.md)","tokens":300,"sizeBytes":253},{"name":"troubleshooting.md","path":".claude/skills/kernel-cute-writing/references/troubleshooting.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/troubleshooting.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/troubleshooting.md)","tokens":300,"sizeBytes":247},{"name":"SKILL.md","path":".claude/skills/kernel-cute-writing/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/SKILL.md","title":"kernel-cute-writing","category":"anthropic-skill","format":"markdown","content":"---\nname: kernel-cute-writing\ndescription: >\n  Write and implement GPU kernels using NVIDIA CuTe DSL (CUTLASS 4.x Python\n  API) — NOT for Triton, CUDA C++, or conceptual explanations.\n  Trigger only when the user wants to write or implement a kernel, not when\n  asking questions about CuTe DSL concepts or layouts.\n  CuTe DSL uses cute.jit/cute.kernel decorators and cutlass.cute imports.\n  Covers element-wise kernels, GEMM patterns, reductions, memory hierarchy\n  (global/shared/register/TMA), MMA tensor core operations, software\n  pipelining, and framework integration.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# CuTe DSL\n\nCuTe DSL is a Python-based domain-specific language for GPU kernel development,\npart of CUTLASS 4.x. It provides Python abstractions over CUTLASS C++ templates\nwith JIT compilation to optimized CUDA kernels via MLIR and ptxas.\n\n## When to Use\n\n**Triggers:**\n- Writing CUDA kernels in Python (element-wise, GEMM, custom ops)\n- Optimizing GPU memory access patterns (vectorized loads, TMA, shared memory)\n- Building tensor core (MMA) kernels for Ampere/Hopper/Blackwell\n- Integrating custom GPU kernels with PyTorch or JAX\n- Prototyping high-performance kernels without C++ metaprogramming\n\n**Symptoms (wrong tool otherwise):**\n- Need shared memory coordination or tensor core MMA → use CuTe DSL (not Triton for complex patterns)\n- Need simple element-wise ops with no shared memory → CuTe DSL or Triton both work\n- Need to call existing CUTLASS C++ kernels → use CUTLASS C++ APIs instead\n- Need reductions, scans, or non-GEMM collective ops → consider CUB/Thrust\n\n**Keywords:** cute, cutlass, cute.jit, cute.kernel, from_dlpack, zipped_divide,\nTiledMMA, TiledCopy, TMA, WGMMA, tcgen05, pipeline, mbarrier\n\n## Requirements\n\n| Requirement | Detail |\n|-------------|--------|\n| Platform | Linux x86_64 only |\n| Python | 3.10–3.13 |\n| GPU | NVIDIA Ampere+ (SM80, SM90, SM100) |\n| CUDA Driver | ≥ 575.51.03 (Toolkit 12.9 compat) |\n| Install | `pip install nvidia-cutlass-dsl` |\n| Optional | `apache-tvm-ffi`, `torch-c-dlpack-ext` |\n\n## Workflows\n\n### Workflow 0: Starting from Examples (Recommended)\n\nFor any non-trivial kernel (GEMM, attention, pipelined, fused ops), start by\nfinding the most similar existing example to use as a **starting point** — study\nits structure, then rework it for your use case. Do not copy examples verbatim;\nthey target specific dtypes, architectures, and problem shapes that likely differ.\n\n1. **Pick the closest example** from the index below.\n   **Prefer examples matching the target GPU architecture** (check with\n   `torch.cuda.get_device_capability()`) when the operation is similar.\n\n   Fetch via `web_fetch` with base URL\n   `https://raw.githubusercontent.com/NVIDIA/cutlass/main/examples/python/CuTeDSL`\n\n   | Operation | Arch | Example path (append to base URL) |\n   |-----------|------|-----------------------------------|\n   | Element-wise add | SM80 | `ampere/elementwise_add.py` |\n   | Element-wise + autotune | SM80 | `ampere/elementwise_add_autotune.py` |\n   | Element-wise apply | SM80 | `ampere/elementwise_apply.py` |\n   | SGEMM (scalar) | SM80 | `ampere/sgemm.py` |\n   | Tensor-core GEMM | SM80 | `ampere/tensorop_gemm.py` |\n   | Flash Attention v2 | SM80 | `ampere/flash_attention_v2.py` |\n   | HSTU Attention | SM80 | `ampere/hstu_attention.py` |\n   | Shared memory allocator | SM80 | `ampere/smem_allocator.py` |\n   | CTA norm (LayerNorm) | SM90 | `hopper/cta_norm.py` |\n   | Dense GEMM | SM90 | `hopper/dense_gemm.py` |\n   | Dense GEMM persistent | SM90 | `hopper/dense_gemm_persistent.py` |\n   | Flash MHA | SM90 | `hopper/fmha.py` |\n   | Dense GEMM | SM100 | `blackwell/dense_gemm.py` |\n   | Dense GEMM persistent | SM100 | `blackwell/dense_gemm_persistent.py` |\n   | Dense GEMM + alpha/beta | SM100 | `blackwell/dense_gemm_alpha_beta_persistent.py` |\n   | RMSNorm | SM100 | `blackwell/rmsnorm.py` |\n   | Reduce | SM100 | `blackwell/reduce.py` |\n   | Flash MHA | SM100 | `blackwell/fmha.py` |\n   | Grouped GEMM | SM100 | `blackwell/grouped_gemm.py` |\n   | Mamba2 SSD | SM100 | `blackwell/mamba2_ssd/` |\n   | GEMM tutorial (notebook) | SM100 | `notebooks/tour_to_sol_gemm.ipynb` |\n\n   **Example:** To fetch the Hopper dense GEMM:\n   ```bash\n   web_fetch https://raw.githubusercontent.com/NVIDIA/cutlass/main/examples/python/CuTeDSL/hopper/dense_gemm.py\n   ```\n\n2. **Read reference materials first** — before diving into example code, read\n   the relevant `references/` docs to understand the patterns and APIs:\n   - For GEMM: `references/patterns-gemm.md` (3-level tiling, epilogue fusion,\n     `cute.compile` with `mark_layout_dynamic`, shared memory layouts)\n   - For reductions: `references/patterns-reduction.md` (warp reductions,\n     `cute.compile` cache pattern)\n   - For element-wise: `references/patterns-elementwise.md` (variations A–E)\n   - Always: `references/api-arch.md` (available APIs, arch-specific caveats)\n\n   This gives you the conceptual foundation so you can rework the example\n   intelligently rather than trying to copy-paste complex pipelines.\n\n3. **Fetch and study the example source** — read for structure, not to copy:\n   - Identify: decorators, tiling strategy, shared memory usage, mainloop flow\n   - Note which dtype/arch it targets (many examples are fp16/bf16-specific)\n   - Check if it uses APIs tied to a specific arch (TMA → SM90+, tcgen05 → SM100)\n\n4. **Rework for the user's workload** (do not copy-paste):\n   - Change shapes, data types, tile sizes to match requirements\n   - Replace compute logic (epilogue, activation fusion) as needed\n   - If dtype differs (e.g., example is fp16, need fp32), expect vectorization\n     and layout changes — the scalar-loop patterns in `references/` may be a\n     better starting point than adapting a vectorized example\n   - **Runtime wrapper must be lightweight**: `kernel_fn()` should only call\n     `from_dlpack()` + the compiled kernel. Never allocate intermediate tensors,\n     copy data, or re-compile per call — these belong in one-time setup\n   - Apply optimizations from this skill's reference docs\n\n   **⛔ Blackwell/Hopper GEMM + extra tensors — STOP:**\n   If the target GPU is SM90+ (Hopper/Blackwell) **and** the GEMM requires\n   extra tensors beyond A, B, C in the epilogue (e.g., bias vector, activation\n   inputs), **do not attempt it**. These examples use TMA descriptors for all\n   data movement — adding tensors requires modifying TMA descriptor setup,\n   which is prohibitively complex. Instead, tell the user this limitation and\n   suggest a **two-kernel approach**: run the GEMM kernel as-is, then apply\n   bias + activation in a separate element-wise kernel (Workflow 1).\n   Plain GEMM (just A×B→C with scalar alpha/beta) on Hopper/Blackwell is fine.\n\n5. **Validate and benchmark** using companion scripts:\n   ```bash\n   python scripts/verify_kernel.py kernel.py --rtol 1e-3 --atol 1e-3\n   python scripts/benchmark_kernel.py kernel.py\n   ```\n   The kernel file must export `kernel_fn`, `reference_fn`, and `get_inputs()`.\n\n**When to skip examples:** Pure element-wise operations (Workflow 1) have\ncomplete patterns in `references/patterns-elementwise.md` — no need to fetch\nexternal examples.\n\n**Reduction kernels** (softmax, layernorm, RMSNorm): Use\n`references/patterns-reduction.md` which provides complete, proven patterns\nfor float32 reductions using scalar loops + butterfly shuffle + shared memory.\n\n### Workflow 1: Element-wise Kernel\n\nFor unary/binary/in-place operations that map inputs to outputs 1:1.\n\n1. **Determine kernel structure**: inputs/outputs count, tensor rank, target arch\n2. **Select pattern** from `references/patterns-elementwise.md` (Variations A–E)\n3. **Write kernel** applying all four invariant principles:\n   - P1: `from_dlpack(tensor, assumed_align=16)` for vector loads\n   - P2: Derive `vec_size` from `element_type.width`\n   - P3: `cute.zipped_divide(mA, tiler)` for coalesced access\n   - P4: `cutlass.dynamic_expr(thread_idx < total)` for bounds\n4. **Critical rules**: No early return, no `a * 2` (use `a + a`), no `cute.math.sigmoid`\n5. **Pre-compile with `cute.compile()`**: Always pre-compile the kernel once\n   using `cute.compile()` so that `kernel_fn` calls the compiled object, not\n   `@cute.jit` directly.  Without pre-compilation, every call recompiles\n   (~20-50ms overhead).  Use `.mark_layout_dynamic()` so a single compiled\n   kernel handles arbitrary input shapes without recompilation:\n   ```python\n   # Compile once with dynamic layouts — works for any shape\n   fake_x = from_dlpack(torch.empty(1, 1, dtype=torch.float16, device=\"cuda\"),\n                         assumed_align=16).mark_layout_dynamic()\n   fake_out = from_dlpack(torch.empty(1, 1, dtype=torch.float16, device=\"cuda\"),\n                           assumed_align=16).mark_layout_dynamic()\n   compiled_kernel = cute.compile(host_fn, fake_x, fake_out)\n\n   def kernel_fn(x):\n       out = torch.empty_like(x)\n       compiled_kernel(from_dlpack(x, assumed_align=16).mark_layout_dynamic(),\n                       from_dlpack(out, assumed_align=16).mark_layout_dynamic())\n       return out\n   ```\n6. **Verify correctness** using companion script:\n   ```bash\n   python scripts/verify_kernel.py kernel.py --rtol 1e-3 --atol 1e-3\n   ```\n   The kernel file must export `kernel_fn`, `reference_fn`, and `get_inputs()`.\n7. **Benchmark** using companion script:\n   ```bash\n   python scripts/benchmark_kernel.py kernel.py\n   ```\n\n### Workflow 2: GEMM Kernel\n\nFor matrix multiplication with tiling, shared memory, and tensor cores.\n\n1. **Define problem**: shapes (M, N, K), data types, target architecture\n2. **Choose tiling**: CTA tile (bM, bN, bK), pipeline stages, cluster shape\n3. **Three-level partitioning** (see `references/patterns-gemm.md`):\n   - Level 1: CTA tiling with `local_tile()`\n   - Level 2: Copy partitioning (global → shared) with `TiledCopy`\n   - Level 3: Compute partitioning (shared → register) with `TiledMMA`\n4. **Shared memory**: Use swizzled layouts (`make_smem_layout_atom`) to avoid bank conflicts\n5. **Mainloop**: K-tile loop with copy → sync → MMA → sync\n6. **Pipeline**: Use `PipelineTmaAsync` (Hopper) or `PipelineTmaUmma` (Blackwell).\n   ⚠️ TMA-based pipelines manage data movement via TMA descriptors — adding\n   extra tensors (bias, activation inputs) to the epilogue requires modifying\n   descriptor setup, which is prohibitively complex. See the stop condition in\n   Workflow 0 step 4.\n7. **Epilogue**: Predicated store with alpha/beta scaling\n8. **Pre-compile with `cute.compile()`**: Always pre-compile the GEMM kernel\n   so `kernel_fn` calls the compiled object, not `@cute.jit` directly.\n   Without pre-compilation, every call recompiles (~20-50ms overhead).\n9. **Autotune**: Search over tile sizes, cluster shapes, pipeline depths\n\n### Workflow 3: Framework Integration\n\nFor wrapping CuTe DSL kernels as PyTorch/JAX custom operators.\n\n1. **Write kernel** using Workflow 1 or 2\n2. **Create wrapper**: Accept `torch.Tensor`, convert via `from_dlpack`, call host fn\n3. **For production**: Compile with TVM FFI for zero-overhead tensor passing:\n   ```python\n   compiled = cute.compile(host_fn, *fake_tensors, options=\"--enable-tvm-ffi\")\n   compiled(torch_a, torch_b)  # Direct torch.Tensor, no from_dlpack\n   ```\n4. **For deployment**: Use AOT compilation → export to `.o` → load at runtime\n\n### Workflow 4: Debugging & Profiling\n\n1. **Set environment**: `CUTE_DSL_PRINT_IR=1`, `CUTE_DSL_KEEP_PTX=1`\n2. **Use `cute.printf()`** for runtime values (not Python `print`)\n3. **Inspect generated code**: `compiled.__ptx__`, `compiled.__mlir__`\n4. **Profile**: Enable `CUTE_DSL_LINEINFO=1`, use Nsight Compute/Systems\n5. **Debug memory**: Run with `compute-sanitizer python script.py`\n\n## Output Formats\n\nA typical CuTe DSL kernel project:\n\n```\nkernel_dir/\n  kernel.py          # @cute.kernel + @cute.jit functions\n  test_kernel.py     # Correctness test vs PyTorch reference\n  bench_kernel.py    # Benchmark with cute.compile() setup\n```\n\n**Success indicators:**\n- Correctness test passes (`torch.testing.assert_close`)\n- Nsight shows vector loads (LDG.128/LDG.256), not scalar loads\n- For GEMM: tensor core utilization > 80% in Nsight Compute\n\n## Companion Script Contract\n\nKernel files used with `scripts/verify_kernel.py` and `scripts/benchmark_kernel.py`\nmust export three names:\n\n- `kernel_fn(*inputs)` — the CuTe DSL kernel wrapper (calls `cute.compile` + runs kernel)\n- `reference_fn(*inputs)` — PyTorch reference implementation (same signature)\n- `get_inputs()` — returns a list of CUDA tensors for testing\n\n```python\n# Example kernel.py contract\nimport torch\nimport cutlass.cute as cute\nfrom cutlass.cute.runtime import from_dlpack\n\ndef kernel_fn(x):\n    out = torch.empty_like(x)\n    # ... call compiled cute kernel ...\n    return out\n\ndef reference_fn(x):\n    return torch.nn.functional.gelu(x)\n\ndef get_inputs():\n    return [torch.randn(1024, 512, dtype=torch.float16, device=\"cuda\")]\n```\n\n## Examples\n\n### Example: 2D Unary Element-wise (ReLU)\n\n```python\nimport torch, cutlass, cutlass.cute as cute\nfrom cutlass.cute.runtime import from_dlpack\n\n@cute.kernel\ndef relu_kernel(gA: cute.Tensor, gC: cute.Tensor):\n    tidx, _, _ = cute.arch.thread_idx()\n    bidx, _, _ = cute.arch.block_idx()\n    bdim, _, _ = cute.arch.block_dim()\n    idx = bidx * bdim + tidx\n    m, n = gA.shape[1]\n    total = m * n\n    if cutlass.dynamic_expr(idx < total):\n        a = gA[(None, (idx // n, idx % n))].load()\n        gC[(None, (idx // n, idx % n))] = cute.where(a > 0, a, 0)\n\n@cute.jit\ndef relu_host(mA: cute.Tensor, mC: cute.Tensor):\n    vec = 16 // (mA.element_type.width // 8)\n    gA = cute.zipped_divide(mA, (1, vec))\n    gC = cute.zipped_divide(mC, (1, vec))\n    T = 256\n    N = cute.size(gA.shape[1])\n    relu_kernel(gA, gC).launch(grid=((N+T-1)//T,1,1), block=(T,1,1))\n\nx = torch.randn(1024, 512, dtype=torch.float16, device=\"cuda\")\nout = torch.empty_like(x)\nrelu_host(from_dlpack(x, assumed_align=16), from_dlpack(out, assumed_align=16))\n```\n\n## Error Handling\n\n| Error | Cause | Fix |\n|-------|-------|-----|\n| `MLIR function requires a Context` | Called @kernel from Python | Launch via @cute.jit host function |\n| `DSLAstPreprocessorError` on return | Early return in @kernel | Use `if cutlass.dynamic_expr(cond):` |\n| Type mismatch on store | `a * 2` promotes FP16→FP32 | Use `a + a` or `.to(cutlass.Float16)` |\n| `could not get source code` | Kernel in `exec()` context | Write to file and import |\n| Scalar loads in Nsight | Missing alignment hint | Add `assumed_align=16` to `from_dlpack` |\n| `Missing required argument` | Not all @jit params passed | Pass ALL declared parameters |\n| `AttributeError: sigmoid` | No `cute.math.sigmoid` | Use `1.0/(1.0+cute.math.exp(-x))` |\n\nSee `references/troubleshooting.md` for the full error table and limitations.\n\n**Debugging rule:** Never delete kernel.py during debugging. Use `backup_file`\nto save a checkpoint, then `edit_file` to iterate. If stuck, `revert_file` to\nrestore the backup. A partially-working kernel is always better than no kernel.\n\n## Finding More Information\n\n### Tier 1: This File (SKILL.md)\n\nWorkflows above cover element-wise kernels, GEMM, framework integration, and\ndebugging. Search this file first for procedural questions.\n\n### Tier 2: references/ Directory\n\nGrep for keywords across `references/`. Headers are grep-friendly.\n\n| File | Content |\n|------|---------|\n| `concepts-architecture.md` | Core abstractions, terminology, compilation pipeline |\n| `concepts-layouts.md` | Layout algebra: composition, complement, divide, swizzle |\n| `concepts-tensors.md` | Tensor types, partitioning, tiling, predication |\n| `concepts-mma.md` | MMA atoms, TiledMMA, per-architecture tensor core ops |\n| `patterns-getting-started.md` | Installation, decorators, first kernel walkthrough |\n| `patterns-elementwise.md` | Invariant principles, pattern variations, reference impl |\n| `patterns-gemm.md` | 3-level tiling, shared memory, pipelining, autotuning |\n| `patterns-memory.md` | from_dlpack, TMA, cp.async, TMEM, copy atoms |\n| `patterns-compilation.md` | Control flow, JIT caching, TVM FFI, AOT compilation |\n| `patterns-pipeline.md` | Producer-consumer, pipeline classes, barriers, warp specialization |\n| `api-core.md` | cute module: layouts, tensors, math, copy, gemm, printing |\n| `api-arch.md` | cute.arch: thread indexing, sync, atomics, memory ops |\n| `api-nvgpu.md` | cute.nvgpu: warp/warpgroup/cpasync/tcgen05 MMA and copy |\n| `api-runtime-utils.md` | Runtime: from_dlpack, fake tensors, utils, schedulers |\n| `troubleshooting.md` | Debugging, env vars, common errors, limitations, FAQ |\n\n**How to search:** Grep for your keyword across `references/`. Read only the\nfile and section that Grep points to.\n\n### Tier 3: Original Documentation\n\nIf Tiers 1–2 don't answer, consult the source:\n- **Web**: https://docs.nvidia.com/cutlass/latest/\n- **GitHub**: https://github.com/NVIDIA/cutlass\n- Fetch specific doc pages or search for \"CUTLASS CuTe DSL <topic>\"\n- Consider distilling the answer back into references/\n","frontmatter":{"name":"kernel-cute-writing","description":"Write and implement GPU kernels using NVIDIA CuTe DSL (CUTLASS 4.x Python API) — NOT for Triton, CUDA C++, or conceptual explanations. Trigger only when the user wants to write or implement a kernel, not when asking questions about CuTe DSL concepts or layouts. CuTe DSL uses cute.jit/cute.kernel decorators and cutlass.cute imports. Covers element-wise kernels, GEMM patterns, reductions, memory hierarchy (global/shared/register/TMA), MMA tensor core operations, software pipelining, and framework integration.\n","license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":4584,"sizeBytes":17091},{"name":"config-templates.md","path":".claude/skills/kernel-tileir-optimization/references/config-templates.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/config-templates.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-tileir-optimization/references/config-templates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/config-templates.md)","tokens":300,"sizeBytes":263},{"name":"tma-conversion.md","path":".claude/skills/kernel-tileir-optimization/references/tma-conversion.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/tma-conversion.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-tileir-optimization/references/tma-conversion.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/tma-conversion.md)","tokens":300,"sizeBytes":259},{"name":"SKILL.md","path":".claude/skills/kernel-tileir-optimization/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/SKILL.md","title":"kernel-tileir-optimization","category":"anthropic-skill","format":"markdown","content":"---\nname: kernel-tileir-optimization\ndescription: >\n  Optimize existing Triton kernels for NVIDIA TileIR backend on Blackwell GPUs (sm_100+).\n  Adds TileIR-specific autotune configs: occupancy, num_ctas, TMA descriptors. Covers\n  kernel classification (dot-related, norm-like, elementwise, reduction), type-specific\n  transformations, and PTX-vs-TileIR benchmarking. Triggered by: \"optimize for TileIR\",\n  \"add TileIR configs\", \"Blackwell optimization\", \"TMA descriptors\", \"2CTA mode\",\n  \"occupancy tuning\". Kernels use standard `import triton`; TileIR activates via\n  ENABLE_TILE=1 when nvtriton is installed.\ncompatibility: >\n  Requires Blackwell GPU (sm_100+) for TileIR execution. Supports development on any GPU.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Triton TileIR Optimization\n\nOptimize EXISTING Triton kernels for NVIDIA's TileIR backend on Blackwell GPUs.\nThis skill does NOT write kernels from scratch -- that is the Triton Specialist's job.\n\n## Principles\n\n### TileIR vs PTX Backend\n\nTileIR is NVIDIA's compiler backend for Triton that generates optimized CUDA code\nusing CGA-level (Cooperative Grid Array) tile representations. Critical differences:\n\n| Parameter | PTX Backend | TileIR Backend |\n|-----------|-------------|----------------|\n| `num_warps` | Strict directive | **Ignored** (compiler decides) |\n| `num_stages` | Strict directive | Cost hint (compiler optimizes) |\n| `occupancy` | Not available | **Critical** tuning param (1-32) |\n| `num_ctas` | Limited | 2CTA mode for Blackwell |\n| Block sizes | Smaller often better | Larger often better |\n| TMA | Not available | Required for dot kernels |\n\n**Key implication**: Do not tune `num_warps` for TileIR -- focus on `occupancy` instead.\n\n### Triton Package Landscape\n\nThree packages share `import triton`:\n\n| Package | Source | Use Case |\n|---------|--------|----------|\n| `pytorch-triton` | PyTorch wheel | `torch.compile`, standard kernels |\n| `triton` | OpenAI PyPI | Official Triton from triton-lang.org |\n| nvtriton | [Triton-to-tile-IR](https://github.com/triton-lang/Triton-to-tile-IR) | TileIR backend for Blackwell |\n\nOnly one triton package should be installed at a time. \"Converting to TileIR\" means\nadding TileIR-specific configs, NOT changing imports. TileIR activates via `ENABLE_TILE=1`.\n\n### When TileIR Applies\n\nTileIR targets Blackwell (sm_100+). Without nvtriton or Blackwell hardware, the\nspecialist still adds TileIR-optimized configs that standard triton safely ignores,\nenabling future deployment.\n\n**Expected speedups** (with nvtriton on Blackwell):\n\n| Kernel Type | Speedup | Key Lever |\n|-------------|---------|-----------|\n| Dot-Related (GEMM, Attention) | 1.2-2.0x | TMA + 2CTA |\n| Norm-Like (LayerNorm, Softmax) | 2.0-5.0x | High occupancy |\n| Element-Wise (ReLU, Add, Exp) | 1.5-3.0x | Occupancy + num_stages |\n| Reduction (Sum, Mean, Max) | 1.8-4.0x | High occupancy |\n\n## Workflow\n\nFive-phase workflow: compatibility, classify, transform, validate, benchmark.\n\n### Phase 1: Compatibility Test (ENABLE_TILE=0)\n\nVerify the kernel works in PTX mode before applying TileIR optimizations.\n\n```bash\npython scripts/tileir_check.py\n```\n\nThen use the **kernel-triton-writing** skill's `verify_kernel.py` to verify with `ENABLE_TILE=0`:\n\n```bash\npython scripts/verify_kernel.py --kernel path/to/kernel.py --reference 'torch reference' --shapes '{\"x\": [32, 512, 4096]}' --dtypes '{\"x\": \"bfloat16\"}'\n```\n\n### Phase 2: Classify Kernel\n\nDetermine kernel type to select the optimization strategy.\n\n```bash\npython scripts/classify_kernel.py --file kernel.py\n```\n\nClassification decision tree:\n\n```\nContains tl.dot()?\n  YES --> dot-related: TMA + 2CTA + occupancy + larger blocks\n  NO  --> Has reduction + normalization?\n            YES --> norm-like: high occupancy (2, 4) + num_warps (4, 8)\n            NO  --> Point-wise only?\n                      YES --> element-wise: occupancy (1-16) + num_stages (2-4)\n                      NO  --> reduction: high occupancy + num_warps\n```\n\n### Phase 3: Apply Transformations\n\nClassify and apply optimizations in one step:\n\n```bash\npython scripts/classify_kernel.py --file kernel.py --apply-optimizations\n```\n\nOutput JSON includes `optimized_code` and `changes_applied` fields.\n\n**Type-specific transformations:**\n\n**Dot-related** (highest priority):\n1. Convert `tl.load`/`tl.store` to TMA descriptors (MANDATORY). See `references/tma-conversion.md`.\n2. Add 2CTA configs (`num_ctas=2`) with SM oversubscription guard in pre-hook.\n3. Add occupancy (1, 2, 4) and extended num_stages (4, 6).\n4. Use larger block sizes (256x256, 256x128).\n\n**Norm-like** (LayerNorm, Softmax, RMSNorm):\n- Add occupancy (2, 4), num_warps (4, 8). No TMA needed.\n\n**Element-wise** (ReLU, GELU, Add, Mul, Exp):\n- Add occupancy (1, 2, 4, 16), num_stages (2, 3, 4). Include extreme configs for small inputs.\n\n**Reduction** (Sum, Mean, Max):\n- Same strategy as norm-like: high occupancy (2, 4), num_warps (4, 8).\n\nGate TileIR-specific configs for sm_100+:\n\n```python\nimport torch\n\ndef get_configs_with_gating(pre_hook=None):\n    configs = get_baseline_configs()\n    if torch.cuda.is_available() and torch.cuda.get_device_capability()[0] >= 10:\n        configs.extend(get_tileir_specific_configs(pre_hook))\n    return configs\n```\n\nSee `references/config-templates.md` for complete config templates per kernel type.\n\n### Phase 4: TileIR Validation (ENABLE_TILE=1)\n\nUse the **kernel-triton-writing** skill's `verify_kernel.py` to verify the optimized kernel with TileIR backend:\n\n```bash\npython scripts/verify_kernel.py --kernel path/to/optimized_kernel.py --reference 'torch reference' --shapes '{\"x\": [32, 512, 4096]}' --dtypes '{\"x\": \"bfloat16\"}'\n```\n\nSet `ENABLE_TILE=1` before running. Check: numerical correctness, no compilation errors,\nTMA/2CTA patterns compile successfully.\n\n### Phase 5: Benchmark\n\nUse `triton.testing.do_bench()` (as documented in the **perf-workload-profiling** skill) to compare PTX (`ENABLE_TILE=0`) vs TileIR (`ENABLE_TILE=1`).\n\nBenchmark across multiple input sizes (128, 1024, 8192) -- performance varies by size.\n\n## Scripts\n\n### tileir_check.py\n\nCheck TileIR availability (nvtriton, ENABLE_TILE, Blackwell GPU):\n\n```bash\npython scripts/tileir_check.py\n```\n\nReturns JSON: `nvtriton_installed`, `tileir_active`, `blackwell_gpu`, `gpu_capability`, `recommendation`.\n\n### classify_kernel.py\n\nClassify kernel type and optionally apply TileIR optimizations:\n\n```bash\n# Classify only\npython scripts/classify_kernel.py --file kernel.py\n\n# Classify + apply optimizations\npython scripts/classify_kernel.py --file kernel.py --apply-optimizations\n\n# From inline code\npython scripts/classify_kernel.py --code '<kernel_code>'\n```\n\nReturns JSON: `classification`, `confidence`, `indicators`, `recommendations`.\nWith `--apply-optimizations`: adds `optimized_code` and `changes_applied`.\n\n## Error Handling\n\n### Common Pitfalls\n\n**TMA descriptor errors** (dot-related kernels):\n- Always pass `pre_hook=tma_set_block_size_hook` to config generation -- without it,\n  TMA descriptors keep dummy block sizes, causing runtime errors or wrong results.\n- For GEMM: pass `b.T.contiguous()` in wrapper and use `tl.dot(a, b.T, accumulator)`\n  in kernel. Transposition mismatch produces incorrect results silently.\n\n**2CTA oversubscription**:\n- Adjust SM count in pre-hook when using `num_ctas=2`:\n  ```python\n  if \"NUM_SMS\" in nargs and \"NUM_CTAS\" in nargs:\n      nargs[\"NUM_SMS\"] = nargs[\"NUM_SMS\"] // nargs[\"NUM_CTAS\"]\n  ```\n\n**Config function signatures**:\n- ALL config helper functions MUST accept `pre_hook=None`, even if unused.\n  Without it: `TypeError: get_autotune_configs() takes 0 positional arguments`.\n\n**Hardware gating**:\n- Gate TileIR configs with `torch.cuda.get_device_capability()[0] >= 10`.\n  TMA/2CTA on pre-Blackwell GPUs causes runtime crashes.\n\n**API availability**:\n- Use `1.0 / (1.0 + tl.exp(-x))` instead of `tl.sigmoid(x)` -- not available in\n  all Triton versions including some nvtriton builds.\n\n**Performance tuning**:\n- Do not over-tune `num_warps` -- TileIR ignores it. Focus on `occupancy`.\n- Use larger block sizes (256x256, 256x128) for TileIR, not PTX-tuned small blocks.\n- Benchmark across small/medium/large inputs; one-size configs underperform.\n- For exp/log heavy kernels, enable approximate math:\n  ```bash\n  export TILEIR_ENABLE_APPROX=1\n  export TILEIR_ENABLE_FTZ=1\n  ```\n\n### When to Abort\n\nStop and report if:\n\n1. **No triton installed** -- cannot proceed.\n2. **Compatibility test fails** -- kernel has syntax/runtime errors before optimization.\n3. **TileIR validation fails** -- optimized kernel produces wrong results.\n4. **No speedup** -- TileIR version is slower than PTX baseline (with nvtriton).\n5. **Not Blackwell GPU** -- still add configs for future deployment, but skip\n   ENABLE_TILE testing and benchmarking.\n\n### Output Format\n\nAfter optimization, return:\n\n```\n## TileIR Optimization: kernel_name\n\n### Classification\n- Kernel type: [dot-related | norm-like | element-wise | reduction]\n- Strategy: [TMA + 2CTA | High occupancy | Occupancy + num_stages]\n\n### Compatibility Check (ENABLE_TILE=0)\n[PASSED | FAILED] — Max difference: X.Xe-Y\n\n### Transformations Applied\n- [List of transformations]\n\n### TileIR Validation (ENABLE_TILE=1)\n[PASSED | FAILED] — Max difference: X.Xe-Y\n\n### Benchmark Comparison\n| Backend | Time (ms) | Speedup |\n|---------|-----------|---------|\n| PTX (ENABLE_TILE=0) | X.XXX | 1.0x |\n| TileIR (ENABLE_TILE=1) | X.XXX | Y.Yx |\n\n### Output\nFile: kernel_name_tileir.py\n```\n","frontmatter":{"name":"kernel-tileir-optimization","description":"Optimize existing Triton kernels for NVIDIA TileIR backend on Blackwell GPUs (sm_100+). Adds TileIR-specific autotune configs: occupancy, num_ctas, TMA descriptors. Covers kernel classification (dot-related, norm-like, elementwise, reduction), type-specific transformations, and PTX-vs-TileIR benchmarking. Triggered by: \"optimize for TileIR\", \"add TileIR configs\", \"Blackwell optimization\", \"TMA descriptors\", \"2CTA mode\", \"occupancy tuning\". Kernels use standard `import triton`; TileIR activates via ENABLE_TILE=1 when nvtriton is installed.\n","compatibility":"Requires Blackwell GPU (sm_100+) for TileIR execution. Supports development on any GPU.\n","license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":2479,"sizeBytes":9493},{"name":"api-core.md","path":".claude/skills/kernel-triton-writing/references/api-core.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/api-core.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/api-core.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/api-core.md)","tokens":300,"sizeBytes":237},{"name":"api-language.md","path":".claude/skills/kernel-triton-writing/references/api-language.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/api-language.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/api-language.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/api-language.md)","tokens":300,"sizeBytes":245},{"name":"concepts-semantics.md","path":".claude/skills/kernel-triton-writing/references/concepts-semantics.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/concepts-semantics.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/concepts-semantics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/concepts-semantics.md)","tokens":300,"sizeBytes":257},{"name":"operator-routing.md","path":".claude/skills/kernel-triton-writing/references/operator-routing.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/operator-routing.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/operator-routing.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/operator-routing.md)","tokens":300,"sizeBytes":253},{"name":"patterns-advanced.md","path":".claude/skills/kernel-triton-writing/references/patterns-advanced.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-advanced.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-advanced.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-advanced.md)","tokens":300,"sizeBytes":255},{"name":"patterns-basic.md","path":".claude/skills/kernel-triton-writing/references/patterns-basic.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-basic.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-basic.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-basic.md)","tokens":300,"sizeBytes":249},{"name":"patterns-fusion.md","path":".claude/skills/kernel-triton-writing/references/patterns-fusion.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-fusion.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-fusion.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-fusion.md)","tokens":300,"sizeBytes":251},{"name":"patterns-gemm.md","path":".claude/skills/kernel-triton-writing/references/patterns-gemm.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-gemm.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/patterns-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/patterns-gemm.md)","tokens":300,"sizeBytes":247},{"name":"troubleshooting.md","path":".claude/skills/kernel-triton-writing/references/troubleshooting.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/troubleshooting.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/kernel-triton-writing/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/references/troubleshooting.md)","tokens":300,"sizeBytes":251},{"name":"SKILL.md","path":".claude/skills/kernel-triton-writing/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/SKILL.md","title":"kernel-triton-writing","category":"anthropic-skill","format":"markdown","content":"---\nname: kernel-triton-writing\ntags: [triton]\ndescription: >\n  ONLY for OpenAI Triton (@triton.jit) kernel development. NEVER use for\n  CUDA C++ kernels, CuTe DSL kernels, TileIR, or profiling tools (ncu, nsys).\n  The user's request must involve Triton explicitly. Covers Triton-specific\n  patterns: fused elementwise, reductions (softmax, LayerNorm, RMSNorm),\n  tiled GEMM with triton.autotune, and flash attention. Workflow:\n  design, write, verify (with fast-path for explicit requests).\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Triton Kernel Writing\n\n## Principles\n\n### Correctness First\n\n1. Never benchmark before verification passes.\n2. Always mask loads and stores for non-divisible shapes.\n3. Include `kernel_fn`, `reference_fn`, and `get_inputs()` exports for companion scripts.\n4. Always run `scripts/verify_kernel.py` to validate against the reference.\n\n### FP16/BF16 Precision Rules (LOW FREEDOM -- follow exactly)\n\nTranscendental functions (`tl.exp`, `tl.log`, `tl.math.erf`, `tl.math.tanh`) require fp32 inputs.\n\n```python\n# WRONG -- compilation error or wrong results with fp16/bf16:\nresult = tl.exp(x)\n\n# CORRECT -- cast to fp32, compute, cast back:\nx_fp32 = x.to(tl.float32)\nresult = tl.exp(x_fp32).to(x.dtype)\n```\n\nRule: any math function beyond basic arithmetic (+, -, *, /) requires fp32 cast in, original dtype cast out.\n\nAdditional precision constraints:\n- `tl.sigmoid()` is unavailable in some Triton versions. Use `1.0 / (1.0 + tl.exp(-x_fp32))`.\n- Always cast back to `x.dtype` before `tl.store` -- mismatches cause \"Type mismatch, store Float32 to Float16\".\n- Unlike PyTorch, Triton does NOT auto-promote fp16/bf16 to fp32 for accumulation. Always use `tl.float32` accumulators for `tl.dot`.\n- **TF32 for matmul:** On Ampere+/Hopper, `tl.dot` uses TF32 by default for fp32 inputs (same as `torch.mm`). Do NOT add `input_precision=\"ieee\"` — it is 3-8x slower. TF32 is the correct default. If verification fails due to TF32 precision (~0.01-0.1 abs diff), ensure `reference_fn` also uses TF32 (plain `torch.mm`, no `allow_tf32=False`).\n\n### CPU-GPU Sync Avoidance (LOW FREEDOM)\n\nNever call `.item()` in kernel wrappers. It forces a CPU-GPU sync (~50-100us per call).\n\n| Pitfall | Fix |\n|---------|-----|\n| `tensor.item()` for seed | `x.data_ptr() % (2**31)` |\n| `torch.randint(...).item()` | Use tensor metadata for pseudo-random seed |\n| Allocating output every call | Accept pre-allocated output as parameter |\n| Python loops calling kernel | Batch operations |\n\n### C Integer Division Semantics (CRITICAL)\n\nTriton uses **C semantics** (round toward zero) for `//` and `%`, NOT Python semantics (round toward negative infinity). This only matters when operands can be negative.\n\n| Expression | Python | Triton/C |\n|------------|--------|----------|\n| `-7 // 2` | `-4` | `-3` |\n| `-7 % 2` | `1` | `-1` |\n\n**Safe pattern:** Ensure all index/offset values are non-negative. If negative values are possible, use `(idx % BLOCK + BLOCK) % BLOCK`.\n\nSee [references/concepts-semantics.md](references/concepts-semantics.md) for full rules and scalar-only exception.\n\n### Kernel Design Mental Model\n\n- **Parallelization axis:** Element-wise kernels parallelize over flattened elements. Row-wise kernels (LayerNorm, softmax) parallelize over rows. Matmul kernels tile in 2D (M, N).\n- **Block size:** Power-of-2 only (256, 512, 1024, 2048). Start with 1024 for H100, 512 for V100.\n- **Memory coalescing:** Adjacent threads must access adjacent memory addresses. The compiler handles this automatically from block-level pointer arithmetic.\n- **Grid:** Use `triton.cdiv(n_elements, BLOCK_SIZE)`. With autotune, grid must be a lambda: `lambda meta: (triton.cdiv(n, meta['BLOCK_SIZE']),)`.\n- **Decorator order:** `@triton.autotune` (outermost) -> `@triton.heuristics` -> `@triton.jit` (innermost).\n- **`reset_to_zero`:** Required for autotune on kernels that accumulate (e.g., matmul output). Without it, later configs see leftover values from earlier trials.\n\n## Workflow\n\n**Fast path:** If the user explicitly requests a Triton kernel (e.g., \"Write a Triton kernel for X\", \"Implement softmax in Triton\"), start at **Phase 2**. Only use Phase 0-1 when the request is ambiguous about whether Triton is appropriate.\n\n### Phase 0: Route the Operator (only for ambiguous requests)\n\nSkip this phase if the user explicitly asks for a Triton kernel. Only use when the request is ambiguous (e.g., \"optimize this operation\").\n\nTriton wins when 2+ operations can share registers instead of writing/reading global memory. Quick rules:\n\n| Pattern | Decision |\n|---------|----------|\n| Single element-wise op (`relu`, `sigmoid`) | SKIP — PyTorch already optimal |\n| Standalone matmul | SKIP — cuBLAS is optimized |\n| Standard attention | SKIP — Use FlashAttention |\n| Element-wise chain (2+ ops), reduction, matmul + epilogue | USE TRITON |\n\nIf SKIP, recommend the alternative and STOP. See [references/operator-routing.md](references/operator-routing.md) for edge cases.\n\n### Phase 1: Analyze the Operator (only for ambiguous requests)\n\nFrom the user's request, identify: (1) operation type, (2) parallelization strategy, (3) input shapes and dtypes.\n\n### Phase 2: Design the Kernel\n\nPick the skeleton below that matches your operation. **These skeletons are sufficient for element-wise, reduction, matmul, and fusion kernels — do NOT read reference files for these common patterns.** Only consult `references/` when implementing uncommon patterns (grouped GEMM, TMA, extern functions) or debugging issues.\n\n**Element-wise skeleton** (GELU, dropout, fused ops on flat tensors):\n\n```python\n@triton.jit\ndef kernel(x_ptr, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):\n    pid = tl.program_id(0)\n    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n    mask = offsets < n_elements\n    x = tl.load(x_ptr + offsets, mask=mask)\n    # ... compute ...\n    tl.store(out_ptr + offsets, result, mask=mask)\n```\n\n**Row-wise skeleton** (softmax, LayerNorm, RMSNorm — one program per row):\n\n```python\n@triton.jit\ndef kernel(x_ptr, out_ptr, n_cols, BLOCK_SIZE: tl.constexpr):\n    row_idx = tl.program_id(0)\n    col_offsets = tl.arange(0, BLOCK_SIZE)\n    mask = col_offsets < n_cols\n    x = tl.load(x_ptr + row_idx * n_cols + col_offsets, mask=mask, other=0.0)\n    # ... reduce / normalize ...\n    tl.store(out_ptr + row_idx * n_cols + col_offsets, result, mask=mask)\n```\n\n**Tiled matmul skeleton** (GEMM with 2D tiling, grouped ordering, and autotune):\n\n```python\n@triton.autotune(\n    configs=[\n        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 256, 'BLOCK_K': 64, 'GROUP_SIZE_M': 8}, num_warps=8, num_stages=3),\n        triton.Config({'BLOCK_M': 64, 'BLOCK_N': 256, 'BLOCK_K': 32, 'GROUP_SIZE_M': 8}, num_warps=4, num_stages=4),\n        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 128, 'BLOCK_K': 32, 'GROUP_SIZE_M': 8}, num_warps=4, num_stages=4),\n        triton.Config({'BLOCK_M': 256, 'BLOCK_N': 64, 'BLOCK_K': 32, 'GROUP_SIZE_M': 8}, num_warps=4, num_stages=4),\n    ],\n    key=['M', 'N', 'K'],\n)\n@triton.jit\ndef matmul_kernel(\n    a_ptr, b_ptr, c_ptr, M, N, K,\n    stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn,\n    BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,\n    GROUP_SIZE_M: tl.constexpr,\n):\n    pid = tl.program_id(0)\n    num_m_blocks = tl.cdiv(M, BLOCK_M)\n    num_n_blocks = tl.cdiv(N, BLOCK_N)\n    # Grouped ordering for L2 cache locality\n    num_pid_in_group = GROUP_SIZE_M * num_n_blocks\n    group_id = pid // num_pid_in_group\n    first_pid_m = group_id * GROUP_SIZE_M\n    group_size_m = min(num_m_blocks - first_pid_m, GROUP_SIZE_M)\n    pid_m = first_pid_m + ((pid % num_pid_in_group) % group_size_m)\n    pid_n = (pid % num_pid_in_group) // group_size_m\n\n    offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)\n    offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)\n    offs_k = tl.arange(0, BLOCK_K)\n\n    a_ptrs = a_ptr + offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak\n    b_ptrs = b_ptr + offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn\n    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)\n\n    for k in range(0, tl.cdiv(K, BLOCK_K)):\n        a_mask = (offs_m[:, None] < M) & (offs_k[None, :] < K)\n        b_mask = (offs_k[:, None] < K) & (offs_n[None, :] < N)\n        a = tl.load(a_ptrs, mask=a_mask, other=0.0)\n        b = tl.load(b_ptrs, mask=b_mask, other=0.0)\n        acc += tl.dot(a, b)\n        a_ptrs += BLOCK_K * stride_ak\n        b_ptrs += BLOCK_K * stride_bk\n        offs_k += BLOCK_K\n\n    c_mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)\n    c_ptrs = c_ptr + offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn\n    tl.store(c_ptrs, acc.to(c_ptr.dtype.element_ty), mask=c_mask)\n```\n\n### Phase 3: Write the Kernel\n\nCreate an output directory, then write the kernel file to `{output_dir}/kernel.py`.\n\nThe kernel file MUST include:\n- `@triton.jit` decorated kernel function\n- `@triton.autotune` for production kernels (see [references/api-core.md](references/api-core.md))\n- Python wrapper function (descriptive name for external import)\n- **Fixed contract exports** (companion scripts rely on these exact names):\n  - `kernel_fn` — alias to the wrapper function\n  - `reference_fn(*args)` — PyTorch reference with identical signature\n  - `get_inputs()` — returns `list` of fresh CUDA tensors for testing/benchmarking\n\nConcise example (fused GELU + dropout):\n\n```python\nimport triton\nimport triton.language as tl\nimport torch\n\n@triton.autotune(\n    configs=[\n        triton.Config({'BLOCK_SIZE': 1024}, num_warps=4),\n        triton.Config({'BLOCK_SIZE': 2048}, num_warps=8),\n    ],\n    key=['n_elements'],\n)\n@triton.jit\ndef fused_gelu_dropout_kernel(\n    x_ptr, out_ptr, n_elements, p, seed,\n    BLOCK_SIZE: tl.constexpr,\n):\n    pid = tl.program_id(0)\n    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n    mask = offsets < n_elements\n\n    x = tl.load(x_ptr + offsets, mask=mask)\n    x_fp32 = x.to(tl.float32)\n    x = (0.5 * x_fp32 * (1.0 + tl.math.erf(x_fp32 * 0.7071067811865476))).to(x.dtype)\n\n    random = tl.rand(seed, offsets)\n    x = tl.where(random > p, x / (1.0 - p), 0.0)\n\n    tl.store(out_ptr + offsets, x, mask=mask)\n\n\ndef fused_gelu_dropout_triton(x: torch.Tensor, p: float = 0.1) -> torch.Tensor:\n    n_elements = x.numel()\n    out = torch.empty_like(x)\n    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)\n    seed = (x.data_ptr() % (2**31)) ^ n_elements  # sync-free seed\n    fused_gelu_dropout_kernel[grid](x, out, n_elements, p, seed)\n    return out\n\n\n# --- Fixed contract (companion scripts rely on these names) ---\nkernel_fn = fused_gelu_dropout_triton\n\ndef reference_fn(x, p=0.1):\n    torch.manual_seed((x.data_ptr() % (2**31)) ^ x.numel())\n    return torch.nn.functional.dropout(\n        torch.nn.functional.gelu(x), p, training=True\n    )\n\ndef get_inputs():\n    return [torch.randn(128 * 1024 * 1024, device=\"cuda\")]\n```\n\nFor more patterns (SiLU+mul, RMSNorm, linear+GELU, add+LayerNorm), see [references/patterns-fusion.md](references/patterns-fusion.md). For GEMM patterns, see [references/patterns-gemm.md](references/patterns-gemm.md).\n\n### Phase 4: Verify Correctness\n\nRun the companion verification script:\n\n```bash\npython scripts/verify_kernel.py {output_dir}/kernel.py --rtol 1e-3 --atol 1e-3\n```\n\nOutput:\n```json\n{\"correct\": true, \"max_abs_diff\": 1.2e-7, \"max_rel_diff\": 3.4e-6, \"details\": \"...\"}\n```\n\n**Stop if `correct: false`.** Fix the kernel before benchmarking.\n\n**Tolerance guide:**\n\n| Dtype | rtol | atol | Notes |\n|-------|------|------|-------|\n| float16 | 1e-3 | 1e-3 | |\n| bfloat16 | 1e-2 | 1e-2 | |\n| float32 | 1e-5 | 1e-5 | Element-wise ops |\n| float32 (matmul) | 1e-2 | 1e-1 | TF32 accumulation order differs between Triton tiles and cuBLAS |\n\n### Phase 5: Benchmark Performance (optional)\n\nOnly benchmark if the user explicitly requests performance numbers. Skip this phase for correctness-focused requests.\n\n```bash\npython scripts/benchmark_kernel.py {output_dir}/kernel.py\n```\n\nOutput:\n```json\n{\"kernel_time_ms\": 0.45, \"reference_time_ms\": 1.23, \"speedup\": 2.73, \"warmup_iters\": 10, \"benchmark_iters\": 40}\n```\n\n## References (consult only when stuck)\n\nThe skeletons and principles above cover element-wise, reduction, matmul, and fusion kernels. **Do NOT read reference files for these common patterns.**\n\nOnly consult `references/` when:\n- Implementing **uncommon patterns** (grouped GEMM, TMA, persistent matmul, extern functions)\n- **Debugging** a compile error or incorrect result not covered by the error table below\n- Needing **API details** for an unfamiliar `tl.*` operation\n\n**How to search:** Grep for your keyword across `references/`. Read only the file Grep points to.\n\n| File | When to use |\n|---|---|\n| `references/api-core.md` | Unfamiliar `triton.autotune` / `triton.Config` options |\n| `references/api-language.md` | Unfamiliar `tl.*` operations |\n| `references/patterns-gemm.md` | Grouped GEMM, persistent matmul, TMA, MX formats |\n| `references/patterns-advanced.md` | Flash attention details, backward passes, libdevice |\n| `references/troubleshooting.md` | Debug ops, interpreter mode, env vars |\n\n## Error Handling and Troubleshooting\n\n### Common Errors\n\n| Error / Symptom | Cause | Fix |\n|---------|-------|-----|\n| \"Type mismatch, store Float32 to Float16\" | Missing `.to(x.dtype)` before store | Cast fp32 result back |\n| `BLOCK_SIZE is not a constexpr` | Block size passed as runtime value | Add `: tl.constexpr` annotation |\n| `shape mismatch` in binary op | Tensor shapes don't broadcast | Check with `tl.static_print`; use `[:, None]` / `[None, :]` |\n| Large diffs everywhere | Wrong dtype in `tl.load` | Check load dtype matches input |\n| Matmul 3-8x slower than expected | `input_precision=\"ieee\"` on `tl.dot` | Remove it; use TF32 default. Ensure `reference_fn` also uses TF32 |\n| Matmul ~0.01-0.1 abs diff vs reference | TF32 vs IEEE mismatch | Use same precision in both kernel and reference (TF32 for both) |\n| Diffs at boundaries | Missing mask | Add mask to all load/store ops |\n| Random diffs | Race condition | Check atomics and ordering |\n| NaN/Inf | Division by zero or fp16 overflow | Guard with epsilon; use `tl.float32` accumulator |\n| `grid must be a tuple` | Grid lambda returns int, not tuple | Return `(value,)` with trailing comma |\n| `expected constexpr` in `tl.arange` | Non-constexpr argument | Both args of `tl.arange(start, end)` must be constexpr |\n| `triton.OutOfResources` | Register/shared memory pressure | Reduce BLOCK_SIZE or `num_stages` |\n| Kernel not updating after edit | Stale compilation cache | `rm -rf ~/.triton/cache/` |\n| Mismatched results vs PyTorch | C integer division semantics | Triton uses truncation; see `references/concepts-semantics.md` |\n\nFor extended error table, interpreter mode issues, and environment variables, see [references/troubleshooting.md](references/troubleshooting.md).\n\n### When to Abort\n\nStop and report failure if:\n1. **Not a good fit** -- Pure matmul or complex control flow (Phase 0 should catch this).\n2. **Verification fails after 3 attempts** -- Numerical issues too severe to fix.\n3. **No speedup** -- Reference is already well-optimized (cuBLAS, cuDNN).\n4. **Hardware mismatch** -- Target GPU not available for testing.\n","frontmatter":{"name":"kernel-triton-writing","tags":["triton"],"description":"ONLY for OpenAI Triton (@triton.jit) kernel development. NEVER use for CUDA C++ kernels, CuTe DSL kernels, TileIR, or profiling tools (ncu, nsys). The user's request must involve Triton explicitly. Covers Triton-specific patterns: fused elementwise, reductions (softmax, LayerNorm, RMSNorm), tiled GEMM with triton.autotune, and flash attention. Workflow: design, write, verify (with fast-path for explicit requests).\n","license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":4209,"sizeBytes":15233},{"name":"SKILL.md","path":".claude/skills/perf-analysis/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-analysis/SKILL.md","title":"perf-analysis","category":"anthropic-skill","format":"markdown","content":"---\nname: perf-analysis\ndescription: >\n  Performance analysis coordination workflow. Guides profiling delegation,\n  bottleneck classification (compute/memory/launch/communication/sync),\n  and structured report generation. Use when the user asks to analyze\n  performance, profile a workload, check MFU/SOL, or diagnose bottlenecks.\ntags:\n  - analysis\n  - profiling\n  - bottleneck\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Performance Analysis\n\n## Principles\n\n1. **Delegate profiling, own analysis.** You coordinate the analysis workflow\n   but do not run profiling tools directly. Delegate all profiling and\n   measurement tasks to **perf-profiling-specialist** or other domain specialists.\n2. **Metrics from tools, never invented.** All performance numbers must come\n   from profiling tool output. Never fabricate metrics.\n3. **Classify before recommending.** Identify the bottleneck type before\n   suggesting optimizations. The wrong classification leads to wasted effort.\n4. **Structured reports.** Every analysis produces a report with Summary,\n   Metrics, Findings, and Recommendations.\n\n## Key Performance Metrics\n\n- **Throughput**: samples/sec, tokens/sec, iterations/sec\n- **Latency**: end-to-end time, kernel time, communication time\n- **MFU (Model FLOPs Utilization)**: actual FLOPs / theoretical peak FLOPs\n- **% of SOL (Speed of Light)**: current perf / hardware peak perf\n- **GPU Utilization**: SM occupancy, tensor core usage\n- **Memory Bandwidth**: DRAM bandwidth utilization vs peak\n\n## Analysis Workflow\n\n1. **Understand**: Clarify what metrics the user needs (MFU, SOL, latency, etc.)\n2. **Plan**: Plan your profiling and analysis steps before starting\n3. **Profile**: Delegate to **perf-profiling-specialist** for actual measurements\n4. **Measure**: Extract requested metrics from profiling results\n5. **Classify**: If diagnosing issues, determine primary bottleneck type\n6. **Report**: Generate performance analysis report with findings\n\n## Bottleneck Classification\n\nWhen diagnosing performance issues, classify the primary bottleneck:\n\n| Type | Indicator | Description |\n|------|-----------|-------------|\n| **Compute-bound** | High GPU utilization, low memory bandwidth usage | Limited by compute capacity (FLOPs) |\n| **Memory-bound** | High memory bandwidth, low compute utilization | Limited by DRAM throughput |\n| **Launch-overhead** | Many small kernels, high CPU time | CPU becoming bottleneck from kernel launch overhead |\n| **Communication-bound** | Significant time in collective operations | Limited by inter-GPU or inter-node communication |\n| **Sync-bound** | Excessive CPU-GPU synchronization points | Stalls from unnecessary synchronization |\n\nAfter classifying, consult the **perf-optimization-casebook** skill for\ncandidate optimizations (decision precedents) matching the bottleneck class,\nmodel, and hardware. Use it to populate the Recommendations section with\nproven, adapted approaches rather than inventing them — but report only\nmeasured numbers, never a case's expected effect as an achieved result.\n\n## Delegation Guidelines\n\nWhen delegating to specialists, describe the **desired outcome** -- not the\ntool methodology.\n\n**DO include:**\n- The workload: file path, code snippet, or command to profile\n- Problem context: dimensions, dtypes, FLOPs calculations, batch sizes\n- Desired metrics: SOL%, MFU, throughput, occupancy, bottleneck classification\n- Any constraints: specific kernel to target, profiling region markers\n\n**DO NOT include:**\n- Specific tool flags or command patterns (e.g., `--set=full`, `--section SpeedOfLight`)\n- Step-by-step tool usage instructions\n- Fallback strategies for tool failures\n- Example commands\n- Output file paths or artifact locations (specialists create their own workspace artifacts)\n\nSpecialists have their own skills that encode best practices for tool usage\nand their own workspace artifacts for output. Prescribing commands in the\ndelegation overrides their skills and may lead to suboptimal profiling\nstrategies (e.g., collecting 8000+ metrics with `--set=full` when a targeted\nsection analysis would be faster and more surgical).\n\n### Good Example\n\n```\nProfile the batched GEMM kernel in bmm_workload.py with NCU.\nThe workload uses cudaProfilerStart/Stop markers to isolate the region of interest.\nCollect kernel-level metrics: SOL%, compute/memory throughput, DRAM bandwidth,\ntensor core utilization, occupancy, warp stall reasons, and roofline classification.\nThe batched GEMM performs 68.72 GFLOP per call (B=32, M=512, N=1024, K=2048, FP16).\nCalculate MFU against the GPU's peak FP16 tensor core TFLOP/s.\n```\n\n### Bad Example\n\n```\nRun NCU with --set=full --profile-from-start off --target-processes all.\nIf --set=full fails, try --set=detailed. Parse the CSV output for\nsm__throughput.avg.pct_of_peak_sustained_elapsed.\nSave raw NCU output to ncu_output.txt.\n```\n\n### Remote Profiling\n\nWhen profiling on a remote SLURM cluster, include the\n**Remote Execution Context** block in the delegation prompt with the SSH+srun\nwrapper for the target cluster. The perf-profiling-specialist will prefix its\ncommands (nsys, ncu, nvidia-smi) with this wrapper.\n\nThe perf-profiling-specialist does not need the `exec-remote-slurm` skill — the\ncontext block provides everything it needs to execute remotely.\n\n## Available Specialists\n\nDelegate profiling and domain-specific analysis to these specialists:\n\n- **perf-profiling-specialist**: Runs nvidia-smi, nsys, ncu, torch.profiler. Use for ALL profiling tasks.\n- **perf-torch-cuda-graph-specialist**: Analyzes CUDA Graph compatibility and applies capture workflows\n\n## Report Format\n\nStructure every analysis report with these four sections:\n\n1. **Summary**: High-level performance status or bottleneck classification\n2. **Metrics**: Key performance numbers from profiling\n3. **Findings**: Detailed observations with evidence\n4. **Recommendations**: Prioritized list of optimizations (if applicable)\n\n### Example Report\n\n```\n## Summary\nTraining at 42% MFU, memory-bound due to large attention tensors.\n\n## Metrics\n- Throughput: 1,247 samples/sec\n- MFU: 42% (vs 65% theoretical for this model)\n- % of SOL: 58% (room for 1.7x improvement)\n- GPU Utilization: 45%\n- Memory Bandwidth: 850 GB/s (89% of peak)\n- Kernel Count: 1,247 per iteration\n\n## Findings\n1. Self-attention consumes 60% of memory bandwidth\n2. Optimizer step has 3 unnecessary synchronizations\n3. Batch size could be increased by 2x\n\n## Recommendations\n1. Enable FlashAttention (expected: +15% MFU)\n2. Remove synchronizations in optimizer (expected: +5% throughput)\n3. Increase batch size to improve GPU utilization\n```\n","frontmatter":{"name":"perf-analysis","description":"Performance analysis coordination workflow. Guides profiling delegation, bottleneck classification (compute/memory/launch/communication/sync), and structured report generation. Use when the user asks to analyze performance, profile a workload, check MFU/SOL, or diagnose bottlenecks.\n","tags":["analysis","profiling","bottleneck"],"license":"Apache-2.0","metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":1490,"sizeBytes":6640},{"name":"examples.md","path":".claude/skills/perf-host-analysis/references/examples.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/examples.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/examples.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/examples.md)","tokens":300,"sizeBytes":231},{"name":"iteration-isolation-techniques.md","path":".claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/iteration-isolation-techniques.md)","tokens":300,"sizeBytes":275},{"name":"kernel-level-analysis.md","path":".claude/skills/perf-host-analysis/references/kernel-level-analysis.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/kernel-level-analysis.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/kernel-level-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/kernel-level-analysis.md)","tokens":300,"sizeBytes":257},{"name":"metrics.md","path":".claude/skills/perf-host-analysis/references/metrics.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/metrics.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/metrics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/metrics.md)","tokens":300,"sizeBytes":229},{"name":"nsys-schema.md","path":".claude/skills/perf-host-analysis/references/nsys-schema.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/nsys-schema.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/nsys-schema.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/nsys-schema.md)","tokens":300,"sizeBytes":237},{"name":"output-format.md","path":".claude/skills/perf-host-analysis/references/output-format.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/output-format.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/output-format.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/output-format.md)","tokens":300,"sizeBytes":241},{"name":"phase-classification.md","path":".claude/skills/perf-host-analysis/references/phase-classification.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/phase-classification.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/phase-classification.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/phase-classification.md)","tokens":300,"sizeBytes":255},{"name":"thresholds.md","path":".claude/skills/perf-host-analysis/references/thresholds.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/thresholds.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/thresholds.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/thresholds.md)","tokens":300,"sizeBytes":235},{"name":"trtllm-nvtx-ranges.md","path":".claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/references/trtllm-nvtx-ranges.md)","tokens":300,"sizeBytes":251},{"name":"SKILL.md","path":".claude/skills/perf-host-analysis/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/SKILL.md","title":"perf-host-analysis","category":"anthropic-skill","format":"markdown","content":"---\nname: perf-host-analysis\ndescription: >\n  Analyze host/CPU overhead in TensorRT-LLM inference from nsys traces.\n  Detect whether host overhead is the bottleneck using GPU idle ratio, host prep\n  exposed ratio, and per-phase evidence. For regressions, isolate forward steps\n  via allreduce/NVTX patterns, compare host operation breakdowns across versions,\n  and identify scheduling or request-management overhead. Supports optional\n  inter-kernel gap, eager-vs-graph, pattern mapping, and multi-rank straggler\n  drill-down. Use standalone or within perf-analysis. Triggers: host overhead,\n  inter-step gap, scheduling overhead, forward step isolation, nsys iteration\n  analysis, NVTX breakdown, request management overhead, GPU idle, host bottleneck,\n  host prep exposed, inter-kernel gap, bubble analysis, graph coverage, eager kernel,\n  rank imbalance, straggler detection.\nlicense: Apache-2.0\ntags:\n  - analysis\n  - profiling\n  - host-overhead\n  - nsys\n  - inference\ndependencies:\n  - workload-instrumentation\n  - trace-interpretation\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Host Performance Analysis\n\nAnalyze host/CPU overhead in TensorRT-LLM inference workloads from nsys traces. This skill operates in two phases:\n\n| Phase | Question | Input | Output |\n|-------|----------|-------|--------|\n| **Detection** | Is host overhead the bottleneck? | Single nsys trace | YES/NO verdict with metric evidence |\n| **Root Cause** | What specifically regressed? | One or two nsys traces | NVTX per-step breakdown, regression sources, optional kernel-level drill-down |\n\n## When to Use\n\n- Before starting host optimization work -- confirms the bottleneck is real (Detection)\n- As a sub-step of `perf-analysis` for bottleneck classification (Detection)\n- When GPU utilization is suspiciously low and you need to know why (Detection)\n- When throughput regressed but GPU kernel execution times are unchanged (Root Cause)\n- When the gap between forward step iterations has increased (Root Cause)\n- To compare inter-iteration overhead between two versions of the inference engine (Root Cause)\n- When you need sub-operation granularity on inter-kernel gaps or graph coverage (Root Cause, kernel-level drill-down)\n- When piecewise CUDA graph coverage is unexpectedly low (Root Cause, kernel-level drill-down)\n- When multi-rank inference shows unexplained performance asymmetry (Root Cause, kernel-level drill-down)\n\nDo NOT use when:\n- The regression is in individual kernel performance (use `perf-nsight-compute-analysis`)\n- You need to profile a workload from scratch (use `workload-instrumentation` first)\n- The issue is NCCL communication (use distributed analysis)\n\n## Prerequisites\n\n- An nsys trace file (`.sqlite` or `.nsys-rep`) from a TRT-LLM benchmark run\n- For Root Cause comparison: two traces (baseline and target)\n- Python 3 with sqlite3 support\n\n---\n\n## Key Concepts\n\n### Host Overhead in LLM Inference\n\nIn an LLM inference loop, each iteration consists of:\n```\n[inter-step gap] -> [_forward_step] -> [inter-step gap] -> [_forward_step] -> ...\n```\n\nThe **forward step** includes GPU kernel execution (GEMM, attention, normalization, allreduce) plus host-side preparation. The **inter-step gap** includes host-side work between forward steps (scheduling, request fetching, broadcasting, sampling, response handling).\n\nSee [references/trtllm-nvtx-ranges.md](references/trtllm-nvtx-ranges.md) for the full per-operation breakdown and timing ranges.\n\n### Hidden vs Exposed Host Overhead\n\nHost overhead only hurts performance when it is **exposed** -- the GPU is idle waiting for work. When host prep overlaps with GPU execution, it is **hidden** and free. See [references/metrics.md](references/metrics.md) (M3 section) for diagrams and the exposed/hidden computation.\n\n### Forward Step Isolation\n\nIn TP configurations, forward steps are isolated via allreduce kernel grouping (deterministic count per transformer layer). For TP=1, NVTX `_forward_step` ranges are used directly. See [references/iteration-isolation-techniques.md](references/iteration-isolation-techniques.md) for the full algorithm.\n\n### Phase Classification (Context vs Generation)\n\nIterations are classified by NVTX marker text into **context** (eager, no CUDA graphs) and **generation** (CUDA graph replay). Per-phase analysis is critical because aggregate metrics can mask phase-specific bottlenecks. See [references/phase-classification.md](references/phase-classification.md).\n\n---\n\n## Phase 1: Detection (YES/NO Verdict)\n\nDetermine whether host overhead is the primary bottleneck.\n\n### Detection Metrics\n\nSix metrics in four categories. See [references/metrics.md](references/metrics.md) for full definitions, formulas, and SQL queries.\n\n| # | Metric | Threshold | What it answers |\n|---|--------|-----------|-----------------|\n| M1 | GPU idle ratio | > 0.30 | Is the GPU starved for work? |\n| M2 | Launch overhead ratio | > 0.10 | Is kernel launch itself expensive? |\n| M3a | Host prep exposed ratio | > 0.50 | How well is host prep pipelined? |\n| M3b | Host prep perf impact | > 0.05 | How much throughput does exposed prep cost? |\n| M3c | Host prep idle attribution | > 0.50 | Is host prep the main cause of GPU idle? |\n| M4 | GPU utilization | < 0.60 | Is GPU utilization too low? |\n| M5 | NCCL ratio (caveat) | > 0.20 | Is communication a confounding factor? |\n\n**Host prep confirmation rule**: Host prep is a confirmed bottleneck only when **both** M3b AND M3c cross their thresholds.\n\nThresholds are configurable with per-phase variants. See [references/thresholds.md](references/thresholds.md).\n\n### Detection Workflow\n\n#### Step 1: Input Validation\n\n```bash\n# Accept .sqlite or .nsys-rep\nls -la <trace_file>\n\n# If .nsys-rep, export to SQLite first\nnsys export -t sqlite -o <output.sqlite> <input.nsys-rep>\n```\n\n#### Step 2: Run Detection Script\n\n```bash\npython scripts/detect_host_overhead.py \\\n  --trace /path/to/trace.sqlite \\\n  --output /path/to/verdict.json\n```\n\nThe script computes M1, M2, M4, M5 from SQL, optionally M3 via range intersection, applies the verdict logic, and outputs structured JSON. See [references/output-format.md](references/output-format.md) for the output schema.\n\nFor manual metric extraction via SQL, see [references/nsys-schema.md](references/nsys-schema.md).\n\n#### Step 3: Interpret Verdict\n\n**Overall Verdict:**\n```\nif aggregate_verdict == YES or context_verdict == YES or generation_verdict == YES:\n    overall_verdict = YES\n```\n\nPer-phase analysis can **elevate** the verdict but never **demote** it.\n\nFormat using the template in [references/output-format.md](references/output-format.md).\n\n**Next Steps:**\n- **If YES** -> Proceed to Phase 2 (Root Cause) below, then use `perf-host-optimization` skill\n- **If NO** -> Use `perf-nsight-compute-analysis` for kernel SOL% or `perf-analysis` for full classification\n\n---\n\n## Phase 2: Root Cause Analysis\n\nIdentify which specific host operations regressed and by how much. Works with a single trace (breakdown) or two traces (comparison).\n\n### Principles\n\n1. **Isolate forward steps, not the full trace.** nsys traces contain warmup, JIT, model loading, and teardown.\n2. **Use structural kernel patterns for iteration detection.** Allreduce grouping is more robust than kernel density.\n3. **Compare steady-state iterations.** Filter to identical workload (same batch size, same ctx/gen mix).\n4. **Per-step metrics, not totals.** Always compare per-step averages.\n\n### Root Cause Workflow\n\n#### Step 1: Collect nsys Traces\n\nProfile both versions (if comparing) with identical settings:\n\n```bash\nnsys profile -o /path/to/trace \\\n  -t cuda,nvtx,osrt \\\n  --force-overwrite=true \\\n  --cuda-memory-usage=true \\\n  -w true \\\n  <benchmark_command> --num_requests 500\n```\n\n#### Step 2: Export to SQLite\n\n```bash\nnsys export --type=sqlite --force-overwrite=true -o trace.sqlite trace.nsys-rep\n```\n\n#### Step 3: Run Host Overhead Analysis\n\n```bash\n# Two-trace comparison\npython scripts/analyze_host_overhead.py \\\n  --baseline /path/to/baseline/trace.sqlite \\\n  --target /path/to/target/trace.sqlite \\\n  --baseline-label \"v1.1\" \\\n  --target-label \"main\" \\\n  --output /path/to/output/analysis.txt\n\n# Single-trace breakdown\npython scripts/analyze_host_overhead.py \\\n  --baseline /path/to/trace.sqlite \\\n  --baseline-label \"current\"\n```\n\n#### Step 4: Interpret Results\n\nThe script produces:\n1. **Allreduce-based iteration detection** -- confirms forward step boundaries\n2. **Per-step wall time comparison** -- quantifies the regression\n3. **NVTX per-step breakdown** -- identifies which host operations regressed\n4. **GPU kernel comparison** -- confirms GPU execution is unchanged\n5. **CUDA API comparison** -- detects kernel launch overhead changes\n\n### Reading the Output\n\n**Per-Step Wall Time:**\n```\nAvg wall time per step: 3,317 us (baseline) vs 3,978 us (target)  +19.9%\n```\nThis is the primary regression metric.\n\n**NVTX Breakdown:**\n```\nOperation           | baseline (us/step) | target (us/step) | Delta    | Status\n_fetch_new_requests |               36   |             270  | +234     | REGRESSION\nbroadcast_requests  |                -   |             250  | +250     | NEW\n_update_requests    |              413   |             723  | +310     | REGRESSION\n```\nFocus on operations with large absolute deltas.\n\n**GPU Kernel Comparison:**\n```\nKernels per step (launched): 6.2 (baseline) vs 21.9 (target)  +253%\n```\nMore individual launches = more host-side launch overhead.\n\n### Step 5: Kernel-Level Drill-Down (Optional)\n\nWhen the NVTX breakdown identifies a regressing operation but does not reveal *why* (the overhead is inside the GPU dispatch, not between NVTX ranges), drill below NVTX operations into individual GPU kernel launches.\n\nSee [references/kernel-level-analysis.md](references/kernel-level-analysis.md) for full technique details, SQL queries, and examples.\n\n**When to drill down:**\n- An operation has high wall time but the overhead is inside GPU dispatch, not between NVTX ranges\n- You need to understand how much of the forward pass is graph-captured vs eager\n- Per-layer overhead is significant and you need to map kernels to functional groups\n- Multi-rank inference shows unexplained performance asymmetry\n\n#### Kernel-Level Techniques\n\n| Technique | Question | Key Output |\n|-----------|----------|------------|\n| **Inter-Kernel Gap Analysis** | Where is the GPU idle between kernels? | Gap bucket distribution, top-N largest gaps with source mapping |\n| **Eager vs Graph Classification** | What fraction of kernels are graph-captured? | Graph coverage ratio, list of eager kernels with source attribution |\n| **Repeating-Pattern Mapping** | Which functional group within a layer has the most overhead? | Per-group gap totals, priority ranking |\n| **Straggler Detection** | Is one rank consistently slower? | Straggler rank ID, root cause (extra host work, queue depth feedback loop) |\n\n#### Workflow\n\n1. **Start with Inter-Kernel Gap Analysis** — bucket the gap distribution to understand the dominant overhead type (graph dispatch, Python interpreter, host-device sync)\n2. **If piecewise graph is in use**, run Eager vs Graph Classification to measure graph coverage and identify unnecessary eager kernels\n3. **For per-layer overhead**, use Repeating-Pattern Mapping to isolate the highest-overhead functional group within a single layer\n4. **For multi-rank setups**, run Straggler Detection if per-step wall time varies across ranks\n\n#### Kernel-Level Findings to Optimization Patterns\n\n| Finding | Optimization Pattern |\n|---------|---------------------|\n| Large gaps from Python tensor view chains | CUSTOM_OP — replace with C++ custom op |\n| Graph-capturable kernels running eagerly | GRAPH_EXPAND — fix partition poisoning |\n| Monolithic custom op blocking graph capture | GRAPH_SPLIT — split into capturable + eager parts |\n| Host-device sync (`.item()`) in per-layer code | SYNC (Pattern 1: pre-compute on CPU) + HOIST (Variant B: pass from step level) |\n| Per-layer buffer allocation | ALLOC — pre-allocate at init |\n| Straggler rank with extra host work | Apply targeted optimization to coordinator-only code paths |\n\n---\n\n## Common Patterns and Root Causes\n\n### Pattern 1: Request Management Refactor\n**Symptom**: `_fetch_new_requests` regressed 5-10x, new `broadcast_requests` operation.\n**Cause**: Request fetching refactored for multi-rank broadcasting in TP.\n**Mitigation**: Optimize broadcast path; batch request state updates.\n\n### Pattern 2: Increased Kernel Launch Count\n**Symptom**: 3-5x more `cudaLaunchKernel` calls per step, similar GPU time.\n**Cause**: Operations that were fused or graph-captured are now individual launches.\n**Mitigation**: Re-fuse kernels; extend CUDA graph capture scope.\n\n### Pattern 3: New Bookkeeping Operations\n**Symptom**: New NVTX ranges like `_write_finish_reasons`, `handle_additional_outputs`.\n**Cause**: New features added to the inference loop without overhead budgeting.\n**Mitigation**: Defer non-critical bookkeeping to async paths; batch updates.\n\n### Pattern 4: Flashinfer JIT Warmup Masquerading as Inference\n**Symptom**: Massive elementwise/reduce kernel counts in \"steady state\" analysis.\n**Cause**: Analysis window includes flashinfer JIT compilation phase.\n**Fix**: Use allreduce-based iteration isolation, not kernel density or time windows.\n\n### Pattern 5: Context-Only Bottleneck (Masked by Aggregate)\n**Symptom**: Aggregate metrics below threshold, but context iterations have 50% GPU idle.\n**Cause**: Generation iterations dilute the context-phase bottleneck.\n**Fix**: Per-phase analysis in Detection phase catches this.\n\n---\n\n## Pitfalls\n\n### 1. shortName is an Integer ID\nIn `CUPTI_ACTIVITY_KIND_KERNEL`, `shortName` is an integer referencing `StringIds.id`. Always join. See [references/nsys-schema.md](references/nsys-schema.md).\n\n### 2. NVTX textId vs text\nMost NVTX events have `textId` (integer) but NULL `text`. Join with StringIds. See [references/nsys-schema.md](references/nsys-schema.md).\n\n### 3. Duplicate NVTX Ranges from TP Ranks\nIn TP configurations, each rank reports NVTX ranges independently. De-duplicate by grouping entries within 100us of each other.\n\n### 4. Negative Inter-Step Gaps\nWhen TP ranks report overlapping NVTX ranges, `gap = next_start - prev_end` can be negative. Use the maximum end time when de-duplicating.\n\n### 5. Benchmark Window Selection\nThe allreduce-based window captures context+generation phases; steady-state NVTX filtering captures generation-only. Both are valid; use the appropriate one for your comparison goal.\n\n---\n\n## Handoff to Optimization\n\nWhen analysis is complete and the verdict is **YES**, hand off to the `perf-host-optimization` skill with:\n\n1. **Detection verdict and evidence**: Which metrics crossed thresholds (M1-M5), whether host prep was confirmed (M3b+M3c), and per-phase breakdown.\n\n2. **NVTX-based triage** (from Root Cause): Top regressing operations by absolute delta (us/step). Map NVTX range names to source functions -- see [references/trtllm-nvtx-ranges.md](references/trtllm-nvtx-ranges.md).\n\n3. **Handoff data block**: Include structured data from [references/output-format.md](references/output-format.md) (see \"Handoff to Optimization\" section).\n\n4. **Kernel-level findings** (from drill-down, if performed): Inter-kernel gap distribution, graph coverage ratio, per-group overhead map, and straggler rank identification. Map findings to optimization patterns using the table in the Root Cause kernel-level drill-down section above.\n\n---\n\n## Reference\n\n| File | Contents |\n|------|----------|\n| [references/metrics.md](references/metrics.md) | Full metric definitions, formulas, SQL queries, M3 sub-metric analysis |\n| [references/thresholds.md](references/thresholds.md) | Aggregate and per-phase threshold tables |\n| [references/phase-classification.md](references/phase-classification.md) | NVTX marker parsing, iteration classification, per-phase aggregation |\n| [references/output-format.md](references/output-format.md) | Report template and integration JSON schema |\n| [references/examples.md](references/examples.md) | Worked scenarios (aggregate, phase-specific, and case study) |\n| [references/iteration-isolation-techniques.md](references/iteration-isolation-techniques.md) | Allreduce, NVTX, and kernel-density iteration isolation techniques |\n| [references/trtllm-nvtx-ranges.md](references/trtllm-nvtx-ranges.md) | TRT-LLM NVTX range reference with per-operation timings |\n| [references/kernel-level-analysis.md](references/kernel-level-analysis.md) | Kernel-level drill-down techniques: gap analysis, graph classification, pattern mapping, straggler detection |\n| [references/nsys-schema.md](references/nsys-schema.md) | nsys SQLite schema reference and useful queries |\n| [scripts/analyze_host_overhead.py](scripts/analyze_host_overhead.py) | Python script for Phase 2 root cause analysis |\n| [scripts/detect_host_overhead.py](scripts/detect_host_overhead.py) | Python script for Phase 1 detection verdict |\n","frontmatter":{"name":"perf-host-analysis","description":"Analyze host/CPU overhead in TensorRT-LLM inference from nsys traces. Detect whether host overhead is the bottleneck using GPU idle ratio, host prep exposed ratio, and per-phase evidence. For regressions, isolate forward steps via allreduce/NVTX patterns, compare host operation breakdowns across versions, and identify scheduling or request-management overhead. Supports optional inter-kernel gap, eager-vs-graph, pattern mapping, and multi-rank straggler drill-down. Use standalone or within perf-analysis. Triggers: host overhead, inter-step gap, scheduling overhead, forward step isolation, nsys iteration analysis, NVTX breakdown, request management overhead, GPU idle, host bottleneck, host prep exposed, inter-kernel gap, bubble analysis, graph coverage, eager kernel, rank imbalance, straggler detection.\n","license":"Apache-2.0","tags":["analysis","profiling","host-overhead","nsys","inference"],"dependencies":["workload-instrumentation","trace-interpretation"],"metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":3950,"sizeBytes":16972},{"name":"communication-patterns.md","path":".claude/skills/perf-host-optimization/references/communication-patterns.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/communication-patterns.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/communication-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/communication-patterns.md)","tokens":300,"sizeBytes":267},{"name":"examples.md","path":".claude/skills/perf-host-optimization/references/examples.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/examples.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/examples.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/examples.md)","tokens":300,"sizeBytes":239},{"name":"hot-path-files.md","path":".claude/skills/perf-host-optimization/references/hot-path-files.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/hot-path-files.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/hot-path-files.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/hot-path-files.md)","tokens":300,"sizeBytes":251},{"name":"hotspot-classification.md","path":".claude/skills/perf-host-optimization/references/hotspot-classification.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/hotspot-classification.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/hotspot-classification.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/hotspot-classification.md)","tokens":300,"sizeBytes":267},{"name":"optimization-patterns.md","path":".claude/skills/perf-host-optimization/references/optimization-patterns.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/optimization-patterns.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/optimization-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/optimization-patterns.md)","tokens":300,"sizeBytes":265},{"name":"optimization-strategy.md","path":".claude/skills/perf-host-optimization/references/optimization-strategy.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/optimization-strategy.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-host-optimization/references/optimization-strategy.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/optimization-strategy.md)","tokens":300,"sizeBytes":265},{"name":"compound-pitfalls.md","path":".claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md","title":"Patterns Skill","category":"anthropic-skill","format":"markdown","content":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/compound-pitfalls.md)","tokens":300,"sizeBytes":273},{"name":"gpu-graph.md","path":".claude/skills/perf-host-optimization/references/patterns/gpu-graph.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/gpu-graph.md","title":"Patterns Skill","category":"anthropic-skill","format":"markdown","content":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/gpu-graph.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/gpu-graph.md)","tokens":300,"sizeBytes":257},{"name":"loop-iteration.md","path":".claude/skills/perf-host-optimization/references/patterns/loop-iteration.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/loop-iteration.md","title":"Patterns Skill","category":"anthropic-skill","format":"markdown","content":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/loop-iteration.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/loop-iteration.md)","tokens":300,"sizeBytes":267},{"name":"python-overhead.md","path":".claude/skills/perf-host-optimization/references/patterns/python-overhead.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/python-overhead.md","title":"Patterns Skill","category":"anthropic-skill","format":"markdown","content":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/python-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/python-overhead.md)","tokens":300,"sizeBytes":269},{"name":"sync-alloc.md","path":".claude/skills/perf-host-optimization/references/patterns/sync-alloc.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/sync-alloc.md","title":"Patterns Skill","category":"anthropic-skill","format":"markdown","content":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/sync-alloc.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/sync-alloc.md)","tokens":300,"sizeBytes":259},{"name":"system.md","path":".claude/skills/perf-host-optimization/references/patterns/system.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/system.md","title":"Patterns Skill","category":"anthropic-skill","format":"markdown","content":"# Patterns Skill\n\nPath: `.claude/skills/perf-host-optimization/references/patterns/system.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/references/patterns/system.md)","tokens":300,"sizeBytes":251},{"name":"SKILL.md","path":".claude/skills/perf-host-optimization/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-optimization/SKILL.md","title":"perf-host-optimization","category":"anthropic-skill","format":"markdown","content":"---\nname: perf-host-optimization\ndescription: Profiles and optimizes TensorRT-LLM host/CPU overhead using line_profiler (with nsys support planned). Runs iterative profile-analyze-optimize-validate rounds. Use when GPU utilization is low or optimizing PyExecutor throughput.\nlicense: Apache-2.0\ntags:\n  - optimization\n  - profiling\n  - host-overhead\n  - line-profiler\n  - inference\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Host Performance Optimization Skill\n\nAutomates detection and optimization of host-side (CPU) overhead in TensorRT-LLM's PyTorch backend.\n\n## When to Use\n\n- GPU utilization is low during inference (CPU bottleneck suspected)\n- User asks to reduce host overhead or CPU latency\n- Optimizing PyExecutor throughput (requests/sec)\n- Need line-by-line profiling of specific Python functions\n\n### Confirming the Bottleneck\n\nline_profiler measures *where* CPU time is spent but not *whether* CPU is the bottleneck.\nIf you need to confirm CPU is the limiting factor, run the `perf-host-analysis` skill first -- it provides a YES/NO verdict with metric evidence.\n\nAs a rough heuristic without nsys: if doubling the batch size does not proportionally increase GPU utilization or throughput, CPU overhead is likely the bottleneck.\n\n### Using Analysis Skill Results\n\nIf the `perf-host-analysis` skill has already been run, use its output to skip the confirmation step and prioritize targets:\n\n1. **Detection verdict**: If YES with host_prep_confirmed, start with `_prepare_tp_inputs`.\n2. **NVTX triage** (from Root Cause): The `top_regressing_ops` in the handoff data block maps NVTX range names to source functions. Profile the function with the largest absolute delta first.\n3. **Cross-function triage**: When the top NVTX regression is NOT in `_prepare_tp_inputs` (e.g., `_fetch_new_requests`, `broadcast_requests`, `_update_requests`), target that function's source file directly instead of defaulting to `_prepare_tp_inputs`. See [references/trtllm-nvtx-ranges.md](../perf-host-analysis/references/trtllm-nvtx-ranges.md) for the NVTX-to-source mapping.\n\n---\n\n## Profiling Setup\n\n### line_profiler (Primary Method)\n\n**Environment Variables:**\n- `TLLM_LINE_PROFILER_ENABLED=True` — Enable the profiler\n- `TLLM_LINE_PROFILER_PATH` — Output file path\n- `TLLM_LINE_PROFILER_FUNCTIONS` — Additional functions to profile (comma-separated)\n\n**Function specification format:**\n```bash\n# Class methods: module.path.ClassName.method_name\nTLLM_LINE_PROFILER_FUNCTIONS=\"tensorrt_llm._torch.pyexecutor.model_engine.PyTorchModelEngine._prepare_tp_inputs\"\n\n# Standalone functions: module.path::function_name\nTLLM_LINE_PROFILER_FUNCTIONS=\"tensorrt_llm._torch.pyexecutor.sampler::_group_requests_by_strategy_key\"\n\n# Multiple functions (comma-separated)\nTLLM_LINE_PROFILER_FUNCTIONS=\"module.Class.method1,module.Class.method2\"\n```\n\n### CPU Affinity (Environment Factor)\n\nCPU core affinity can significantly affect host overhead measurements,\nespecially on multi-socket systems (e.g., B300). Pinning processes to cores\nnear the GPU's NUMA node reduces cross-socket memory access latency.\n\n- Check current affinity: `taskset -p <pid>` or `numactl --show`\n- Pin to local NUMA node: `numactl --cpunodebind=<node> --membind=<node>`\n- **Impact**: Up to 2x difference in host overhead on B300 systems\n\nWhen comparing profiling results across runs, ensure CPU affinity is consistent.\nDo not externally modify the affinity, unless user requires to do this to examine the affects of this part.\nDocument the affinity setting in each round's report if it varies.\n\n### Workspace & Suffix Management\n\nEach profiling run should have a unique suffix to track progress across rounds:\n```bash\nEXTRA_SUFFIX=round0_baseline bash profile.sh\nEXTRA_SUFFIX=round1_eliminate_redundant_iter bash profile.sh\n```\n\n---\n\n## Autonomous Optimization Loop\n\nBefore starting the loop, review [references/optimization-strategy.md](references/optimization-strategy.md) for strategic guidance on ordering (zero-risk-first), measurement traps, and overhead scoping.\n\n**Key insight**: Optimizations are NOT independent. Fixing a 50ms bottleneck may reveal a 30ms bottleneck that was previously masked (hidden behind the larger one). Always re-profile after each significant change — the bottleneck landscape shifts.\n\n**Ordering principle**: Within each round, prefer zero-risk optimizations (caching, pre-allocation, hoisting invariants) over medium/high-risk ones (graph partition changes, algorithm fusion). Zero-risk changes provide free gains and make subsequent profiling cleaner.\n\nRun N rounds (default 3) of the following cycle:\n\n```\nFOR round = 1 to MAX_ROUNDS:\n\n  1. PROFILE (with Drill-Down)\n  2. ANALYZE (Multi-Option)\n  3. OPTIMIZE (Apply Change — prefer zero-risk first)\n  4. TEST (Unit Test Validation)\n  5. VALIDATE (Re-Profile — expect bottleneck landscape to shift)\n  6. REPORT\n\nEND FOR → FINAL SUMMARY\n```\n\n### Phase 1: PROFILE (with Drill-Down)\n\n- Run workload with profiler enabled\n- Parse output: identify functions with highest Total time and lines with highest % Time\n- **CRITICAL: Drill down into sub-functions that are not yet profiled** (see below)\n\n#### Drill-Down Profiling\n\nThe default profiler covers top-level executor functions but **not all sub-functions**. When a profiled function shows most time in a single sub-call, you must drill down.\n\n**When:** A single line consumes >80% of a function's time calling an unprofiled sub-function:\n```\nLine #      Hits         Time    Per Hit   % Time  Line Contents\n==============================================================\n  2848      4100  59200000000.0  14439024.4   98.7      output = self.model_engine.forward(...)\n```\n\n**How:**\n1. Identify the sub-function's full qualified path (e.g., `tensorrt_llm._torch.pyexecutor.model_engine.PyTorchModelEngine._prepare_tp_inputs`)\n2. Add it to `TLLM_LINE_PROFILER_FUNCTIONS`\n3. Re-profile to get line-level data inside it\n4. Now analyze the **inner** hotspots\n\nFor common drill-down targets, see [references/hot-path-files.md](references/hot-path-files.md).\n\n### Phase 2: ANALYZE (Multi-Option)\n\nFor the chosen hotspot:\n\n1. **Identify** the top hotspots by **absolute time** (not just %) within the target function\n2. **Classify** each hotspot by type. Summary table:\n\n| Type | Indicators | Severity |\n|------|------------|----------|\n| **HOST_SYNC** | `.item()`, `.cpu()` in per-layer forward path | Critical |\n| **SYNC** | `.item()`, `.cpu()`, `synchronize()` in step-level code | Critical |\n| **CUSTOM_OP** | Chain of Python tensor ops (view/slice/cast) before kernel launch | Critical |\n| **GRAPH_BREAK** | Op that prevents CUDA graph capture of surrounding code (fix via GRAPH_EXPAND / GRAPH_SPLIT) | High |\n| **ALLOC** | `torch.zeros/empty/tensor()` in loops, `.clone()` | High |\n| **HOIST** | Per-layer recomputation of step-invariant values | High |\n| **PYLOOP** | `for x in collection:` with many iterations | High |\n| **REDUNDANT_ITER** | Multiple passes over the same collection | High |\n| **DEAD_WORK** | Object construction whose results are always discarded | High |\n| **CONTAINER** | Dict/set lookups in hot loops | Medium |\n| **FUNCALL** | Repeated method/property calls | Medium |\n| **COMM** | `dist.all_reduce`, `dist.barrier`, NCCL overhead in TP/PP paths | Medium |\n| **GIL** | Lock/queue contention | Medium |\n| **SERIALIZE** | `pickle.dumps/loads`, `json.dumps/loads` in request processing | Medium |\n| **GC** | Periodic latency spikes, non-deterministic pauses (tail latency) | Low |\n| **COMPUTE** | Actual computation (may not be optimizable) | Low |\n\nFor detailed classification with code examples, see [references/hotspot-classification.md](references/hotspot-classification.md).\n\n3. **Propose 2-4 optimization options** in a table:\n\n| Option | Description | Estimated Savings | Risk | Complexity |\n|--------|-------------|-------------------|------|------------|\n| A | ... | ... | Low/Med/High | ... |\n| B | ... | ... | ... | ... |\n\n4. **Select the best option** and explain reasoning (prefer high-savings + low-risk; follow zero-risk-first ordering from [references/optimization-strategy.md](references/optimization-strategy.md))\n\nFor optimization patterns by type, see [references/optimization-patterns.md](references/optimization-patterns.md) (index) — it links to the relevant sub-file for each hotspot type. For GPU-specific patterns (CUSTOM_OP, GRAPH_SPLIT, GRAPH_EXPAND), see [references/patterns/gpu-graph.md](references/patterns/gpu-graph.md).\n\n### Phase 3: OPTIMIZE (Apply Change)\n\n- Apply the selected code change with Edit tool\n- **One optimization per round** — keep changes minimal and targeted\n- Record the exact change (file, line range, before/after) for potential rollback\n\n### Phase 4: TEST (Unit Test Validation)\n\n**Mandatory** after each optimization. Find and run related UTs to verify correctness.\n\n**Finding related tests:**\n```bash\n# Search by modified file name\ngrep -rl \"model_engine\\|PyTorchModelEngine\" tests/unittest/_torch/executor/\n\n# Search by modified function name\ngrep -rl \"_prepare_tp_inputs\\|prepare_inputs\" tests/\n```\n\n**Running tests:**\n```bash\n# Run specific test file with stop-on-first-failure\npytest tests/unittest/_torch/executor/test_pytorch_model_engine.py -v -x --timeout=120\n\n# Run specific test method\npytest tests/unittest/_torch/executor/test_pytorch_model_engine.py::PyTorchModelEngineTestCase::test_position_id_preparation -v -x\n```\n\nFor the full UT-to-file mapping, see [references/hot-path-files.md](references/hot-path-files.md).\n\n**If tests fail:**\n1. Read the failure message\n2. Rollback immediately (`git checkout -- <file>`)\n3. Analyze why the optimization broke correctness\n4. Try the next-best option from Phase 2\n\n### Phase 5: VALIDATE (Re-Profile)\n\n- Re-run profiler with identical workload, using suffix `round<N>_<description>`\n- Compare three things:\n  1. Did the **target hotspot** time decrease?\n  2. Did the **overall function** Total time decrease?\n  3. Did **benchmark metrics** (TPOT, throughput) improve?\n\n**If regression detected** (function time increased or metrics worsened):\n- The \"optimization\" may have triggered a CPython pitfall — see [references/patterns/compound-pitfalls.md](references/patterns/compound-pitfalls.md) (CPython Pitfalls section)\n- Rollback and try the next-best option from Phase 2\n\n### Phase 6: REPORT\n\nLog for this round:\n- Round number\n- Hotspot location (file:line) and classification\n- Optimization applied (with before/after code summary)\n- Time delta: function Total time before → after\n- Benchmark delta: TPOT, throughput before → after\n\n---\n\n## Reading Profile Output\n\n```\nTimer unit: 1e-06 s\nTotal time: 1.234 s\nFile: /path/to/file.py\nFunction: my_function at line 100\n\nLine #      Hits         Time  Per Hit   % Time  Line Contents\n==============================================================\n   100                                           def my_function(self):\n   101       500    890000.0   1780.0     72.1       result = tensor.item()\n   102       500    234567.0    469.1     19.0       return result\n```\n\n**How to read effectively:**\n1. Start with **Total time** for each function — this is the overall budget\n2. Sort lines mentally by **absolute Time**, not just % Time (3% of a 60s function = 1.8s)\n3. Check **Hits count** to understand iteration patterns:\n   - Hits = 2 × expected count → `for x in range(1):` loop overhead (2 hits = enter + exit check)\n   - Hits ≫ expected → the line is inside a nested loop\n4. Look for **repeated patterns**: if 10 lines each take 3% in a loop body, the loop itself costs 30%\n\n---\n\n## Stopping Criteria\n\nStop the optimization loop when:\n1. **Iteration limit reached**: Completed N rounds (default 3)\n2. **No actionable hotspots**: Top hotspots are pure GPU compute (COMPUTE type)\n3. **Diminishing returns**: < 5% improvement in last 2 rounds\n4. **Risk threshold**: Further optimizations require architectural changes (e.g., Cython, struct-of-arrays)\n5. **Test failures**: Cannot find an optimization that passes UTs\n\n**Primary success metric**: Benchmark throughput (requests/sec or tokens/sec) as measured by the profiling script. line_profiler time reductions are leading indicators, but throughput is the ground truth — a function-level speedup that doesn't improve throughput is not a real win.\n\n---\n\n## Final Summary Output\n\nThe final report should include:\n- **Rounds executed**: Number of profile-optimize cycles completed\n- **Cumulative improvement**: Total host time reduction (percentage and absolute)\n- **Benchmark metrics**: Before/after comparison table (TPOT, throughput, ITL, E2EL)\n- **Optimizations applied**: List of changes with file:line locations and classification\n- **Failed attempts**: Any optimizations tried and reverted (with why)\n- **Remaining hotspots**: Top bottlenecks that couldn't be optimized (with classification)\n- **Recommendations**: Suggested follow-up for architectural changes if needed\n\nFor a concrete multi-round example, see [references/examples.md](references/examples.md).\n\n---\n\n## Reference Files\n\n| File | Contents |\n|------|----------|\n| [references/optimization-patterns.md](references/optimization-patterns.md) | Pattern index — links to 6 sub-files: [sync-alloc](references/patterns/sync-alloc.md), [loop-iteration](references/patterns/loop-iteration.md), [python-overhead](references/patterns/python-overhead.md), [gpu-graph](references/patterns/gpu-graph.md), [system](references/patterns/system.md), [compound-pitfalls](references/patterns/compound-pitfalls.md) |\n| [references/optimization-strategy.md](references/optimization-strategy.md) | Zero-risk-first ordering, metric traps, three scopes of host overhead, pattern selection guide |\n| [references/hotspot-classification.md](references/hotspot-classification.md) | Extended per-type indicators and code examples (including CUSTOM_OP, GRAPH_BREAK, HOST_SYNC) |\n| [references/communication-patterns.md](references/communication-patterns.md) | Communication overhead patterns (NCCL batching, barrier removal, async overlap, reduce_scatter) |\n| [references/hot-path-files.md](references/hot-path-files.md) | Key file tables, drill-down targets, UT mapping |\n| [references/examples.md](references/examples.md) | Usage examples and multi-round walkthrough |\n| [trtllm-nvtx-ranges.md](../perf-host-analysis/references/trtllm-nvtx-ranges.md) | TRT-LLM NVTX range reference (from analysis skill) — maps range names to source functions |\n","frontmatter":{"name":"perf-host-optimization","description":"Profiles and optimizes TensorRT-LLM host/CPU overhead using line_profiler (with nsys support planned). Runs iterative profile-analyze-optimize-validate rounds. Use when GPU utilization is low or optimizing PyExecutor throughput.","license":"Apache-2.0","tags":["optimization","profiling","host-overhead","line-profiler","inference"],"metadata":{"author":"NVIDIA Corporation"}},"isInternal":false,"tokens":3427,"sizeBytes":14450},{"name":"aliases.yaml","path":".claude/skills/perf-instability-cookbook/data/aliases.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/aliases.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/aliases.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/aliases.yaml)","tokens":300,"sizeBytes":229},{"name":"modules.yaml","path":".claude/skills/perf-instability-cookbook/data/modules.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/modules.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/modules.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/modules.yaml)","tokens":300,"sizeBytes":229},{"name":"patterns.yaml","path":".claude/skills/perf-instability-cookbook/data/patterns.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/patterns.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/patterns.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/patterns.yaml)","tokens":300,"sizeBytes":231},{"name":"pending.yaml","path":".claude/skills/perf-instability-cookbook/data/pending.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/pending.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/pending.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/pending.yaml)","tokens":300,"sizeBytes":229},{"name":"removed.yaml","path":".claude/skills/perf-instability-cookbook/data/removed.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/removed.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/removed.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/removed.yaml)","tokens":300,"sizeBytes":229},{"name":"tags.yaml","path":".claude/skills/perf-instability-cookbook/data/tags.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/tags.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-instability-cookbook/data/tags.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/data/tags.yaml)","tokens":300,"sizeBytes":223},{"name":"case-template.md","path":".claude/skills/perf-instability-cookbook/references/case-template.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/case-template.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/case-template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/case-template.md)","tokens":300,"sizeBytes":255},{"name":"index.md","path":".claude/skills/perf-instability-cookbook/references/communication/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/index.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/index.md)","tokens":300,"sizeBytes":270},{"name":"nccl-symmetric-graceful-fallback.md","path":".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-graceful-fallback.md)","tokens":300,"sizeBytes":324},{"name":"nccl-symmetric-off-on-gb10-spark.md","path":".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-off-on-gb10-spark.md)","tokens":300,"sizeBytes":324},{"name":"nccl-symmetric-preallocation-for-autotuning.md","path":".claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/communication/nccl-symmetric-preallocation-for-autotuning.md)","tokens":300,"sizeBytes":346},{"name":"deepgemm-paged-mqa-logits-prewarm.md","path":".claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/deepgemm-paged-mqa-logits-prewarm.md)","tokens":300,"sizeBytes":329},{"name":"flashinfer-sampling-jit-on-served-path.md","path":".claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/flashinfer-sampling-jit-on-served-path.md)","tokens":300,"sizeBytes":339},{"name":"fmha-cubin-drop-align-kernel-selection.md","path":".claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/fmha-cubin-drop-align-kernel-selection.md)","tokens":300,"sizeBytes":339},{"name":"index.md","path":".claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/index.md)","tokens":300,"sizeBytes":273},{"name":"kv-aware-adp-cold-start.md","path":".claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/kv-aware-adp-cold-start.md)","tokens":300,"sizeBytes":309},{"name":"mamba-hybrid-warmup-gap.md","path":".claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/mamba-hybrid-warmup-gap.md)","tokens":300,"sizeBytes":309},{"name":"visual-gen-i2v-warmup.md","path":".claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-i2v-warmup.md)","tokens":300,"sizeBytes":305},{"name":"visual-gen-warmup-shapes-configurable.md","path":".claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/jit-and-warmup/visual-gen-warmup-shapes-configurable.md)","tokens":300,"sizeBytes":337},{"name":"force-num-accepted-tokens-in-spec-perf-test.md","path":".claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/force-num-accepted-tokens-in-spec-perf-test.md)","tokens":300,"sizeBytes":367},{"name":"fractional-synthetic-acceptance-rates.md","path":".claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/fractional-synthetic-acceptance-rates.md)","tokens":300,"sizeBytes":355},{"name":"gen-only-log-flush-sentinel.md","path":".claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-log-flush-sentinel.md)","tokens":300,"sizeBytes":335},{"name":"gen-only-prev-device-step-time-race.md","path":".claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/gen-only-prev-device-step-time-race.md)","tokens":300,"sizeBytes":351},{"name":"ignore-eos-with-spec-decoding.md","path":".claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/ignore-eos-with-spec-decoding.md)","tokens":300,"sizeBytes":339},{"name":"index.md","path":".claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/measurement-and-test/index.md)","tokens":300,"sizeBytes":291},{"name":"index.md","path":".claude/skills/perf-instability-cookbook/references/moe/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/moe/index.md","title":"Moe Skill","category":"anthropic-skill","format":"markdown","content":"# Moe Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/moe/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/moe/index.md)","tokens":300,"sizeBytes":240},{"name":"moe-deepep-gc-hang.md","path":".claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md","title":"Moe Skill","category":"anthropic-skill","format":"markdown","content":"# Moe Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/moe/moe-deepep-gc-hang.md)","tokens":300,"sizeBytes":266},{"name":"autodeploy-gc-gen0-threshold-unforwarded.md","path":".claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md","title":"Scheduler-and-executor Skill","category":"anthropic-skill","format":"markdown","content":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/autodeploy-gc-gen0-threshold-unforwarded.md)","tokens":300,"sizeBytes":367},{"name":"index.md","path":".claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md","title":"Scheduler-and-executor Skill","category":"anthropic-skill","format":"markdown","content":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/references/scheduler-and-executor/index.md)","tokens":300,"sizeBytes":297},{"name":"SKILL.md","path":".claude/skills/perf-instability-cookbook/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/SKILL.md","title":"perf-instability-cookbook","category":"anthropic-skill","format":"markdown","content":"# perf-instability-cookbook\n\nPath: `.claude/skills/perf-instability-cookbook/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-instability-cookbook/SKILL.md)","tokens":300,"sizeBytes":226},{"name":"advanced-profiling.md","path":".claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/advanced-profiling.md)","tokens":300,"sizeBytes":271},{"name":"bottleneck-guide.md","path":".claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/bottleneck-guide.md)","tokens":300,"sizeBytes":267},{"name":"cli-reference.md","path":".claude/skills/perf-nsight-compute-analysis/references/cli-reference.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/cli-reference.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/cli-reference.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/cli-reference.md)","tokens":300,"sizeBytes":261},{"name":"memory-analysis.md","path":".claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/memory-analysis.md)","tokens":300,"sizeBytes":265},{"name":"metrics-guide.md","path":".claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/metrics-guide.md)","tokens":300,"sizeBytes":261},{"name":"python-report-api.md","path":".claude/skills/perf-nsight-compute-analysis/references/python-report-api.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/python-report-api.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/python-report-api.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/python-report-api.md)","tokens":300,"sizeBytes":269},{"name":"roofline-analysis.md","path":".claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/roofline-analysis.md)","tokens":300,"sizeBytes":269},{"name":"sections-guide.md","path":".claude/skills/perf-nsight-compute-analysis/references/sections-guide.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/sections-guide.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-compute-analysis/references/sections-guide.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/references/sections-guide.md)","tokens":300,"sizeBytes":263},{"name":"SKILL.md","path":".claude/skills/perf-nsight-compute-analysis/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/SKILL.md","title":"perf-nsight-compute-analysis","category":"anthropic-skill","format":"markdown","content":"# perf-nsight-compute-analysis\n\nPath: `.claude/skills/perf-nsight-compute-analysis/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-compute-analysis/SKILL.md)","tokens":300,"sizeBytes":235},{"name":"app-preparation.md","path":".claude/skills/perf-nsight-systems/references/app-preparation.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/app-preparation.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/app-preparation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/app-preparation.md)","tokens":300,"sizeBytes":247},{"name":"cli-post-collection.md","path":".claude/skills/perf-nsight-systems/references/cli-post-collection.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/cli-post-collection.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/cli-post-collection.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/cli-post-collection.md)","tokens":300,"sizeBytes":255},{"name":"cli-profiling.md","path":".claude/skills/perf-nsight-systems/references/cli-profiling.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/cli-profiling.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/cli-profiling.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/cli-profiling.md)","tokens":300,"sizeBytes":243},{"name":"expert-systems.md","path":".claude/skills/perf-nsight-systems/references/expert-systems.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/expert-systems.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/expert-systems.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/expert-systems.md)","tokens":300,"sizeBytes":245},{"name":"nvtx-analysis.md","path":".claude/skills/perf-nsight-systems/references/nvtx-analysis.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/nvtx-analysis.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/nvtx-analysis.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/nvtx-analysis.md)","tokens":300,"sizeBytes":243},{"name":"recipes-dl.md","path":".claude/skills/perf-nsight-systems/references/recipes-dl.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/recipes-dl.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/recipes-dl.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/recipes-dl.md)","tokens":300,"sizeBytes":237},{"name":"stats-reports.md","path":".claude/skills/perf-nsight-systems/references/stats-reports.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/stats-reports.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-nsight-systems/references/stats-reports.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/references/stats-reports.md)","tokens":300,"sizeBytes":243},{"name":"SKILL.md","path":".claude/skills/perf-nsight-systems/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/SKILL.md","title":"perf-nsight-systems","category":"anthropic-skill","format":"markdown","content":"# perf-nsight-systems\n\nPath: `.claude/skills/perf-nsight-systems/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-nsight-systems/SKILL.md)","tokens":300,"sizeBytes":208},{"name":"aliases.yaml","path":".claude/skills/perf-optimization-casebook/data/aliases.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/aliases.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-optimization-casebook/data/aliases.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/aliases.yaml)","tokens":300,"sizeBytes":231},{"name":"patterns.yaml","path":".claude/skills/perf-optimization-casebook/data/patterns.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/patterns.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-optimization-casebook/data/patterns.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/patterns.yaml)","tokens":300,"sizeBytes":233},{"name":"tags.yaml","path":".claude/skills/perf-optimization-casebook/data/tags.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/tags.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-optimization-casebook/data/tags.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/data/tags.yaml)","tokens":300,"sizeBytes":225},{"name":"case-template.md","path":".claude/skills/perf-optimization-casebook/references/case-template.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/case-template.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/case-template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/case-template.md)","tokens":300,"sizeBytes":257},{"name":"deepep.md","path":".claude/skills/perf-optimization-casebook/references/communication/deepep.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/deepep.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/deepep.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/deepep.md)","tokens":300,"sizeBytes":274},{"name":"index.md","path":".claude/skills/perf-optimization-casebook/references/communication/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/index.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/index.md)","tokens":300,"sizeBytes":272},{"name":"low-precision-dispatch.md","path":".claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/low-precision-dispatch.md)","tokens":300,"sizeBytes":306},{"name":"mnnvl-twoshot-allreduce.md","path":".claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/mnnvl-twoshot-allreduce.md)","tokens":300,"sizeBytes":308},{"name":"shape-aware-allreduce-autotune.md","path":".claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/shape-aware-allreduce-autotune.md)","tokens":300,"sizeBytes":322},{"name":"userbuffers-symmetric-memory.md","path":".claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/communication/userbuffers-symmetric-memory.md)","tokens":300,"sizeBytes":318},{"name":"fold-scale-swizzle-into-kernel.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fold-scale-swizzle-into-kernel.md)","tokens":300,"sizeBytes":334},{"name":"fp8-mla-kv-cache.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fp8-mla-kv-cache.md)","tokens":300,"sizeBytes":306},{"name":"fuse-add-norm-quant.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-add-norm-quant.md)","tokens":300,"sizeBytes":312},{"name":"fuse-ar-epilogue.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-ar-epilogue.md)","tokens":300,"sizeBytes":306},{"name":"fuse-datamovement-into-quantize.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-datamovement-into-quantize.md)","tokens":300,"sizeBytes":336},{"name":"fuse-moe-routing-kernel.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-moe-routing-kernel.md)","tokens":300,"sizeBytes":320},{"name":"fuse-qk-norm-rope.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/fuse-qk-norm-rope.md)","tokens":300,"sizeBytes":308},{"name":"hw-matched-lowprec-moe-gemm.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/hw-matched-lowprec-moe-gemm.md)","tokens":300,"sizeBytes":328},{"name":"index.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/index.md)","tokens":300,"sizeBytes":284},{"name":"mega-fuse-moe-deepgemm.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/mega-fuse-moe-deepgemm.md)","tokens":300,"sizeBytes":318},{"name":"ranking-only-precision-tf32.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/ranking-only-precision-tf32.md)","tokens":300,"sizeBytes":328},{"name":"reevaluate-fusion-boundary-per-dtype.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/reevaluate-fusion-boundary-per-dtype.md)","tokens":300,"sizeBytes":346},{"name":"relax-tl-constexpr.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/relax-tl-constexpr.md)","tokens":300,"sizeBytes":310},{"name":"sparse-mla-topk-attention.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/sparse-mla-topk-attention.md)","tokens":300,"sizeBytes":324},{"name":"specialize-topk-selection-kernel.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/specialize-topk-selection-kernel.md)","tokens":300,"sizeBytes":338},{"name":"split-mla-reduction-kernel.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/split-mla-reduction-kernel.md)","tokens":300,"sizeBytes":326},{"name":"triton-to-cpp-op.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/triton-to-cpp-op.md)","tokens":300,"sizeBytes":306},{"name":"trtllm-gen-fp4-moe-backend.md","path":".claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md","title":"Kernel-and-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-and-fusion Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/kernel-and-fusion/trtllm-gen-fp4-moe-backend.md)","tokens":300,"sizeBytes":326},{"name":"attention-dp-padding.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/attention-dp-padding.md)","tokens":300,"sizeBytes":314},{"name":"auxiliary-cache-in-kv-manager.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/auxiliary-cache-in-kv-manager.md)","tokens":300,"sizeBytes":332},{"name":"cache-step-invariant-per-layer.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/cache-step-invariant-per-layer.md)","tokens":300,"sizeBytes":334},{"name":"chunked-prefill-aligned-auxiliary.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/chunked-prefill-aligned-auxiliary.md)","tokens":300,"sizeBytes":340},{"name":"cuda-graph-padding.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/cuda-graph-padding.md)","tokens":300,"sizeBytes":310},{"name":"free-mla-intermediates.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/free-mla-intermediates.md)","tokens":300,"sizeBytes":318},{"name":"hoist-torch-compile-closures.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/hoist-torch-compile-closures.md)","tokens":300,"sizeBytes":330},{"name":"index.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/index.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/index.md)","tokens":300,"sizeBytes":284},{"name":"mla-kv-cache-reuse.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/mla-kv-cache-reuse.md)","tokens":300,"sizeBytes":310},{"name":"move-bookkeeping-into-cpp-op.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/move-bookkeeping-into-cpp-op.md)","tokens":300,"sizeBytes":330},{"name":"multi-stream-shared-routed-expert.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/multi-stream-shared-routed-expert.md)","tokens":300,"sizeBytes":340},{"name":"overlap-mla-rope-uk-bgemm.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-mla-rope-uk-bgemm.md)","tokens":300,"sizeBytes":324},{"name":"overlap-online-eplb.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-online-eplb.md)","tokens":300,"sizeBytes":312},{"name":"overlap-scheduler.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/overlap-scheduler.md)","tokens":300,"sizeBytes":308},{"name":"pdl.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/pdl.md)","tokens":300,"sizeBytes":280},{"name":"piecewise-cuda-graph.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/piecewise-cuda-graph.md)","tokens":300,"sizeBytes":314},{"name":"pybind-wrapper-pure-python.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/pybind-wrapper-pure-python.md)","tokens":300,"sizeBytes":326},{"name":"relaxed-mtp-acceptance.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/relaxed-mtp-acceptance.md)","tokens":300,"sizeBytes":318},{"name":"skip-sparse-path-when-degenerate.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/skip-sparse-path-when-degenerate.md)","tokens":300,"sizeBytes":338},{"name":"split-custom-op-for-piecewise-capture.md","path":".claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md","title":"Runtime-execution Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-execution Skill\n\nPath: `.claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/references/runtime-execution/split-custom-op-for-piecewise-capture.md)","tokens":300,"sizeBytes":348},{"name":"SKILL.md","path":".claude/skills/perf-optimization-casebook/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/SKILL.md","title":"perf-optimization-casebook","category":"anthropic-skill","format":"markdown","content":"# perf-optimization-casebook\n\nPath: `.claude/skills/perf-optimization-casebook/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization-casebook/SKILL.md)","tokens":300,"sizeBytes":229},{"name":"SKILL.md","path":".claude/skills/perf-optimization/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization/SKILL.md","title":"perf-optimization","category":"anthropic-skill","format":"markdown","content":"# perf-optimization\n\nPath: `.claude/skills/perf-optimization/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-optimization/SKILL.md)","tokens":300,"sizeBytes":202},{"name":"aliases.yaml","path":".claude/skills/perf-regression-cookbook/data/aliases.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/aliases.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/aliases.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/aliases.yaml)","tokens":300,"sizeBytes":227},{"name":"modules.yaml","path":".claude/skills/perf-regression-cookbook/data/modules.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/modules.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/modules.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/modules.yaml)","tokens":300,"sizeBytes":227},{"name":"patterns.yaml","path":".claude/skills/perf-regression-cookbook/data/patterns.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/patterns.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/patterns.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/patterns.yaml)","tokens":300,"sizeBytes":229},{"name":"pending.yaml","path":".claude/skills/perf-regression-cookbook/data/pending.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/pending.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/pending.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/pending.yaml)","tokens":300,"sizeBytes":227},{"name":"removed.yaml","path":".claude/skills/perf-regression-cookbook/data/removed.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/removed.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/removed.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/removed.yaml)","tokens":300,"sizeBytes":227},{"name":"tags.yaml","path":".claude/skills/perf-regression-cookbook/data/tags.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/tags.yaml","title":"Data Skill","category":"anthropic-skill","format":"yaml","content":"# Data Skill\n\nPath: `.claude/skills/perf-regression-cookbook/data/tags.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/data/tags.yaml)","tokens":300,"sizeBytes":221},{"name":"dsa-indexer-host-overhead.md","path":".claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md","title":"Attention-fmha Skill","category":"anthropic-skill","format":"markdown","content":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-host-overhead.md)","tokens":300,"sizeBytes":311},{"name":"dsa-indexer-tp-duplicated-prefill.md","path":".claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md","title":"Attention-fmha Skill","category":"anthropic-skill","format":"markdown","content":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsa-indexer-tp-duplicated-prefill.md)","tokens":300,"sizeBytes":327},{"name":"dsmla-sparse-ctx-2cta-fallback.md","path":".claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md","title":"Attention-fmha Skill","category":"anthropic-skill","format":"markdown","content":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/dsmla-sparse-ctx-2cta-fallback.md)","tokens":300,"sizeBytes":321},{"name":"fmha-dispatcher-sm-count-query.md","path":".claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md","title":"Attention-fmha Skill","category":"anthropic-skill","format":"markdown","content":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/fmha-dispatcher-sm-count-query.md)","tokens":300,"sizeBytes":321},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/attention-fmha/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/index.md","title":"Attention-fmha Skill","category":"anthropic-skill","format":"markdown","content":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/index.md)","tokens":300,"sizeBytes":271},{"name":"trtllmgen-mla-decode-unwaived-slower.md","path":".claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md","title":"Attention-fmha Skill","category":"anthropic-skill","format":"markdown","content":"# Attention-fmha Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/attention-fmha/trtllmgen-mla-decode-unwaived-slower.md)","tokens":300,"sizeBytes":333},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/autotuner/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/autotuner/index.md","title":"Autotuner Skill","category":"anthropic-skill","format":"markdown","content":"# Autotuner Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/autotuner/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/autotuner/index.md)","tokens":300,"sizeBytes":256},{"name":"nested-nvfp4-gemm-autotune-dispatch.md","path":".claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md","title":"Autotuner Skill","category":"anthropic-skill","format":"markdown","content":"# Autotuner Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/autotuner/nested-nvfp4-gemm-autotune-dispatch.md)","tokens":300,"sizeBytes":316},{"name":"case-template.md","path":".claude/skills/perf-regression-cookbook/references/case-template.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/case-template.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/case-template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/case-template.md)","tokens":300,"sizeBytes":253},{"name":"allreduce-host-overhead-small-model-tp.md","path":".claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/allreduce-host-overhead-small-model-tp.md)","tokens":300,"sizeBytes":334},{"name":"allreduce-strategy-lut.md","path":".claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/allreduce-strategy-lut.md)","tokens":300,"sizeBytes":302},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/communication/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/index.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/index.md)","tokens":300,"sizeBytes":268},{"name":"nemotronh-mnnvl-nccl-fallback.md","path":".claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md","title":"Communication Skill","category":"anthropic-skill","format":"markdown","content":"# Communication Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/communication/nemotronh-mnnvl-nccl-fallback.md)","tokens":300,"sizeBytes":316},{"name":"cudagraph-default-batch-sizes.md","path":".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md","title":"Cuda-graph-and-compile Skill","category":"anthropic-skill","format":"markdown","content":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/cudagraph-default-batch-sizes.md)","tokens":300,"sizeBytes":343},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md","title":"Cuda-graph-and-compile Skill","category":"anthropic-skill","format":"markdown","content":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/index.md)","tokens":300,"sizeBytes":295},{"name":"mrope-graph-gate-text-only-eager.md","path":".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md","title":"Cuda-graph-and-compile Skill","category":"anthropic-skill","format":"markdown","content":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/mrope-graph-gate-text-only-eager.md)","tokens":300,"sizeBytes":349},{"name":"piecewise-attention-torch-compile-host-overhead.md","path":".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md","title":"Cuda-graph-and-compile Skill","category":"anthropic-skill","format":"markdown","content":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-attention-torch-compile-host-overhead.md)","tokens":300,"sizeBytes":379},{"name":"piecewise-cudagraph-capture-coverage.md","path":".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md","title":"Cuda-graph-and-compile Skill","category":"anthropic-skill","format":"markdown","content":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-capture-coverage.md)","tokens":300,"sizeBytes":357},{"name":"piecewise-cudagraph-far-ceiling-append.md","path":".claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md","title":"Cuda-graph-and-compile Skill","category":"anthropic-skill","format":"markdown","content":"# Cuda-graph-and-compile Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/cuda-graph-and-compile/piecewise-cudagraph-far-ceiling-append.md)","tokens":300,"sizeBytes":361},{"name":"disagg-fill-throttle-slow-start.md","path":".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md","title":"Disagg-orchestrator Skill","category":"anthropic-skill","format":"markdown","content":"# Disagg-orchestrator Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-fill-throttle-slow-start.md)","tokens":300,"sizeBytes":338},{"name":"disagg-server-gc-pauses.md","path":".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md","title":"Disagg-orchestrator Skill","category":"anthropic-skill","format":"markdown","content":"# Disagg-orchestrator Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/disagg-server-gc-pauses.md)","tokens":300,"sizeBytes":322},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md","title":"Disagg-orchestrator Skill","category":"anthropic-skill","format":"markdown","content":"# Disagg-orchestrator Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/disagg-orchestrator/index.md)","tokens":300,"sizeBytes":286},{"name":"deepgemm-pdl-import-time-cuda-context.md","path":".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md","title":"Gemm-and-quantization Skill","category":"anthropic-skill","format":"markdown","content":"# Gemm-and-quantization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/deepgemm-pdl-import-time-cuda-context.md)","tokens":300,"sizeBytes":356},{"name":"fp8-pertensor-quant-slower-than-te.md","path":".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md","title":"Gemm-and-quantization Skill","category":"anthropic-skill","format":"markdown","content":"# Gemm-and-quantization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/fp8-pertensor-quant-slower-than-te.md)","tokens":300,"sizeBytes":350},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md","title":"Gemm-and-quantization Skill","category":"anthropic-skill","format":"markdown","content":"# Gemm-and-quantization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/gemm-and-quantization/index.md)","tokens":300,"sizeBytes":292},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/index.md)","tokens":300,"sizeBytes":271},{"name":"maybe-compile-recompiles-every-call.md","path":".claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/maybe-compile-recompiles-every-call.md)","tokens":300,"sizeBytes":331},{"name":"mla-chunked-prefill-maybe-compiled-cat-warmup.md","path":".claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/mla-chunked-prefill-maybe-compiled-cat-warmup.md)","tokens":300,"sizeBytes":351},{"name":"model-engine-warmup-orchestration-gaps.md","path":".claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/model-engine-warmup-orchestration-gaps.md)","tokens":300,"sizeBytes":337},{"name":"trtllmgen-fmha-densify-grid.md","path":".claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-densify-grid.md)","tokens":300,"sizeBytes":315},{"name":"trtllmgen-fmha-jit-warmup.md","path":".claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/trtllmgen-fmha-jit-warmup.md)","tokens":300,"sizeBytes":311},{"name":"warmup-token-cap-revert.md","path":".claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md","title":"Jit-and-warmup Skill","category":"anthropic-skill","format":"markdown","content":"# Jit-and-warmup Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/jit-and-warmup/warmup-token-cap-revert.md)","tokens":300,"sizeBytes":307},{"name":"flinear-bias-copy-sm100.md","path":".claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md","title":"Kernel-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-fusion Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/flinear-bias-copy-sm100.md)","tokens":300,"sizeBytes":304},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md","title":"Kernel-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-fusion Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/index.md)","tokens":300,"sizeBytes":268},{"name":"qwen35-fused-ops-disabled-under-torch-compile.md","path":".claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md","title":"Kernel-fusion Skill","category":"anthropic-skill","format":"markdown","content":"# Kernel-fusion Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kernel-fusion/qwen35-fused-ops-disabled-under-torch-compile.md)","tokens":300,"sizeBytes":348},{"name":"dsa-kcache-dtype-size-estimate.md","path":".claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md","title":"Kv-cache-manager Skill","category":"anthropic-skill","format":"markdown","content":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/dsa-kcache-dtype-size-estimate.md)","tokens":300,"sizeBytes":327},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md","title":"Kv-cache-manager Skill","category":"anthropic-skill","format":"markdown","content":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/index.md)","tokens":300,"sizeBytes":277},{"name":"mamba-hybrid-recurrent-state-d2h-syncs.md","path":".claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md","title":"Kv-cache-manager Skill","category":"anthropic-skill","format":"markdown","content":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/mamba-hybrid-recurrent-state-d2h-syncs.md)","tokens":300,"sizeBytes":343},{"name":"max-utilization-reuse-token-budget.md","path":".claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md","title":"Kv-cache-manager Skill","category":"anthropic-skill","format":"markdown","content":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/max-utilization-reuse-token-budget.md)","tokens":300,"sizeBytes":335},{"name":"qwen35-ssm-cache-fp32-fallback.md","path":".claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md","title":"Kv-cache-manager Skill","category":"anthropic-skill","format":"markdown","content":"# Kv-cache-manager Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-manager/qwen35-ssm-cache-fp32-fallback.md)","tokens":300,"sizeBytes":327},{"name":"disagg-gen-retokenize-and-memcpy.md","path":".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md","title":"Kv-cache-transceiver Skill","category":"anthropic-skill","format":"markdown","content":"# Kv-cache-transceiver Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-gen-retokenize-and-memcpy.md)","tokens":300,"sizeBytes":343},{"name":"disagg-quorum-vote-world-scope.md","path":".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md","title":"Kv-cache-transceiver Skill","category":"anthropic-skill","format":"markdown","content":"# Kv-cache-transceiver Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/disagg-quorum-vote-world-scope.md)","tokens":300,"sizeBytes":339},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md","title":"Kv-cache-transceiver Skill","category":"anthropic-skill","format":"markdown","content":"# Kv-cache-transceiver Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/kv-cache-transceiver/index.md)","tokens":300,"sizeBytes":289},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/index.md)","tokens":300,"sizeBytes":289},{"name":"ipc-hmac-key-via-fd-breaks-bench.md","path":".claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/ipc-hmac-key-via-fd-breaks-bench.md)","tokens":300,"sizeBytes":343},{"name":"multi-gpu-bench-warmup-too-few.md","path":".claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/multi-gpu-bench-warmup-too-few.md)","tokens":300,"sizeBytes":339},{"name":"nixl-ctx-only-gap-not-reproduced-home-mount.md","path":".claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/nixl-ctx-only-gap-not-reproduced-home-mount.md)","tokens":300,"sizeBytes":365},{"name":"perf-sanity-sampler-options-applied-to-all-cases.md","path":".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-sanity-sampler-options-applied-to-all-cases.md)","tokens":300,"sizeBytes":375},{"name":"perf-test-gpt-oss-20b-moe-backend-pin.md","path":".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-gpt-oss-20b-moe-backend-pin.md)","tokens":300,"sizeBytes":353},{"name":"perf-test-wrong-moe-backend.md","path":".claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md","title":"Measurement-and-test Skill","category":"anthropic-skill","format":"markdown","content":"# Measurement-and-test Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/measurement-and-test/perf-test-wrong-moe-backend.md)","tokens":300,"sizeBytes":333},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/model-definition/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/index.md","title":"Model-definition Skill","category":"anthropic-skill","format":"markdown","content":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/index.md)","tokens":300,"sizeBytes":277},{"name":"ltx2-bf16-lora-restore.md","path":".claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md","title":"Model-definition Skill","category":"anthropic-skill","format":"markdown","content":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/ltx2-bf16-lora-restore.md)","tokens":300,"sizeBytes":311},{"name":"mamba2-flashinfer-head-group-ratio-gate.md","path":".claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md","title":"Model-definition Skill","category":"anthropic-skill","format":"markdown","content":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/mamba2-flashinfer-head-group-ratio-gate.md)","tokens":300,"sizeBytes":345},{"name":"vlm-vision-tower-loaded-for-text-only-bench.md","path":".claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md","title":"Model-definition Skill","category":"anthropic-skill","format":"markdown","content":"# Model-definition Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/model-definition/vlm-vision-tower-loaded-for-text-only-bench.md)","tokens":300,"sizeBytes":353},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/moe/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/index.md","title":"Moe Skill","category":"anthropic-skill","format":"markdown","content":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/index.md)","tokens":300,"sizeBytes":238},{"name":"moe-a2a-rank-mask-hot-path.md","path":".claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md","title":"Moe Skill","category":"anthropic-skill","format":"markdown","content":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/moe-a2a-rank-mask-hot-path.md)","tokens":300,"sizeBytes":280},{"name":"moe-doactivation-optimization-side-effect.md","path":".claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md","title":"Moe Skill","category":"anthropic-skill","format":"markdown","content":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/moe-doactivation-optimization-side-effect.md)","tokens":300,"sizeBytes":310},{"name":"router-gemm-missing-instantiation.md","path":".claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md","title":"Moe Skill","category":"anthropic-skill","format":"markdown","content":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/router-gemm-missing-instantiation.md)","tokens":300,"sizeBytes":294},{"name":"triton-kernels-36-uplift-and-routing-port.md","path":".claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md","title":"Moe Skill","category":"anthropic-skill","format":"markdown","content":"# Moe Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/moe/triton-kernels-36-uplift-and-routing-port.md)","tokens":300,"sizeBytes":310},{"name":"buffer-pool-zero-fill-fillfunctor.md","path":".claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md","title":"Runtime-and-serving Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/buffer-pool-zero-fill-fillfunctor.md)","tokens":300,"sizeBytes":342},{"name":"custom-op-dispatch-tax-ltx2.md","path":".claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md","title":"Runtime-and-serving Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/custom-op-dispatch-tax-ltx2.md)","tokens":300,"sizeBytes":330},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md","title":"Runtime-and-serving Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/index.md)","tokens":300,"sizeBytes":286},{"name":"nvle-only-treated-as-cc.md","path":".claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md","title":"Runtime-and-serving Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/nvle-only-treated-as-cc.md)","tokens":300,"sizeBytes":322},{"name":"png-encode-on-serving-hot-path.md","path":".claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md","title":"Runtime-and-serving Skill","category":"anthropic-skill","format":"markdown","content":"# Runtime-and-serving Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/runtime-and-serving/png-encode-on-serving-hot-path.md)","tokens":300,"sizeBytes":336},{"name":"beam-search-host-handoff.md","path":".claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md","title":"Sampler Skill","category":"anthropic-skill","format":"markdown","content":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/beam-search-host-handoff.md)","tokens":300,"sizeBytes":288},{"name":"cc-blocking-d2h-copies-worker-thread.md","path":".claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md","title":"Sampler Skill","category":"anthropic-skill","format":"markdown","content":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/cc-blocking-d2h-copies-worker-thread.md)","tokens":300,"sizeBytes":312},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/sampler/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/index.md","title":"Sampler Skill","category":"anthropic-skill","format":"markdown","content":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/index.md)","tokens":300,"sizeBytes":250},{"name":"sampler-stop-words-implicit-d2h-sync.md","path":".claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md","title":"Sampler Skill","category":"anthropic-skill","format":"markdown","content":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/sampler-stop-words-implicit-d2h-sync.md)","tokens":300,"sizeBytes":312},{"name":"torchsampler-logprobs0-overhead.md","path":".claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md","title":"Sampler Skill","category":"anthropic-skill","format":"markdown","content":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-logprobs0-overhead.md)","tokens":300,"sizeBytes":302},{"name":"torchsampler-optimizations-lost-in-branch.md","path":".claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md","title":"Sampler Skill","category":"anthropic-skill","format":"markdown","content":"# Sampler Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/sampler/torchsampler-optimizations-lost-in-branch.md)","tokens":300,"sizeBytes":322},{"name":"canceled-req-ids-unbounded-growth.md","path":".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md","title":"Scheduler-and-executor Skill","category":"anthropic-skill","format":"markdown","content":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/canceled-req-ids-unbounded-growth.md)","tokens":300,"sizeBytes":351},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md","title":"Scheduler-and-executor Skill","category":"anthropic-skill","format":"markdown","content":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/index.md)","tokens":300,"sizeBytes":295},{"name":"model-engine-per-request-attr-overhead.md","path":".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md","title":"Scheduler-and-executor Skill","category":"anthropic-skill","format":"markdown","content":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/model-engine-per-request-attr-overhead.md)","tokens":300,"sizeBytes":361},{"name":"overlap-scheduler-first-token-delay.md","path":".claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md","title":"Scheduler-and-executor Skill","category":"anthropic-skill","format":"markdown","content":"# Scheduler-and-executor Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/scheduler-and-executor/overlap-scheduler-first-token-delay.md)","tokens":300,"sizeBytes":355},{"name":"cutedsl-argmax-revert.md","path":".claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md","title":"Spec-decode Skill","category":"anthropic-skill","format":"markdown","content":"# Spec-decode Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/cutedsl-argmax-revert.md)","tokens":300,"sizeBytes":294},{"name":"dsa-mtp-stale-token-to-request-map.md","path":".claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md","title":"Spec-decode Skill","category":"anthropic-skill","format":"markdown","content":"# Spec-decode Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/dsa-mtp-stale-token-to-request-map.md)","tokens":300,"sizeBytes":320},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/spec-decode/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/index.md","title":"Spec-decode Skill","category":"anthropic-skill","format":"markdown","content":"# Spec-decode Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/spec-decode/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/spec-decode/index.md)","tokens":300,"sizeBytes":262},{"name":"index.md","path":".claude/skills/perf-regression-cookbook/references/tokenization/index.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/tokenization/index.md","title":"Tokenization Skill","category":"anthropic-skill","format":"markdown","content":"# Tokenization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/tokenization/index.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/tokenization/index.md)","tokens":300,"sizeBytes":265},{"name":"kimi-tokenizer-slow-path-on-hot-path.md","path":".claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md","title":"Tokenization Skill","category":"anthropic-skill","format":"markdown","content":"# Tokenization Skill\n\nPath: `.claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/references/tokenization/kimi-tokenizer-slow-path-on-hot-path.md)","tokens":300,"sizeBytes":327},{"name":"SKILL.md","path":".claude/skills/perf-regression-cookbook/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/SKILL.md","title":"perf-regression-cookbook","category":"anthropic-skill","format":"markdown","content":"# perf-regression-cookbook\n\nPath: `.claude/skills/perf-regression-cookbook/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-regression-cookbook/SKILL.md)","tokens":300,"sizeBytes":223},{"name":"api-pytorch.md","path":".claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/api-pytorch.md)","tokens":300,"sizeBytes":245},{"name":"api-te-megatron.md","path":".claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/api-te-megatron.md)","tokens":300,"sizeBytes":253},{"name":"patterns-compatibility.md","path":".claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/patterns-compatibility.md)","tokens":300,"sizeBytes":267},{"name":"patterns-dynamic.md","path":".claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/patterns-dynamic.md)","tokens":300,"sizeBytes":255},{"name":"troubleshooting.md","path":".claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/references/troubleshooting.md)","tokens":300,"sizeBytes":253},{"name":"SKILL.md","path":".claude/skills/perf-torch-cuda-graphs/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/SKILL.md","title":"perf-torch-cuda-graphs","category":"anthropic-skill","format":"markdown","content":"# perf-torch-cuda-graphs\n\nPath: `.claude/skills/perf-torch-cuda-graphs/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/SKILL.md)","tokens":300,"sizeBytes":217},{"name":"sync-patterns.md","path":".claude/skills/perf-torch-sync-free/references/sync-patterns.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/references/sync-patterns.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-torch-sync-free/references/sync-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/references/sync-patterns.md)","tokens":300,"sizeBytes":245},{"name":"SKILL.md","path":".claude/skills/perf-torch-sync-free/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/SKILL.md","title":"perf-torch-sync-free","category":"anthropic-skill","format":"markdown","content":"# perf-torch-sync-free\n\nPath: `.claude/skills/perf-torch-sync-free/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/SKILL.md)","tokens":300,"sizeBytes":211},{"name":"benchmarking-patterns.md","path":".claude/skills/perf-workload-profiling/references/benchmarking-patterns.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/benchmarking-patterns.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-workload-profiling/references/benchmarking-patterns.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/benchmarking-patterns.md)","tokens":300,"sizeBytes":267},{"name":"nvtx-api.md","path":".claude/skills/perf-workload-profiling/references/nvtx-api.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/nvtx-api.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-workload-profiling/references/nvtx-api.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/nvtx-api.md)","tokens":300,"sizeBytes":241},{"name":"pytorch-profiler-api.md","path":".claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/references/pytorch-profiler-api.md)","tokens":300,"sizeBytes":265},{"name":"SKILL.md","path":".claude/skills/perf-workload-profiling/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/SKILL.md","title":"perf-workload-profiling","category":"anthropic-skill","format":"markdown","content":"# perf-workload-profiling\n\nPath: `.claude/skills/perf-workload-profiling/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-workload-profiling/SKILL.md)","tokens":300,"sizeBytes":220},{"name":"SKILL.md","path":".claude/skills/trtllm-case-executor/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-case-executor/SKILL.md","title":"trtllm-case-executor","category":"anthropic-skill","format":"markdown","content":"# trtllm-case-executor\n\nPath: `.claude/skills/trtllm-case-executor/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-case-executor/SKILL.md)","tokens":300,"sizeBytes":211},{"name":"SKILL.md","path":".claude/skills/trtllm-code-contribution/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-code-contribution/SKILL.md","title":"trtllm-code-contribution","category":"anthropic-skill","format":"markdown","content":"# trtllm-code-contribution\n\nPath: `.claude/skills/trtllm-code-contribution/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-code-contribution/SKILL.md)","tokens":300,"sizeBytes":223},{"name":"SKILL.md","path":".claude/skills/trtllm-codebase-exploration/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-codebase-exploration/SKILL.md","title":"trtllm-codebase-exploration","category":"anthropic-skill","format":"markdown","content":"# trtllm-codebase-exploration\n\nPath: `.claude/skills/trtllm-codebase-exploration/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-codebase-exploration/SKILL.md)","tokens":300,"sizeBytes":232},{"name":"SKILL.md","path":".claude/skills/trtllm-flashinfer-upgrade/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-flashinfer-upgrade/SKILL.md","title":"trtllm-flashinfer-upgrade","category":"anthropic-skill","format":"markdown","content":"# trtllm-flashinfer-upgrade\n\nPath: `.claude/skills/trtllm-flashinfer-upgrade/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-flashinfer-upgrade/SKILL.md)","tokens":300,"sizeBytes":226},{"name":"SKILL.md","path":".claude/skills/trtllm-model-onboard-multimodal/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-model-onboard-multimodal/SKILL.md","title":"trtllm-model-onboard-multimodal","category":"anthropic-skill","format":"markdown","content":"# trtllm-model-onboard-multimodal\n\nPath: `.claude/skills/trtllm-model-onboard-multimodal/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-model-onboard-multimodal/SKILL.md)","tokens":300,"sizeBytes":244},{"name":"moe-canonical-code-examples.md","path":".claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-moe-develop/references/moe-canonical-code-examples.md)","tokens":300,"sizeBytes":269},{"name":"SKILL.md","path":".claude/skills/trtllm-moe-develop/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-moe-develop/SKILL.md","title":"trtllm-moe-develop","category":"anthropic-skill","format":"markdown","content":"# trtllm-moe-develop\n\nPath: `.claude/skills/trtllm-moe-develop/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-moe-develop/SKILL.md)","tokens":300,"sizeBytes":205},{"name":"knob-heuristics.md","path":".claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-serve-config-guide/references/knob-heuristics.md)","tokens":300,"sizeBytes":259},{"name":"SKILL.md","path":".claude/skills/trtllm-serve-config-guide/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-serve-config-guide/SKILL.md","title":"trtllm-serve-config-guide","category":"anthropic-skill","format":"markdown","content":"# trtllm-serve-config-guide\n\nPath: `.claude/skills/trtllm-serve-config-guide/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-serve-config-guide/SKILL.md)","tokens":300,"sizeBytes":226},{"name":"trtllm_test_template.md","path":".claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/references/trtllm_test_template.md)","tokens":300,"sizeBytes":271},{"name":"SKILL.md","path":".claude/skills/trtllm-test-script-builder/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/SKILL.md","title":"trtllm-test-script-builder","category":"anthropic-skill","format":"markdown","content":"# trtllm-test-script-builder\n\nPath: `.claude/skills/trtllm-test-script-builder/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/SKILL.md)","tokens":300,"sizeBytes":229},{"name":"agg_config_template.yaml","path":".claude/skills/trtllm-test-specialist/references/agg_config_template.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/agg_config_template.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/agg_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/agg_config_template.yaml)","tokens":300,"sizeBytes":265},{"name":"benchmark_config_template.yaml","path":".claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/benchmark_config_template.yaml)","tokens":300,"sizeBytes":277},{"name":"disagg_config_template.yaml","path":".claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/disagg_config_template.yaml)","tokens":300,"sizeBytes":271},{"name":"test_config_template.yaml","path":".claude/skills/trtllm-test-specialist/references/test_config_template.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/test_config_template.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/test_config_template.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/test_config_template.yaml)","tokens":300,"sizeBytes":267},{"name":"trtllm_test_fix_recommendations.md","path":".claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# References Skill\n\nPath: `.claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/trtllm_test_fix_recommendations.md)","tokens":300,"sizeBytes":285},{"name":"SKILL.md","path":".claude/skills/trtllm-test-specialist/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/SKILL.md","title":"trtllm-test-specialist","category":"anthropic-skill","format":"markdown","content":"# trtllm-test-specialist\n\nPath: `.claude/skills/trtllm-test-specialist/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/SKILL.md)","tokens":300,"sizeBytes":217},{"name":"cosmos3-i2v.json","path":".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json","title":"Prompts Skill","category":"anthropic-skill","format":"json","content":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-i2v.json)","tokens":300,"sizeBytes":240},{"name":"cosmos3-t2av.json","path":".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json","title":"Prompts Skill","category":"anthropic-skill","format":"json","content":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2av.json)","tokens":300,"sizeBytes":242},{"name":"cosmos3-t2i.json","path":".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json","title":"Prompts Skill","category":"anthropic-skill","format":"json","content":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2i.json)","tokens":300,"sizeBytes":240},{"name":"cosmos3-t2v.json","path":".claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json","title":"Prompts Skill","category":"anthropic-skill","format":"json","content":"# Prompts Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/prompts/cosmos3-t2v.json)","tokens":300,"sizeBytes":240},{"name":"wan22-ti2v-5b-t2v.json","path":".claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json","title":"Reference Skill","category":"anthropic-skill","format":"json","content":"# Reference Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/reference/wan22-ti2v-5b-t2v.json)","tokens":300,"sizeBytes":258},{"name":"README.md","path":".claude/skills/trtllm-visualgen-bench/serve/README.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/serve/README.md","title":"Serve Skill","category":"anthropic-skill","format":"markdown","content":"# Serve Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/serve/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/serve/README.md)","tokens":300,"sizeBytes":220},{"name":"SKILL.md","path":".claude/skills/trtllm-visualgen-bench/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/SKILL.md","title":"trtllm-visualgen-bench","category":"anthropic-skill","format":"markdown","content":"# trtllm-visualgen-bench\n\nPath: `.claude/skills/trtllm-visualgen-bench/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/SKILL.md)","tokens":300,"sizeBytes":217},{"name":"cosmos3-nano-t2i.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-nano-t2i.yaml)","tokens":300,"sizeBytes":256},{"name":"cosmos3-super-i2v.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-i2v.yaml)","tokens":300,"sizeBytes":258},{"name":"cosmos3-super-t2av.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2av.yaml)","tokens":300,"sizeBytes":260},{"name":"cosmos3-super-t2v.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/cosmos3-super-t2v.yaml)","tokens":300,"sizeBytes":258},{"name":"flux2-dev-edit.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-edit.yaml)","tokens":300,"sizeBytes":252},{"name":"flux2-dev-t2i.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/flux2-dev-t2i.yaml)","tokens":300,"sizeBytes":250},{"name":"ltx2-i2v.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/ltx2-i2v.yaml)","tokens":300,"sizeBytes":240},{"name":"ltx2-t2v.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/ltx2-t2v.yaml)","tokens":300,"sizeBytes":240},{"name":"qwen-image-edit-2511.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-edit-2511.yaml)","tokens":300,"sizeBytes":264},{"name":"qwen-image-layered.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image-layered.yaml)","tokens":300,"sizeBytes":260},{"name":"qwen-image.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/qwen-image.yaml)","tokens":300,"sizeBytes":244},{"name":"README.md","path":".claude/skills/trtllm-visualgen-bench/workloads/README.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/README.md","title":"Workloads Skill","category":"anthropic-skill","format":"markdown","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/README.md)","tokens":300,"sizeBytes":232},{"name":"wan22-i2v-a14b.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-i2v-a14b.yaml)","tokens":300,"sizeBytes":252},{"name":"wan22-t2v-a14b.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-t2v-a14b.yaml)","tokens":300,"sizeBytes":252},{"name":"wan22-ti2v-5b-i2v.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-i2v.yaml)","tokens":300,"sizeBytes":258},{"name":"wan22-ti2v-5b-t2v.yaml","path":".claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml","title":"Workloads Skill","category":"anthropic-skill","format":"yaml","content":"# Workloads Skill\n\nPath: `.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/workloads/wan22-ti2v-5b-t2v.yaml)","tokens":300,"sizeBytes":258},{"name":"SKILL.md","path":".claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md","title":"trtllm-visualgen-lpips-golden-refresh","category":"anthropic-skill","format":"markdown","content":"# trtllm-visualgen-lpips-golden-refresh\n\nPath: `.claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-lpips-golden-refresh/SKILL.md)","tokens":300,"sizeBytes":262},{"name":"SKILL.md","path":".codex/skills/visual-gen-component-test/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/skills/visual-gen-component-test/SKILL.md","title":"visual-gen-component-test","category":"anthropic-skill","format":"markdown","content":"# visual-gen-component-test\n\nPath: `.codex/skills/visual-gen-component-test/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/skills/visual-gen-component-test/SKILL.md)","tokens":300,"sizeBytes":224},{"name":"SKILL.md","path":"agent-flow/.claude/skills/perf-analyze/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/.claude/skills/perf-analyze/SKILL.md","title":"perf-analyze","category":"anthropic-skill","format":"markdown","content":"# perf-analyze\n\nPath: `agent-flow/.claude/skills/perf-analyze/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/.claude/skills/perf-analyze/SKILL.md)","tokens":300,"sizeBytes":209},{"name":"SKILL.md","path":"agent-flow/.claude/skills/perf-optimize/SKILL.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/.claude/skills/perf-optimize/SKILL.md","title":"perf-optimize","category":"anthropic-skill","format":"markdown","content":"# perf-optimize\n\nPath: `agent-flow/.claude/skills/perf-optimize/SKILL.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/.claude/skills/perf-optimize/SKILL.md)","tokens":300,"sizeBytes":212},{"name":"README.md","path":"jenkins/scripts/cbts/cbts/rules/README.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/jenkins/scripts/cbts/cbts/rules/README.md","title":"Rules Skill","category":"anthropic-skill","format":"markdown","content":"# Rules Skill\n\nPath: `jenkins/scripts/cbts/cbts/rules/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/jenkins/scripts/cbts/cbts/rules/README.md)","tokens":300,"sizeBytes":196},{"name":"trtllm_engine_model.json","path":"tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json","title":"Nsys_profile_tools Skill","category":"anthropic-skill","format":"json","content":"# Nsys_profile_tools Skill\n\nPath: `tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/nsys_profile_tools/trtllm_engine_model.json)","tokens":300,"sizeBytes":269},{"name":"acceptance_length.yaml","path":"tests/integration/defs/accuracy/references/acceptance_length.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/acceptance_length.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/acceptance_length.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/acceptance_length.yaml)","tokens":300,"sizeBytes":249},{"name":"cnn_dailymail.yaml","path":"tests/integration/defs/accuracy/references/cnn_dailymail.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/cnn_dailymail.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/cnn_dailymail.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/cnn_dailymail.yaml)","tokens":300,"sizeBytes":241},{"name":"gpqa_diamond.yaml","path":"tests/integration/defs/accuracy/references/gpqa_diamond.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gpqa_diamond.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/gpqa_diamond.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gpqa_diamond.yaml)","tokens":300,"sizeBytes":239},{"name":"gsm8k_inferencex.yaml","path":"tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gsm8k_inferencex.yaml)","tokens":300,"sizeBytes":247},{"name":"gsm8k.yaml","path":"tests/integration/defs/accuracy/references/gsm8k.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gsm8k.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/gsm8k.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/gsm8k.yaml)","tokens":300,"sizeBytes":225},{"name":"humaneval.yaml","path":"tests/integration/defs/accuracy/references/humaneval.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/humaneval.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/humaneval.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/humaneval.yaml)","tokens":300,"sizeBytes":233},{"name":"json_mode_eval.yaml","path":"tests/integration/defs/accuracy/references/json_mode_eval.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/json_mode_eval.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/json_mode_eval.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/json_mode_eval.yaml)","tokens":300,"sizeBytes":243},{"name":"longbench_v1.yaml","path":"tests/integration/defs/accuracy/references/longbench_v1.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/longbench_v1.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/longbench_v1.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/longbench_v1.yaml)","tokens":300,"sizeBytes":239},{"name":"longbench_v2.yaml","path":"tests/integration/defs/accuracy/references/longbench_v2.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/longbench_v2.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/longbench_v2.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/longbench_v2.yaml)","tokens":300,"sizeBytes":239},{"name":"mmlu.yaml","path":"tests/integration/defs/accuracy/references/mmlu.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/mmlu.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/mmlu.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/mmlu.yaml)","tokens":300,"sizeBytes":223},{"name":"mmmu.yaml","path":"tests/integration/defs/accuracy/references/mmmu.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/mmmu.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/mmmu.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/mmmu.yaml)","tokens":300,"sizeBytes":223},{"name":"passkey_retrieval_64k.yaml","path":"tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/passkey_retrieval_64k.yaml)","tokens":300,"sizeBytes":257},{"name":"videomme.yaml","path":"tests/integration/defs/accuracy/references/videomme.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/videomme.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/videomme.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/videomme.yaml)","tokens":300,"sizeBytes":231},{"name":"voxpopuli.yaml","path":"tests/integration/defs/accuracy/references/voxpopuli.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/voxpopuli.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/voxpopuli.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/voxpopuli.yaml)","tokens":300,"sizeBytes":233},{"name":"zero_scrolls.yaml","path":"tests/integration/defs/accuracy/references/zero_scrolls.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/zero_scrolls.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/integration/defs/accuracy/references/zero_scrolls.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/integration/defs/accuracy/references/zero_scrolls.yaml)","tokens":300,"sizeBytes":239},{"name":"batched_logits_processor.yaml","path":"tests/unittest/api_stability/references/batched_logits_processor.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/batched_logits_processor.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/batched_logits_processor.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/batched_logits_processor.yaml)","tokens":300,"sizeBytes":257},{"name":"calib_config.yaml","path":"tests/unittest/api_stability/references/calib_config.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/calib_config.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/calib_config.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/calib_config.yaml)","tokens":300,"sizeBytes":233},{"name":"completion_output.yaml","path":"tests/unittest/api_stability/references/completion_output.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/completion_output.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/completion_output.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/completion_output.yaml)","tokens":300,"sizeBytes":243},{"name":"conversation_params.yaml","path":"tests/unittest/api_stability/references/conversation_params.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/conversation_params.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/conversation_params.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/conversation_params.yaml)","tokens":300,"sizeBytes":247},{"name":"decoding_base_config.yaml","path":"tests/unittest/api_stability/references/decoding_base_config.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/decoding_base_config.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/decoding_base_config.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/decoding_base_config.yaml)","tokens":300,"sizeBytes":249},{"name":"guided_decoding_params.yaml","path":"tests/unittest/api_stability/references/guided_decoding_params.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/guided_decoding_params.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/guided_decoding_params.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/guided_decoding_params.yaml)","tokens":300,"sizeBytes":253},{"name":"llm.yaml","path":"tests/unittest/api_stability/references/llm.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/llm.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/llm.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/llm.yaml)","tokens":300,"sizeBytes":215},{"name":"logits_processor.yaml","path":"tests/unittest/api_stability/references/logits_processor.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/logits_processor.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/logits_processor.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/logits_processor.yaml)","tokens":300,"sizeBytes":241},{"name":"quant_config.yaml","path":"tests/unittest/api_stability/references/quant_config.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/quant_config.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/quant_config.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/quant_config.yaml)","tokens":300,"sizeBytes":233},{"name":"request_output.yaml","path":"tests/unittest/api_stability/references/request_output.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/request_output.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/request_output.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/request_output.yaml)","tokens":300,"sizeBytes":237},{"name":"sampling_params.yaml","path":"tests/unittest/api_stability/references/sampling_params.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/sampling_params.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/sampling_params.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/sampling_params.yaml)","tokens":300,"sizeBytes":239},{"name":"trtllm_serve_api.yaml","path":"tests/unittest/api_stability/references/trtllm_serve_api.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/trtllm_serve_api.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/trtllm_serve_api.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/trtllm_serve_api.yaml)","tokens":300,"sizeBytes":241},{"name":"trtllm_serve_cli.yaml","path":"tests/unittest/api_stability/references/trtllm_serve_cli.yaml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/trtllm_serve_cli.yaml","title":"References Skill","category":"anthropic-skill","format":"yaml","content":"# References Skill\n\nPath: `tests/unittest/api_stability/references/trtllm_serve_cli.yaml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/unittest/api_stability/references/trtllm_serve_cli.yaml)","tokens":300,"sizeBytes":241},{"name":"mini_cnn_eval_spec_decoding.json","path":"triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json","title":"Dataset Skill","category":"anthropic-skill","format":"json","content":"# Dataset Skill\n\nPath: `triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/triton_backend/tools/dataset/mini_cnn_eval_spec_decoding.json)","tokens":300,"sizeBytes":238},{"name":"mini_cnn_eval.json","path":"triton_backend/tools/dataset/mini_cnn_eval.json","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/triton_backend/tools/dataset/mini_cnn_eval.json","title":"Dataset Skill","category":"anthropic-skill","format":"json","content":"# Dataset Skill\n\nPath: `triton_backend/tools/dataset/mini_cnn_eval.json`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/triton_backend/tools/dataset/mini_cnn_eval.json)","tokens":300,"sizeBytes":210},{"name":"build.sh","path":".claude/skills/exec-remote-slurm/scripts/build.sh","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-remote-slurm/scripts/build.sh","title":"Claude Directive: build.sh","category":"claude-rule","format":"markdown","content":"# Claude Directive: build.sh\n\nPath: `.claude/skills/exec-remote-slurm/scripts/build.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-remote-slurm/scripts/build.sh)","tokens":300,"sizeBytes":227},{"name":"build.slurm","path":".claude/skills/exec-remote-slurm/scripts/build.slurm","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-remote-slurm/scripts/build.slurm","title":"Claude Directive: build.slurm","category":"claude-rule","format":"markdown","content":"# Claude Directive: build.slurm\n\nPath: `.claude/skills/exec-remote-slurm/scripts/build.slurm`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-remote-slurm/scripts/build.slurm)","tokens":300,"sizeBytes":236},{"name":"compile.sh","path":".claude/skills/exec-slurm-compile/scripts/compile.sh","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/compile.sh","title":"Claude Directive: compile.sh","category":"claude-rule","format":"markdown","content":"# Claude Directive: compile.sh\n\nPath: `.claude/skills/exec-slurm-compile/scripts/compile.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/compile.sh)","tokens":300,"sizeBytes":235},{"name":"compile.slurm","path":".claude/skills/exec-slurm-compile/scripts/compile.slurm","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/compile.slurm","title":"Claude Directive: compile.slurm","category":"claude-rule","format":"markdown","content":"# Claude Directive: compile.slurm\n\nPath: `.claude/skills/exec-slurm-compile/scripts/compile.slurm`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/compile.slurm)","tokens":300,"sizeBytes":244},{"name":"enroot-import","path":".claude/skills/exec-slurm-compile/scripts/enroot-import","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/enroot-import","title":"Claude Directive: enroot-import","category":"claude-rule","format":"markdown","content":"# Claude Directive: enroot-import\n\nPath: `.claude/skills/exec-slurm-compile/scripts/enroot-import`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/enroot-import)","tokens":300,"sizeBytes":244},{"name":"submit_compile.sh","path":".claude/skills/exec-slurm-compile/scripts/submit_compile.sh","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/submit_compile.sh","title":"Claude Directive: submit_compile.sh","category":"claude-rule","format":"markdown","content":"# Claude Directive: submit_compile.sh\n\nPath: `.claude/skills/exec-slurm-compile/scripts/submit_compile.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/exec-slurm-compile/scripts/submit_compile.sh)","tokens":300,"sizeBytes":256},{"name":"__init__.py","path":".claude/skills/kernel-cute-writing/scripts/__init__.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/__init__.py","title":"Claude Directive: __init__.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: __init__.py\n\nPath: `.claude/skills/kernel-cute-writing/scripts/__init__.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/__init__.py)","tokens":300,"sizeBytes":240},{"name":"benchmark_kernel.py","path":".claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py","title":"Claude Directive: benchmark_kernel.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: benchmark_kernel.py\n\nPath: `.claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/benchmark_kernel.py)","tokens":300,"sizeBytes":264},{"name":"verify_kernel.py","path":".claude/skills/kernel-cute-writing/scripts/verify_kernel.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/verify_kernel.py","title":"Claude Directive: verify_kernel.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: verify_kernel.py\n\nPath: `.claude/skills/kernel-cute-writing/scripts/verify_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/scripts/verify_kernel.py)","tokens":300,"sizeBytes":255},{"name":"classify_kernel.py","path":".claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py","title":"Claude Directive: classify_kernel.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: classify_kernel.py\n\nPath: `.claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/scripts/classify_kernel.py)","tokens":300,"sizeBytes":275},{"name":"tileir_check.py","path":".claude/skills/kernel-tileir-optimization/scripts/tileir_check.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/scripts/tileir_check.py","title":"Claude Directive: tileir_check.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: tileir_check.py\n\nPath: `.claude/skills/kernel-tileir-optimization/scripts/tileir_check.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/scripts/tileir_check.py)","tokens":300,"sizeBytes":266},{"name":"__init__.py","path":".claude/skills/kernel-triton-writing/scripts/__init__.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/__init__.py","title":"Claude Directive: __init__.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: __init__.py\n\nPath: `.claude/skills/kernel-triton-writing/scripts/__init__.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/__init__.py)","tokens":300,"sizeBytes":244},{"name":"benchmark_kernel.py","path":".claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py","title":"Claude Directive: benchmark_kernel.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: benchmark_kernel.py\n\nPath: `.claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/benchmark_kernel.py)","tokens":300,"sizeBytes":268},{"name":"verify_kernel.py","path":".claude/skills/kernel-triton-writing/scripts/verify_kernel.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/verify_kernel.py","title":"Claude Directive: verify_kernel.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: verify_kernel.py\n\nPath: `.claude/skills/kernel-triton-writing/scripts/verify_kernel.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-triton-writing/scripts/verify_kernel.py)","tokens":300,"sizeBytes":259},{"name":"analyze_host_overhead.py","path":".claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py","title":"Claude Directive: analyze_host_overhead.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: analyze_host_overhead.py\n\nPath: `.claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/scripts/analyze_host_overhead.py)","tokens":300,"sizeBytes":277},{"name":"detect_host_overhead.py","path":".claude/skills/perf-host-analysis/scripts/detect_host_overhead.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/scripts/detect_host_overhead.py","title":"Claude Directive: detect_host_overhead.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: detect_host_overhead.py\n\nPath: `.claude/skills/perf-host-analysis/scripts/detect_host_overhead.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-host-analysis/scripts/detect_host_overhead.py)","tokens":300,"sizeBytes":274},{"name":"verify_workload.py","path":".claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py","title":"Claude Directive: verify_workload.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: verify_workload.py\n\nPath: `.claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-cuda-graphs/scripts/verify_workload.py)","tokens":300,"sizeBytes":267},{"name":"verify_workload.py","path":".claude/skills/perf-torch-sync-free/scripts/verify_workload.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/scripts/verify_workload.py","title":"Claude Directive: verify_workload.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: verify_workload.py\n\nPath: `.claude/skills/perf-torch-sync-free/scripts/verify_workload.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/perf-torch-sync-free/scripts/verify_workload.py)","tokens":300,"sizeBytes":263},{"name":"detect_slurm_env.sh","path":".claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh","title":"Claude Directive: detect_slurm_env.sh","category":"claude-rule","format":"markdown","content":"# Claude Directive: detect_slurm_env.sh\n\nPath: `.claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-case-executor/scripts/detect_slurm_env.sh)","tokens":300,"sizeBytes":266},{"name":"build_slurm_script.py","path":".claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py","title":"Claude Directive: build_slurm_script.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: build_slurm_script.py\n\nPath: `.claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/scripts/build_slurm_script.py)","tokens":300,"sizeBytes":284},{"name":"slurm_run_custom.sh","path":".claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh","title":"Claude Directive: slurm_run_custom.sh","category":"claude-rule","format":"markdown","content":"# Claude Directive: slurm_run_custom.sh\n\nPath: `.claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-script-builder/scripts/slurm_run_custom.sh)","tokens":300,"sizeBytes":278},{"name":"smoke_test_config_template.yml","path":".claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml","title":"Claude Directive: smoke_test_config_template.yml","category":"claude-rule","format":"markdown","content":"# Claude Directive: smoke_test_config_template.yml\n\nPath: `.claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/references/smoke_test_config_template.yml)","tokens":300,"sizeBytes":309},{"name":"build_test_command.py","path":".claude/skills/trtllm-test-specialist/scripts/build_test_command.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/build_test_command.py","title":"Claude Directive: build_test_command.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: build_test_command.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/build_test_command.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/build_test_command.py)","tokens":300,"sizeBytes":276},{"name":"extract_test_markers.py","path":".claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py","title":"Claude Directive: extract_test_markers.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: extract_test_markers.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/extract_test_markers.py)","tokens":300,"sizeBytes":282},{"name":"generate_benchmark_config.py","path":".claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py","title":"Claude Directive: generate_benchmark_config.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: generate_benchmark_config.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/generate_benchmark_config.py)","tokens":300,"sizeBytes":297},{"name":"generate_report.py","path":".claude/skills/trtllm-test-specialist/scripts/generate_report.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/generate_report.py","title":"Claude Directive: generate_report.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: generate_report.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/generate_report.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/generate_report.py)","tokens":300,"sizeBytes":267},{"name":"parse_config.py","path":".claude/skills/trtllm-test-specialist/scripts/parse_config.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/parse_config.py","title":"Claude Directive: parse_config.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: parse_config.py\n\nPath: `.claude/skills/trtllm-test-specialist/scripts/parse_config.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-test-specialist/scripts/parse_config.py)","tokens":300,"sizeBytes":258},{"name":"cat_piano.png","path":".claude/skills/trtllm-visualgen-bench/media/cat_piano.png","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/cat_piano.png","title":"Claude Directive: cat_piano.png","category":"claude-rule","format":"markdown","content":"# Claude Directive: cat_piano.png\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/cat_piano.png`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/cat_piano.png)","tokens":300,"sizeBytes":248},{"name":"image_with_red_background.png","path":".claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png","title":"Claude Directive: image_with_red_background.png","category":"claude-rule","format":"markdown","content":"# Claude Directive: image_with_red_background.png\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/image_with_red_background.png)","tokens":300,"sizeBytes":296},{"name":"landscape.png","path":".claude/skills/trtllm-visualgen-bench/media/landscape.png","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/landscape.png","title":"Claude Directive: landscape.png","category":"claude-rule","format":"markdown","content":"# Claude Directive: landscape.png\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/landscape.png`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/landscape.png)","tokens":300,"sizeBytes":248},{"name":"robot_153.jpg","path":".claude/skills/trtllm-visualgen-bench/media/robot_153.jpg","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/robot_153.jpg","title":"Claude Directive: robot_153.jpg","category":"claude-rule","format":"markdown","content":"# Claude Directive: robot_153.jpg\n\nPath: `.claude/skills/trtllm-visualgen-bench/media/robot_153.jpg`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/media/robot_153.jpg)","tokens":300,"sizeBytes":248},{"name":"serve_check.py","path":".claude/skills/trtllm-visualgen-bench/scripts/serve_check.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/scripts/serve_check.py","title":"Claude Directive: serve_check.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: serve_check.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/scripts/serve_check.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/scripts/serve_check.py)","tokens":300,"sizeBytes":255},{"name":"serve_wait.py","path":".claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py","title":"Claude Directive: serve_wait.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: serve_wait.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/scripts/serve_wait.py)","tokens":300,"sizeBytes":252},{"name":"test_discovery.py","path":".claude/skills/trtllm-visualgen-bench/tests/test_discovery.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_discovery.py","title":"Claude Directive: test_discovery.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: test_discovery.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/tests/test_discovery.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_discovery.py)","tokens":300,"sizeBytes":260},{"name":"test_serve_wait.py","path":".claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py","title":"Claude Directive: test_serve_wait.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: test_serve_wait.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_serve_wait.py)","tokens":300,"sizeBytes":263},{"name":"test_workloads.py","path":".claude/skills/trtllm-visualgen-bench/tests/test_workloads.py","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_workloads.py","title":"Claude Directive: test_workloads.py","category":"claude-rule","format":"markdown","content":"# Claude Directive: test_workloads.py\n\nPath: `.claude/skills/trtllm-visualgen-bench/tests/test_workloads.py`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/trtllm-visualgen-bench/tests/test_workloads.py)","tokens":300,"sizeBytes":260},{"name":"CLAUDE.md","path":"agent-flow/CLAUDE.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/CLAUDE.md","title":"agent-flow - Claude Agent Guidelines & System Prompt","category":"claude-rule","format":"markdown","content":"# agent-flow - Claude Agent Guidelines & System Prompt\n\nPath: `agent-flow/CLAUDE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/CLAUDE.md)","tokens":300,"sizeBytes":195},{"name":"CLAUDE.md","path":"CLAUDE.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/CLAUDE.md","title":"Claude Agent Guidelines & System Prompt","category":"claude-rule","format":"markdown","content":"# Claude Agent Guidelines & System Prompt\n\nPath: `CLAUDE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/CLAUDE.md)","tokens":300,"sizeBytes":160},{"name":"CLAUDE.md","path":"tests/CLAUDE.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/CLAUDE.md","title":"tests - Claude Agent Guidelines & System Prompt","category":"claude-rule","format":"markdown","content":"# tests - Claude Agent Guidelines & System Prompt\n\nPath: `tests/CLAUDE.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/CLAUDE.md)","tokens":300,"sizeBytes":180},{"name":"README.md","path":"tensorrt_llm/tools/profiler/host_profile_tools/README.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/host_profile_tools/README.md","title":"host_profile_tools Documentation","category":"plugin-manifest","format":"markdown","content":"# host_profile_tools Documentation\n\nPath: `tensorrt_llm/tools/profiler/host_profile_tools/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/host_profile_tools/README.md)","tokens":300,"sizeBytes":247},{"name":"README.md","path":"tensorrt_llm/tools/profiler/nsys_profile_tools/README.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/nsys_profile_tools/README.md","title":"nsys_profile_tools Documentation","category":"plugin-manifest","format":"markdown","content":"# nsys_profile_tools Documentation\n\nPath: `tensorrt_llm/tools/profiler/nsys_profile_tools/README.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tensorrt_llm/tools/profiler/nsys_profile_tools/README.md)","tokens":300,"sizeBytes":247},{"name":"exec-compile-specialist.md","path":".claude/agents/exec-compile-specialist.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-compile-specialist.md","title":"Subagent: exec-compile-specialist","category":"subagent-persona","format":"markdown","content":"# Subagent: exec-compile-specialist\n\nPath: `.claude/agents/exec-compile-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-compile-specialist.md)","tokens":300,"sizeBytes":218},{"name":"exec-local-slurm.md","path":".claude/agents/exec-local-slurm.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-local-slurm.md","title":"Subagent: exec-local-slurm","category":"subagent-persona","format":"markdown","content":"# Subagent: exec-local-slurm\n\nPath: `.claude/agents/exec-local-slurm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-local-slurm.md)","tokens":300,"sizeBytes":197},{"name":"exec-remote-slurm.md","path":".claude/agents/exec-remote-slurm.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-remote-slurm.md","title":"Subagent: exec-remote-slurm","category":"subagent-persona","format":"markdown","content":"# Subagent: exec-remote-slurm\n\nPath: `.claude/agents/exec-remote-slurm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/exec-remote-slurm.md)","tokens":300,"sizeBytes":200},{"name":"kernel-cuda-specialist.md","path":".claude/agents/kernel-cuda-specialist.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-cuda-specialist.md","title":"Subagent: kernel-cuda-specialist","category":"subagent-persona","format":"markdown","content":"# Subagent: kernel-cuda-specialist\n\nPath: `.claude/agents/kernel-cuda-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-cuda-specialist.md)","tokens":300,"sizeBytes":215},{"name":"kernel-cute-specialist.md","path":".claude/agents/kernel-cute-specialist.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-cute-specialist.md","title":"Subagent: kernel-cute-specialist","category":"subagent-persona","format":"markdown","content":"# Subagent: kernel-cute-specialist\n\nPath: `.claude/agents/kernel-cute-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-cute-specialist.md)","tokens":300,"sizeBytes":215},{"name":"kernel-tileir-specialist.md","path":".claude/agents/kernel-tileir-specialist.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-tileir-specialist.md","title":"Subagent: kernel-tileir-specialist","category":"subagent-persona","format":"markdown","content":"# Subagent: kernel-tileir-specialist\n\nPath: `.claude/agents/kernel-tileir-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-tileir-specialist.md)","tokens":300,"sizeBytes":221},{"name":"kernel-triton-specialist.md","path":".claude/agents/kernel-triton-specialist.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-triton-specialist.md","title":"Subagent: kernel-triton-specialist","category":"subagent-persona","format":"markdown","content":"# Subagent: kernel-triton-specialist\n\nPath: `.claude/agents/kernel-triton-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/kernel-triton-specialist.md)","tokens":300,"sizeBytes":221},{"name":"modeling-test-specialist.md","path":".claude/agents/modeling-test-specialist.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/modeling-test-specialist.md","title":"Subagent: modeling-test-specialist","category":"subagent-persona","format":"markdown","content":"# Subagent: modeling-test-specialist\n\nPath: `.claude/agents/modeling-test-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/modeling-test-specialist.md)","tokens":300,"sizeBytes":221},{"name":"perf-profiling-specialist.md","path":".claude/agents/perf-profiling-specialist.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-profiling-specialist.md","title":"Subagent: perf-profiling-specialist","category":"subagent-persona","format":"markdown","content":"# Subagent: perf-profiling-specialist\n\nPath: `.claude/agents/perf-profiling-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-profiling-specialist.md)","tokens":300,"sizeBytes":224},{"name":"perf-test-sync.md","path":".claude/agents/perf-test-sync.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-test-sync.md","title":"Subagent: perf-test-sync","category":"subagent-persona","format":"markdown","content":"# Subagent: perf-test-sync\n\nPath: `.claude/agents/perf-test-sync.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-test-sync.md)","tokens":300,"sizeBytes":191},{"name":"perf-torch-cuda-graph-specialist.md","path":".claude/agents/perf-torch-cuda-graph-specialist.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-torch-cuda-graph-specialist.md","title":"Subagent: perf-torch-cuda-graph-specialist","category":"subagent-persona","format":"markdown","content":"# Subagent: perf-torch-cuda-graph-specialist\n\nPath: `.claude/agents/perf-torch-cuda-graph-specialist.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/agents/perf-torch-cuda-graph-specialist.md)","tokens":300,"sizeBytes":245},{"name":"onboard_update_reviewer.toml","path":".codex/agents/onboard_update_reviewer.toml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/agents/onboard_update_reviewer.toml","title":"Codex Directive: onboard_update_reviewer.toml","category":"codex-rule","format":"markdown","content":"# Codex Directive: onboard_update_reviewer.toml\n\nPath: `.codex/agents/onboard_update_reviewer.toml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/agents/onboard_update_reviewer.toml)","tokens":300,"sizeBytes":232},{"name":"config.toml","path":".codex/config.toml","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/config.toml","title":"Codex Directive: config.toml","category":"codex-rule","format":"markdown","content":"# Codex Directive: config.toml\n\nPath: `.codex/config.toml`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/config.toml)","tokens":300,"sizeBytes":167},{"name":"render-endpoints.md","path":"docs/source/commands/trtllm-serve/render-endpoints.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/docs/source/commands/trtllm-serve/render-endpoints.md","title":"Command: /render-endpoints","category":"command-prompt","format":"markdown","content":"# Command: /render-endpoints\n\nPath: `docs/source/commands/trtllm-serve/render-endpoints.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/docs/source/commands/trtllm-serve/render-endpoints.md)","tokens":300,"sizeBytes":235},{"name":"run-benchmark-with-trtllm-serve.md","path":"docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md","rawUrl":"https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md","title":"Command: /run-benchmark-with-trtllm-serve","category":"command-prompt","format":"markdown","content":"# Command: /run-benchmark-with-trtllm-serve\n\nPath: `docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/docs/source/commands/trtllm-serve/run-benchmark-with-trtllm-serve.md)","tokens":300,"sizeBytes":280}],"systemPromptSnippet":"<yakaai_skills repo=\"NVIDIA/TensorRT-LLM\">\n<!-- File: .codex/AGENTS.md (Tokens: ~300 | Category: root-instruction) -->\n# .codex - AI Agent Protocol & Instructions\n\nPath: `.codex/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.codex/AGENTS.md)\n\n<!-- File: agent-flow/AGENTS.md (Tokens: ~300 | Category: root-instruction) -->\n# agent-flow - AI Agent Protocol & Instructions\n\nPath: `agent-flow/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/agent-flow/AGENTS.md)\n\n<!-- File: AGENTS.md (Tokens: ~300 | Category: root-instruction) -->\n# AI Agent Protocol & Instructions\n\nPath: `AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/AGENTS.md)\n\n<!-- File: tests/AGENTS.md (Tokens: ~300 | Category: root-instruction) -->\n# tests - AI Agent Protocol & Instructions\n\nPath: `tests/AGENTS.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/tests/AGENTS.md)\n\n<!-- File: .claude/skills/exec-env-check/SKILL.md (Tokens: ~1570 | Category: anthropic-skill) -->\n---\nname: exec-env-check\ndescription: >-\n  Check the local execution environment for GPU availability, Docker support,\n  and Slurm access. Returns the execution scenario (`satisfied, local, docker`,\n  `satisfied, local, direct`, `satisfied, slurm, local`, or `not_satisfied`),\n  the number of available GPUs, and the GPU type. On Slurm login nodes\n  without local GPUs, the cluster is identified by delegating the hostname\n  to internal-env-info (hostname-based mode), which owns the\n  hostname → cluster_name patterns; GPU type and gpus_per_node then come\n  from that skill's reference files. If internal-env-info is not\n  installed, the scenario falls back to `not_satisfied` without probing\n  compute nodes via srun.\ntags: [infrastructure, slurm, environment]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# TensorRT-LLM Environment Check\n\nDetect whether the current machine can run a GPU workload locally (Docker) or via Slurm, and report hardware details.\n\n## Input\n\n| Field | Description | Required |\n|-------|-------------|----------|\n| `required_devices` | Minimum number of GPUs needed | Yes |\n| `account` | Slurm account (unused for GPU probing, kept for compatibility) | No |\n\n## Procedure\n\n### 1. Check local GPUs\n\n```bash\ntimeout 5 nvidia-smi --query-gpu=name,memory.total --format=csv,noheader 2>/dev/null\n```\n\nIf this succeeds:\n- Count the number of GPU lines → `available_gpus`\n- Extract the GPU name from the first line → `device_type` (e.g., `NVIDIA B200`, `NVIDIA H100 80GB HBM3`)\n- Normalize `device_type`: strip `NVIDIA ` prefix and trailing memory info to get the short name (e.g., `B200`, `H100`, `A100`, `L40S`, `RTX 6000`)\n\nIf `nvidia-smi` fails or returns no GPUs → `available_gpus = 0`, `device_type = null`.\n\n### 2. Check if local GPUs are sufficient\n\nIf `available_gpus >= required_devices`, continue to step 2a to determine whether Docker is available on this host.\n\n### 2a. Check Docker availability\n\n```bash\ncommand -v docker >/dev/null 2>&1 && timeout 3 docker info >/dev/null 2>&1\n```\n\n- If the command succeeds (Docker CLI exists **and** the daemon responds) → **Result**: `satisfied, local, docker`\n- Otherwise (Docker CLI missing, daemon not running, or permission denied) → **Result**: `satisfied, local, direct`\n\nInclude `available_gpus` in the result. Do NOT include `device_type` — local execution does not need it.\n\n### 3. Check Slurm availability\n\nIf local GPUs are insufficient (or as additional detection), check for Slurm:\n\n```bash\nwhich squeue 2>/dev/null && squeue --version 2>/dev/null\n```\n\nIf Slurm is NOT available → go to step 5.\n\n### 4. Resolve GPU type\n\n**Optional dependency.** Before doing anything else in this step, check whether `skills/internal-env-info/` exists in the toolkit. If it does **not**, skip this step entirely and return `scenario: not_satisfied` with `available_gpus = 0`, `device_type = null`, `gpus_per_node = null`, `cluster_name = null`, `default_models_repo = null`, and `default_user_root_dir = null`. Do **not** report this as an error — the skill is an internal-only dependency.\n\nWhen Slurm is available but local `nvidia-smi` returned no GPUs or `device_type` is null (login nodes typically have no GPUs), capture the hostname and delegate cluster identification to `internal-env-info`:\n\n```bash\nhostname -f 2>/dev/null || hostname\n```\n\n- Pass the captured hostname to `internal-env-info` in **hostname-based mode**. That skill owns the NVIDIA-internal login-host patterns and the hostname → `<cluster>` mapping; do **not** parse the hostname or hard-code any cluster identifier here.\n- It returns the standard output template (`cluster_name`, `device_type`, `gpus_per_node`, `default_models_repo`, `default_user_root_dir`) plus supplementary field (`mfa_style`).\n- Set `available_gpus` = `gpus_per_node` (if resolved).\n- Set `cluster_name` = the value returned (a placeholder `<cluster>` token in this skill). The orchestrator uses this to fetch per-cluster info (`mfa_style`, `default_models_repo`, `default_user_root_dir`, `gpus_per_node`) from `internal-env-info`; connection fields (`mounts`, `ssh_host`, `partition`, etc.) come from caller-supplied inputs in `job_spec.json` (with `internal-env-info` default values / ask-the-user fallbacks).\n- Set `default_user_root_dir` = the user root directory returned (with `<user_name>` substituted with the actual SLURM username). Set to `null` if not found.\n- If `internal-env-info` returns `null` for `cluster_name` (no pattern matched), set `device_type = null`, `cluster_name = null`, `default_user_root_dir = null` and fall through to Step 5 (`not_satisfied`).\n\n**Result**: `satisfied, slurm, local` with `device_type`, `gpus_per_node`, `cluster_name`, and `default_user_root_dir`.\n\n### 5. Not satisfied\n\nIf neither local GPUs nor Slurm is available:\n- **Result**: `not_satisfied`\n\n## Output\n\nReturn a single structured result:\n\n```\nscenario: <satisfied, local, docker | satisfied, local, direct | satisfied, slurm, local | not_satisfied>\navailable_gpus: <N>\ndevice_type: <short GPU/device name or null>\ngpus_per_node: <N or null>\ncluster_name: <cluster name or null>\ndefault_models_repo: <host path to llm-models directory, or null>\ndefault_user_root_dir: <per-user root directory on cluster storage, or null>\n```\n\n**Examples:**\n\n```\nscenario: satisfied, local, docker\navailable_gpus: 4\n```\n\n```\nscenario: satisfied, local, direct\navailable_gpus: 4\n```\n\n```\nscenario: satisfied, slurm, local\navailable_gpus: 4\ndevice_type: B200\ngpus_per_node: 4\ncluster_name: <cluster>\ndefault_user_root_dir: <default_user_root_dir>/<user_name>\n```\n\n```\nscenario: not_satisfied\navailable_gpus: 0\ndevice_type: null\ncluster_name: null\n```\n\n## Rules\n\n- Never install drivers or modify the system\n- If `nvidia-smi` hangs, use a 5-second timeout: `timeout 5 nvidia-smi ...`\n- GPU type is derived from the hostname by matching the cluster name via the `internal-env-info` skill when it is installed — no `srun` allocation needed. If that skill is absent, the scenario falls back to `not_satisfied` (see Step 4); do not error.\n- Report the GPU type exactly as found in the mapping (do not guess or fabricate)\n\n<!-- File: .claude/skills/exec-local-compile/SKILL.md (Tokens: ~786 | Category: anthropic-skill) -->\n---\nname: exec-local-compile\ndescription: Compile TensorRT-LLM on a compute node inside a Docker container. Use this when already on a compute node with GPUs visible.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Compile TensorRT-LLM (Local / Compute Node)\n\nCompile TensorRT-LLM from source on a compute node inside a Docker container.\n\n## When to Use\n\n| Scenario | Use This Skill? |\n|----------|----------------|\n| On a compute node with GPUs visible (`nvidia-smi` works) | Yes |\n| On a SLURM login node (no GPUs) | No — use `exec-slurm-compile` instead |\n\n## Prerequisites\n\n- You are inside a Docker/enroot container on a compute node\n- `nvidia-smi` succeeds (GPUs visible)\n- `/usr/local/tensorrt` exists (TensorRT installation in the container)\n\n## Instructions\n\n### Step 1: Verify Environment\n\nRun `nvidia-smi` to confirm you are on a compute node with GPU access.\n\n### Step 2: Locate the Codebase\n\n`cd` to the TensorRT-LLM repository. If the path is not provided by the user, ask for it.\n\n### Step 3: (Optional) Checkout Branch\n\nIf the user specifies a branch (e.g., \"compile ToT\"), checkout and pull:\n```bash\ngit checkout main && git pull\n```\n\n### Step 4: Build\n\nRun the build command (**incremental by default** — omit `-c`/`--clean` unless explicitly requested or the incremental build fails):\n\n```bash\n./scripts/build_wheel.py --use_ccache -a \"<arch>\" -f --nvtx\n```\n\nReplace `<arch>` with the target GPU architecture (see Architecture Reference below). If not specified by the user, auto-detect from `nvidia-smi`.\n\n### Step 5: Install\n\n```bash\npip install -e .[devel]\n```\n\n### Step 6: Verify\n\n```bash\npython3 -c \"import tensorrt_llm; print(tensorrt_llm.__version__)\"\n```\n\n## Build Flags\n\n| Flag | Description |\n|------|-------------|\n| `-a \"<arch>\"` | Target GPU architecture(s) |\n| `--nvtx` | Enable NVTX markers for profiling |\n| `--use_ccache` | Use ccache for faster recompilation |\n| `-f` / `--fast_build` | Skip some kernels for faster dev compilation. **Always use for dev builds.** |\n| `-c` / `--clean` | Clean build directory before building. Only when needed (see below). |\n| `--skip_building_wheel` | Build in-place without creating a wheel file |\n| `--no-venv` | Skip virtual environment creation |\n\n## Architecture Reference\n\n| Value | GPU Family |\n|-------|-----------|\n| `\"100-real\"` | Blackwell (B200, GB200) |\n| `\"90-real\"` | Hopper (H100, H200) |\n| `\"89-real\"` | Ada Lovelace (L40S) |\n| `\"80-real\"` | Ampere (A100) |\n| `\"90;100-real\"` | Multiple architectures |\n\n## Incremental vs. Clean Builds\n\n**Default to incremental builds** — CMake only recompiles changed files, saving significant time.\n\nUse a **clean build** (`-c`) only when:\n- The user explicitly requests a clean/fresh build\n- An incremental build fails with linker errors, stale object files, or CMake cache issues\n- Major branch changes (e.g., rebasing across many commits) that may invalidate the build cache\n- Build system files changed (`CMakeLists.txt`, `*.cmake`)\n\n<!-- File: .claude/skills/exec-local-docker/SKILL.md (Tokens: ~1307 | Category: anthropic-skill) -->\n---\nname: exec-local-docker\ndescription: >-\n  Execute a TensorRT-LLM workload locally in Docker. Runs a fully-resolved\n  Docker command in background, monitors completion, reads logs, and reports\n  results. Workflow-agnostic — does not need to know if the workload is pytest,\n  eval, benchmark, or a custom script.\ntags: [docker, execution, infrastructure]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Local Docker Executor\n\nRun a Docker command locally, monitor it, and report results.\n\n## Input (from orchestrator prompt)\n\nThe orchestrator passes these fields in the skill prompt:\n\n| Field | Description |\n|-------|-------------|\n| `docker_cmd` | Complete `docker run` command string, ready to execute |\n| `work_dir` | Local work directory for logs and artifacts |\n| `log_file` | Full path to the log file (output redirected here) |\n| `model_name` | Short model name for reporting |\n| `workflow_type` | `pytest`, `eval`, `custom`, or `benchmark` — for output parsing hints |\n| `success_patterns` | Comma-separated patterns indicating success (e.g., `passed,accuracy:`) |\n| `failure_patterns` | Comma-separated patterns indicating failure (e.g., `FAILED,Error,AssertionError`) |\n\n## Procedure\n\n### Step 0: Resolve Image and Build (when `build_project=true`)\n\nThis executor owns image selection and the build for the local Docker target. Skip the entire step when `build_project=false`.\n\n1. **Detect the target GPU type.** Use `gpu_type` from `job_spec.json` if upstream env-check resolved it; otherwise probe locally with `nvidia-smi --query-gpu=name --format=csv,noheader | head -1`.\n2. **Detect the host CPU arch** with `uname -m` (`x86_64` or `aarch64`).\n3. **Resolve the container image.** Read `<repo_root>/jenkins/current_image_tags.properties` and pick the tag whose CPU-arch flavor matches the host. If the orchestrator already passed a `container_image` field in the job spec, use that and skip the lookup.\n4. **Map GPU → build arch (`-a` flag):** `H100`/`H200` → `90-real`; `B200`/`GB200`/`B300`/`GB300` → `100-real`; `A100` → `80-real`; `L40S` → `89-real`. Default `100-real` when the GPU is unknown.\n5. **Compile.** Invoke the `exec-local-compile` skill with `repo_dir=<repo_root>`, `image=<resolved tag>`, `arch=<arch>`. Wait for completion.\n6. **On failure**, do not launch the workload. Report `BUILD_FAILED` with the last 100 lines of the compile log.\n\n`build_project`, `gpu_type`, `repo_root`, and (optionally) `container_image` come from `job_spec.json`.\n\n### Step 1: Launch\n\nRun the Docker command in background using `run_in_background`:\n\n```bash\n<docker_cmd> 2>&1 | tee <log_file>\n```\n\nReport to the orchestrator: \"Launched locally, log at `<log_file>`\"\n\n### Step 2: Monitor for Hangs\n\nWhile waiting for the background process to complete, actively monitor the log\nfile for hang indicators. Launch a monitoring loop using `run_in_background`:\n\n```bash\nwhile true; do\n  sleep 60\n  if [ -f \"<log_file>\" ] && grep -qi \"hang detected\" \"<log_file>\"; then\n    echo \"HANG_DETECTED: Found 'hang detected' in log file\"\n    docker ps --filter \"ancestor=<image>\" -q | xargs -r docker kill 2>/dev/null\n    exit 1\n  fi\ndone\n```\n\n- If the monitor detects a hang, it kills the Docker container and exits with\n  code 1. The main background process will also terminate.\n- When the main process completes normally (background notification received),\n  kill the monitoring loop (it is no longer needed).\n- If a hang is detected, skip to Step 4 and report `HANG_DETECTED` status\n  instead of proceeding to normal result collection.\n\n### Step 3: Wait for Completion\n\nThe Bash tool's `run_in_background` will notify when the process finishes\n(either normally or because the container was killed by the hang monitor).\n\n### Step 4: Read Results\n\nOn completion:\n\n1. **Read exit code** from the background command result.\n2. **Read the last 100 lines** of `<log_file>` using the Read tool.\n3. **If exit code != 0**, also read the first 50 lines to catch early errors (import failures, setup crashes).\n4. **Search for patterns**:\n   - Grep `<log_file>` for each `success_patterns` entry\n   - Grep `<log_file>` for each `failure_patterns` entry\n\n### Step 5: Report\n\nReturn a structured result:\n\n```\nStatus: PASSED | FAILED | ERROR | HANG_DETECTED\nExit code: <N>\nLog file: <log_file>\nWork directory: <work_dir>\nSummary: <relevant output lines — pytest summary, accuracy score, or last few lines>\nErrors: <if failed, the relevant error output>\n```\n\n### Output Parsing by Workflow Type\n\n- **pytest**: Look for `X passed, Y failed in Zs` summary line\n- **eval**: Look for `accuracy:` or `score:` lines; check for `Expected accuracy >= X, but got Y` assertion\n- **custom**: No specific patterns — report last 10 lines of output\n- **benchmark**: Look for throughput/latency numbers\n\n## Rules\n\n- Never run the Docker command in foreground — always use `run_in_background`\n- Never `cat` the full log file — use Read with offset/limit or tail\n- If the Docker command fails immediately (exit code within seconds), check if the image exists locally\n- Report results even if the log file is empty (container may have failed to start)\n\n<!-- File: .claude/skills/exec-local-slurm/SKILL.md (Tokens: ~5160 | Category: anthropic-skill) -->\n---\nname: exec-local-slurm\ndescription: >-\n  Submit and monitor a Slurm job on a local cluster. Supports two modes:\n  (1) Persistent allocation (default) — allocates nodes once via nohup salloc,\n  imports the container once, installs once, and reuses across runs by setting\n  SLURM env vars and running the sbatch script via bash. (2) One-shot sbatch —\n  submits a fully-generated Slurm script via sbatch, polls job status, reads\n  logs on completion, and reports results. Workflow-agnostic — handles pytest,\n  eval, benchmark, and custom scripts identically.\ntags: [slurm, execution, infrastructure]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Local Slurm Executor\n\nSubmit a Slurm job locally, monitor it, and report results. Uses persistent\nallocation by default to eliminate queue wait, container import, and install\noverhead on repeated runs.\n\n## Input (from orchestrator prompt)\n\nThe orchestrator passes these fields in the skill prompt:\n\n| Field | Description |\n|-------|-------------|\n| `script_path` | Full local path to the generated `.slurm` script |\n| `work_dir` | Local work directory for logs and artifacts |\n| `model_name` | Short model name for reporting |\n| `workflow_type` | `pytest`, `eval`, `custom`, or `benchmark` — for output parsing hints |\n| `success_patterns` | Comma-separated patterns indicating success |\n| `failure_patterns` | Comma-separated patterns indicating failure |\n| `log_file_pattern` | Log filename pattern with `%j` placeholder (e.g., `llama_auto_test_%j.out`) |\n| `persistent_mode` | Default: `true` for all local slurm workflows. Set to `false` to force one-shot sbatch (opt-out). |\n| `release_allocation` | `true` to release the current allocation and stop. Only set when the user explicitly says no more jobs are needed. Default: `false`. Never auto-release. |\n| `alloc_time_limit` | Walltime for the persistent allocation. Default: `04:00:00`. |\n| `docker_image` | Container image for persistent container import. From `job_spec.json`. |\n| `container_name` | Container name used in the `.slurm` script (e.g., `llama_auto_test`). Must match exactly. From `job_spec.json`. |\n| `container_mounts` | Comma-separated mount mappings. From `job_spec.json`. |\n| `repo_root` | Repo root path (for locating state file and project path). |\n| `slurm_params` | Slurm parameters object: `partition`, `account`, `nodes`, `ntasks`, `ntasks_per_node`, `gpus_per_node`. From `job_spec.json`. |\n\n## Procedure\n\n### Pre-step: Resolve Image and Build (when `build_project=true`)\n\nThis executor owns image selection and the build for the local SLURM cluster. Skip the entire pre-step when `build_project=false`.\n\n1. **Read the cluster GPU type** from `device_type` in `job_spec.json` (resolved upstream by env-check / internal-env-info). If absent, probe from any compute node via `srun -p <partition> --ntasks-per-node=1 nvidia-smi --query-gpu=name --format=csv,noheader | head -1`. If `skills/internal-env-info/` is not installed, skip the upstream lookup silently and rely on the srun probe — do not report the missing skill as an error.\n2. **Determine the partition's CPU arch.** Look up `job_spec.slurm_env.partitions[]` for the entry where `name == partition` and read its `arch` (`x86_64` / `aarch64`). Case-executor's Step 2.5 already detected this — do **not** run `scontrol show node` here. If `slurm_env` is absent or the entry's `arch` is null (case-executor ran without SLURM tools), fall back to inferring from `device_type`: Grace-based parts (`GB*` / `GH*`) → `aarch64`; everything else → `x86_64`.\n3. **Resolve the container image.** If `docker_image` is already set in `job_spec.json`, use it directly. Otherwise read `<repo_root>/jenkins/current_image_tags.properties` and pick the tag matching the partition's CPU arch.\n4. **Map GPU → build arch (`-a` flag):** `H100`/`H200` → `90-real`; `B200`/`GB200`/`B300`/`GB300` → `100-real`; `A100` → `80-real`; `L40S` → `89-real`. Default `100-real`.\n5. **Compile.** Invoke the `exec-slurm-compile` skill with `repo_dir=<repo_root>`, `partition`, `account`, `container_image=<resolved tag>`, `user_root_dir`, `arch`. Wait for completion.\n6. **On failure**, do not allocate or run the workload. Report `BUILD_FAILED` with the last 100 lines of the build log.\n\n`build_project`, `device_type`, `repo_root`, `partition`, `account`, `user_root_dir`, and (optionally) `docker_image` come from `job_spec.json`.\n\n### Step 0: Allocation Management\n\nThis step runs before any execution. It determines whether to reuse an existing\npersistent allocation, create a new one, or fall through to one-shot sbatch.\n\n#### Step 0A — Release mode\n\nIf `release_allocation=true`:\n\n1. Read `<repo_root>/work_dirs/.slurm_alloc.json`\n2. If file exists and `job_id` is present:\n   ```bash\n   scancel <job_id>\n   ```\n3. Delete the state file and `.salloc.log`\n4. Report: \"Allocation <job_id> released.\" Stop.\n\n#### Step 0B — One-shot mode\n\nIf `persistent_mode=false`, skip entirely to Step 1 (One-Shot Path).\n\n#### Step 0C — Validate existing allocation\n\nRead `<repo_root>/work_dirs/.slurm_alloc.json`.\n\n- If state file **does not exist** → go to Step 0D.\n- If state file exists, validate the allocation:\n  ```bash\n  squeue -j <job_id> -h -o \"%T %L\"\n  ```\n  - **RUNNING + remaining > 5 min + params compatible** → **reuse** (skip to Step 0F)\n  - **RUNNING + remaining <= 5 min** → warn user (\"Allocation expiring soon\"), `scancel <job_id>`, delete state file, go to Step 0D\n  - **RUNNING + params incompatible** → `scancel <job_id>`, delete state file, go to Step 0D\n  - **Empty output / error** (job gone) → stale state file, delete it, go to Step 0D\n\n**Params compatibility check:**\n- `partition` must match\n- `nodes` in state must be **>=** requested nodes (a 2-node allocation can serve 1-node jobs)\n- `docker_image` must match\n\n#### Step 0D — Justify and prepare allocation\n\nBefore allocating, ensure no orphaned allocations exist and log the reason:\n\n1. Check for existing persistent jobs:\n   ```bash\n   squeue -u $(whoami) -h -o \"%i %T %j\" --name=<account>-trtllm.persistent\n   ```\n2. If found → orphan (state file was missing/corrupt). Cancel it:\n   ```bash\n   scancel <orphan_job_id>\n   ```\n   Log: \"Released orphaned allocation <orphan_job_id> — state file was missing.\"\n3. Log justification: \"No reusable allocation found. Allocating <nodes> node(s) on partition <partition> for <alloc_time_limit>.\"\n\n#### Step 0E — Allocate new\n\nConvert `alloc_time_limit` from `HH:MM:SS` to seconds (e.g., `04:00:00` → `14400`).\n\n```bash\nmkdir -p <repo_root>/work_dirs\nnohup salloc --partition=<partition> --account=<account> \\\n  --nodes=<nodes> --time=<alloc_time_limit> \\\n  --job-name=<account>-trtllm.persistent \\\n  sleep <alloc_time_limit_seconds> \\\n  > <repo_root>/work_dirs/.salloc.log 2>&1 &\n```\n\nRetrieve job ID:\n```bash\nsqueue -u $(whoami) -h -o \"%i\" --name=<account>-trtllm.persistent\n```\n\nIf no job appears after 10 seconds, read `<repo_root>/work_dirs/.salloc.log`\nfor errors (bad partition, invalid account, etc.) and report to user.\n\nOtherwise, poll until RUNNING (every 10s, max 60 polls). Once RUNNING, get\nthe nodelist:\n```bash\nsqueue -j <JOB_ID> -h -o \"%N\"\n```\n\n**Import container** on all nodes (using the **same container name** as the\n`.slurm` script — critical for pyxis reuse):\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-image=<docker_image> \\\n  --container-name=<container_name> true\n```\n\n**Warm up filesystem mounts** on all nodes — `ls` each mounted path so\nLustre/NFS metadata is cached for later use:\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-name=<container_name> \\\n  --container-mounts=<container_mounts> \\\n  bash -c 'for p in <mount_target_1> <mount_target_2> ...; do ls \"$p\" > /dev/null 2>&1; done'\n```\nParse mount targets from `container_mounts` — the right-hand side of each\n`host:container` pair.\n\n**Check GPU status** on all nodes (no container needed — `nvidia-smi` is on\nthe host):\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  bash -c 'echo \"=== $(hostname) ===\" && nvidia-smi --query-compute-apps=pid,name,used_memory --format=csv,noheader'\n```\n- If output shows no processes → GPUs are clean, proceed\n- If unexpected processes found → warn user: \"GPU processes found on\n  <node>: <process_list>. These may interfere with the job.\"\n\n**Run install** on all nodes with `--container-writable` (so packages persist\nin the named container across srun calls). Install all common requirements\nupfront so that subsequent jobs of any workflow type can skip install:\n```bash\nsrun --jobid=<JOB_ID> -N <nodes> --ntasks-per-node=1 \\\n  --container-name=<container_name> --container-writable \\\n  --container-mounts=<container_mounts> \\\n  bash -c 'cd <repo_root> && pip install -e . && pip install -r requirements-dev.txt && \\\n    if [ -f examples/trtllm-eval/requirements.txt ]; then pip install -r examples/trtllm-eval/requirements.txt; fi'\n```\nFor custom workflow with `skip_install=true`, skip this step entirely.\n\n**Write state file** `<repo_root>/work_dirs/.slurm_alloc.json`:\n```json\n{\n  \"job_id\": \"<JOB_ID>\",\n  \"container_name\": \"<container_name>\",\n  \"nodelist\": \"<NODELIST>\",\n  \"partition\": \"<partition>\",\n  \"account\": \"<account>\",\n  \"nodes\": <nodes>,\n  \"gpus_per_node\": <gpus_per_node>,\n  \"docker_image\": \"<docker_image>\",\n  \"container_mounts\": \"<container_mounts>\",\n  \"allocated_at\": \"<ISO timestamp>\",\n  \"time_limit\": \"<alloc_time_limit>\",\n  \"installed\": true\n}\n```\n\nProceed to Step 1 (Persistent Path).\n\n#### Step 0F — Reuse path validation\n\nAllocation is valid. Check if container name matches current request:\n\n- If `container_name` in state file **matches** `container_name` from\n  `job_spec.json` → proceed to Step 1 (Persistent Path). Container and install\n  are already set up.\n- If **different** (model changed) → import the new container, warm up mounts,\n  and install:\n  ```bash\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-image=<docker_image> \\\n    --container-name=<new_container_name> true\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-name=<new_container_name> \\\n    --container-mounts=<container_mounts> \\\n    bash -c 'for p in <mount_targets>; do ls \"$p\" > /dev/null 2>&1; done'\n  srun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n    --container-name=<new_container_name> --container-writable \\\n    --container-mounts=<container_mounts> \\\n    bash -c 'cd <repo_root> && pip install -e . && pip install -r requirements-dev.txt'\n  ```\n  Update `container_name` in state file. Proceed to Step 1 (Persistent Path).\n\n---\n\n### Step 1: Execute\n\nTwo execution paths depending on `persistent_mode`.\n\n#### Persistent Path (persistent_mode=true)\n\n**Pre-flight time check** — verify the allocation has enough remaining time:\n```bash\nsqueue -j <job_id> -h -o \"%L\"\n```\nIf remaining time < job's `time_limit` from `slurm_params`, warn user:\n\"Allocation has <remaining> left but job expects <time_limit>. The job may\nbe killed early.\"\n\n**Pre-flight GPU check** — verify GPUs are not occupied by leftover processes:\n```bash\nsrun --jobid=<job_id> -N <nodes> --ntasks-per-node=1 \\\n  bash -c 'procs=$(nvidia-smi --query-compute-apps=pid,name,used_memory --format=csv,noheader 2>/dev/null); [ -n \"$procs\" ] && echo \"WARNING: GPU processes on $(hostname): $procs\"'\n```\nIf unexpected processes found, warn user before proceeding.\n\n**Run the sbatch script** with SLURM env vars — the `#SBATCH` directives are\ncomments when run via `bash`; the inner `srun` inherits the env vars and uses\nthe persistent allocation:\n```bash\nexport SLURM_JOB_ID=<job_id>\nexport SLURM_JOB_NUM_NODES=<nodes>\nexport SLURM_NNODES=<nodes>\nexport SLURM_NTASKS=<ntasks>\nexport SLURM_NTASKS_PER_NODE=<ntasks_per_node>\nexport SLURM_NODELIST=<nodelist>\nexport SLURM_JOB_NODELIST=<nodelist>\nbash <script_path> 2>&1 | tee <work_dir>/<model_name>_<workflow_type>.log\n```\n\nRun with `run_in_background`. Container reuse is automatic (pyxis skips\nre-import when `--container-name` already exists on the node). The install\nstep (Step 1 srun) always runs to keep the container up to date.\n\n**Step cancellation:** If a job hangs or the user wants to abort, cancel just\nthe srun step without killing the allocation:\n```bash\n# List active steps:\nsqueue -s -j <job_id>\n# Cancel a specific step (e.g., step 0):\nscancel <job_id>.0\n```\nThe allocation stays RUNNING — new jobs can run immediately. The job name\n`<account>-trtllm.persistent` uniquely identifies allocations created by this\nskill via `squeue --name=<account>-trtllm.persistent`.\n\n#### One-Shot Path (persistent_mode=false)\n\n```bash\nsbatch <script_path>\n```\n\nParse the job ID from output: `Submitted batch job <JOB_ID>`.\n\nIf sbatch fails, report the error immediately and stop.\n\n### Step 2: Report Submission\n\n**One-shot mode:**\n```\nJob ID: <JOB_ID>\nScript: <script_path>\nWork directory: <work_dir>\nLog files: <work_dir>/<log_file_pattern with %j replaced>\nMonitor: squeue -j <JOB_ID>\n```\n\n**Persistent mode:**\n```\nAllocation: <job_id> (persistent, remaining: <time>)\nContainer: <container_name>\nScript: <script_path>\nWork directory: <work_dir>\nLog file: <work_dir>/<model_name>_<workflow_type>.log\n```\n\n### Step 3: Poll / Wait for Completion (with Hang Monitoring)\n\nBoth polling modes must also monitor the log file for hang indicators. Every\n60 seconds, check the log file for the pattern `hang detected`\n(case-insensitive). If found, terminate the job and skip to Step 5 with\n`HANG_DETECTED` status.\n\n**One-shot mode:** Poll every 30 seconds, max 60 polls (30 minutes):\n\n```bash\nsqueue -j <JOB_ID> -h -o \"%T %M %R\"\n```\n\n- Report state transitions: `PENDING` → `RUNNING` → `COMPLETING` → done\n- If `PENDING` for > 5 minutes, check the reason:\n  ```bash\n  squeue -j <JOB_ID> -o \"%i %T %r %S\"\n  ```\n- The job is done when `squeue -j <JOB_ID>` returns empty (no rows)\n- If max polls exceeded, report the job is still running and provide manual\n  check commands\n- **Hang check** (every other poll, i.e., every 60s): Once the job is RUNNING,\n  resolve the log file path (`<work_dir>/<log_file_pattern>` with `%j` replaced\n  by `<JOB_ID>`) and check:\n  ```bash\n  grep -qi \"hang detected\" <log_file> 2>/dev/null && echo \"HANG_DETECTED\"\n  ```\n  If `HANG_DETECTED` is printed, immediately cancel the job:\n  ```bash\n  scancel <JOB_ID>\n  ```\n  Then read the last 200 lines of the log file and skip to Step 5.\n\n**Persistent mode:** While waiting for `run_in_background` completion\nnotification, launch a hang monitoring loop using `run_in_background`:\n\n```bash\nwhile true; do\n  sleep 60\n  LOG_FILE=\"<work_dir>/<model_name>_<workflow_type>.log\"\n  if [ -f \"$LOG_FILE\" ] && grep -qi \"hang detected\" \"$LOG_FILE\"; then\n    echo \"HANG_DETECTED: Found 'hang detected' in log file\"\n    squeue -s -j <job_id> -h -o \"%i\" | head -1 | xargs -r scancel 2>/dev/null\n    exit 1\n  fi\ndone\n```\n\nWhen the main process completes normally, kill the monitoring loop. If the\nmonitor fires first, it cancels the active Slurm step (not the allocation\nitself — the allocation stays for future runs).\n\n### Step 4: Collect Results\n\nOn completion:\n\n1. **Get exit code:**\n   - **One-shot:** `sacct -j <JOB_ID> --format=JobID,State,ExitCode,Elapsed,MaxRSS -n`\n   - **Persistent:** From the `run_in_background` result exit code\n\n2. **Resolve log file paths:**\n   - **One-shot:** Replace `%j` in `log_file_pattern` with `<JOB_ID>`:\n     - Output: `<work_dir>/<MODEL_NAME>_..._<JOB_ID>.out`\n     - Error: `<work_dir>/<MODEL_NAME>_..._<JOB_ID>.err`\n   - **Persistent:** Single log at `<work_dir>/<model_name>_<workflow_type>.log`\n\n3. **Read log files:**\n   - Read the last 100 lines of the output log\n   - If job failed, also read the first 50 lines to catch early errors\n   - **One-shot only:** If job failed, read the last 50 lines of the `.err` file\n\n4. **Search for patterns:**\n   - Grep the log file for each `success_patterns` entry\n   - Grep the log file for each `failure_patterns` entry\n\n### Step 5: Report\n\nReturn a structured result:\n\n```\nJob ID: <JOB_ID> (or allocation <job_id> for persistent mode)\nStatus: COMPLETED | FAILED | TIMEOUT | CANCELLED | OUT_OF_MEMORY | HANG_DETECTED\nExit code: <exit code>\nElapsed: <time>\nLog files:\n  stdout: <path>\n  stderr: <path> (one-shot only)\nWork directory: <work_dir>\nSummary: <relevant output — pytest summary, accuracy score, throughput, or last lines>\nErrors: <if failed, relevant error output>\n```\n\n### Output Parsing by Workflow Type\n\n- **pytest**: Look for `X passed, Y failed in Zs` summary line. Status maps:\n  `COMPLETED` + 0 exit → `PASSED`; non-zero → `FAILED`\n- **eval**: Look for `accuracy:` lines. If `--check_accuracy` was used,\n  assertion error means threshold not met\n- **benchmark**: Look for throughput/latency numbers. Check for\n  `8_done_*.txt` completion marker\n- **custom**: No specific patterns — report last 10 lines of output\n\n### Failure Diagnosis\n\n| Symptom | Likely Cause | Action |\n|---------|-------------|--------|\n| `TIMEOUT` | Job exceeded time limit | Report and suggest increasing time |\n| `OUT_OF_MEMORY` | OOM on node | Report and suggest reducing batch size |\n| Exit `1:0` + `FAILED` | Non-zero exit from srun | Read `.err` file for details |\n| `CANCELLED` | Preempted or user cancelled | Check `sacct --format=JobID,State,Reason` |\n| Empty output log | Container failed to start | Check `.err` / `.salloc.log` for mount or image errors |\n| `HANG_DETECTED` in log | Process/GPU hang (deadlock, NCCL timeout) | Job auto-cancelled; check last 200 log lines for root cause |\n| `srun: error: Unable to create step` | Allocation expired | Delete state file, re-run (will re-allocate) |\n| `srun: error: ... does not exist` | Named container missing | Delete state file, re-run (will re-import) |\n\n## Rules\n\n- Never use `squeue` with the `-uall` flag\n- Always use `sacct` (not `squeue`) for final job status in one-shot mode\n- Poll with `run_in_background` + `sleep 30 &&` to avoid blocking (one-shot)\n- Never read full log files — use Read with offset/limit\n- If the Slurm script path doesn't exist, report the error immediately\n- For cluster hardware info (GPU types, partitions, constraints, storage\n  paths), use the `trtllm-agent-toolkit:internal-env-info` skill\n  when it is installed. If the skill is absent, skip the lookup silently and\n  use the caller-supplied fields or the srun probe in Step 1 — do not report\n  the missing skill as an error.\n- Never auto-release an allocation after job completion — it should persist\n  for potential future runs until walltime expires or user explicitly releases\n- Must pass `-N <nodes>` to every direct `srun --jobid` call — without it,\n  srun defaults to 1 node even in a multi-node allocation\n- The container name used in setup steps must **exactly match** the\n  `--container-name` in the `.slurm` script, otherwise pyxis will re-import\n  the image and lose `--container-writable` state\n\n<!-- File: .claude/skills/exec-remote-slurm/SKILL.md (Tokens: ~10397 | Category: anthropic-skill) -->\n---\nname: exec-remote-slurm\ndescription: >-\n  Remote SLURM cluster development via SSH. Use when running jobs,\n  profiling, or developing on a remote SLURM cluster with pyxis/enroot\n  containers. Covers SSH connection management, srun/sbatch/salloc job\n  patterns, tmux-based allocation persistence, file transfer, and safe\n  remote file access. Works with any SLURM cluster accessible via SSH.\ntags: [infrastructure, remote, slurm]\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Remote SLURM Development\n\nDevelop and execute on remote SLURM clusters via SSH commands. No filesystem\nsync needed — all remote operations go through SSH.\n\n## Step 1: Resolve Cluster\n\n**Optional dependency — `internal-env-info`.** Before doing\nanything else in Step 1, check whether `skills/internal-env-info/`\nexists in the toolkit. If it does **not**, skip every sub-step that invokes\nthe skill or reads its `references/` files and proceed using only the\ncluster fields the orchestrator wrote into `job_spec.json` (`ssh_host`,\n`slurm_user`, `remote_cwd`, `partition`, `account`, `container_image`,\n`mounts`, `gpus_per_node`, `mfa_style`). If a required field is also\nmissing, stop and ask the user to supply it — do **not** emit an error\nabout the missing skill.\n\nThe orchestrator supplies one of two input shapes; pick the right path:\n\n- **Explicit cluster** — `slurm_cluster` name is provided. Skip\n  auto-selection and go straight to **Load the per-cluster info**.\n- **Auto-select cluster** — only hardware constraints are provided\n  (`device_type`, `total_required_devices`, `required_devices_per_node`).\n  Resolve a cluster from the constraints, then load its per-cluster info.\n\n### 1a. Auto-select (only when no explicit cluster was provided)\n\nInvoke the `internal-env-info` skill in **constraint-based mode**, passing\n`device_type`, `required_devices_per_node`, and (optionally)\n`total_required_devices`. That skill owns the filter rules and either\nreturns a single resolved cluster entry or stops and asks the user. Do not\nimplement the filter here.\n\nUse `node_count` from `job_spec.json` directly as `--nodes` and\n`required_devices_per_node` as `--ntasks-per-node`. **Do not recompute\n`node_count`** — `trtllm-case-executor` is the single source of truth and\nwrites the resolved value into `job_spec.json` for every executor to read.\n\nThe returned entry's `cluster_name`, `mfa_style`, `default_user_root_dir`,\n`default_models_repo`, and `gpus_per_node` feed Step 1b.\n\n### 1b. Load the per-cluster info\n\nConnection fields are read **only** from `job_spec.json` (written by\n`trtllm-case-executor`): `ssh_host`, `slurm_user`, `remote_cwd`, `account`,\n`partition`, `mounts`, `container_image`, `gpus_per_node`, and `packages`.\nThere is no per-cluster connection-config file to parse — if a required\nfield is absent from `job_spec.json`, stop and ask the user to supply it.\nDo not silently default.\n\nInvoke the `internal-env-info` skill in single-cluster mode with\nthe resolved `cluster_name` to retrieve per-cluster info: partition limits,\ntopology-aware scheduling guidance, `default_models_repo`,\n`default_user_root_dir`, `gpus_per_node`, `mfa_style` (used by Step 2 to\nchoose the SSH path — see below), and shared constraint and tuning\npatterns. The runtime `container_image` always comes from `job_spec.json`,\nnever from this skill.\n\n## Step 2: SSH Preflight\n\nThe cluster's `mfa_style` (returned by `internal-env-info` in Step 1b) is\nthe **single source of truth for which SSH path to use**. Switch on the\nvalue — do **not** probe-and-fall-back on the SSH error string.\n\n| `mfa_style` | Path | Why |\n|-------------|------|-----|\n| `false` | Direct (Step 2a) | Cluster accepts key-based / passwordless SSH; no special handling needed. |\n| `true`  | MFA (Step 2b)    | Cluster requires interactive MFA (Duo / device-code / one-time PIN). |\n| `null`  | Unknown (Step 2c)| `mfa_style` not declared in the cluster reference (or `internal-env-info` is not installed). Probe direct first, fall back to MFA. When the reference exists, also ask the user to update it so future runs skip the probe; otherwise just continue silently. |\n\n### 2.0. Build `<ssh_host>` from `slurm_user` + `cluster_name`\n\nBefore any SSH attempt, derive the SSH destination. Priority order:\n\n1. **Caller-supplied `ssh_host`** (passed in through case-executor) → use\n   verbatim, skip everything below. This is the canonical escape hatch.\n2. **`internal-env-info` skill installed** → delegate for the\n   `ssh_host_pattern` key (defined in the skill's `Default Environment Values`\n   section). It returns a single-line snippet that concatenates `slurm_user`\n   and `cluster_name` into `SSH_HOST`. Do **not** re-derive or re-suffix\n   `slurm_user` here — `internal-env-info` already applied whatever its\n   `slurm_user` construction rule requires for that cluster.\n\n   Inputs to substitute:\n   - `slurm_user` — pre-resolved from `internal-env-info` (already in its\n     final, SSH-ready form for that cluster). Read it from\n     `job_spec.json`.\n   - `cluster_name` — from `internal-env-info` Step 1b.\n3. **Skill not installed** → stop and ask the user to supply `ssh_host`\n   explicitly. Do **not** hard-code a login-host pattern in this skill —\n   the NVIDIA-internal convention lives only in `internal-env-info`.\n\nThe resolved value is the canonical `<ssh_host>` placeholder used everywhere\nin the rest of this skill (Steps 2a/2b/2c, Step 3, recipes, file-transfer,\nhang detection).\n\n### 2a. Direct path — `mfa_style: false`\n\nVerify connectivity using the user's existing SSH config:\n\n```bash\nssh -o ConnectTimeout=5 <ssh_host> \"echo ok\"\n```\n\nIf this succeeds, **skip to Step 3**. Do not create a trtllm-managed\nControlMaster — reuse whatever the user already has.\n\nIf it fails, do **not** silently switch to the MFA path. Stop and ask the\nuser to verify their `ssh_host` value or their key-based SSH setup.\nA direct-path cluster failing here is a configuration issue, not an MFA\nprompt to interpret.\n\n### 2b. MFA path — `mfa_style: true`\n\nFirst check whether a ControlMaster is already live (the user may have set\none up earlier in this session):\n\n```bash\nssh -O check <ssh_host> 2>&1 || ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -O check <ssh_host> 2>&1\n```\n\nIf either succeeds, **skip to Step 3** — reuse it.\n\nOtherwise the user must complete an interactive MFA handshake. MFA uses\nbrowser-based device-code verification that needs an interactive SSH\nsession — the skill cannot complete this on the user's behalf because the\nsession must stay alive while the user authenticates in their browser and\npresses ENTER.\n\nTell the user to run the ControlMaster command:\n\n> Run this in your terminal:\n>\n>     ssh -fNM -o ControlMaster=auto -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -o ControlPersist=24h <ssh_host>\n>\n> If the cluster uses a device-code flow you will see a PIN and a URL like:\n>\n>     Authenticate with PIN XXXXXX at <your-identity-provider-device-url>\n>\n> Open that URL, enter the PIN, approve, then press ENTER. The SSH session\n> will background itself as a persistent ControlMaster. Let me know when done.\n\nAfter the user confirms, verify with\n`ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p -O check <ssh_host>` and\nproceed to Step 3.\n\n### 2c. Unknown — `mfa_style: null`\n\nThe cluster reference (`references/<cluster>.md`) does not declare an\n`mfa_style` value. Try the direct path (Step 2a) first. If it fails with\n`Permission denied (keyboard-interactive)` — and only that — switch to\nthe MFA path (Step 2b). Any other failure stops here and is reported to\nthe user.\n\nAfter connecting successfully, **ask the user to update the cluster\nreference file** with the discovered value (`mfa_style: true` or\n`false`, in the `## Authentication` table) so that future runs skip\nthis probe and use the deterministic switch in Step 2a / 2b.\n\n## SSH Command Shorthand\n\nEvery SSH and SCP invocation in the rest of this skill uses the placeholders below. Both pin `ControlPath` to the trtllm-managed socket so they reuse the ControlMaster set up in Step 2 — without this, OpenSSH starts a fresh handshake on every call, which on MFA clusters re-prompts for keyboard-interactive auth (and defeats the whole point of multiplexing):\n\n| Placeholder | Expansion |\n|-------------|-----------|\n| `<ssh_cmd>` | `ssh -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p` |\n| `<scp_cmd>` | `scp -o ControlPath=/tmp/ssh-trtllm-%r@%h:%p` |\n\nWhen the user has their own ControlMaster configured in `~/.ssh/config` (no trtllm-managed socket needed), the explicit `-o ControlPath` is harmless: if the path has no live master, OpenSSH falls back to a fresh connection. The connectivity probe in Step 2 Check 2 (`ssh -o ConnectTimeout=5 <ssh_host> \"echo ok\"`) deliberately uses bare `ssh` to test un-multiplexed access; everywhere else use `<ssh_cmd>` / `<scp_cmd>`.\n\n## Step 3: Remote Environment Check\n\nVerify the remote environment is ready:\n\n```bash\n<ssh_cmd> <ssh_host> \"test -d <remote_cwd> && echo ok\"\n<ssh_cmd> <ssh_host> \"which tmux\"\n<ssh_cmd> <ssh_host> \"squeue --version\"\n```\n\nIf `remote_cwd` does not exist, create it or ask the user. If tmux or SLURM\ncommands are not available, stop and inform the user.\n\n## Step 4: Prepare Remote Environment\n\nThis step sets up the user's working directory and clones the TensorRT-LLM\nrepository on the remote cluster. It runs once before any job submission.\n\n> **MANDATORY: two-mechanism sync — both must run.**\n> - **Tracked changes** → temp branch (**4c–4f**): commit on a session-scoped\n>   branch, push to `origin`, fetch + checkout on remote, delete the ref both\n>   sides.\n> - **Untracked files** (`git ls-files --others --exclude-standard`) →\n>   `tar | ssh` pipe (**4e.1**), right after the remote checkout. Never touches\n>   `origin`.\n>\n> **Do not** shortcut with `git fetch origin <repo_branch> && git checkout <repo_branch>` —\n> `repo_branch` may not exist on `origin`, may be stale, or may lack uncommitted\n> edits. Always run 4c–4f **and** 4e.1.\n\n**Inputs received from `trtllm-case-executor`:**\n\n| Field | Description |\n|-------|-------------|\n| `default_user_root_dir` | Per-user root directory on the remote cluster (from `exec-env-check` / `internal-env-info`) |\n| `repo_url` | Git remote URL of the local TensorRT-LLM repo (from `git remote get-url origin`). **Used to clone the bare repository on the remote (Step 4b) and to push the temp branch to (Step 4d).** |\n| `repo_branch` | **Informational only** — the local current branch name (from `git rev-parse --abbrev-ref HEAD`). Used in Step 4c to remember which branch to restore locally after the temp branch is deleted. **Never used as the ref the remote checks out** — the remote always checks out the temp branch from Step 4c. |\n| `device_type` | Required short GPU name (e.g., `B200`, `B300`). Used by Step 1a auto-selection. Ignored when an explicit cluster was provided. |\n| `total_required_devices` | Total GPU processes across all nodes. Used by Step 1a auto-selection only. SLURM job sizing reads `node_count` from `job_spec.json` (precomputed by `trtllm-case-executor`) — do not recompute it here. |\n| `required_devices_per_node` | Minimum GPUs per node. Used by Step 1a auto-selection (cluster must have `gpus_per_node >= required_devices_per_node`) and as `--ntasks-per-node`. |\n\n**Procedure:**\n\n### 4a. Create user root directory\n\n```bash\n<ssh_cmd> <ssh_host> \"mkdir -p <default_user_root_dir>\"\n```\n\nIf `default_user_root_dir` is null or empty, skip this step and use `remote_cwd` directly.\n\n### 4b. Clone or update the repository\n\nDetermine the repo directory name from `repo_url` (strip `.git` suffix if present):\n\n```bash\nREPO_DIR_NAME=$(basename <repo_url> .git)\nREMOTE_REPO_PATH=\"<default_user_root_dir>/${REPO_DIR_NAME}\"\n```\n\nCheck if it already exists:\n\n```bash\n<ssh_cmd> <ssh_host> \"test -d <remote_repo_path>/.git && echo exists\"\n```\n\nIf the repository does **not** exist, clone it:\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <default_user_root_dir> && git clone <repo_url>\"\n```\n\nIf the repository **already exists**, refresh remote refs only — do **not**\ncheckout or reset to any branch here:\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <remote_repo_path> && git fetch origin --prune\"\n```\n\nThe actual ref the remote will check out is the temp branch created in 4c\nand pushed in 4d; the remote checkout happens in 4e. **Do not** add a\n`git checkout <repo_branch>` or `git reset --hard origin/<repo_branch>`\nhere as a shortcut — that path is forbidden (see the MANDATORY callout\nat the top of Step 4).\n\n### 4c. Create a temporary local branch capturing all local changes\n\nGenerate a unique temp branch name tied to this session. The branch name\n**must start with `<user_name>_<model_name>`** so it is easy to identify\nthe owner and the workload it was created for. A timestamp suffix keeps\nit unique across sessions:\n\n```bash\n# Resolve the user name (prefer the slurm_user input over whoami,\n# so the temp branch is attributed to the cluster user rather than the\n# host that happens to dispatch the job).\nUSER_NAME=\"${slurm_user:-$(whoami)}\"\n\n# Resolve the workload name. job_spec.json's `model_name` is authoritative;\n# fall back to the repo basename if it is missing.\nMODEL_NAME=\"${model_name:-$(basename \"<REPO_ROOT>\")}\"\n\n# Sanitize to a git-ref-safe identifier (alphanum + `_`/`-` only).\n# Any other character is replaced with `_`.\nsanitize() { printf '%s' \"$1\" | tr -c 'A-Za-z0-9_-' '_'; }\nUSER_NAME_SAFE=$(sanitize \"$USER_NAME\")\nMODEL_NAME_SAFE=$(sanitize \"$MODEL_NAME\")\n\nTEMP_BRANCH=\"${USER_NAME_SAFE}_${MODEL_NAME_SAFE}_$(date +%Y%m%d%H%M%S)\"\n```\n\nExample resolved value: `<user_name>_dwdp_deepseek_v3_lite_20260430145200`.\n\nStash the current branch name so it can be restored later:\n\n```bash\nORIGINAL_BRANCH=$(git -C \"<REPO_ROOT>\" rev-parse --abbrev-ref HEAD)\n```\n\nCreate the temp branch at HEAD and stage tracked changes only\n(`-u`, not `-A` — untracked files ship in 4e.1):\n\n```bash\ngit -C \"<REPO_ROOT>\" checkout -b \"$TEMP_BRANCH\"\ngit -C \"<REPO_ROOT>\" add -u\ngit -C \"<REPO_ROOT>\" commit --allow-empty -m \"sync: temp snapshot for remote execution\"\n```\n\n`--allow-empty` covers the no-tracked-changes case; the branch still exists\nat HEAD and 4e.1 still runs.\n\n### 4d. Push temp branch to the remote origin\n\nPush the temp branch so the remote cluster can fetch it:\n\n```bash\ngit -C \"<REPO_ROOT>\" push origin \"$TEMP_BRANCH\"\n```\n\n### 4e. Fetch and check out temp branch on the remote cluster\n\nOn the remote: stash prior-session leftovers, fetch, force-checkout the temp\nbranch, and re-materialize `cpp/` from HEAD (stale `cpp/` produces\nmissing-symbol `ImportError`s in Step 5):\n\n```bash\n<ssh_cmd> <ssh_host> \"cd <remote_repo_path> && \\\n    git stash push -u -m 'auto-stash before remote sync' || true && \\\n    git fetch origin --prune && \\\n    git checkout -B $TEMP_BRANCH origin/$TEMP_BRANCH && \\\n    rm -rf cpp/ && \\\n    git checkout HEAD -- cpp/\"\n```\n\n- `stash -u` sweeps both tracked edits and untracked files left from a prior\n  session; `|| true` covers the no-op.\n- `checkout -B` force-resets even if a same-named branch already exists.\n- `rm -rf cpp/ && git checkout HEAD -- cpp/` always runs, even when checkout\n  reports no diffs.\n\n`remote_cwd` is updated to `<remote_repo_path>`. Untracked files for this\nsession arrive next in 4e.1.\n\n### 4e.1. Transfer untracked files via tar pipe over SSH\n\nShip untracked, non-gitignored files on top of the freshly checked-out tree:\n\n```bash\nUNTRACKED_COUNT=$(git -C \"<REPO_ROOT>\" ls-files --others --exclude-standard | wc -l)\nif [ \"$UNTRACKED_COUNT\" -gt 0 ]; then\n    git -C \"<REPO_ROOT>\" ls-files --others --exclude-standard -z | \\\n        tar -C \"<REPO_ROOT>\" --null -T - -cf - | \\\n        <ssh_cmd> <ssh_host> \"tar -C <remote_repo_path> -xf -\"\nfi\n```\n\n`-z` / `--null` keeps filenames with spaces or newlines intact; `tar -x`\nauto-creates parent dirs; the guard skips the round-trip when nothing to\nship. Do **not** substitute `rsync` or `sshfs` (not reliably available on HPC).\n\n### 4f. Delete the temp branch (local and remote)\n\nAfter the remote checkout succeeds, immediately delete the temp branch from\nthe local repo and the remote origin, then restore the original local branch:\n\n```bash\ngit -C \"<REPO_ROOT>\" checkout \"$ORIGINAL_BRANCH\"\ngit -C \"<REPO_ROOT>\" branch -D \"$TEMP_BRANCH\"\ngit -C \"<REPO_ROOT>\" push origin --delete \"$TEMP_BRANCH\"\n```\n\nThe remote cluster keeps the checked-out working tree from the deleted branch\n— only the ref is removed, not the files.\n\n### 4f.1 Create remote work_dir under the cloned repo\n\nThe remote work_dir must mirror the local layout (`<REPO_ROOT>/work_dirs/<name>`)\nso that all per-run artefacts — slurm script, stdout/stderr logs, intermediate\nfiles — live inside the cloned repo and are co-located with the code that\nproduced them.\n\n`remote_work_dir` is precomputed by `trtllm-case-executor` (Step 1) and read\nverbatim from `job_spec.json` — it is `<remote_repo_path>/work_dirs/<basename(local_work_dir)>`.\n**Do not** fall back to `<remote_cwd>/work_dirs/...` or any other root; that\nolder layout is deprecated.\n\n```bash\n<ssh_cmd> <ssh_host> \"mkdir -p <remote_work_dir>\"\n```\n\nAll subsequent steps (file uploads, sbatch submission, log paths) reference\n`<remote_work_dir>` from `job_spec.json`.\n\n### 4g. Copy build scripts to remote node\n\nThe skill ships two build scripts in `scripts/` alongside this SKILL.md.\nCopy them to a dedicated directory on the remote so they are accessible\nfrom inside the SLURM container:\n\n```bash\nREMOTE_SCRIPTS_DIR=\"<default_user_root_dir>/trtllm-build-scripts\"\nSKILL_SCRIPTS_DIR=\"<LOCAL_SKILL_DIR>/scripts\"   # local path to this skill's scripts/\n\n<ssh_cmd> <ssh_host> \"mkdir -p ${REMOTE_SCRIPTS_DIR}/log\"\n\n<ssh_cmd> <ssh_host> \"cat > ${REMOTE_SCRIPTS_DIR}/build.slurm\" < \"${SKILL_SCRIPTS_DIR}/build.slurm\"\n<ssh_cmd> <ssh_host> \"cat > ${REMOTE_SCRIPTS_DIR}/build.sh\"    < \"${SKILL_SCRIPTS_DIR}/build.sh\"\n<ssh_cmd> <ssh_host> \"chmod +x ${REMOTE_SCRIPTS_DIR}/build.slurm ${REMOTE_SCRIPTS_DIR}/build.sh\"\n```\n\n`<LOCAL_SKILL_DIR>` is the directory containing this SKILL.md file\n(`.claude/skills/exec-remote-slurm/`).\n\n---\n\n## Step 5: Build on Remote SLURM Node\n\n> **Mandatory unless `build_project: false`.** Step 4e wipes `cpp/`, so\n> the remote has no compiled C++ extensions until Step 5 succeeds (exit\n> code `0:0`). Submitting any test/bench job before then fails on the\n> first `import tensorrt_llm` with missing-symbol errors.\n>\n> Do **not** replace this with `pip install -e .` inside the user's\n> `.slurm` script — it does not reliably recompile `cpp/`, and when it\n> does it blocks the GPU allocation behind a 30+ minute compile. Step 5\n> runs the build as a 1-node side job and installs into the persistent\n> container volume.\n>\n> If the caller sets `build_project: false`, skip Step 5 (the caller is\n> asserting the install is current). Do not silently downgrade `true` →\n> `false` because Step 5 looks expensive.\n\nSubmit a SLURM batch job that compiles TensorRT-LLM on the remote cluster\ninside the official container. Wait for it to complete before proceeding to\njob submission.\n\n### 5a. Read the container image\n\nThe container image is resolved by `trtllm-case-executor` (Step 3) and\nforwarded in `job_spec.json` as `container_image`. Read it from there —\n**do not grep `current_image_tags.properties` and do not re-derive from\nthe cluster config**. case-executor is the single source of truth for the\nimage URI; this skill only consumes it.\n\nApply the enroot URI rewrite at use-time so the registry path is not\nmistaken for a Docker-style credential lookup. The rewrite recipe lives\nin the `internal-env-info` skill's `Default Environment Values` section\n(`enroot_uri_rewrite` key):\n\n- **Skill installed** → delegate to it, passing the canonical\n  `<container_image>` from `job_spec.json`. The skill returns the bash\n  snippet (an `sed` over the registry-host prefix); run it here to\n  produce `<IMAGE_PATH>`.\n- **Skill not installed** → use `<container_image>` verbatim as `<IMAGE_PATH>`.\n\nUse `<IMAGE_PATH>` as the value passed to `--container-image=` in the\nbuild sbatch below and in every `srun`/`sbatch` recipe later in this\nskill. The canonical (un-rewritten) URI in `job_spec.json` is unchanged;\nthe rewrite is purely a transport adaptation for enroot.\n\n### 5b. Derive build architecture\n\nMap `device_type` (from the env-check output) to the `build_wheel.py` `-a` flag:\n\n| `device_type` | `-a` value |\n|------------|------------|\n| B200, GB200, GB300, B300 | `100-real` |\n| H100, H200 | `90-real` |\n| L40S | `89-real` |\n| A100 | `80-real` |\n\nDefault to `100-real` if `device_type` is unknown.\n\n### 5c. Submit the build job\n\n```bash\n<ssh_cmd> <ssh_host> \"sbatch \\\n    --nodes=1 --ntasks=1 --ntasks-per-node=1 \\\n    --partition=<partition> \\\n    --account=<account> \\\n    --job-name=<account>-trtllm.build \\\n    --time=02:00:00 \\\n    ${REMOTE_SCRIPTS_DIR}/build.slurm \\\n    <IMAGE_PATH> <default_user_root_dir> ${REMOTE_SCRIPTS_DIR} <remote_repo_path> \\\n    -a \\\"<arch>\\\" --nvtx\"\n```\n\nCapture the job ID from `sbatch` stdout (format: `Submitted batch job <id>`).\n\nIf `gpus_per_node` is set in the cluster config, add\n`--gpus-per-node=<gpus_per_node>` to the `sbatch` flags.\n\n### 5d. Monitor the build job\n\nPoll every 30–60 seconds until the job leaves the queue:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <job_id> -o '%.18i %.9P %.30j %.8u %.2t %.10M %R'\"\n```\n\nOnce running, tail the log to track compile progress:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -30 ${REMOTE_SCRIPTS_DIR}/log/build_<job_id>.srun.log\"\n```\n\nProgress indicators in the log:\n- `[XX%] Building CXX object ...` — compilation in progress\n- `Linking CXX ...` — link phase\n- `FAILED:` / `error:` / `fatal error:` — build failure\n- `Successfully built` — success\n\nIf the job disappears from `squeue`, proceed to Step 5e.\n\n### 5e. Verify the build\n\n```bash\n<ssh_cmd> <ssh_host> \"sacct -j <job_id> --format=JobID,State,ExitCode,Elapsed -n\"\n<ssh_cmd> <ssh_host> \"tail -50 ${REMOTE_SCRIPTS_DIR}/log/build_<job_id>.srun.log\"\n```\n\nA successful build shows exit code `0:0` and ends with\n`Successfully built tensorrt_llm` or no error lines.\n\nIf the build fails, report the last 100 lines of the log to the user and\nstop — do not proceed to job submission.\n\n## Build Scripts Reference\n\n| Script | Location | Purpose |\n|--------|----------|---------|\n| `build.slurm` | `scripts/build.slurm` | SLURM batch script — launches container, calls `build.sh` |\n| `build.sh` | `scripts/build.sh` | Runs inside container — calls `build_wheel.py` with defaults |\n\nBoth scripts accept the same positional arguments:\n`<container_image> <user_root_dir> <scripts_dir> <repo_dir> [build_wheel_args...]`\n\n---\n\n## Recipes\n\nPick the recipe that matches your workflow. Each recipe is self-contained —\ncopy it, substitute the `<placeholders>` from config, and run.\n\n**Placeholders** (all from `job_spec.json` unless noted):\n\n| Placeholder | Source |\n|-------------|--------|\n| `<ssh_host>`, `<account>`, `<partition>`, `<mounts>` | `job_spec.json` |\n| `<image>` | `container_image` from `job_spec.json` (resolved by `trtllm-case-executor` Step 3). Apply enroot rewrite (`/` → `#`) at use-time per Step 5a — substitute the rewritten `<IMAGE_PATH>` here. Do not re-resolve from the cluster config or `current_image_tags.properties`. |\n| `<job_name>` | Pre-resolved by `trtllm-case-executor` and forwarded as `job_name` in `job_spec.json`. Substitute directly into `-J <job_name>`; do not reconstruct from `account` / `subproject` / `detail`. |\n| `<nodes>` | `node_count` from `job_spec.json` (computed by `trtllm-case-executor`). Do not recompute. |\n| `<gpus>` | `required_devices_per_node` from `job_spec.json` — probe only if absent (see bottom of this section) |\n| `<remote_repo_path>` | `remote_repo_path` from `job_spec.json` — `<default_user_root_dir>/<repo_dir_name>` precomputed by `trtllm-case-executor` Step 1. The cloned repo and `--container-workdir` for srun. Do not recompute. |\n| `<remote_work_dir>` | `remote_work_dir` from `job_spec.json` — `<remote_repo_path>/work_dirs/<basename(local_work_dir)>` precomputed by `trtllm-case-executor` Step 1. Used for `--output`/`--error` slurm log paths and any per-run artefact uploads. Do not recompute or fall back to `<remote_cwd>/work_dirs/...`. |\n| `<time>` | Walltime in HH:MM:SS |\n\n**Job naming:** Every recipe includes `-J <job_name>` using the\npre-resolved `job_name` value from `job_spec.json` (built by\n`trtllm-case-executor`). This is required — without `-J`, SLURM emits\nwarnings and jobs are hard to identify in `squeue`. Do not reconstruct\nthe name from `account` / `subproject` / `detail` here; that convention\nlives in `trtllm-case-executor`.\n\n**GPU allocation:** Cluster behavior varies:\n\n- **Exclusive-access clusters** (most NVIDIA internal): Nodes are allocated\n  exclusively with all GPUs available. No `--gpus-per-node` needed.\n- **Shared/QOS clusters**: Require `--gpus-per-node=<gpus_per_node>`. Without\n  it, you may get a CPU-only node (no GPUs at all), a \"Cannot find GPU\n  specification\" error, or a \"QOSMinGRES\" error (below minimum).\n\nIf `gpus_per_node` is set in the cluster config, add `--gpus-per-node=<gpus_per_node>`\nto every `srun`, `salloc`, and `sbatch` command. If not set, omit it.\n\nIf a job fails with GPU-related errors and `gpus_per_node` is not configured,\nprobe the node's GPU count, set `gpus_per_node` in the cluster config, and retry.\n\nDo not set `CUDA_VISIBLE_DEVICES` for MPI workloads — the MPI runtime handles\nGPU assignment.\n\n---\n\n### Recipe 1: Single-process command (srun)\n\nOne-off commands on a single node: profiling, inspecting, lightweight scripts.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N 1 --ntasks-per-node=1 \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <command>\"\n```\n\n---\n\n### Recipe 2: MPI workload (srun)\n\nMPI applications across nodes: NCCL tests, HPCX benchmarks. The binary must\nbe MPI-aware (e.g., `all_reduce_perf_mpi`, not `all_reduce_perf`).\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=<gpus> --mpi=pmix \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <command>\"\n```\n\n---\n\n### Recipe 3: Distributed training (srun)\n\nFramework-launched distributed training (PyTorch DDP, DeepSpeed). SLURM starts\n`<nodes> × <gpus>` processes. Pyxis/enroot sets `RANK`, `LOCAL_RANK`, and\n`WORLD_SIZE`. Each process runs the training script directly.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=<gpus> \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  python train.py\"\n```\n\n---\n\n### Recipe 4: Launcher-based training (srun)\n\nOne process per node, an in-container launcher forks locally (torchrun,\ndeepspeed launcher).\n\n```bash\n<ssh_cmd> <ssh_host> \"srun -A <account> -p <partition> -N <nodes> \\\n  --ntasks-per-node=1 \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  torchrun --nproc-per-node=<gpus> train.py\"\n```\n\n---\n\n### Recipe 5: Batch job (sbatch)\n\nFire-and-forget background job:\n\n```bash\n<ssh_cmd> <ssh_host> \"sbatch -A <account> -p <partition> -N <nodes> \\\n  -J <job_name> -t <time> \\\n  --container-image=<image> --container-mounts=<mounts> \\\n  <script-path>\"\n```\n\nCheck status:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <jobid>\"\n<ssh_cmd> <ssh_host> \"sacct -j <jobid> --format=JobID,State,ExitCode,Elapsed -n\"\n```\n\n---\n\n### Recipe 6: Multi-command session (salloc + srun)\n\nInteractive development with persistent container state. Use this when you need\nmultiple commands on the same allocation — installing packages, compiling,\nrunning experiments iteratively.\n\n> **Relationship to `persistent_mode`:** The `persistent_mode` field in\n> `job_spec.json` is consumed only by `exec-local-slurm`. This\n> skill does **not** read `persistent_mode`, `container_name`, or\n> `container_mounts` from `job_spec.json` for allocation reuse — remote\n> persistence is driven by this recipe (tmux + named `salloc` + named\n> container) and controlled by inputs to this skill directly (e.g.,\n> `release_allocation`, `alloc_time_limit`). If a caller forwards a\n> `persistent_mode=true` job spec into a `remote_slurm` run, treat the field\n> as informational only and use this recipe's flow.\n\n**Step 1: Allocate via tmux**\n\ntmux keeps the allocation alive on the frontend — invisible plumbing.\n\n```bash\n<ssh_cmd> <ssh_host> \"tmux new-session -d -s trtllm-<id> \\\n  'salloc -A <account> -p <partition> -N <nodes> \\\n  -J <job_name> -t <time>'\"\n```\n\n**Step 2: Check allocation status**\n\nWait ~30 seconds for the scheduler, then get job ID and status:\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -u <user> -h -o '%i %T %S' --state=RUNNING,PENDING\"\n```\n\nIf RUNNING → proceed to Step 3.\n\nIf PENDING → follow the **walltime negotiation flow** (see Walltime strategy\nsection). Only proceed to Step 3 once the job is RUNNING.\n\n**Step 3: Create a persistent container**\n\nImport the image once. This creates a named container on each node that\npersists across `srun` calls (container persistence):\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-image=<image> --container-name=trtllm-<jobid> true\"\n```\n\n**Container packages:** If `packages` is defined in config, install them\nusing the exact names listed — do not guess or substitute package names:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> --container-writable \\\n  pip install <packages>\"\n```\n\nOnly install when the workload actually needs them — not for simple\ncommands like `nvidia-smi` or `hostname`.\n\n**Step 4: Run commands**\n\nUse `--container-name` (without `--container-image`) to reuse the persistent\ncontainer. All in-container state is preserved between commands.\n\nSingle-process command:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nPer-node operation (mkdir, pip install) — use `--ntasks-per-node=1`:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nMPI workload within the allocation:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --ntasks-per-node=<gpus> --mpi=pmix \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  <command>\"\n```\n\nDistributed training within the allocation:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> \\\n  --ntasks-per-node=<gpus> \\\n  --container-name=trtllm-<jobid> --container-mounts=<mounts> \\\n  python train.py\"\n```\n\nAdd `--container-writable` if the workload writes to container paths outside\nof mounted directories (e.g., `pip install` into site-packages).\n\n**Step 5: Release the allocation**\n\nClean up the named container, then kill tmux. If the allocation has already\nexpired (walltime exceeded), the `srun` will fail — that is fine, the container\nis already gone. Just kill the tmux session.\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  bash -c 'enroot remove -f trtllm-<jobid> 2>/dev/null; true'\" 2>/dev/null\n<ssh_cmd> <ssh_host> \"tmux kill-session -t trtllm-<id>\"\n```\n\n---\n\n### Probing GPUs per node\n\nIf you do not know `<gpus>`, probe once and record in cluster notes:\n\n```bash\n<ssh_cmd> <ssh_host> \"srun --jobid=<jobid> --ntasks-per-node=1 \\\n  --container-name=trtllm-<jobid> bash -c 'nvidia-smi -L | wc -l'\"\n```\n\n## Cluster-Specific Flags\n\nAfter substituting recipe placeholders, apply the cluster-specific flags\nreturned by the `internal-env-info` skill in Step 1:\n\n- **SHARP**: Add `--network=sharp` if the cluster supports it\n- **Topology constraints**: Add `--constraint=<feature>` for topology-aware\n  placement (e.g., `--constraint=nvlblk01` for a single rack)\n- **GPU clocks**: Run `sudo nvidia-smi -ac <mem>,<graphics>` as a separate\n  `srun --ntasks-per-node=1` pre-step in sbatch scripts\n- **Comments**: Add `--comment=metrics` for Prometheus/Grafana collection\n\n## File Transfer\n\nUse SSH pipes for file transfer — no extra tools needed, works on every cluster.\n\n### Local → Remote (writing files)\n\nPipe file content directly via SSH heredoc. Never write locally then scp:\n\n```bash\n<ssh_cmd> <ssh_host> \"cat > <remote_cwd>/script.py\" << 'PYEOF'\nimport torch\n# ... file content ...\nPYEOF\n```\n\nFor multiple files, repeat the pattern — each file is one SSH command.\n\n### Remote → Local (reading results)\n\nRead remote files via SSH stdout:\n\n```bash\n<ssh_cmd> <ssh_host> \"cat <remote_cwd>/output.json\"\n```\n\nFor large files, read only what you need:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/slurm-<jobid>.out\"\n<ssh_cmd> <ssh_host> \"head -n 50 <remote_cwd>/results.csv\"\n```\n\n### When to use scp\n\nFall back to `scp` only for binary files or large directories that cannot be\npiped through SSH:\n\n```bash\n<scp_cmd> <ssh_host>:<remote_cwd>/profile.nsys-rep .    # binary profile\n<scp_cmd> -r local_dir/ <ssh_host>:<remote_cwd>/        # whole directory\n```\n\n**Decision rule:** text files → SSH pipe. Binary files or directories → scp.\nNever use rsync or sshfs — not reliably available on HPC clusters.\n\n## Rules\n\n### Never read full remote files\n\nRemote output can be gigabytes. Always use `tail`, `head`, or `grep`:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/output.log\"\n<ssh_cmd> <ssh_host> \"head -n 50 <remote_cwd>/script.py\"\n<ssh_cmd> <ssh_host> \"grep -n 'error\\|Error\\|ERROR' <remote_cwd>/output.log\"\n```\n\nNever use `cat` on remote files unless you are certain the file is small.\n\n### Never run heavy computation on the frontend\n\nThe frontend node is shared. Only use it for SLURM commands, lightweight file\noperations, and tmux management.\n\n### Check before creating\n\nBefore creating resources, check if they already exist:\n- SSH ControlMaster: `ssh -O check <ssh_host> 2>&1 || <ssh_cmd> -O check <ssh_host>` (probe both the user's default socket and the trtllm-managed socket from Step 2)\n- tmux sessions: `<ssh_cmd> <ssh_host> \"tmux has-session -t trtllm-<id> 2>/dev/null && echo exists\"`\n- Existing allocations: `<ssh_cmd> <ssh_host> \"squeue -u <user> -h -o '%i %j' --state=RUNNING\"`\n\n### Clean up when done\n\n- Remove named containers: `enroot remove -f` (see Recipe 6, Step 5)\n- Kill tmux sessions: `<ssh_cmd> <ssh_host> \"tmux kill-session -t trtllm-<id>\"`\n- Cancel unneeded jobs: `<ssh_cmd> <ssh_host> \"scancel <jobid>\"`\n\n### Walltime strategy\n\n- **Production (sbatch, srun):** Tight walltime with ~20% margin — gets\n  scheduled sooner.\n- **Dev/debug (salloc):** Use **walltime negotiation** — submit with max\n  walltime, then reduce if the queue wait is too long:\n\n**Walltime negotiation flow:**\n\n1. Submit `salloc` with **max walltime** for the partition.\n2. Wait ~30 seconds for the scheduler to process, then check:\n   ```bash\n   <ssh_cmd> <ssh_host> \"squeue -j <jobid> -h -o '%i %T %S'\"\n   ```\n   - `%S` = estimated start time. Parse it to calculate wait.\n3. **If RUNNING** → proceed to next step.\n4. **If PENDING and wait ≤ 10 minutes** → wait for it, proceed.\n5. **If PENDING and wait > 10 minutes** (or no start time shown):\n   a. Infer how long you actually need from context (profiling → ~30 min,\n      iterative dev → ~2h, training run → match expected duration).\n      Round up to the next hour boundary.\n   b. Reduce walltime in-place:\n      ```bash\n      <ssh_cmd> <ssh_host> \"scontrol update JobId=<jobid> TimeLimit=<new_time>\"\n      ```\n   c. Tell the user: \"Reduced walltime from Xh to Yh for faster scheduling —\n      you'll have less time on the allocation.\"\n   d. Wait ~15 seconds, re-check estimated start time.\n   e. **If still > 10 minutes:** inform user with the ETA and suggest\n      alternatives — different partition (e.g., backfill), fewer nodes.\n\n## Hang Detection Monitoring\n\nWhen monitoring a running job (polling via `squeue`), also check the remote log\nfile every 60 seconds for the pattern `hang detected` (case-insensitive). This\ncatches GPU hangs, NCCL deadlocks, and other stalls that keep the job running\nindefinitely without making progress.\n\n**Check for hang (run alongside each `squeue` poll once job is RUNNING):**\n\n```bash\n<ssh_cmd> <ssh_host> \"grep -qi 'hang detected' <remote_work_dir>/slurm-<jobid>.out 2>/dev/null && echo HANG_DETECTED\"\n```\n\nAlso check the `.err` file if the output file has no match:\n\n```bash\n<ssh_cmd> <ssh_host> \"grep -qi 'hang detected' <remote_work_dir>/slurm-<jobid>.err 2>/dev/null && echo HANG_DETECTED\"\n```\n\n**If `HANG_DETECTED` is printed:**\n\n1. Immediately cancel the job:\n   ```bash\n   <ssh_cmd> <ssh_host> \"scancel <jobid>\"\n   ```\n2. Read the last 200 lines of the log file for context:\n   ```bash\n   <ssh_cmd> <ssh_host> \"tail -n 200 <remote_work_dir>/slurm-<jobid>.out\"\n   ```\n3. Report status as `HANG_DETECTED` with the relevant log lines. Do not\n   continue polling — the job is terminated.\n\n## Recovering from Failures\n\n**SSH connection drops:** Re-run SSH preflight (Step 2). tmux-held allocations\nsurvive disconnects.\n\n**Job stuck in PENDING:**\n\n```bash\n<ssh_cmd> <ssh_host> \"squeue -j <jobid> -o '%i %T %r %S'\"\n```\n\nCheck the reason (`%r`) and estimated start (`%S`). If wait > 10 minutes,\nfollow the **walltime negotiation flow** (see Walltime strategy). If already\nnegotiated and still stuck, suggest alternatives: backfill partition, fewer\nnodes, or a different cluster.\n\n**Job fails:** Read last lines, never dump full output:\n\n```bash\n<ssh_cmd> <ssh_host> \"tail -n 100 <remote_cwd>/slurm-<jobid>.out\"\n<ssh_cmd> <ssh_host> \"tail -n 50 <remote_cwd>/slurm-<jobid>.err\"\n```\n\n**Hang detected:** If the log file contains `hang detected`, the job is\nauto-cancelled (see Hang Detection Monitoring section above). Report\n`HANG_DETECTED` with the last 200 lines of log context.\n\n<!-- File: .claude/skills/exec-slurm-compile/SKILL.md (Tokens: ~2733 | Category: anthropic-skill) -->\n---\nname: exec-slurm-compile\ndescription: Compile TensorRT-LLM on a SLURM cluster. Covers submitting a batch job with a container image, monitoring the job, and verifying the build. Use when the user wants to compile TRT-LLM remotely via SLURM rather than on a local compute node.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Compile TensorRT-LLM on SLURM Cluster\n\nSubmit, monitor, and verify a TensorRT-LLM compilation job on a SLURM cluster using enroot containers.\n\n## When to Use\n\n| Scenario | Use This Skill? |\n|----------|----------------|\n| User wants to compile TRT-LLM on a SLURM cluster | Yes |\n| User is already on a compute node and wants to compile | No — use `exec-local-compile` skill instead |\n\n## Finding the Docker Image\n\nThe official Docker image tag for a given TensorRT-LLM version is recorded in the repo itself:\n\n```\n<repo_dir>/jenkins/current_image_tags.properties\n```\n\nRead this file to find the current image URL. The format is `<registry-host>/<namespace>/tensorrt-llm:<tag>`, for example `<your-registry>/<your-namespace>/tensorrt-llm:pytorch-25.12-py3-aarch64-ubuntu24.04-trt10.14.1.48-skip-tritondevel-202602011118-10901`. Substitute `<your-registry>` and `<your-namespace>` with your own container registry coordinates.\n\n\n## Pre-dumping the Container Image (enroot import)\n\nSLURM clusters using enroot/pyxis require a `.sqsh` container image. To avoid download overhead at compile time, **pre-dump the image in advance** using the `enroot-import` companion script:\n\n```bash\n# Basic usage — submits a SLURM job on a CPU partition to import the image\nenroot-import --partition cpu_datamover --debug <docker_image_url>\n```\n\nThe script submits an `sbatch` job that runs `enroot import docker://<image_url>` and produces a `.sqsh` file in the current directory. The output on stdout is the SLURM job ID.\n\n### enroot-import flags\n\n| Flag | Description |\n|------|-------------|\n| `-p, --partition` | SLURM partition for the import job (use a CPU partition like `cpu_datamover`) |\n| `-d, --debug` | Enable debug output and preserve the SLURM log (recommended) |\n| `-o, --output` | Custom output path for the `.sqsh` file |\n| `-A, --account` | SLURM account (defaults to user's first account) |\n| `-t, --time` | Time limit for the import job (default: 1 hour) |\n| `-n, --just-print` | Print the sbatch command without executing |\n| `-J, --job-name` | Custom job name |\n\n### enroot-import workflow\n\n1. Read the image tag from `jenkins/current_image_tags.properties` in the TRT-LLM repo.\n2. Run `enroot-import` to submit the import job:\n   ```bash\n   cd <directory_where_sqsh_should_be_stored>\n   <path_to>/enroot-import --partition cpu_datamover --debug <image_url>\n   ```\n   **IMPORTANT:** Replace the first `/` after the registry host with `#` to avoid credential issues. For example, `<your-registry>/<your-namespace>/tensorrt-llm:xxx` becomes `<your-registry>#<your-namespace>/tensorrt-llm:xxx`.\n3. Wait for the import job to complete (`squeue -j <job_id>`).\n4. The resulting `.sqsh` file is the `container_image` used in the compile step.\n\n\n## Prerequisites\n\nThe user must provide (or you must ask for) these values:\n\n| Parameter | Description | Example |\n|-----------|-------------|---------|\n| `container_image` | Path to `.sqsh` container image (see enroot import above) | `/path/to/pytorch.sqsh` |\n| `repo_dir` | Path to the TensorRT-LLM repository | `/path/to/TensorRT-LLM` |\n| `user_root_dir` | Top-level directory to bind-mount into the container | `/shared/users` |\n| `partition` | SLURM partition | `batch` |\n| `account` | SLURM account | `my_account` |\n\nOptional parameters:\n\n| Parameter | Description | Default |\n|-----------|-------------|---------|\n| `jobname` | SLURM job name | `trtllm-compile.<username>` |\n| `gpu_count` | Number of GPUs to request | `4` |\n| `time_limit` | Job time limit | `02:00:00` |\n| `arch` | GPU architecture(s) for `-a` flag | `100-real` |\n| `extra_build_args` | Extra flags for `build_wheel.py` | (none) |\n\n## Companion Scripts\n\nThis skill includes three companion scripts in `scripts/`:\n\n| Script | Purpose |\n|--------|---------|\n| `enroot-import` | Pre-dump a Docker image to `.sqsh` via a SLURM batch job |\n| `submit_compile.sh` | Template for submitting the SLURM job — copy and customize |\n| `compile.slurm` | SLURM batch script — launches the container and calls `compile.sh` |\n| `compile.sh` | Runs inside the container — executes `build_wheel.py` |\n\nScripts directory: `skills/exec-slurm-compile/scripts/`\n\n## Instructions\n\nFollow these steps in order:\n\n### Step 0: Resolve the Container Image (if needed)\n\nIf the user does not already have a `.sqsh` container image:\n\n1. Read the Docker image tag from `<repo_dir>/jenkins/current_image_tags.properties`.\n2. Use `enroot-import` to pre-dump it:\n   ```bash\n   cd <directory_for_sqsh_files>\n   <scripts_dir>/enroot-import --partition cpu_datamover --debug <image_url>\n   ```\n3. Monitor the import job with `squeue -j <job_id>`.\n4. Once complete, the `.sqsh` file path becomes the `container_image` parameter.\n\nIf the user already has a `.sqsh` file, skip this step.\n\n### Step 1: Gather Information\n\nAsk the user for any missing prerequisite values listed above. At minimum you need:\n- `container_image` (or the Docker image URL — then run Step 0 first)\n- `repo_dir`\n- `user_root_dir`\n- `partition` and `account`\n\nIf the user has used this workflow before, check if previous values are stored in memory files.\n\n### Step 2: Prepare the Scripts Directory\n\nThe compile scripts must be accessible from inside the container (i.e., under `user_root_dir`). Either:\n\n**Option A** — Copy companion scripts to a location under `user_root_dir`:\n```bash\nscripts_dir=<user_root_dir>/<username>/workspace/tensorrt_llm_scripts\nmkdir -p ${scripts_dir}/log\ncp skills/exec-slurm-compile/scripts/compile.sh ${scripts_dir}/\ncp skills/exec-slurm-compile/scripts/compile.slurm ${scripts_dir}/\nchmod +x ${scripts_dir}/compile.sh ${scripts_dir}/compile.slurm\n```\n\n**Option B** — If the user already has scripts at a known location, use those directly.\n\n### Step 3: Submit the Job\n\nRun `sbatch` from the login node (or a node with SLURM client access):\n\n```bash\nsbatch \\\n    --nodes=1 --ntasks=1 --ntasks-per-node=1 \\\n    --gres=gpu:<gpu_count> \\\n    --partition=<partition> \\\n    --account=<account> \\\n    --job-name=<jobname> \\\n    --time=<time_limit> \\\n    <scripts_dir>/compile.slurm \\\n    <container_image> <user_root_dir> <scripts_dir> <repo_dir>\n```\n\nCapture and report the job ID from the `sbatch` output.\n\n### Step 4: Monitor the Job (Proactive — Do NOT Wait for User)\n\n**You MUST actively poll the job until it completes.** Do not submit and walk away.\n\n```bash\n# Check job status (repeat every 30-60 seconds)\nsqueue -j <job_id> -o \"%.18i %.9P %.30j %.8u %.2t %.10M %.6D %R\"\n\n# Once running, periodically tail the log (do NOT use tail -f, use tail -30 instead)\ntail -30 <scripts_dir>/log/compile_<job_id>.srun.log\n```\n\n**Monitoring loop:**\n1. Poll `squeue -j <job_id>` to check state\n2. If `PD` (pending) — report the reason, keep polling every 30-60s\n3. If `R` (running) — tail the build log every 30-60s; look for `[XX%] Building`, errors, or completion\n4. If the job disappears from `squeue`, it has finished — proceed to Step 5\n5. If `F` (failed) — immediately read the full log and report the error\n\n**Progress indicators to look for in the log:**\n- `[XX%] Building CXX object...` — compilation progress\n- `Linking CXX...` — link phase\n- `FAILED:`, `error:`, `fatal error:` — build failure\n- `Successfully built` — success\n\n### Step 5: Verify the Build\n\nOnce the job completes, check for success:\n\n```bash\n# Check SLURM exit code\nsacct -j <job_id> --format=JobID,State,ExitCode,Elapsed\n\n# Check the build log for errors\ntail -50 <scripts_dir>/log/compile_<job_id>.srun.log\n```\n\nA successful build ends with a message like `Successfully built tensorrt_llm` or completes without error.\n\n## Common Build Flags Reference\n\n| Flag | Description |\n|------|-------------|\n| `-a \"100-real\"` | Target architecture — `100` for Blackwell, `90` for Hopper, etc. |\n| `--nvtx` | Enable NVTX markers for profiling |\n| `--no-venv` | Skip virtual environment creation |\n| `--use_ccache` | Use ccache to speed up recompilation |\n| `--skip_building_wheel` | Build in-place without creating a wheel file |\n| `-f` | Fast build — skip some kernels for faster dev compilation |\n| `-c` | Clean build — wipe build directory before building |\n\nCommon architecture values:\n- `\"100-real\"` — Blackwell (B200, GB200)\n- `\"90-real\"` — Hopper (H100, H200)\n- `\"89-real\"` — Ada Lovelace (L40S)\n- `\"80-real\"` — Ampere (A100)\n- `\"90;100-real\"` — Multiple architectures\n\n## Troubleshooting\n\n| Issue | Solution |\n|-------|----------|\n| `sbatch: error: invalid partition` | Verify partition name with `sinfo -s` |\n| `sbatch: error: invalid account` | Check available accounts with `sacctmgr show assoc user=$USER` |\n| Container image not found | Verify the `.sqsh` path exists and is readable |\n| Build OOM (out of memory) | Reduce parallelism with `-j <N>` flag to `build_wheel.py` |\n| `srun: error: Unable to create step` | The node may lack enroot/pyxis — check with cluster admin |\n| Job stuck in `PD` state | Check `squeue -j <id> -o %R` for the reason (e.g., resource limits, priority) |\n| `enroot import` fails with auth error | Check `~/.config/enroot/.credentials` has the correct registry credentials |\n| `enroot import` produces empty/corrupt `.sqsh` | Re-run with `--debug` and check the SLURM log; verify the image URL has no `https://` prefix |\n| Weird compile issues | Retry with a clean build (`-c` flag) |\n| `QOSGrpNodeLimit` shown in `NODELIST(REASON)` | Not a blocker, just wait for the job to get scheduled |\n\n## Example Interaction\n\n**User**: \"Compile TRT-LLM on the OCI cluster\"\n\n**Agent actions**:\n1. Ask for container image path, repo path, mount dir (if not known)\n2. Confirm partition/account for OCI cluster\n3. Copy scripts to accessible location under user_root_dir\n4. Submit with `sbatch`\n5. Report job ID\n6. Monitor with `squeue` until complete\n7. Check logs and report success/failure\n\n<!-- File: .claude/skills/kernel-cute-writing/references/api-arch.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-arch.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-arch.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/api-core.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-core.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-core.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/api-nvgpu.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-nvgpu.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-nvgpu.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/api-runtime-utils.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/api-runtime-utils.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/api-runtime-utils.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/concepts-architecture.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-architecture.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-architecture.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/concepts-layouts.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-layouts.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-layouts.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/concepts-mma.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-mma.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-mma.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/concepts-tensors.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/concepts-tensors.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/concepts-tensors.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/patterns-compilation.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-compilation.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-compilation.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/patterns-elementwise.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-elementwise.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-elementwise.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/patterns-gemm.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-gemm.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-gemm.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/patterns-getting-started.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-getting-started.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-getting-started.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/patterns-memory.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-memory.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-memory.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/patterns-pipeline.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-pipeline.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-pipeline.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/patterns-reduction.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/patterns-reduction.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/patterns-reduction.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/references/troubleshooting.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-cute-writing/references/troubleshooting.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-cute-writing/references/troubleshooting.md)\n\n<!-- File: .claude/skills/kernel-cute-writing/SKILL.md (Tokens: ~4584 | Category: anthropic-skill) -->\n---\nname: kernel-cute-writing\ndescription: >\n  Write and implement GPU kernels using NVIDIA CuTe DSL (CUTLASS 4.x Python\n  API) — NOT for Triton, CUDA C++, or conceptual explanations.\n  Trigger only when the user wants to write or implement a kernel, not when\n  asking questions about CuTe DSL concepts or layouts.\n  CuTe DSL uses cute.jit/cute.kernel decorators and cutlass.cute imports.\n  Covers element-wise kernels, GEMM patterns, reductions, memory hierarchy\n  (global/shared/register/TMA), MMA tensor core operations, software\n  pipelining, and framework integration.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# CuTe DSL\n\nCuTe DSL is a Python-based domain-specific language for GPU kernel development,\npart of CUTLASS 4.x. It provides Python abstractions over CUTLASS C++ templates\nwith JIT compilation to optimized CUDA kernels via MLIR and ptxas.\n\n## When to Use\n\n**Triggers:**\n- Writing CUDA kernels in Python (element-wise, GEMM, custom ops)\n- Optimizing GPU memory access patterns (vectorized loads, TMA, shared memory)\n- Building tensor core (MMA) kernels for Ampere/Hopper/Blackwell\n- Integrating custom GPU kernels with PyTorch or JAX\n- Prototyping high-performance kernels without C++ metaprogramming\n\n**Symptoms (wrong tool otherwise):**\n- Need shared memory coordination or tensor core MMA → use CuTe DSL (not Triton for complex patterns)\n- Need simple element-wise ops with no shared memory → CuTe DSL or Triton both work\n- Need to call existing CUTLASS C++ kernels → use CUTLASS C++ APIs instead\n- Need reductions, scans, or non-GEMM collective ops → consider CUB/Thrust\n\n**Keywords:** cute, cutlass, cute.jit, cute.kernel, from_dlpack, zipped_divide,\nTiledMMA, TiledCopy, TMA, WGMMA, tcgen05, pipeline, mbarrier\n\n## Requirements\n\n| Requirement | Detail |\n|-------------|--------|\n| Platform | Linux x86_64 only |\n| Python | 3.10–3.13 |\n| GPU | NVIDIA Ampere+ (SM80, SM90, SM100) |\n| CUDA Driver | ≥ 575.51.03 (Toolkit 12.9 compat) |\n| Install | `pip install nvidia-cutlass-dsl` |\n| Optional | `apache-tvm-ffi`, `torch-c-dlpack-ext` |\n\n## Workflows\n\n### Workflow 0: Starting from Examples (Recommended)\n\nFor any non-trivial kernel (GEMM, attention, pipelined, fused ops), start by\nfinding the most similar existing example to use as a **starting point** — study\nits structure, then rework it for your use case. Do not copy examples verbatim;\nthey target specific dtypes, architectures, and problem shapes that likely differ.\n\n1. **Pick the closest example** from the index below.\n   **Prefer examples matching the target GPU architecture** (check with\n   `torch.cuda.get_device_capability()`) when the operation is similar.\n\n   Fetch via `web_fetch` with base URL\n   `https://raw.githubusercontent.com/NVIDIA/cutlass/main/examples/python/CuTeDSL`\n\n   | Operation | Arch | Example path (append to base URL) |\n   |-----------|------|-----------------------------------|\n   | Element-wise add | SM80 | `ampere/elementwise_add.py` |\n   | Element-wise + autotune | SM80 | `ampere/elementwise_add_autotune.py` |\n   | Element-wise apply | SM80 | `ampere/elementwise_apply.py` |\n   | SGEMM (scalar) | SM80 | `ampere/sgemm.py` |\n   | Tensor-core GEMM | SM80 | `ampere/tensorop_gemm.py` |\n   | Flash Attention v2 | SM80 | `ampere/flash_attention_v2.py` |\n   | HSTU Attention | SM80 | `ampere/hstu_attention.py` |\n   | Shared memory allocator | SM80 | `ampere/smem_allocator.py` |\n   | CTA norm (LayerNorm) | SM90 | `hopper/cta_norm.py` |\n   | Dense GEMM | SM90 | `hopper/dense_gemm.py` |\n   | Dense GEMM persistent | SM90 | `hopper/dense_gemm_persistent.py` |\n   | Flash MHA | SM90 | `hopper/fmha.py` |\n   | Dense GEMM | SM100 | `blackwell/dense_gemm.py` |\n   | Dense GEMM persistent | SM100 | `blackwell/dense_gemm_persistent.py` |\n   | Dense GEMM + alpha/beta | SM100 | `blackwell/dense_gemm_alpha_beta_persistent.py` |\n   | RMSNorm | SM100 | `blackwell/rmsnorm.py` |\n   | Reduce | SM100 | `blackwell/reduce.py` |\n   | Flash MHA | SM100 | `blackwell/fmha.py` |\n   | Grouped GEMM | SM100 | `blackwell/grouped_gemm.py` |\n   | Mamba2 SSD | SM100 | `blackwell/mamba2_ssd/` |\n   | GEMM tutorial (notebook) | SM100 | `notebooks/tour_to_sol_gemm.ipynb` |\n\n   **Example:** To fetch the Hopper dense GEMM:\n   ```bash\n   web_fetch https://raw.githubusercontent.com/NVIDIA/cutlass/main/examples/python/CuTeDSL/hopper/dense_gemm.py\n   ```\n\n2. **Read reference materials first** — before diving into example code, read\n   the relevant `references/` docs to understand the patterns and APIs:\n   - For GEMM: `references/patterns-gemm.md` (3-level tiling, epilogue fusion,\n     `cute.compile` with `mark_layout_dynamic`, shared memory layouts)\n   - For reductions: `references/patterns-reduction.md` (warp reductions,\n     `cute.compile` cache pattern)\n   - For element-wise: `references/patterns-elementwise.md` (variations A–E)\n   - Always: `references/api-arch.md` (available APIs, arch-specific caveats)\n\n   This gives you the conceptual foundation so you can rework the example\n   intelligently rather than trying to copy-paste complex pipelines.\n\n3. **Fetch and study the example source** — read for structure, not to copy:\n   - Identify: decorators, tiling strategy, shared memory usage, mainloop flow\n   - Note which dtype/arch it targets (many examples are fp16/bf16-specific)\n   - Check if it uses APIs tied to a specific arch (TMA → SM90+, tcgen05 → SM100)\n\n4. **Rework for the user's workload** (do not copy-paste):\n   - Change shapes, data types, tile sizes to match requirements\n   - Replace compute logic (epilogue, activation fusion) as needed\n   - If dtype differs (e.g., example is fp16, need fp32), expect vectorization\n     and layout changes — the scalar-loop patterns in `references/` may be a\n     better starting point than adapting a vectorized example\n   - **Runtime wrapper must be lightweight**: `kernel_fn()` should only call\n     `from_dlpack()` + the compiled kernel. Never allocate intermediate tensors,\n     copy data, or re-compile per call — these belong in one-time setup\n   - Apply optimizations from this skill's reference docs\n\n   **⛔ Blackwell/Hopper GEMM + extra tensors — STOP:**\n   If the target GPU is SM90+ (Hopper/Blackwell) **and** the GEMM requires\n   extra tensors beyond A, B, C in the epilogue (e.g., bias vector, activation\n   inputs), **do not attempt it**. These examples use TMA descriptors for all\n   data movement — adding tensors requires modifying TMA descriptor setup,\n   which is prohibitively complex. Instead, tell the user this limitation and\n   suggest a **two-kernel approach**: run the GEMM kernel as-is, then apply\n   bias + activation in a separate element-wise kernel (Workflow 1).\n   Plain GEMM (just A×B→C with scalar alpha/beta) on Hopper/Blackwell is fine.\n\n5. **Validate and benchmark** using companion scripts:\n   ```bash\n   python scripts/verify_kernel.py kernel.py --rtol 1e-3 --atol 1e-3\n   python scripts/benchmark_kernel.py kernel.py\n   ```\n   The kernel file must export `kernel_fn`, `reference_fn`, and `get_inputs()`.\n\n**When to skip examples:** Pure element-wise operations (Workflow 1) have\ncomplete patterns in `references/patterns-elementwise.md` — no need to fetch\nexternal examples.\n\n**Reduction kernels** (softmax, layernorm, RMSNorm): Use\n`references/patterns-reduction.md` which provides complete, proven patterns\nfor float32 reductions using scalar loops + butterfly shuffle + shared memory.\n\n### Workflow 1: Element-wise Kernel\n\nFor unary/binary/in-place operations that map inputs to outputs 1:1.\n\n1. **Determine kernel structure**: inputs/outputs count, tensor rank, target arch\n2. **Select pattern** from `references/patterns-elementwise.md` (Variations A–E)\n3. **Write kernel** applying all four invariant principles:\n   - P1: `from_dlpack(tensor, assumed_align=16)` for vector loads\n   - P2: Derive `vec_size` from `element_type.width`\n   - P3: `cute.zipped_divide(mA, tiler)` for coalesced access\n   - P4: `cutlass.dynamic_expr(thread_idx < total)` for bounds\n4. **Critical rules**: No early return, no `a * 2` (use `a + a`), no `cute.math.sigmoid`\n5. **Pre-compile with `cute.compile()`**: Always pre-compile the kernel once\n   using `cute.compile()` so that `kernel_fn` calls the compiled object, not\n   `@cute.jit` directly.  Without pre-compilation, every call recompiles\n   (~20-50ms overhead).  Use `.mark_layout_dynamic()` so a single compiled\n   kernel handles arbitrary input shapes without recompilation:\n   ```python\n   # Compile once with dynamic layouts — works for any shape\n   fake_x = from_dlpack(torch.empty(1, 1, dtype=torch.float16, device=\"cuda\"),\n                         assumed_align=16).mark_layout_dynamic()\n   fake_out = from_dlpack(torch.empty(1, 1, dtype=torch.float16, device=\"cuda\"),\n                           assumed_align=16).mark_layout_dynamic()\n   compiled_kernel = cute.compile(host_fn, fake_x, fake_out)\n\n   def kernel_fn(x):\n       out = torch.empty_like(x)\n       compiled_kernel(from_dlpack(x, assumed_align=16).mark_layout_dynamic(),\n                       from_dlpack(out, assumed_align=16).mark_layout_dynamic())\n       return out\n   ```\n6. **Verify correctness** using companion script:\n   ```bash\n   python scripts/verify_kernel.py kernel.py --rtol 1e-3 --atol 1e-3\n   ```\n   The kernel file must export `kernel_fn`, `reference_fn`, and `get_inputs()`.\n7. **Benchmark** using companion script:\n   ```bash\n   python scripts/benchmark_kernel.py kernel.py\n   ```\n\n### Workflow 2: GEMM Kernel\n\nFor matrix multiplication with tiling, shared memory, and tensor cores.\n\n1. **Define problem**: shapes (M, N, K), data types, target architecture\n2. **Choose tiling**: CTA tile (bM, bN, bK), pipeline stages, cluster shape\n3. **Three-level partitioning** (see `references/patterns-gemm.md`):\n   - Level 1: CTA tiling with `local_tile()`\n   - Level 2: Copy partitioning (global → shared) with `TiledCopy`\n   - Level 3: Compute partitioning (shared → register) with `TiledMMA`\n4. **Shared memory**: Use swizzled layouts (`make_smem_layout_atom`) to avoid bank conflicts\n5. **Mainloop**: K-tile loop with copy → sync → MMA → sync\n6. **Pipeline**: Use `PipelineTmaAsync` (Hopper) or `PipelineTmaUmma` (Blackwell).\n   ⚠️ TMA-based pipelines manage data movement via TMA descriptors — adding\n   extra tensors (bias, activation inputs) to the epilogue requires modifying\n   descriptor setup, which is prohibitively complex. See the stop condition in\n   Workflow 0 step 4.\n7. **Epilogue**: Predicated store with alpha/beta scaling\n8. **Pre-compile with `cute.compile()`**: Always pre-compile the GEMM kernel\n   so `kernel_fn` calls the compiled object, not `@cute.jit` directly.\n   Without pre-compilation, every call recompiles (~20-50ms overhead).\n9. **Autotune**: Search over tile sizes, cluster shapes, pipeline depths\n\n### Workflow 3: Framework Integration\n\nFor wrapping CuTe DSL kernels as PyTorch/JAX custom operators.\n\n1. **Write kernel** using Workflow 1 or 2\n2. **Create wrapper**: Accept `torch.Tensor`, convert via `from_dlpack`, call host fn\n3. **For production**: Compile with TVM FFI for zero-overhead tensor passing:\n   ```python\n   compiled = cute.compile(host_fn, *fake_tensors, options=\"--enable-tvm-ffi\")\n   compiled(torch_a, torch_b)  # Direct torch.Tensor, no from_dlpack\n   ```\n4. **For deployment**: Use AOT compilation → export to `.o` → load at runtime\n\n### Workflow 4: Debugging & Profiling\n\n1. **Set environment**: `CUTE_DSL_PRINT_IR=1`, `CUTE_DSL_KEEP_PTX=1`\n2. **Use `cute.printf()`** for runtime values (not Python `print`)\n3. **Inspect generated code**: `compiled.__ptx__`, `compiled.__mlir__`\n4. **Profile**: Enable `CUTE_DSL_LINEINFO=1`, use Nsight Compute/Systems\n5. **Debug memory**: Run with `compute-sanitizer python script.py`\n\n## Output Formats\n\nA typical CuTe DSL kernel project:\n\n```\nkernel_dir/\n  kernel.py          # @cute.kernel + @cute.jit functions\n  test_kernel.py     # Correctness test vs PyTorch reference\n  bench_kernel.py    # Benchmark with cute.compile() setup\n```\n\n**Success indicators:**\n- Correctness test passes (`torch.testing.assert_close`)\n- Nsight shows vector loads (LDG.128/LDG.256), not scalar loads\n- For GEMM: tensor core utilization > 80% in Nsight Compute\n\n## Companion Script Contract\n\nKernel files used with `scripts/verify_kernel.py` and `scripts/benchmark_kernel.py`\nmust export three names:\n\n- `kernel_fn(*inputs)` — the CuTe DSL kernel wrapper (calls `cute.compile` + runs kernel)\n- `reference_fn(*inputs)` — PyTorch reference implementation (same signature)\n- `get_inputs()` — returns a list of CUDA tensors for testing\n\n```python\n# Example kernel.py contract\nimport torch\nimport cutlass.cute as cute\nfrom cutlass.cute.runtime import from_dlpack\n\ndef kernel_fn(x):\n    out = torch.empty_like(x)\n    # ... call compiled cute kernel ...\n    return out\n\ndef reference_fn(x):\n    return torch.nn.functional.gelu(x)\n\ndef get_inputs():\n    return [torch.randn(1024, 512, dtype=torch.float16, device=\"cuda\")]\n```\n\n## Examples\n\n### Example: 2D Unary Element-wise (ReLU)\n\n```python\nimport torch, cutlass, cutlass.cute as cute\nfrom cutlass.cute.runtime import from_dlpack\n\n@cute.kernel\ndef relu_kernel(gA: cute.Tensor, gC: cute.Tensor):\n    tidx, _, _ = cute.arch.thread_idx()\n    bidx, _, _ = cute.arch.block_idx()\n    bdim, _, _ = cute.arch.block_dim()\n    idx = bidx * bdim + tidx\n    m, n = gA.shape[1]\n    total = m * n\n    if cutlass.dynamic_expr(idx < total):\n        a = gA[(None, (idx // n, idx % n))].load()\n        gC[(None, (idx // n, idx % n))] = cute.where(a > 0, a, 0)\n\n@cute.jit\ndef relu_host(mA: cute.Tensor, mC: cute.Tensor):\n    vec = 16 // (mA.element_type.width // 8)\n    gA = cute.zipped_divide(mA, (1, vec))\n    gC = cute.zipped_divide(mC, (1, vec))\n    T = 256\n    N = cute.size(gA.shape[1])\n    relu_kernel(gA, gC).launch(grid=((N+T-1)//T,1,1), block=(T,1,1))\n\nx = torch.randn(1024, 512, dtype=torch.float16, device=\"cuda\")\nout = torch.empty_like(x)\nrelu_host(from_dlpack(x, assumed_align=16), from_dlpack(out, assumed_align=16))\n```\n\n## Error Handling\n\n| Error | Cause | Fix |\n|-------|-------|-----|\n| `MLIR function requires a Context` | Called @kernel from Python | Launch via @cute.jit host function |\n| `DSLAstPreprocessorError` on return | Early return in @kernel | Use `if cutlass.dynamic_expr(cond):` |\n| Type mismatch on store | `a * 2` promotes FP16→FP32 | Use `a + a` or `.to(cutlass.Float16)` |\n| `could not get source code` | Kernel in `exec()` context | Write to file and import |\n| Scalar loads in Nsight | Missing alignment hint | Add `assumed_align=16` to `from_dlpack` |\n| `Missing required argument` | Not all @jit params passed | Pass ALL declared parameters |\n| `AttributeError: sigmoid` | No `cute.math.sigmoid` | Use `1.0/(1.0+cute.math.exp(-x))` |\n\nSee `references/troubleshooting.md` for the full error table and limitations.\n\n**Debugging rule:** Never delete kernel.py during debugging. Use `backup_file`\nto save a checkpoint, then `edit_file` to iterate. If stuck, `revert_file` to\nrestore the backup. A partially-working kernel is always better than no kernel.\n\n## Finding More Information\n\n### Tier 1: This File (SKILL.md)\n\nWorkflows above cover element-wise kernels, GEMM, framework integration, and\ndebugging. Search this file first for procedural questions.\n\n### Tier 2: references/ Directory\n\nGrep for keywords across `references/`. Headers are grep-friendly.\n\n| File | Content |\n|------|---------|\n| `concepts-architecture.md` | Core abstractions, terminology, compilation pipeline |\n| `concepts-layouts.md` | Layout algebra: composition, complement, divide, swizzle |\n| `concepts-tensors.md` | Tensor types, partitioning, tiling, predication |\n| `concepts-mma.md` | MMA atoms, TiledMMA, per-architecture tensor core ops |\n| `patterns-getting-started.md` | Installation, decorators, first kernel walkthrough |\n| `patterns-elementwise.md` | Invariant principles, pattern variations, reference impl |\n| `patterns-gemm.md` | 3-level tiling, shared memory, pipelining, autotuning |\n| `patterns-memory.md` | from_dlpack, TMA, cp.async, TMEM, copy atoms |\n| `patterns-compilation.md` | Control flow, JIT caching, TVM FFI, AOT compilation |\n| `patterns-pipeline.md` | Producer-consumer, pipeline classes, barriers, warp specialization |\n| `api-core.md` | cute module: layouts, tensors, math, copy, gemm, printing |\n| `api-arch.md` | cute.arch: thread indexing, sync, atomics, memory ops |\n| `api-nvgpu.md` | cute.nvgpu: warp/warpgroup/cpasync/tcgen05 MMA and copy |\n| `api-runtime-utils.md` | Runtime: from_dlpack, fake tensors, utils, schedulers |\n| `troubleshooting.md` | Debugging, env vars, common errors, limitations, FAQ |\n\n**How to search:** Grep for your keyword across `references/`. Read only the\nfile and section that Grep points to.\n\n### Tier 3: Original Documentation\n\nIf Tiers 1–2 don't answer, consult the source:\n- **Web**: https://docs.nvidia.com/cutlass/latest/\n- **GitHub**: https://github.com/NVIDIA/cutlass\n- Fetch specific doc pages or search for \"CUTLASS CuTe DSL <topic>\"\n- Consider distilling the answer back into references/\n\n<!-- File: .claude/skills/kernel-tileir-optimization/references/config-templates.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-tileir-optimization/references/config-templates.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/config-templates.md)\n\n<!-- File: .claude/skills/kernel-tileir-optimization/references/tma-conversion.md (Tokens: ~300 | Category: anthropic-skill) -->\n# References Skill\n\nPath: `.claude/skills/kernel-tileir-optimization/references/tma-conversion.md`\n\n[View Raw Content on GitHub](https://raw.githubusercontent.com/NVIDIA/TensorRT-LLM/HEAD/.claude/skills/kernel-tileir-optimization/references/tma-conversion.md)\n\n<!-- File: .claude/skills/kernel-tileir-optimization/SKILL.md (Tokens: ~2479 | Category: anthropic-skill) -->\n---\nname: kernel-tileir-optimization\ndescription: >\n  Optimize existing Triton kernels for NVIDIA TileIR backend on Blackwell GPUs (sm_100+).\n  Adds TileIR-specific autotune configs: occupancy, num_ctas, TMA descriptors. Covers\n  kernel classification (dot-related, norm-like, elementwise, reduction), type-specific\n  transformations, and PTX-vs-TileIR benchmarking. Triggered by: \"optimize for TileIR\",\n  \"add TileIR configs\", \"Blackwell optimization\", \"TMA descriptors\", \"2CTA mode\",\n  \"occupancy tuning\". Kernels use standard `import triton`; TileIR activates via\n  ENABLE_TILE=1 when nvtriton is installed.\ncompatibility: >\n  Requires Blackwell GPU (sm_100+) for TileIR execution. Supports development on any GPU.\nlicense: Apache-2.0\nmetadata:\n  author: NVIDIA Corporation\n---\n\n# Triton TileIR Optimization\n\nOptimize EXISTING Triton kernels for NVIDIA's TileIR backend on Blackwell GPUs.\nThis skill does NOT write kernels from scratch -- that is the Triton Specialist's job.\n\n## Principles\n\n### TileIR vs PTX Backend\n\nTileIR is NVIDIA's compiler backend for Triton that generates optimized CUDA code\nusing CGA-level (Cooperative Grid Array) tile representations. Critical differences:\n\n| Parameter | PTX Backend | TileIR Backend |\n|-----------|-------------|----------------|\n| `num_warps` | Strict directive | **Ignored** (compiler decides) |\n| `num_stages` | Strict directive | Cost hint (compiler optimizes) |\n| `occupancy` | Not available | **Critical** tuning param (1-32) |\n| `num_ctas` | Limited | 2CTA mode for Blackwell |\n| Block sizes | Smaller often better | Larger often better |\n| TMA | Not available | Required for dot kernels |\n\n**Key implication**: Do not tune `num_warps` for TileIR -- focus on `occupancy` instead.\n\n### Triton Package Landscape\n\nThree packages share `import triton`:\n\n| Package | Source | Use Case |\n|---------|--------|----------|\n| `pytorch-triton` | PyTorch wheel | `torch.compile`, standard kernels |\n| `triton` | OpenAI PyPI | Official Triton from triton-lang.org |\n| nvtriton | [Triton-to-tile-IR](https://github.com/triton-lang/Triton-to-tile-IR) | TileIR backend for Blackwell |\n\nOnly one triton package should be installed at a time. \"Converting to TileIR\" means\nadding TileIR-specific configs, NOT changing imports. TileIR activates via `ENABLE_TILE=1`.\n\n### When TileIR Applies\n\nTileIR targets Blackwell (sm_100+). Without nvtriton or Blackwell hardware, the\nspecialist still adds TileIR-optimized configs that standard triton safely ignores,\nenabling future deployment.\n\n**Expected speedups** (with nvtriton on Blackwell):\n\n| Kernel Type | Speedup | Key Lever |\n|-------------|---------|-----------|\n| Dot-Related (GEMM, Attention) | 1.2-2.0x | TMA + 2CTA |\n| Norm-Like (LayerNorm, Softmax) | 2.0-5.0x | High occupancy |\n| Element-Wise (ReLU, Add, Exp) | 1.5-3.0x | Occupancy + num_stages |\n| Reduction (Sum, Mean, Max) | 1.8-4.0x | High occupancy |\n\n## Workflow\n\nFive-phase workflow: compatibility, classify, transform, validate, benchmark.\n\n### Phase 1: Compatibility Test (ENABLE_TILE=0)\n\nVerify the kernel works in PTX mode before applying TileIR optimizations.\n\n```bash\npython scripts/tileir_check.py\n```\n\nThen use the **kernel-triton-writing** skill's `verify_kernel.py` to verify with `ENABLE_TILE=0`:\n\n```bash\npython scripts/verify_kernel.py --kernel path/to/kernel.py --reference 'torch reference' --shapes '{\"x\": [32, 512, 4096]}' --dtypes '{\"x\": \"bfloat16\"}'\n```\n\n### Phase 2: Classify Kernel\n\nDetermine kernel type to select the optimization strategy.\n\n```bash\npython scripts/classify_kernel.py --file kernel.py\n```\n\nClassification decision tree:\n\n```\nContains tl.dot()?\n  YES --> dot-related: TMA + 2CTA + occupancy + larger blocks\n  NO  --> Has reduction + normalization?\n            YES --> norm-like: high occupancy (2, 4) + num_warps (4, 8)\n            NO  --> Point-wise only?\n                      YES --> element-wise: occupancy (1-16) + num_stages (2-4)\n                      NO  --> reduction: high occupancy + num_warps\n```\n\n### Phase 3: Apply Transformations\n\nClassify and apply optimizations in one step:\n\n```bash\npython scripts/classify_kernel.py --file kernel.py --apply-optimizations\n```\n\nOutput JSON includes `optimized_code` and `changes_applied` fields.\n\n**Type-specific transformations:**\n\n**Dot-related** (highest priority):\n1. Convert `tl.load`/`tl.store` to TMA descriptors (MANDATORY). See `references/tma-conversion.md`.\n2. Add 2CTA configs (`num_ctas=2`) with SM oversubscription guard in pre-hook.\n3. Add occupancy (1, 2, 4) and extended num_stages (4, 6).\n4. Use larger block sizes (256x256, 256x128).\n\n**Norm-like** (LayerNorm, Softmax, RMSNorm):\n- Add occupancy (2, 4), num_warps (4, 8). No TMA needed.\n\n**Element-wise** (ReLU, GELU, Add, Mul, Exp):\n- Add occupancy (1, 2, 4, 16), num_stages (2, 3, 4). Include extreme configs for small inputs.\n\n**Reduction** (Sum, Mean, Max):\n- Same strategy as norm-like: high occupancy (2, 4), num_warps (4, 8).\n\nGate TileIR-specific configs for sm_100+:\n\n```python\nimport torch\n\ndef get_configs_with_gating(pre_hook=None):\n    configs = get_baseline_configs()\n    if torch.cuda.is_available() and torch.cuda.get_device_capability()[0] >= 10:\n        configs.extend(get_tileir_specific_configs(pre_hook))\n    return configs\n```\n\nSee `references/config-templates.md` for complete config templates per kernel type.\n\n### Phase 4: TileIR Validation (ENABLE_TILE=1)\n\nUse the **kernel-triton-writing** skill's `verify_kernel.py` to verify the optimized kernel with TileIR backend:\n\n```bash\npython scripts/verify_kernel.py --kernel path/to/optimized_kernel.py --reference 'torch reference' --shapes '{\"x\": [32, 512, 4096]}' --dtypes '{\"x\": \"bfloat16\"}'\n```\n\nSet `ENABLE_TILE=1` before running. Check: numerical correctness, no compilation errors,\nTMA/2CTA patterns compile successfully.\n\n### Phase 5: Benchmark\n\nUse `triton.testing.do_bench()` (as documented in the **perf-workload-profiling** skill) to compare PTX (`ENABLE_TILE=0`) vs TileIR (`ENABLE_TILE=1`).\n\nBenchmark across multiple input sizes (128, 1024, 8192) -- performance varies by size.\n\n## Scripts\n\n### tileir_check.py\n\nCheck TileIR availability (nvtriton, ENABLE_TILE, Blackwell GPU):\n\n```bash\npython scripts/tileir_check.py\n```\n\nReturns JSON: `nvtriton_installed`, `tileir_active`, `blackwell_gpu`, `gpu_capability`, `recommendation`.\n\n### classify_kernel.py\n\nClassify kernel type and optionally apply TileIR optimizations:\n\n```bash\n# Classify only\npython scripts/classify_kernel.py --file kernel.py\n\n# Classify + apply optimizations\npython scripts/classify_kernel.py --file kernel.py --apply-optimizations\n\n# From inline code\npython scripts/classify_kernel.py --code '<kernel_code>'\n```\n\nReturns JSON: `classification`, `confidence`, `indicators`, `recommendations`.\nWith `--apply-optimizations`: adds `optimized_code` and `changes_applied`.\n\n## Error Handling\n\n### Common Pitfalls\n\n**TMA descriptor errors** (dot-related kernels):\n- Always pass `pre_hook=tma_set_block_size_hook` to config generation -- without it,\n  TMA descriptors keep dummy block sizes, causing runtime errors or wrong results.\n- For GEMM: pass `b.T.contiguous()` in wrapper and use `tl.dot(a, b.T, accumulator)`\n  in kernel. Transposition mismatch produces incorrect results silently.\n\n**2CTA oversubscription**:\n- Adjust SM count in pre-hook when using `num_ctas=2`:\n  ```python\n  if \"NUM_SMS\" in nargs and \"NUM_CTAS\" in nargs:\n      nargs[\"NUM_SMS\"] = nargs[\"NUM_SMS\"] // nargs[\"NUM_CTAS\"]\n  ```\n\n**Config function signatures**:\n- ALL config helper functions MUST accept `pre_hook=None`, even if unused.\n  Without it: `TypeError: get_autotune_configs() takes 0 positional arguments`.\n\n**Hardware gating**:\n- Gate TileIR configs with `torch.cuda.get_device_capability()[0] >= 10`.\n  TMA/2CTA on pre-Blackwell GPUs causes runtime crashes.\n\n**API availability**:\n- Use `1.0 / (1.0 + tl.exp(-x))` instead of `tl.sigmoid(x)` -- not available in\n  all Triton versions including some nvtriton builds.\n\n**Performance tuning**:\n- Do not over-tune `num_warps` -- TileIR ignores it. Focus on `occupancy`.\n- Use larger block sizes (256x256, 256x128) for TileIR, not PTX-tuned small blocks.\n- Benchmark across small/medium/large inputs; one-size configs underperform.\n- For exp/log heavy kernels, enable approximate math:\n  ```bash\n  export TILEIR_ENABLE_APPROX=1\n  export TILEIR_ENABLE_FTZ=1\n  ```\n\n### When to Abort\n\nStop and report if:\n\n1. **No triton installed** -- cannot proceed.\n2. **Compatibility test fails** -- kernel has syntax/runtime errors before optimization.\n3. **TileIR validation fails** -- optimized kernel produces wrong results.\n4. **No speedup** -- TileIR version is slower than PTX baseline (with nvtriton).\n5. **Not Blackwell GPU** -- still add configs for future deployment, but skip\n   ENABLE_TILE testing and benchmarking.\n\n### Output Format\n\nAfter optimization, return:\n\n```\n## TileIR Optimization: kernel_name\n\n### Classification\n- Kernel type: [dot-related | norm-like | element-wise | reduction]\n- Strategy: [TMA + 2CTA | High occupancy | Occupancy + num_stages]\n\n### Compatibility Check (ENABLE_TILE=0)\n[PASSED | FAILED] — Max difference: X.Xe-Y\n\n### Transformations Applied\n- [List of transformations]\n\n### TileIR Validation (ENABLE_TILE=1)\n[PASSED | FAILED] — Max difference: X.Xe-Y\n\n### Benchmark Comparison\n| Backend | Time (ms) | Speedup |\n|---------|-----------|---------|\n| PTX (ENABLE_TILE=0) | X.XXX | 1.0x |\n| TileIR (ENABLE_TILE=1) | X.XXX | Y.Yx |\n\n### Output\nFile: kernel_name_tileir.py\n```\n\n<!-- ... +361 more modular skills available via /api/skills/NVIDIA/TensorRT-LLM?path={path} -->\n</yakaai_skills>"}