facebookexperimental
- 22 skills
- 0 followers
- 9 hours ago last updated
- ▌ Fabler · facebookexperimentalApply a disciplined read, plan, execute, and adversarially verify workflow. Use for complex research, architecture, implementation, or audit tasks.
- ▌ Progress Rubric · facebookexperimentalCreate evidence-based Hermit progress reports from live measurements. Use when gathering, validating, or writing a project progress report.
- ▌ Hermit Debugging · facebookexperimentalDebug hermit/detcore runs (nondeterminism, hangs, syscall gaps, scheduling) using hermit's built-in logging and log-diff FIRST, before reading source. Use whenever a guest program behaves unexpectedly, --verify fails, or a run hangs under hermit.
- ▌ Post Facto Review · facebookexperimentalLand reviewed, CI-green Hermit changes before human review and mark them for follow-up. Use as the default autonomous landing discipline.
- ▌ Human Review First · facebookexperimentalGate key Hermit changes on adversarial review and explicit human approval. Use only when the user explicitly requests human-review-first mode.
- ▌ Autows Docs · facebookexperimentalConsult and maintain AutoWS documentation. Use BEFORE exploring AutoWS source code — when investigating, planning, or modifying files under WarpSpecialization/, partition scheduling, warp_specialize ops, WSCodePartition, WSDataPartition, WSTaskPartition, WSMemoryPlanner, or related passes. Also use AFTER making non-trivial changes to AutoWS code to keep docs in sync.
- ▌ Ir Debugging · facebookexperimentalDebug Triton compilation by dumping IR at each stage (TTIR, TTGIR, LLVM, PTX). Use when investigating compilation failures, kernel performance, register spills, or when user asks to inspect IR output. Covers TRITON_KERNEL_DUMP, MLIR_ENABLE_DUMP, LLVM_IR_ENABLE_DUMP, TRITON_DUMP_PTXAS_LOG, and related env vars.
- ▌ Autows Testing · facebookexperimentalRun autoWS (automatic warp specialization) correctness tests. Use when working on autoWS compiler code — files under WarpSpecialization/, partition scheduling, warp_specialize ops, WSCodePartition, WSDataPartition, WSTaskPartition, WSMemoryPlanner, or related passes. Do NOT use TLX correctness tests (third_party/tlx/tutorials/testing/test_correctness.py) for autoWS work — those test manual warp specialization via TLX, not the automatic compiler pipeline.
- ▌ Tlx Amd Testing · facebookexperimentalTest and run TLX-AMD tutorial kernels (gfx950/CDNA4 and gfx1250) and understand their CI. Use when working on AMD TLX tutorial kernels — GEMM (warp-pipeline, LDS-pipelined, TDM, MXFP), Flash Attention (simple, prefetch, persistent), addmm+GLU, or IKBO (FA, LCE) — running their correctness or perf, checking arch gating (gfx950 vs gfx1250), or the MI350 CI workflow. Covers the standardized layout (one correctness file, one perf file per op×arch).
- ▌ Autows Authoring · facebookexperimentalAuthor Triton kernels with automatic warp specialization (AutoWS). Use when writing new AutoWS kernels, adding warp_specialize=True to tl.range loops, choosing tl.range kwargs and JIT options, debugging why WS was not applied, or structuring a kernel to work with both Meta WS and upstream OAI Triton. Covers GEMM and Flash Attention patterns on Hopper and Blackwell.
- ▌ Compute Sanitizer · facebookexperimentalRun NVIDIA compute-sanitizer (memcheck, racecheck, initcheck, synccheck) against a Triton/TLX kernel to find runtime memory and synchronization bugs. Use when a kernel produces wrong results, crashes with an illegal/misaligned access, or is suspected of a shared-memory data race or invalid barrier usage — especially warp-specialized (WS) kernels using mbarriers, named barriers, TMA copies, or MMA accumulators. This is a runtime check: it runs the real kernel via its reproduce command, so it needs a working GPU and is 10-100x slower than a normal run.
- ▌ Debug Failing Gpu · facebookexperimentalRecover from GPU-busy / GPU-unavailable failures. Use when a command (pytest, python, a TLX/Triton kernel run, a benchmark) fails with errors indicating the GPU is busy, out of memory, or unavailable — e.g. "CUDA error: out of memory", "all CUDA-capable devices are busy or unavailable", "CUDA-capable device(s) is/are busy or unavailable", "RuntimeError: No CUDA GPUs are available", "device-side assert", or a hang on the first CUDA call. Runs find_working_gpu.sh to locate a healthy GPU and re-runs the failed command pinned to it via CUDA_VISIBLE_DEVICES.
- ▌ Running With Buck · facebookexperimentalHow to build and run GPU targets under Buck in fbcode. Use when invoking buck2 run / buck2 build for any GPU benchmark, test, or kernel — selecting the GPU architecture and CUDA version, using @mode/opt and the beta Triton modifier, passing environment variables through, and running from the right directory. Covers the general requirements plus the B200/GB200 (b200a, CUDA >= 12.8) and GB300 (b300a, CUDA >= 13.0) hardware requirements.
- ▌ Tlx API Reference · facebookexperimentalTLX DSL API reference for low-level GPU primitives. Use when writing or modifying TLX kernel code that uses barriers (mbarrier, named barriers), memory allocation (local_alloc, SMEM, TMEM), TMA operations, warp specialization (async_tasks, async_task), CLC (cluster launch control), or wgmma instructions. Covers Hopper and Blackwell hardware differences.
- ▌ Kernel Perf Testing · facebookexperimentalRun TLX kernel performance benchmarks on Hopper, Blackwell, and AMD (gfx950/CDNA4, gfx1250) GPUs. Use when user asks to benchmark, profile, or measure performance of any TLX kernel (GEMM, Flash Attention, addmm+GLU, IKBO variants). Handles GPU selection, denoise wrapping (NVIDIA only), and version flags. Never run unless explicitly asked.
- ▌ Ir Override Ablation · facebookexperimental bundleDesign and run Triton TTGIR debugging ablations using ir_override. Use when reducing a provided or dumped TTGIR, trying user-provided or agent-generated ablation/oblation ideas, updating a test harness around ir_override, or preserving a compile/runtime failure while simplifying IR to expose a fundamental compiler or lowering gap.
- ▌ Barrier Visualization · facebookexperimental bundleProduce a structured barrier report for AutoWS (automatic warp specialization) IR. Use when the user wants to visualize, audit, or debug barrier usage across warp-specialized partitions, or when debugging a GPU kernel hang (deadlock). For hangs, first dump IR using the ir-debugging skill, then run this barrier analysis to find the barrier that actually deadlocks -- reasoning with the mbarrier phase model (NOT raw arrive/wait counts, which give false positives), plus missing backward barriers and other synchronization issues. Covers mbarriers, named barriers, tcgen05 commit, TMA-implicit arrives, Aref-based synchronization, and producer/consumer barrier patterns.
- ▌ Proxy Fence Insertion · facebookexperimentalUse when working on fence-related compiler passes, TMA store lowering, proxy fence insertion, investigating missing or spurious fences, or debugging correctness issues in TLX kernels that use tlx.async_descriptor_store or MMA operations.
- ▌ Sched2tlx Perf Testing · facebookexperimentalRun the sched2tlx perf/correctness harness over the modulo-scheduling example corpus (case1-9: GEMM, persistent GEMM, FA fwd/bwd, addmm+bias, LayerNorm, wgrad+bias, multiphase GEMM, scaled_mm). Use when the user asks to benchmark generated-vs-handwritten kernels, check corpus correctness, compare emitter revisions, or regenerate schedule_graph.json fixtures. Never run perf unless explicitly asked.
- ▌ Tma Illegal Instruction · facebookexperimentalDiagnose CUDA "illegal instruction" / kernel crashes on Triton kernels that reference to TMA loads or stores (`make_tensor_descriptor`, `TensorDescriptor`, `descriptor.load`, `descriptor.store`, `tl.async_descriptor_load`, async TMA copies) as the source code line. Use when the user reports CUDA error 716, "an illegal instruction was encountered", segfault inside a TMA op, kernel hang followed by an illegal instruction trap, or a crash that only fires on the first or last tile of a launch. Covers the pattern where a TMA store/load is issued at an offset entirely past a tensor's shape — TMA does NOT silently mask out-of-bounds tile accesses; it traps. The root cause is almost never "missing in-kernel mask" — it is commonly a structural launcher / tile-mapping bug.
- ▌ Amd Att Trace · facebookexperimental bundleCollect, validate, package, and inspect rocprofv3 Advanced Thread Trace bundles for AMD GPU kernels. Use for ATT collection, viewer-ready artifacts, or loop cycle windows; do not trigger for ordinary performance benchmarks.
- ▌ Tlx Kernel Optimization Agent · facebookexperimental bundleExecute the TLX Kernel Optimization Agent CLI on a Triton or TLX kernel. Use this skill whenever the user says "use the TLX agent", "use the kernel optimization agent", "用 TLX agent 优化", or asks Carl/Claude to optimize a kernel with the repository agent. The required outcome is an actual agent CLI invocation and its measured result, not a walkthrough or reimplementation.