New: 98,000+ skills indexed from GitHub repos with 100+ stars · every skill safety-checkedBrowse top skillsTop skills
Esc to closeBrowse all skills
Submit a skillBrowse skills
Submit a skillBrowse skills

Agent skill

B200 Mbarrier Protocol Auditor

by mirage-project in mirage-project/mirage

A skill your agent uses when a Blackwell/B200 asynchronous kernel deadlocks, fails intermittently, reads stale data, reuses a stage too early, or when the arrival, tx-count, phase, and wait of…

Apache-2.0Auto-check passedDevelopment

Install B200 Mbarrier Protocol Auditor

skills CLI
$ npx skills add mirage-project/mirage --skill b200-mbarrier-protocol-auditor -a claude-code

Project install by default; add -g for ~/.claude/skills/.

GitHub CLI
$ gh skill install mirage-project/mirage b200-mbarrier-protocol-auditor --agent claude-code

Project scope by default; add --scope user for a personal install. Needs GitHub CLI 2.90.0 or later (public preview).

Manual copy
$ git clone --depth 1 https://github.com/mirage-project/mirage.git skills-src && mkdir -p .claude/skills && cp -r skills-src/.claude/skills/b200-mbarrier-protocol-auditor .claude/skills/b200-mbarrier-protocol-auditor && rm -rf skills-src

Use ~/.claude/skills/ instead of .claude/skills for a personal install. The folder must contain SKILL.md.

Claude Code skills documentation · loads skills from .claude/skills/

Facts

Skill name
b200-mbarrier-protocol-auditor
GitHub stars
2.5k
Token cost
~1.8k tokens
SKILL.md length
853 words
Files
2
Skills in repo
24
Repo updated
First seen
Licence
Apache-2.0

At a glance

