Agent skill

B200 Cluster Persistent Scheduler

by mirage-project in mirage-project/mirage

A skill your agent uses when the user wants to use Thread Block Cluster, DSMEM, 2-CTA cooperative MMA, persistent kernels, a tile scheduler, or Cluster Launch Control on B200 to handle uneven tails.

Apache-2.0Auto-check passed

Install B200 Cluster Persistent Scheduler

skills CLI
$ npx skills add mirage-project/mirage --skill b200-cluster-persistent-scheduler -a claude-code

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

GitHub CLI
$ gh skill install mirage-project/mirage b200-cluster-persistent-scheduler --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-cluster-persistent-scheduler .claude/skills/b200-cluster-persistent-scheduler && 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-cluster-persistent-scheduler
GitHub stars
2.5k
Token cost
~1.8k tokens
SKILL.md length
831 words
Files
2
Skills in repo
24
Repo updated
First seen
Licence
Apache-2.0

At a glance

A skill your agent uses when the user wants to use Thread Block Cluster, DSMEM, 2-CTA cooperative MMA, persistent kernels, a tile scheduler, or Cluster Launch Control on B200 to handle uneven tails.

  • Works in 3 steps: "Is a 2-CTA cluster worth it for this… → "Help me design a persistent tile… → "How do I reason about DSMEM access,…
  • The user wants to use Thread Block Cluster
  • 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 Cluster Persistent Scheduler is an agent skill from mirage-project/mirage. Use when the user wants to use Thread Block Cluster, DSMEM, 2-CTA cooperative MMA, persistent kernels, a tile scheduler, or Cluster Launch Control on B200 to handle uneven tails. Outputs the cluster tile, occupancy, sharing/multicast, and static or dynamic scheduling plan. Not for kernels where independent small CTAs already saturate the machine and the workload is uniform.

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`).

The repository describes itself as: Mirage Persistent Kernel: Compiling LLMs into a MegaKernel. The licence is Apache-2.0.

When your agent uses it

  • The user wants to use Thread Block Cluster
  • 2-CTA cooperative MMA
  • Persistent kernels
  • A tile scheduler

Example prompts

  • “/b200-cluster-persistent-scheduler”

Workflow steps

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

  1. "Is a 2-CTA cluster worth it for this GEMM?"
  2. "Help me design a persistent tile scheduler and CLC tail stealing."
  3. "How do I reason about DSMEM access, cluster size, and occupancy together?"

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 Cluster Persistent Scheduler loads about 1.8k tokens when it runs. Until then it costs about 103 tokens; SKILL.md has 831 words of instructions outside code blocks.

Always · name and description, kept in context so the agent knows when to use it
~103
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). 831 words, ~1,789 tokens.

Download SKILL.mdSave it as .claude/skills/b200-cluster-persistent-scheduler/SKILL.md (or your agent's skills folder). This skill also uses 1 other file; get the full folder from GitHub.
name
b200-cluster-persistent-scheduler
description
Use when the user wants to use Thread Block Cluster, DSMEM, 2-CTA cooperative MMA, persistent kernels, a tile scheduler, or Cluster Launch Control on B200 to handle uneven tails. Outputs the cluster tile, occupancy, sharing/multicast, and static or dynamic scheduling plan. Not for kernels where independent small CTAs already saturate the machine and the workload is uniform.
source_book
Modern GPU Programming For MLSys (MLC Community) + NVIDIA Blackwell Tuning/Compatibility Guides
source_chapter
S2; S10; S13; S16
tags
b200, blackwell
related_skills
b200-scope-layout-dispatch, b200-mbarrier-protocol-auditor, b200-tcgen05-mma-contract-builder, b200-gemm-optimization-ladder, b200-kernel-roofline-triage
version
0.1.0
<!-- Distilled from "Modern GPU Programming for MLSys" — https://mlc.ai/modern-gpu-programming-for-mlsys/ -->

B200 Cluster & Persistent Scheduler

R — Source evidence (Reading, paraphrased)

  • [S2/S16] CTAs within a cluster can access each other's SMEM (DSMEM); the B200 portable cluster size is 8, and an explicit opt-in reaches the non-portable 16, but this may reduce the number of active blocks.
  • [S13] A 2-CTA cluster can jointly compute a larger MMA tile and share operands through cluster-scope handoff.
  • [S10] A persistent kernel keeps a fixed set of CTAs/clusters resident to process multiple tiles; CLC allows "stealing" work at runtime from not-yet-launched cluster coordinates, improving the tail.

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)

Cluster and persistent scheduling solve two different problems:

  • cluster enlarges the spatial extent and data-sharing scope of a single cooperative tile;
  • persistent scheduling reduces the static one-tile-one-CTA binding, letting a limited set of work owners fetch tasks in a loop;
  • CLC on Blackwell further turns task fetching from static grid-stride into hardware-assisted dynamic tail scheduling.

Before using them, you must prove the gain comes from higher reuse, larger MMAs, or a better tail — not from "Blackwell has this feature, so use it".


A1 — Applications in the source (Past Application)

Case 1: 2-CTA GEMM
  • The two CTAs each hold part of A/B/accumulator.
  • DSMEM or multicast gives the pair the operands they both need.
  • The barrier's CTA mask changes from single-CTA to cluster remote notification.
Case 2: uneven tile tail
  • Static grid-stride may leave some SMs idle early during the final phase.
  • CLC allows a finished resident cluster to cancel a not-yet-launched cluster and take over its coordinates to continue working.

A2 — Trigger scenarios (Future Trigger) ★

In what situations will the user need this skill?
  1. "Is a 2-CTA cluster worth it for this GEMM?"
  2. "Help me design a persistent tile scheduler and CLC tail stealing."
  3. "How do I reason about DSMEM access, cluster size, and occupancy together?"
Language signals
  • "Is a 2-CTA cluster worth it for this GEMM?"
  • "Help me design a persistent tile scheduler and CLC tail stealing."
  • "How do I reason about DSMEM access, cluster size, and occupancy together?"
Distinction from adjacent skills

Difference from b200-tcgen05-mma-contract-builder: this skill covers cluster-level cooperation and task scheduling; the MMA builder covers the contract of a single Tensor Core operation. Difference from b200-gemm-optimization-ladder: this skill also applies to non-GEMM persistent/cluster kernels.


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

E — Executable steps (Execution)

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

  1. Prove the necessity of the cluster
    • Is a single CTA limited by SMEM/TMEM/operand reuse?
    • Does a 2-CTA tile raise Tensor Core utilization or reduce HBM/L2 traffic?
  2. Define the cluster shape and tile mapping
    • clusterDim, each CTA's logical coordinates, the M/N/K slices, owners, and shared operands.
  3. Plan DSMEM/multicast
    • Make accesses as coalesced as possible and 32B-aligned; avoid non-unit strides.
    • Write out the producers/consumers and fences for local SMEM and remote SMEM.
  4. Compute residency
    • Use cudaOccupancyMaxActiveClusters or an equivalent tool.
    • Compare portable size ≤8 with the B200 opt-in 16; if using 16, set the nonportable attribute and flag the portability.
  5. Upgrade the barrier scope
    • Update the CTA mask, remote arrival, and pair completion; check the even/leader CTA responsibilities.
  6. Choose the scheduling mode
    • Equal-cost, uniform work: static tile formula/grid-stride.
    • Uneven cost or a pronounced tail: persistent owner + CLC work stealing.
  7. Implement the CLC protocol
    • async try-cancel → mbarrier completion → query predicate → read the cluster coordinate only on success.
    • No numeric sentinel; termination is decided by the predicate.
  8. Evaluate the tail and fairness
    • Test the boundary cases where the tile count is fewer than, equal to, and slightly more than the resident clusters.
  9. Output the fallback
    • Provide the portable single-CTA/static-scheduler path and the conditions for enabling it.
Required outputs
  1. Conclusion: the current choice/diagnosis; do not use a vague "it could be any of them".
  2. Evidence or assumptions: which items 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 requirements are not met.

B — Boundaries (Boundary) ★

Do not use when
  • A single CTA already saturates and tile costs are uniform; a cluster would only add synchronization and resource pressure.
  • It must run fully identically on non-B200/non-Blackwell platforms, yet there is no fallback.
Failure modes
  • Assuming a larger cluster size is always faster.
  • DSMEM accesses that are not coalesced or not correctly fenced.
  • Still reading an invalid coordinate when CLC fails.
  • A persistent pool that is too large or too small, ignoring resource residency.
Limitations
  • CLC and non-portable cluster sizes are architecture-dependent features; production code must have device-capability checks and a fallback.

  • depends-on: b200-scope-layout-dispatch, b200-mbarrier-protocol-auditor
  • contrasts-with: none
  • composes-with: b200-tcgen05-mma-contract-builder, b200-gemm-optimization-ladder, b200-kernel-roofline-triage

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-cluster-persistent-scheduler of mirage-project/mirage.

  • SKILL.md
  • test-prompts.json

Open the folder on GitHubat commit f9eb70c

Compare with similar skills

B200 Cluster Persistent Scheduler 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 Cluster Persistent Scheduler compared with similar skills
SkillStarsUsed inTokensAuto-checkLicenceRepo updated
B200 Cluster Persistent Scheduler this skillmirage-project/mirage2.5k—~1.8kAutomated safety check: PassApache-2.0
Hunting For Scheduled Task Persistencemukul975/Anthropic-Cybersecurity-Skills34k—~907Automated safety check: PassApache-2.0
Scheduleasgeirtj/system_prompts_leaks69k—~2.9kAutomated safety check: PassCC0-1.0
Scheduleasgeirtj/system_prompts_leaks69k—~597Automated safety check: PassCC0-1.0
Time Blocking Scheduleriflytek/skillhub5.2k—~1.1kAutomated safety check: PassMIT
Cron Scheduleruvnet/ruflo74k—~257Automated safety check: PassMIT

Similar skills

  • Hunting For Scheduled Task Persistence

    mukul975/Anthropic-Cybersecurity-Skills

    Runs a hypothesis-driven threat hunt for Windows Scheduled Task persistence (T1053), guiding SIEM/EDR queries against task creation events (e.g.

    34k GitHub stars~907 tokensUpdated 1 mo ago
    Productivity & AutomationAuto-check passed
  • Schedule

    asgeirtj/system_prompts_leaks

    Create, update, list, or run scheduled cloud agents (routines) that execute on a cron schedule.

    69k GitHub stars~2.9k tokensUpdated yesterday
    Productivity & AutomationAuto-check passed
  • Schedule

    asgeirtj/system_prompts_leaks

    Create or update a scheduled task that runs automatically. An agent skill from asgeirtj/system_prompts_leaks.

    69k GitHub stars~597 tokensUpdated yesterday
    Productivity & AutomationAuto-check passed
  • Time Blocking Scheduler

    iflytek/skillhub

    Draft flexible daily or weekly schedules around a user's priorities, availability, energy patterns, and fixed commitments.

    5.2k GitHub stars~1.1k tokensUpdated today
    Productivity & AutomationAuto-check passed
  • Cron Schedule

    ruvnet/ruflo

    Schedule persistent background workers via CronCreate. An agent skill from ruvnet/ruflo.

    74k GitHub stars~257 tokensUpdated yesterday
    Productivity & AutomationAuto-check passed
  • Thread-safe data persistence in Swift using actors — in-memory cache with file-backed storage, eliminating data races by design.

    276k GitHub starsUsed in 4 repos~1.2k tokens
    MobileAuto-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 2 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 2 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 2 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 2 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 2 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 2 days ago
    Auto-check passed

Questions about B200 Cluster Persistent Scheduler

What does B200 Cluster Persistent Scheduler do?

A skill your agent uses when the user wants to use Thread Block Cluster, DSMEM, 2-CTA cooperative MMA, persistent kernels, a tile scheduler, or Cluster Launch Control on B200 to handle uneven tails. B200 Cluster Persistent Scheduler is an agent skill from mirage-project/mirage. Use when the user wants to use Thread Block Cluster, DSMEM, 2-CTA cooperative MMA, persistent kernels, a tile scheduler, or Cluster Launch Control on B200 to handle uneven tails.

When should I use B200 Cluster Persistent Scheduler?

B200 Cluster Persistent Scheduler fits situations like: the user wants to use Thread Block Cluster; 2-CTA cooperative MMA; persistent kernels; A tile scheduler.

How do I install B200 Cluster Persistent Scheduler in Claude Code?

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

How do I install B200 Cluster Persistent Scheduler in Codex?

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

Can I use B200 Cluster Persistent Scheduler 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-cluster-persistent-scheduler -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-cluster-persistent-scheduler, .gemini/skills/b200-cluster-persistent-scheduler, .github/skills/b200-cluster-persistent-scheduler and .opencode/skills/b200-cluster-persistent-scheduler in your project.

What does B200 Cluster Persistent Scheduler need to run?

SKILL.md names no scripts, command-line tools or credentials: B200 Cluster Persistent Scheduler is instructions for the agent only.

Does B200 Cluster Persistent Scheduler 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 Cluster Persistent Scheduler 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 Cluster Persistent Scheduler use?

B200 Cluster Persistent Scheduler 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 Cluster Persistent Scheduler use?

About 1.8k tokens (SKILL.md is roughly 7.2k 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 Cluster Persistent Scheduler?

Skills that share tags, products or a category with B200 Cluster Persistent Scheduler: Hunting For Scheduled Task Persistence (mukul975/Anthropic-Cybersecurity-Skills, 34k stars), Schedule (asgeirtj/system_prompts_leaks, 69k stars), Schedule (asgeirtj/system_prompts_leaks, 69k stars) and Time Blocking Scheduler (iflytek/skillhub, 5.2k stars). The comparison table on this page puts their stars, adoption, token cost, safety result and licence side by side.

Who maintains B200 Cluster Persistent Scheduler?

mirage-project (a GitHub organization) maintains it in mirage-project/mirage, which has 2,543 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.