A skill your agent uses when a Blackwell/B200 asynchronous kernel deadlocks, fails intermittently, reads stale data, reuses a stage too early, or when the arrival, tx-count, phase, and wait of…

  • Works in 3 steps: "The kernel is stuck in mbarrier.wait —… → "My double buffering occasionally reads… → "How should tcgen05 commit be wired to…
  • A Blackwell/B200 asynchronous kernel deadlocks
  • SKILL.md covers R — Source evidence (Reading,…, I — Methodology skeleton…, A1 — Applications in the… and A2 — Trigger scenarios (Future…, plus 4 more sections
  • Instructions only: no scripts, shell commands, URLs or credentials in SKILL.md

What it does

B200 Mbarrier Protocol Auditor is an agent skill from mirage-project/mirage. Use when a Blackwell/B200 asynchronous kernel deadlocks, fails intermittently, reads stale data, reuses a stage too early, or when the arrival, tx-count, phase, and wait of TMA/tcgen05/CLC need auditing. Produces a per-barrier protocol ledger and fix points. Not for cases already known to be out-of-bounds accesses, layout mismatches, or ordinary host-side synchronization problems.

Its SKILL.md is about 1.8k tokens, which your agent loads only when the skill is triggered. The skill folder holds 1 other file (for example `test-prompts.json`).

It sits in Development, covering Async programming. The repository describes itself as: Mirage Persistent Kernel: Compiling LLMs into a MegaKernel. The licence is Apache-2.0.

When your agent uses it

  • A Blackwell/B200 asynchronous kernel deadlocks
  • Fails intermittently
  • Reads stale data
  • Reuses a stage too early

Example prompts

  • “/b200-mbarrier-protocol-auditor”

Workflow steps

3 steps, taken from the first numbered list in SKILL.md.

  1. "The kernel is stuck in mbarrier.wait — help me check the arrival/phase."
  2. "My double buffering occasionally reads the previous round's tile."
  3. "How should tcgen05 commit be wired to the barrier?"

What it can do on your machine

Read from SKILL.md and the folder at commit f9eb70c. It shows what the files ask for, not the result of running them.

  • Tool permissions

    Pre-approves nothing: there is no allowed-tools line, so your agent's usual permission prompts apply.

    From allowed-tools in the SKILL.md frontmatter.

  • Runs code

    No scripts in the folder and no shell commands in SKILL.md.

    From the folder's file list and the shell code blocks in SKILL.md.

  • Network

    Links to these hosts (documentation or services it may open):

    • mlc.ai

    From URLs in SKILL.md, links to its own repository left out.

  • Credentials

    Names no API keys, tokens, secrets or passwords.

    From names ending in _API_KEY, _TOKEN, _SECRET, _KEY or _PASSWORD in SKILL.md.

Context cost

B200 Mbarrier Protocol Auditor loads about 1.8k tokens when it runs. Until then it costs about 104 tokens; SKILL.md has 853 words of instructions outside code blocks.

Always · name and description, kept in context so the agent knows when to use it
~104
When it runs · the whole SKILL.md, loaded when a task matches
~1.8k

Estimates: characters ÷ 4, the usual rule of thumb; real counts depend on the model's tokenizer. Scripts and assets cost tokens only if the agent reads them.

Safety

Auto-check passed

The automated check found no risky patterns in SKILL.md.

Automated static check — not a guarantee. Review scripts before installing. It scans the text of SKILL.md for risky patterns (piping downloads into a shell, reading credential files, hidden Unicode, destructive commands); files beside SKILL.md are not scanned.

SKILL.md

The full file from mirage-project/mirage at commit f9eb70c, republished under its Apache-2.0 licence (© mirage-project). 853 words, ~1,815 tokens.

Download SKILL.mdSave it as .claude/skills/b200-mbarrier-protocol-auditor/SKILL.md (or your agent's skills folder). This skill also uses 1 other file; get the full folder from GitHub.
name
b200-mbarrier-protocol-auditor
description
Use when a Blackwell/B200 asynchronous kernel deadlocks, fails intermittently, reads stale data, reuses a stage too early, or when the arrival, tx-count, phase, and wait of TMA/tcgen05/CLC need auditing. Produces a per-barrier protocol ledger and fix points. Not for cases already known to be out-of-bounds accesses, layout mismatches, or ordinary host-side synchronization problems.
source_book
Modern GPU Programming For MLSys (MLC Community) + NVIDIA Blackwell Tuning/Compatibility Guides
source_chapter
S9; S6; S7; S13; S14
tags
b200, blackwell
related_skills
b200-scope-layout-dispatch, b200-layout-contract-auditor, b200-tma-pipeline-designer, b200-tmem-lifecycle-planner, b200-tcgen05-mma-contract-builder…
version
0.1.0
<!-- Distilled from "Modern GPU Programming for MLSys" — https://mlc.ai/modern-gpu-programming-for-mlsys/ -->

B200 mbarrier Protocol Auditor

R — Source evidence (Reading, paraphrased)

  • [S9] An mbarrier keeps an arrival counter and a phase in SMEM; completion of asynchronous operations cannot be inferred from program order.
  • [S9] A TMA load's expect_tx registers both the issuing thread's arrival and the expected bytes; the phase flips only when both the arrivals and the pending bytes are complete.
  • [S9] A tcgen05 MMA must explicitly associate a barrier arrival on its commit path; otherwise consumers may wait forever.
  • [S13] Multi-role GEMMs commonly carry two barrier sets, forward ready and backward release; a wrong initial phase deadlocks or silently corrupts.

Source: distilled from "Modern GPU Programming for MLSys" (https://mlc.ai/modern-gpu-programming-for-mlsys/) and the NVIDIA Blackwell tuning/compatibility guides. Short paraphrases only; no long passages are reproduced.

I — Methodology skeleton (Interpretation)

Treat every barrier as a bidirectional contract, not "a wait somewhere":

  • Who initializes it, and how many arrivals are expected?
  • Who produces the data, and does the completion signal come from threads, a TMA byte count, an MMA commit, or a CLC response?
  • Which consumers wait on which phase?
  • On completion, does it release data-ready or buffer-free?
  • When does the phase flip, and who holds the local expected phase for the next round?

The audit's goal is a complete barrier ledger; any field without a unique answer is a potential bug.

A1 — Applications in the source (Past Application)

Case 1: producer/consumer initial phases
  • With an empty buffer, the producer's first round should start filling immediately; the consumer's first round should block until the data is ready.
  • If the two sides share the same initial phase, both may block on the first round, or an old completion may be mistaken for a new one.
Case 2: the four-barrier GEMM
  • tma2mma: SMEM data ready.
  • mma2tma: SMEM stage overwritable.
  • mma2ld: TMEM result ready.
  • ld2mma: TMEM region reusable.
  • Forward data flow and backward resource release must close the loop in pairs.

A2 — Trigger scenarios (Future Trigger) ★

In what situations will the user need this skill?
  1. "The kernel is stuck in mbarrier.wait — help me check the arrival/phase."
  2. "My double buffering occasionally reads the previous round's tile."
  3. "How should tcgen05 commit be wired to the barrier?"
Language signals
  • "The kernel is stuck in mbarrier.wait — help me check the arrival/phase."
  • "My double buffering occasionally reads the previous round's tile."
  • "How should tcgen05 commit be wired to the barrier?"
Distinction from adjacent skills

Versus b200-layout-contract-auditor: this skill only proves "when it is safe", not "whether the address is correct". Versus b200-warp-specialized-debugger: this skill is a dedicated barrier audit; the debugger is a full symptom-driven workflow.

Show full SKILL.md (426 more words)Show less

E — Executable steps (Execution)

Once the skill is activated, the agent must execute the following procedure:

  1. Build the barrier ledger
    • Fields: name, storage address, init scope, expected arrivals, tx bytes, producer, arrival mechanism, consumer, wait phase, released resource.
  2. Verify the initialization sites
    • Barrier init must happen before the role branch and be observed by the correct scope.
    • Check alignment and the number of stages.
  3. Verify the producer side
    • TMA load: is the expected byte count exact, and does the actual copy complete against the same barrier?
    • MMA/tcgen05.cp: is the commit bound to the correct barrier arrival?
    • Plain threads: does the number of threads actually arriving equal the init count?
  4. Verify the consumer side
    • Do the waited-on barrier and the stage index match?
    • Does the local phase flip exactly once, and only once, after a successful consume?
  5. Verify the ready/free closed loop
    • Every SMEM/TMEM slot has both a "when readable" and a "when overwritable/freeable" condition.
  6. Verify scope legality
    • A CTA-wide collective must not hide inside a warpgroup-only branch.
    • Cluster signals must use the correct CTA mask/remote arrival.
  7. Check the first and last rounds
    • Simulate the prologue's first wait, one steady-state lap, and the epilogue's final drain.
    • Hand-compute the 0/1/0/1 phase sequence in a table.
  8. Change only one handoff
    • After changing exactly one of init count, phase, arrival, or fence, run minimal correctness first, then measure performance.
Required outputs
  1. Conclusion: the current choice/diagnosis, without vague "could be any of them" hedging.
  2. Evidence or assumptions: which come from user data, and which are hypotheses awaiting verification.
  3. Contract/table/timeline: the auditable intermediate artifacts corresponding to this skill.
  4. Minimal validation: correctness tests, boundary tests, and one falsifiable experiment.
  5. Risks and fallback: alternative paths when hardware, version, or resource conditions are not met.

B — Boundaries (Boundary) ★

Do not use when
  • There is already conclusive evidence of an out-of-bounds access, an invalid pointer, or a wrong descriptor.
  • Ordinary CPU mutex/condition_variable synchronization problems.
Failure modes
  • Treating expect_tx as a plain arrival and ignoring the byte budget.
  • Forgetting the MMA commit arrival.
  • Stage index correct but the phase belongs to the previous round.
  • Building only the ready barrier and not the buffer-release barrier.
Limitations
  • Compiler lowering may change the surface code structure; ultimately verify the barrier initialization and commit/wait in the generated CUDA/PTX.
  • depends-on: b200-scope-layout-dispatch
  • contrasts-with: b200-layout-contract-auditor
  • composes-with: b200-tma-pipeline-designer, b200-tmem-lifecycle-planner, b200-tcgen05-mma-contract-builder, b200-warp-specialized-debugger

Audit info

  • Validation passed: V1 ✓ / V2 ✓ / V3 ✓
  • Test definitions: 6 (3 should_trigger / 2 should_not_trigger / 1 edge_case)
  • Hardware validation: not performed; must be verified on a target B200
  • Distilled: 2026-06-25

© mirage-project, Apache-2.0. Rendered from Markdown: HTML in the file is shown as text, images as links, and headings moved down two levels. Raw file

Files

SKILL.md and 1 other file in .claude/skills/b200-mbarrier-protocol-auditor of mirage-project/mirage.

  • SKILL.md
  • test-prompts.json

Open the folder on GitHubat commit f9eb70c

Compare with similar skills

B200 Mbarrier Protocol Auditor next to the 5 skills that share the most tags, products or categories with it. Stars are the repository's; “used in” counts other GitHub owners with a copy.

B200 Mbarrier Protocol Auditor compared with similar skills
SkillStarsUsed inTokensAuto-checkLicenceRepo updated
B200 Mbarrier Protocol Auditor this skillmirage-project/mirage2.5k—~1.8kAutomated safety check: PassApache-2.0
YugabyteDB ASH Instrumentationyugabyte/yugabyte-db11k—~4.5kAutomated safety check: PassCustom licence
Mirage VFS Adapter Authoringstrukto-ai/mirage3.7k—~2.5kAutomated safety check: PassApache-2.0
Golang Patternsantoniopaya22/go-rest-template1729 repos~3.5kAutomated safety check: PassNone
Rust Async Patternsdiodeme/Gold-Band14310 repos~3.1kAutomated safety check: PassAGPL-3.0
Swift Concurrencyhenrypldev/react-native-nitro-mlx1003 repos~3.1kAutomated safety check: PassMIT

Similar skills

  • YugabyteDB ASH Instrumentation

    yugabyte/yugabyte-db

    Procedure for adding or changing YugabyteDB Active Session History wait states in TServer and DocDB C++ code, including the macro to use for sync and async paths.

    11k GitHub stars~4.5k tokensUpdated today
    DevelopmentAuto-check passed

  • Mirage VFS Adapter Authoring

    strukto-ai/mirage

    Builds or extends a custom Mirage virtual filesystem adapter for an API, database, object store or app data, with a working mount configuration and filesystem tests.

    3.7k GitHub stars~2.5k tokensUpdated today
    DevelopmentAuto-check passed

  • Golang Patterns

    antoniopaya22/go-rest-template

    Idiomatic Go patterns, best practices, and conventions for building robust, efficient, and maintainable Go applications.

    172 GitHub starsUsed in 9 repos~3.5k tokens
    DevelopmentAuto-check passed

  • Rust Async Patterns

    diodeme/Gold-Band

    Master Rust async programming with Tokio, async traits, error handling, and concurrent patterns.

    143 GitHub starsUsed in 10 repos~3.1k tokens
    DevelopmentAuto-check passed

  • Swift Concurrency

    henrypldev/react-native-nitro-mlx

    Diagnose Swift Concurrency issues, refactor callback-based code to async/await, and guide Swift 6 migration when working with tasks, actors, @MainActor, Sendable, data races, thread safety, or…

    100 GitHub starsUsed in 3 repos~3.1k tokens
    DevelopmentAuto-check passed

  • Rust Engineer

    farm-fe/farm

    Writes, reviews, and debugs idiomatic Rust code with memory safety and zero-cost abstractions.

    5.6k GitHub starsUsed in 1 repo~1.5k tokens
    DevelopmentAuto-check passed

More from mirage-project/mirage

All 24 skills in this repo

  • V2 Perf Iteration

    mirage-project/mirage

    Runtime-V2 performance-iteration workflow. An agent skill from mirage-project/mirage.

    2.5k GitHub stars~4k tokensUpdated 3 days ago
    Auto-check passed

  • Add Mpk Task

    mirage-project/mirage

    Step-by-step guide for adding a new task implementation to Mirage Persistent Kernel (MPK).

    2.5k GitHub stars~4.5k tokensUpdated 3 days ago
    Auto-check passed

  • B200 Flash Attention4 Planner

    mirage-project/mirage

    A skill your agent uses when the user wants to design or extend a FlashAttention-style forward kernel on B200/Blackwell, involving the two MMAs QKᵀ and PV, online softmax, S/P/O in TMEM, warp roles…

    2.5k GitHub stars~1.9k tokensUpdated 3 days ago
    Auto-check passed

  • Mpk Faithful Gate

    mirage-project/mirage

    Build or run a FAITHFUL in-MPK per-task latency gate (slowCTA at the production grid + cos) for a DeepSeek-V3 MPK decode kernel or shape.

    2.5k GitHub stars~2.6k tokensUpdated 3 days ago
    Auto-check passed

  • Mpk Lever Cleanup

    mirage-project/mirage

    A skill your agent uses when a batch of env-gated (ifdef MPKDSV3 / os.environ-controlled, default-OFF) MPK optimization levers needs to be consolidated into a single clean code path for a PR…

    2.5k GitHub stars~2.2k tokensUpdated 3 days ago
    Auto-check passed

  • Test Mode

    mirage-project/mirage

    Guide for using MPK test mode to unit-test individual layers or multi-layer pipelines through the full compilation pipeline.

    2.5k GitHub stars~4.6k tokensUpdated 3 days ago
    Auto-check passed

Categories

Questions about B200 Mbarrier Protocol Auditor

What does B200 Mbarrier Protocol Auditor do?

A skill your agent uses when a Blackwell/B200 asynchronous kernel deadlocks, fails intermittently, reads stale data, reuses a stage too early, or when the arrival, tx-count, phase, and wait of…. B200 Mbarrier Protocol Auditor is an agent skill from mirage-project/mirage. Use when a Blackwell/B200 asynchronous kernel deadlocks, fails intermittently, reads stale data, reuses a stage too early, or when the arrival, tx-count, phase, and wait of TMA/tcgen05/CLC need auditing.

When should I use B200 Mbarrier Protocol Auditor?

B200 Mbarrier Protocol Auditor fits situations like: A Blackwell/B200 asynchronous kernel deadlocks; fails intermittently; reads stale data; reuses a stage too early.

How do I install B200 Mbarrier Protocol Auditor in Claude Code?

Run `npx skills add mirage-project/mirage --skill b200-mbarrier-protocol-auditor -a claude-code`. Or copy the skill folder (.claude/skills/b200-mbarrier-protocol-auditor in mirage-project/mirage) into .claude/skills/b200-mbarrier-protocol-auditor in your project. Claude Code loads it when a task matches its description.

How do I install B200 Mbarrier Protocol Auditor in Codex?

Run `npx skills add mirage-project/mirage --skill b200-mbarrier-protocol-auditor -a codex`. Or copy the skill folder (.claude/skills/b200-mbarrier-protocol-auditor in mirage-project/mirage) into .agents/skills/b200-mbarrier-protocol-auditor in your project. Codex loads it when a task matches its description.

Can I use B200 Mbarrier Protocol Auditor in Cursor, Gemini CLI or GitHub Copilot?

Cursor, Gemini CLI, GitHub Copilot and OpenCode also load SKILL.md folders. With the skills CLI, run `npx skills add mirage-project/mirage --skill b200-mbarrier-protocol-auditor -a cursor` (or -a gemini-cli, github-copilot or opencode for the others). To copy it by hand, put the folder in .cursor/skills/b200-mbarrier-protocol-auditor, .gemini/skills/b200-mbarrier-protocol-auditor, .github/skills/b200-mbarrier-protocol-auditor and .opencode/skills/b200-mbarrier-protocol-auditor in your project.

What does B200 Mbarrier Protocol Auditor need to run?

SKILL.md names no scripts, command-line tools or credentials: B200 Mbarrier Protocol Auditor is instructions for the agent only.

Does B200 Mbarrier Protocol Auditor access the network?

SKILL.md names 1 domain. As links in the text: mlc.ai. This is read from the text; nothing was executed.

Is B200 Mbarrier Protocol Auditor safe to install?

Our automated static check of SKILL.md found no risky patterns, such as piping downloads into a shell, reading credential files or hidden Unicode. It is not a guarantee. Review the folder before installing.

What licence does B200 Mbarrier Protocol Auditor use?

B200 Mbarrier Protocol Auditor is published under the Apache-2.0 licence (the repository's licence). It allows redistribution, so the full SKILL.md is shown on this page.

How many tokens does B200 Mbarrier Protocol Auditor use?

About 1.8k tokens (SKILL.md is roughly 7.3k characters). Agents keep only the skill's name and description in context until a task matches; then they load SKILL.md in full.

What are the alternatives to B200 Mbarrier Protocol Auditor?

Skills that share tags, products or a category with B200 Mbarrier Protocol Auditor: YugabyteDB ASH Instrumentation (yugabyte/yugabyte-db, 11k stars), Mirage VFS Adapter Authoring (strukto-ai/mirage, 3.7k stars), Golang Patterns (antoniopaya22/go-rest-template, 172 stars) and Rust Async Patterns (diodeme/Gold-Band, 143 stars). The comparison table on this page puts their stars, adoption, token cost, safety result and licence side by side.

Who maintains B200 Mbarrier Protocol Auditor?

mirage-project (a GitHub organization) maintains it in mirage-project/mirage, which has 2,545 GitHub stars. The repository holds 24 skills in this directory. The repository was last updated on October 7, 2026.

Source: mirage-project/mirage on GitHub. Facts on this page come from the repository at the commit we read; the author's words are quoted as theirs.