{"id":"1378ab14-77ca-4a59-8e44-de65d87377f5","entityType":"agent","slug":"clawhub-wbavon-kernelgen-flagos","name":"Kernelgen Flagos","canonicalUrl":"https://www.xpersona.co/agent/clawhub-wbavon-kernelgen-flagos","canonicalPath":"/agent/clawhub-wbavon-kernelgen-flagos","generatedAt":"2026-10-15T03:21:35.702Z","source":"CLAWHUB","claimStatus":"UNCLAIMED","verificationTier":"NONE","summary":{"evidence":{"source":"CLAWHUB","verified":false,"confidence":"medium","updatedAt":"2026-10-11T23:08:28.846Z","emptyReason":null},"description":"Unified GPU kernel operator generation and optimization skill. Automatically detects the target repository type (FlagGems, vLLM, or general Python/Triton) an...","descriptionLabel":"Source description","evidenceSummary":"Capability contract not published. No trust telemetry is available yet. 979 downloads reported by the source. Last updated 10/11/2026.","installCommand":"clawhub skill install s179mcxns07mhashdtzg4tbf6584wt8v:kernelgen-flagos","sourceUrl":"https://clawhub.ai/wbavon/kernelgen-flagos","homepage":"https://clawhub.ai/wbavon/skills/kernelgen-flagos","primaryLinks":[{"label":"View on ClawHub","url":"https://clawhub.ai/wbavon/kernelgen-flagos","kind":"source"},{"label":"Homepage","url":"https://clawhub.ai/wbavon/skills/kernelgen-flagos","kind":"homepage"}],"safetyScore":84,"overallRank":62,"popularityScore":60,"trustScore":null,"claimedByName":null,"isOwner":false,"seoDescription":"Kernelgen Flagos technical dossier on Xpersona with agent coverage, OPENCLEW support, and live trust metadata."},"coverage":{"evidence":{"source":"public-profile","verified":false,"confidence":"medium","updatedAt":"2026-10-11T23:08:28.846Z","emptyReason":null},"protocols":[{"protocol":"OPENCLEW","label":"OpenClaw","status":"self-declared","notes":"Declared in the public agent profile."}],"capabilities":[],"verifiedCount":0,"selfDeclaredCount":1,"capabilityMatrix":{"rows":[{"key":"OPENCLEW","type":"protocol","support":"unknown","confidenceSource":"profile","notes":"Listed on profile"}],"flattenedTokens":"protocol:OPENCLEW|unknown|profile"}},"adoption":{"evidence":{"source":"CLAWHUB","verified":false,"confidence":"medium","updatedAt":"2026-10-11T23:08:28.846Z","emptyReason":null},"stars":null,"forks":null,"downloads":979,"likes":null,"task":null,"library":null,"packageName":null,"latestVersion":"1.1.0","tractionLabel":"979 downloads"},"release":{"evidence":{"source":"CLAWHUB","verified":false,"confidence":"medium","updatedAt":"2026-10-11T23:08:28.833Z","emptyReason":null},"lastUpdatedAt":"2026-10-11T23:08:28.846Z","lastCrawledAt":"2026-10-11T23:08:28.833Z","lastIndexedAt":null,"nextCrawlAt":"2026-10-12T23:08:28.833Z","lastVerifiedAt":null,"highlights":[{"version":"1.1.0","createdAt":"2026-05-19T07:04:17.149Z","changelog":"Sync from flagos-ai/skills latest","fileCount":16,"zipByteSize":115476},{"version":"1.0.0","createdAt":"2026-03-23T04:33:20.901Z","changelog":"- Initial release of kernelgen-flagos as a unified GPU kernel operator generation skill. - Automatically detects repository type (FlagGems, vLLM, or general Python/Triton) and dispatches to the appropriate specialized workflow. - Combines four sub-skills: kernelgen-general, kernelgen-for-flaggems, kernelgen-for-vllm, and kernelgen-submit-feedback, removing the need for separate installations. - Enforces mandatory use of the MCP tool for all operator code generation. - Supports user-invoked feedback and bug report submission via the bundled feedback sub-skill. - Enhances user experience with clear routing, explicit detection rules, and a quick reference guide.","fileCount":9,"zipByteSize":68655}]},"execution":{"evidence":{"source":"CLAWHUB","verified":false,"confidence":"low","updatedAt":null,"emptyReason":"No published capability contract is available yet."},"installCommand":"clawhub skill install s179mcxns07mhashdtzg4tbf6584wt8v:kernelgen-flagos","setupComplexity":"low","setupSteps":["Install using `clawhub skill install s179mcxns07mhashdtzg4tbf6584wt8v:kernelgen-flagos` in an isolated environment before connecting it to live workloads.","No published capability contract is available yet, so validate auth and request/response behavior manually.","Review the upstream CLAWHUB listing at https://clawhub.ai/wbavon/kernelgen-flagos before using production credentials."],"contract":{"contractStatus":"missing","authModes":[],"requires":[],"forbidden":[],"supportsMcp":false,"supportsA2a":false,"supportsStreaming":false,"inputSchemaRef":null,"outputSchemaRef":null,"dataRegion":null,"contractUpdatedAt":null,"sourceUpdatedAt":null,"freshnessSeconds":null},"invocationGuide":{"preferredApi":{"snapshotUrl":"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/snapshot","contractUrl":"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/contract","trustUrl":"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/trust"},"curlExamples":["curl -s \"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/snapshot\"","curl -s \"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/contract\"","curl -s \"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/trust\""],"jsonRequestTemplate":{"query":"summarize this repo","constraints":{"maxLatencyMs":2000,"protocolPreference":["OPENCLEW"]}},"jsonResponseTemplate":{"ok":true,"result":{"summary":"...","confidence":0.9},"meta":{"source":"CLAWHUB","generatedAt":"2026-10-15T03:21:35.700Z"}},"retryPolicy":{"maxAttempts":3,"backoffMs":[500,1500,3500],"retryableConditions":["HTTP_429","HTTP_503","NETWORK_TIMEOUT"]}},"endpoints":{"dossierUrl":"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/dossier","snapshotUrl":"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/snapshot","contractUrl":"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/contract","trustUrl":"https://www.xpersona.co/api/v1/agents/clawhub-wbavon-kernelgen-flagos/trust"}},"reliability":{"evidence":{"source":"runtime-metrics","verified":false,"confidence":"low","updatedAt":null,"emptyReason":"No trust, reliability, or runtime telemetry is available."},"trust":{"status":"unavailable","handshakeStatus":"UNKNOWN","verificationFreshnessHours":null,"reputationScore":null,"p95LatencyMs":null,"successRate30d":null,"fallbackRate":null,"attempts30d":null,"trustUpdatedAt":null,"trustConfidence":"unknown","sourceUpdatedAt":null,"freshnessSeconds":null},"decisionGuardrails":{"doNotUseIf":["Contract metadata is missing or unavailable for deterministic execution."],"safeUseWhen":[],"riskFlags":["missing_or_unavailable_contract","trust_data_unavailable","schema_references_missing"],"operationalConfidence":"low"},"executionMetrics":{"observedLatencyMsP50":null,"observedLatencyMsP95":null,"estimatedCostUsd":null,"uptime30d":null,"rateLimitRpm":null,"rateLimitBurst":null,"lastVerifiedAt":null,"verificationSource":null},"runtimeMetrics":{"successRate":null,"avgLatencyMs":null,"avgCostUsd":null,"hallucinationRate":null,"retryRate":null,"disputeRate":null,"p50Latency":null,"p95Latency":null,"lastUpdated":null}},"benchmarks":{"evidence":{"source":"no-benchmark-data","verified":false,"confidence":"low","updatedAt":null,"emptyReason":"No benchmark suites or observed failure patterns are available."},"suites":[],"failurePatterns":[]},"artifacts":{"evidence":{"source":"CLAWHUB","verified":false,"confidence":"medium","updatedAt":"2026-10-11T23:08:28.846Z","emptyReason":null},"readme":"Skill: Kernelgen Flagos\n\nOwner: wbavon\n\nSummary: Unified GPU kernel operator generation and optimization skill. Automatically detects the target repository type (FlagGems, vLLM, or general Python/Triton) an...\n\nTags: latest:1.1.0\n\nVersion history:\n\nv1.1.0 | 2026-05-19T07:04:17.149Z | user\n\nSync from flagos-ai/skills latest\n\nv1.0.0 | 2026-03-23T04:33:20.901Z | auto\n\n- Initial release of kernelgen-flagos as a unified GPU kernel operator generation skill.\n- Automatically detects repository type (FlagGems, vLLM, or general Python/Triton) and dispatches to the appropriate specialized workflow.\n- Combines four sub-skills: kernelgen-general, kernelgen-for-flaggems, kernelgen-for-vllm, and kernelgen-submit-feedback, removing the need for separate installations.\n- Enforces mandatory use of the MCP tool for all operator code generation.\n- Supports user-invoked feedback and bug report submission via the bundled feedback sub-skill.\n- Enhances user experience with clear routing, explicit detection rules, and a quick reference guide.\n\nArchive index:\n\nArchive v1.1.0: 16 files, 115476 bytes\n\nFiles: kernelgen-generate-for-flaggems.md (49340b), kernelgen-generate-for-vllm.md (47402b), kernelgen-generate.md (50179b), kernelgen-mcp-setup.md (3207b), kernelgen-optimize-for-flaggems.md (29665b), kernelgen-optimize-for-vllm.md (34928b), kernelgen-optimize.md (20770b), kernelgen-specialize-for-flaggems.md (18758b), kernelgen-specialize.md (17974b), kernelgen-submit-feedback.md (11593b), LICENSE.txt (11262b), README_zh.md (9748b), README.md (10949b), skill-card.md (3144b), SKILL.md (8355b), _meta.json (135b)\n\nFile v1.1.0:SKILL.md\n\n---\nname: kernelgen-flagos\ndescription: >\n Unified GPU kernel operator generation and optimization skill. Automatically detects the target\n repository type (FlagGems, vLLM, or general Python/Triton) and dispatches to the appropriate\n specialized sub-skill. Includes operator generation, MCP-based iterative optimization, and\n feedback submission sub-skills. Use this skill when the user wants to generate or optimize a\n GPU kernel operator, create a Triton kernel, or says things like \"generate an operator\",\n \"create a kernel for X\", \"optimize triton kernel\", or \"/kernelgen-flagos\".\nargument-hint: \" [--func-type ]\"\nuser-invokable: true\ncompatibility: \"Python 3.8+, PyTorch with CUDA, Triton\"\nmetadata:\n version: \"1.0.0\"\n author: flagos-ai\n category: gpu-kernel-generation\n tags: [kernelgen, triton, gpu, mcp, operator-generation, operator-optimization, flaggems, vllm, feedback]\nallowed-tools:\n - Bash\n - Bash(gh:*)\n - Bash(python:*)\n - Bash(python3:*)\n - Bash(command:*)\n - Read\n - Write\n - Edit\n - Glob\n - Grep\n - AskUserQuestion\n---\n\n# kernelgen-flagos — Unified GPU Operator Generation Skill\n\nThis is a **unified entry point** that bundles generation and optimization sub-skills into one:\n\n| Sub-skill file | Purpose |\n|---|---|\n| **Generation** | |\n| `kernelgen-generate.md` | Generate GPU kernels for **any** Python/Triton repository |\n| `kernelgen-generate-for-flaggems.md` | Specialized generation for **FlagGems** repositories |\n| `kernelgen-generate-for-vllm.md` | Specialized generation for **vLLM** repositories |\n| **Optimization** | |\n| `kernelgen-optimize.md` | Optimize existing Triton kernels via MCP iterative optimization (general purpose) |\n| `kernelgen-optimize-for-flaggems.md` | Optimize Triton operators and integrate into **FlagGems** (3 modes: built-in/external/experimental) |\n| `kernelgen-optimize-for-vllm.md` | Optimize Triton operators and integrate into **vLLM** (with CustomOp registration) |\n| **Platform Specialization** | |\n| `kernelgen-specialize.md` | Specialize Triton operators to target platforms (e.g., GPU → Ascend NPU) via MCP `specialize_kernel` |\n| `kernelgen-specialize-for-flaggems.md` | Platform specialization + **FlagGems** integration (4 modes: vendor-ops/vendor-fused/override-builtin/experimental) |\n| **MCP Configuration** | |\n| `kernelgen-mcp-setup.md` | Check and auto-configure the `kernelgen-server` MCP service (URL built-in, user only provides Token) |\n| **Feedback** | |\n| `kernelgen-submit-feedback.md` | Submit bug reports and feedback via GitHub or email |\n\nAll sub-skill files are located in the **same directory** as this `SKILL.md` file.\n\n---\n\n## Routing Protocol — Follow This BEFORE Doing Anything Else\n\n### Phase 0: MCP Configuration Check\n\nBefore anything else, ensure the `kernelgen-server` MCP service is configured and ready.\n\nUse the Glob tool to find `kernelgen-mcp-setup.md` in this skill's directory:\n\n```\nGlob: **/skills/kernelgen-flagos/kernelgen-mcp-setup.md\n```\n\nThen use the Read tool to read the matched file and **follow its instructions exactly**.\n\n- If MCP is already configured → proceed to Phase 1.\n- If MCP is not configured → the setup skill will guide the user through configuration.\n Once configuration is written and the user is prompted to restart, **stop here** — do not\n continue to Phase 1.\n\n### Phase 1: Detect Repository Type\n\nUse the Glob tool to check for project identity files in the current working directory:\n\n```\nGlob: pyproject.toml\nGlob: setup.py\nGlob: setup.cfg\n```\n\nThen use the Read tool to read whichever file exists. Determine the **project name** from\nthe file contents (e.g., `name = \"flag_gems\"` in pyproject.toml, or `name='vllm'` in setup.py).\n\nAlso use the Glob tool to check for characteristic directory structures:\n\n**FlagGems indicators** (match ANY):\n- `src/flag_gems/` directory exists\n- Project name is `flag_gems` or `flag-gems` or `FlagGems`\n- `import flag_gems` appears in test files\n\n**vLLM indicators** (match ANY):\n- `vllm/` directory exists at the repo root (with `vllm/__init__.py`)\n- Project name is `vllm`\n- `csrc/` directory exists alongside `vllm/`\n\n### Phase 2: Dispatch to Sub-skill\n\nBased on the detection result, use the **Read tool** to read the appropriate sub-skill file\nfrom this skill's directory, then **follow the instructions in that file exactly**.\n\n**To locate the sub-skill files**: They are in the same directory as this SKILL.md. Use the\nGlob tool to find the path:\n\n```\nGlob: **/skills/kernelgen-flagos/kernelgen-generate.md\n```\n\nThen use the Read tool to read the matched path.\n\n#### Decision Table\n\n**Generation requests** (user wants to create/generate a new operator):\n\n| Detection Result | Action |\n|---|---|\n| FlagGems repository detected | Read `kernelgen-generate-for-flaggems.md` and follow it |\n| vLLM repository detected | Read `kernelgen-generate-for-vllm.md` and follow it |\n| Neither detected (or unknown) | Read `kernelgen-generate.md` and follow it |\n\n**Optimization requests** (user wants to optimize an existing operator, mentions \"optimize\", \"speedup\", \"improve performance\"):\n\n| Detection Result | Action |\n|---|---|\n| FlagGems repository detected | Read `kernelgen-optimize-for-flaggems.md` and follow it |\n| vLLM repository detected | Read `kernelgen-optimize-for-vllm.md` and follow it |\n| Neither detected (or unknown) | Read `kernelgen-optimize.md` and follow it |\n\n**Specialization requests** (user wants to migrate/specialize an operator to a different platform, mentions \"specialize\", \"migrate to Ascend/NPU\", \"platform migration\"):\n\n| Detection Result | Action |\n|---|---|\n| FlagGems repository detected | Read `kernelgen-specialize-for-flaggems.md` and follow it |\n| Neither detected (or unknown) | Read `kernelgen-specialize.md` and follow it |\n\n**Feedback requests**:\n\n| Detection Result | Action |\n|---|---|\n| User reports a bug or requests feedback submission | Read `kernelgen-submit-feedback.md` and follow it |\n\n**Important rules:**\n1. **Always detect first, dispatch second.** Never skip detection.\n2. **Read the entire sub-skill file** before starting execution — do not partially read it.\n3. **Follow the sub-skill instructions exactly** as if they were the main SKILL.md. All steps,\n rules, and protocols in the sub-skill apply fully.\n4. **Do not mix sub-skills.** Once you dispatch to a sub-skill, follow it to completion.\n5. If the user explicitly requests a specific sub-skill (e.g., \"use the FlagGems version\"),\n honor that request regardless of auto-detection results.\n6. **CRITICAL — MCP is mandatory**: ALL operator code generation MUST go through the\n `mcp__kernelgen-mcp__generate_kernel` MCP tool. Optimization uses\n `mcp__kernelgen-mcp__optimize_kernel`, and platform specialization uses\n `mcp__kernelgen-mcp__specialize_kernel`. NEVER generate Triton kernels, PyTorch\n wrappers, or operator implementations yourself. If MCP is not configured, not reachable,\n or fails after all retries, STOP and report the issue — do NOT fall back to writing code\n manually.\n\n### Phase 3: Feedback Handling\n\nAt **any point** during the workflow, if the user reports a bug, says something is broken,\nor asks to submit feedback about the skill:\n\n1. Use the Read tool to read `kernelgen-submit-feedback.md` from this skill's directory.\n2. Follow the feedback submission workflow described in that file.\n3. After feedback is submitted, ask the user if they want to continue with the operator\n generation workflow or stop.\n\n---\n\n## Quick Reference for Users\n\n```bash\n# === Generation ===\n# Generate a kernel operator (auto-detects repo type)\n/kernelgen-flagos relu\n\n# Generate with explicit function type\n/kernelgen-flagos rms_norm --func-type normalization\n\n# === Optimization ===\n# Optimize an existing Triton kernel (auto-detects repo type)\n# Just say \"optimize the relu kernel\" or \"improve kernel performance\"\n# The skill will automatically dispatch to the right optimization sub-skill\n\n# The skill will automatically:\n# - Detect if you're in a FlagGems repo → use FlagGems-specific workflow\n# - Detect if you're in a vLLM repo → use vLLM-specific workflow\n# - Otherwise → use the general-purpose workflow\n```\n\nIf you encounter any issues during generation, just say \"submit feedback\" or \"report a bug\"\nand the skill will guide you through the feedback submission process.\n\nFile v1.1.0:README.md\n\n# kernelgen-flagos: Unified GPU Kernel Generation Skills\n\n[中文版](README_zh.md)\n\n## Overview\n\nThe KernelGen `kernelgen-flagos` skill is a unified AI programming skill launched by FlagOS. By leveraging the `kernelgen-mcp` KernelGen Operator Development MCP Toolkit service, it is compatible with various domestic and international chip platforms to generate high-precision GPU kernel operators. It automatically detects the type of the target code repository and dispatches it to the corresponding dedicated workflow.\n\n## Quick Start\n\nThis section introduces how to quickly start a task, for example, creating, optimizing, specialize Kernels for a specific platform, or submitting feedback. You can use either prompt or the command line to execute these tasks. In this section, we only teach you how to use prompts. For how to use command lines or manual configurations, see [KernelGen Skills User Guide](https://docs.flagos.io/projects/kernelgen/en/latest/skills_user_guide/skills-user-guide.html#).\n\nPerform all of the following steps except for obtaining the KernelGen Token in your agent client, for example, OpenClaw, VSCode with activated Github Copilot, or Claude Code.\n\n### Prerequisites\n\n* Claude Code version 2.1 and later\n\n* OpenClaw version 2026.3.2 and later\n\n* VSCode with Github Copilot activated\n\n### Steps\n\n1. Obtain the **KernelGen Token** from: [https://kernelgen.flagos.io/mcp](https://kernelgen.flagos.io/mcp).\n\n2. In your agent client, send a prompt to connect to the KernelGen Operator Development MCP Toolkit, for example:\n\n ```{code-block} shell\n Based on the Claude Code configuration documentation: https://code.claude.com/docs/en/mcp, connect to the MCP. The MCP URL is https://kernelgen.flagos.io/sse, and the token is . Configure this in the claude.json file.\n ```\n\n or\n\n ```{code-block} shell\n Based on the VSCode documentation: https://code.visualstudio.com/docs/copilot/customization/mcp-servers, configure the kernelgen MCP. The MCP URL is https://kernelgen.flagos.io/sse, and the token is . Configure this in the mcp.json file.\n ```\n\n **Note**: You may need to restart your agent to let the settings take effect. Check this in the documentation of the relevant AI agent.\n\n3. Verify KernelGen Operator Development MCP Toolkit connection, send a prompt:\n\n ```{code-block} shell\n Please verify the kernelgen mcp connection is successful.\n ```\n\n4. Send a prompt to setup the `kernelgen-flagos` unified skill, including all sub-skills, for example:\n\n * **Copilot**:\n \n ```{code-block} shell\n Setup kernelgen-flagos skills in Copilot from .\n ```\n\n * **Other agent clients**:\n\n ```{code-block} shell\n Setup kernelgen-flagos skills from .\n ```\n\n5. Verify skills are successfully installed，send a prompt\n\n ```{code-block} shell\n Please verify if the kernelgen-flagos skills are working correctly.\n ```\n\n6. Send a prompt to run a task by invoking the corresponding sub-skill.\n\n### Send a prompt to run a task by invoking the sub-skill\n\nKernelGen automatically detects whether the project is a FlagGems repository, a vLLM repository, or another type, and executes the corresponding workflow.\n\n* **Use Case 1**: Create an Kernel for the vLLM project：\n\n 1. Preinstall [vLLM](https://docs.vllm.ai/en/latest/getting_started/installation/).\n\n **Note**: Please install from source via git clone.\n\n 2. Send a prompt to generate a Kernel, for example:\n\n ```{code-block} shell\n Use kernelgen-flagos to generate the ReLU operator. Use MetaX. Integrated into vLLM.\n ```\n\n or\n\n ```{code-block} shell\n /kernelgen-flagos Generate the ReLU operator. Use MetaX. Integrated into vLLM.\n ```\n\n* **Use Case 2**: Send a prompt to optimize an Kernel on NVIDIA, for example: \n \n ```{code-block} shell\n Use kernelgen-flagos to optimize the scaled_dot_product_attention_math operator. Optimize 5 iterations. The scaled_dot_product_attention_math operator is located at .\n ```\n\n or\n\n ```{code-block} shell\n /kernelgen-flagos Optimize the scaled_dot_product_attention_math operator. Optimize 5 iterations. The scaled_dot_product_attention_math operator is located at .\n ```\n\n* **Use Case 3**: Send a prompt to specialize the CUDA-implemented Kernel for Huawei Ascend for the FlagGems project：\n\n 1. Preinstall [FlagGems](https://docs.flagos.io/projects/FlagGems/en/latest/getting_started/install.html#).\n\n **Note**: Please install from source via git clone.\n\n 2. Send a prompt to specialize a Kernel, for example:\n\n ```{code-block} shell\n Use kernelgen-flagos to migrate the CUDA-implemented operator to the Ascend chip, with the operator file stored in the FlagGems repository, and the directory is , ensuring that the accuracy verification passes.\n ```\n\n or\n\n ```{code-block} shell\n /kernelgen-flagos Migrate the CUDA-implemented operator to the Ascend chip, with the operator file stored in the FlagGems repository, and the directory is , ensuring that the accuracy verification passes.\n ```\n\n### Output files\n\nThe output files are as follows:\n\n* \\\\_triton.py - Triton kernel implementation\n\n* \\\\_torch.py - PyTorch baseline implementation\n\n* test\\_\\.py - Complete test cases\n\n* benchmark\\_\\.py - Performance benchmarking\n\n**Note**: If you have other files generated, send a prompt to verify the skills and KernelGen Operator Development MCP Toolkit are still working.\n\nFurther steps:\n\n* Pull Requests: If the user is developing within a GitHub repository, the skill will ask if a Pull Request needs to be submitted. It can then automatically submit the PR and provide the link.\n\n* Chat Integration: If the user configures the skill via chat applications (such as WeChat or Feishu), the skill will ask if the modified files should be sent to the chat application.\n\n* Feedback: If you encounter bugs during usage, you can use this skill to submit feedback via email or by opening a GitHub issue. send a prompt to submit feedback, for example: \"Use the kernelgen-flagos skill to submit feedback\" or \"Use the kernelgen-flagos skill to report a bug\".\n\n## Why Choose Skill?\n\nKernelGen Skill doesn't rely solely on LLMs to generate code from scratch. Instead, it follows a structured four-step workflow:\n\n1. **Collect Kernel Information**: Extracts operator parameters (input/output types, shapes, computation logic) from the user's natural language description.\n\n2. **Search for Similar Code Snippets**: Retrieves similar implementations from the existing codebase for reference (users can choose whether to adopt them).\n\n3. **Generate Kernel Code + CUDA Implementation**: Simultaneously generates the Triton Kernel and its corresponding CUDA implementation, with the latter serving as the PyTorch benchmark.\n\n4. **Test**: Uses the CUDA implementation as the ground truth to test the Kernel, outputting Correctness and Speedup Ratio.\n\nOnce generation is complete, the path from generation to contribution is fully automated: if developing within the FlagGems repository, a specialized sub-skill automatically handles format conversion and directory alignment. It directly injects the four core deliverables (kernel, baseline, accuracy test, speedup test) into the `src/flag_gems/experimental_ops/` directory. Developers can generate a compliant PR with one click, without needing to manually invoke conversion scripts.\n\n### Core Features\n\nWriting high-performance GPU operators is complex and error-prone. Different projects (FlagGems, vLLM, custom Triton repositories) each have unique operator implementation specifications, testing patterns, and registration systems. Previously, users needed to install multiple skills separately.\n\nThis unified skill packages ten sub-skills into one, requiring installation only once:\n\n### File Descriptions\n\n### `SKILL.md`\n\nThe unified entry point. Contains routing logic that auto-detects the repository type (FlagGems, vLLM, or generic) and reads the appropriate sub-skill file to execute.\n\n### `kernelgen-mcp-setup.md`\n\nMCP service configuration check and auto-setup. Before any sub-skill executes, SKILL.md dispatches here to verify that the kernelgen-mcp MCP service is configured. If not, guides the user through token acquisition and writes the configuration automatically.\n\n### `kernelgen-generate.md`\n\nFull 10-step workflow for generating GPU kernel operators in any Python/Triton repository. Includes dynamic repo structure discovery, convention detection, and adaptive code placement.\n\n### `kernelgen-generate-for-flaggems.md`\n\n9-step workflow specialized for FlagGems repositories. Handles `pointwise_dynamic` wrappers, promotion methods, `_FULL_CONFIG` registration, categorized test files, and FlagGems-specific conventions.\n\n### `kernelgen-generate-for-vllm.md`\n\n9-step workflow specialized for vLLM repositories. Handles SPDX license headers, `vllm.logger.init_logger`, `@triton.autotune`, custom op registration, and vLLM directory conventions.\n\n### `kernelgen-submit-feedback.md`\n\nFeedback submission workflow. Collects bug reports with auto-detected environment info and submits via GitHub Issues (`gh` CLI) or email fallback.\n\n### `kernelgen-optimize.md`\n\nGeneral-purpose Triton kernel optimization via MCP iterative loop. Analyzes existing kernels, identifies bottlenecks, and applies optimizations through multiple rounds until the target speedup is reached.\n\n### `kernelgen-optimize-for-flaggems.md`\n\nFlagGems-specific kernel optimization with 3 modes: optimize built-in operators in-place, optimize external operators and integrate into experimental\\_ops, or optimize existing experimental operators. Includes accuracy tests and performance benchmarks.\n\n### `kernelgen-optimize-for-vllm.md`\n\nvLLM-specific kernel optimization with CustomOp registration, accuracy tests, and performance benchmark integration. Optimizes Triton operators and automatically integrates them into the vLLM project.\n\n### `kernelgen-specialize.md`\n\nPlatform specialization for Triton operators via MCP `specialize_kernel` tool. Migrates GPU Triton operators to target platforms (e.g., Huawei Ascend NPU), handling architecture differences, Grid configuration, UB overflow, and memory alignment.\n\n### `kernelgen-specialize-for-flaggems.md`\n\nCombines MCP platform specialization with FlagGems framework integration. Supports four integration modes: vendor-ops (default), vendor-fused, override-builtin, and experimental. Includes automated testing and performance benchmarking.\n\n***\n\n## License\n\nThis project is licensed under the Apache 2.0 License. See [LICENSE.txt](https://github.com/flagos-ai/skills/blob/main/skills/kernelgen-flagos/LICENSE.txt) for details.\n\nFile v1.1.0:_meta.json\n\n{\n \"ownerId\": \"kn7c1yzknmhkmxasrtdbxa3fhd81z80k\",\n \"slug\": \"kernelgen-flagos\",\n \"version\": \"1.1.0\",\n \"publishedAt\": 1779174257149\n}\n\nFile v1.1.0:kernelgen-generate-for-flaggems.md\n\n# KernelGen Skill — Generate FlagGems Operators via MCP\n\nYou are an expert at generating GPU kernel operators for the FlagGems project using the `kernelgen-mcp` MCP service.\n\n**Tool usage**: This skill relies on the following capabilities:\n- **Shell**: execute shell commands (python, pip, pytest, etc.) via the Bash/shell tool\n- **Read**: read files from the codebase\n- **Write / Edit**: create or modify files\n- **Grep**: search file contents\n- **Glob**: find files by pattern\n- **MCP tools**: `mcp__kernelgen-mcp__generate_kernel` and `mcp__kernelgen-mcp__optimize_kernel`\n\n> **⚠️ MCP Prerequisite**: MCP configuration is handled by `SKILL.md` Phase 0. If MCP tools fail during execution,\n> stop and tell the user to visit https://kernelgen.flagos.io/mcp 获取 Token 后重新配置。\n\nWhen you need user input or clarification, ask a question directly and wait for their reply.\nAlways use the appropriate built-in tool rather than outputting commands for the user to run manually.\n\n## Execution Rules\n\n1. **Follow the steps in order**. Never jump ahead.\n2. **Never skip Step 2** (operator existence check) — always verify before generating.\n3. **Do not generate or write any code before Step 4** (MCP call). Steps 0–3 are preparation only.\n4. **Always confirm with the user before destructive actions**.\n5. **Report progress to the user when entering each step**.\n6. **Never fabricate repository files or paths**.\n7. **CRITICAL — MCP is mandatory**: ALL operator code generation MUST go through the\n `mcp__kernelgen-mcp__generate_kernel` MCP tool. NEVER generate Triton kernels, PyTorch\n wrappers, or operator implementations yourself — even if the operator seems simple (e.g.,\n relu, abs). If MCP is not configured, not reachable, or fails after all retries in Step 4,\n STOP and report the issue to the user. Do NOT fall back to writing kernel code manually.\n The MCP service produces optimized, tested code that manual writing cannot match.\n\n---\n\n## FlagGems Project Layout Reference\n\nThe FlagGems project has three operator locations and corresponding test/benchmark locations:\n\n### Operator Locations\n| Location | Path | Description |\n|---|---|---|\n| **Core (src)** | `src/flag_gems/ops/.py` | Mature, production-quality operators. Uses `pointwise_dynamic` decorators for pointwise ops. Registered in `_FULL_CONFIG` for aten dispatch. |\n| **Experimental** | `src/flag_gems/experimental_ops/.py` | New/experimental operators. Uses raw Triton pointer-based style (self-contained, no `pointwise_dynamic`). NOT registered in `_FULL_CONFIG`. |\n| **Backend** | `src/flag_gems/runtime/backend/_/ops/.py` | Non-NVIDIA chip-specific operators. Vendor names: `_ascend` (Huawei), `_cambricon`, `_hygon` (DCU), `_mthreads` (Moore), `_iluvatar` (Tianshu), `_metax`, `_amd`, `_kunlunxin`, `_arm`, `_aipu`. |\n\n### Test Locations\n| Operator Location | Accuracy Tests | Performance Tests |\n|---|---|---|\n| **Core (src)** | `tests/test__ops.py` (append to existing file) | `benchmark/test__perf.py` (append to existing file) |\n| **Experimental** | `experimental_tests/_test.py` (standalone file) | Same file as accuracy test (includes benchmark function) |\n| **Backend** | Backend-specific test directory (follow existing pattern per vendor) | Backend-specific benchmark directory |\n\n### Test Style Conventions (CRITICAL)\n\n**`tests/` folder style** — uses `flag_gems.use_gems()` context manager and calls ops via `torch.()`:\n```python\nimport flag_gems\nfrom .accuracy_utils import gems_assert_close, to_reference, POINTWISE_SHAPES, FLOAT_DTYPES\n\ndef test_accuracy_(shape, dtype):\n inp = torch.randn(shape, dtype=dtype, device=flag_gems.device)\n ref_inp = to_reference(inp)\n ref_out = (ref_inp)\n with flag_gems.use_gems():\n res_out = torch.(inp)\n gems_assert_close(res_out, ref_out, dtype)\n```\n\n**`experimental_tests/` folder style** — uses direct import from `flag_gems.experimental_ops`:\n```python\nimport flag_gems\nfrom flag_gems.experimental_ops. import as gems_op\n\ndef test__tensor(shape, dtype):\n x = torch.randn(shape, dtype=dtype, device=flag_gems.device)\n ref_x = to_reference(x)\n ref_out = torch.ops.aten.(ref_x)\n with flag_gems.use_gems():\n act_out = gems_op(x)\n gems_assert_close(act_out, ref_out, dtype=dtype)\n```\n\n**IMPORTANT**: Before writing a test, always read 1-2 existing test files in the **same folder** to match the exact style (imports, utility functions, assertion helpers, parametrize decorators, etc.). Never mix styles between folders.\n\n### Decorator Conventions\n\n**Core ops (`src/flag_gems/ops/`)** use FlagGems decorators:\n- `@pointwise_dynamic(promotion_methods=...)` + `@triton.jit` — for pointwise ops\n- The kernel receives **scalar elements** (not pointers), no `tl.load`/`tl.store`\n- See `src/flag_gems/ops/abs.py`, `src/flag_gems/ops/sigmoid.py` for examples\n\n**Experimental ops (`src/flag_gems/experimental_ops/`)** use raw Triton style:\n- `@triton.jit` only — standard Triton pointer-based kernels\n- The kernel receives **pointers**, uses `tl.load`/`tl.store`, `BLOCK_SIZE`, masks\n- Self-contained, no `flag_gems.utils` imports needed\n- See `src/flag_gems/experimental_ops/abs.py` for example\n\n**When adapting MCP output**: The MCP always generates raw pointer-based Triton code.\n- For **core ops**: MUST rewrite to `pointwise_dynamic` style — strip all pointer logic\n- For **experimental ops**: Keep the MCP Triton code mostly as-is (it's already pointer-based)\n- For **backend ops**: Follow the existing backend vendor patterns\n\n---\n\n## Step 0: Pre-flight — Environment & MCP Check\n\n### 0a. Check Python Environment\n\nRun the following diagnostic command:\n\n```bash\npython -c \"import torch; import triton; import flag_gems; print('torch', torch.__version__); print('triton', triton.__version__); print('flag_gems', flag_gems.__version__); print('device', flag_gems.device)\"\n```\n\n**If any import fails**, identify what's missing and use the shell tool to install it:\n\n| Missing package | Install command |\n|---|---|\n| `torch` | Do NOT auto-install. Ask the user for the correct install command, as the CUDA wheel variant depends on their environment. |\n| `triton` | `pip install triton` |\n| `flag_gems` | `pip install -e .` from the repo root, or `pip install flag_gems` |\n| `pytest` | `pip install pytest` |\n| `numpy` | `pip install numpy` |\n| `scipy` | `pip install scipy` |\n| `pyyaml` | `pip install pyyaml` |\n| `packaging` | `pip install packaging` |\n\nIf `torch` is missing or has no CUDA support, run a separate GPU diagnostic:\n\n```bash\npython -c \"import torch; print('cuda_version:', torch.version.cuda); print('cuda_available:', torch.cuda.is_available())\"\n```\n\nThen diagnose:\n- If `torch.version.cuda` is `None` → torch is a CPU-only build. Ask the user to reinstall with CUDA.\n- If `torch.version.cuda` has a value but `torch.cuda.is_available() == False` → CUDA build but no GPU\n detected. Warn the user that GPU tests will fail and ask how to proceed.\n\nIf `flag_gems` is not installed, first verify `pyproject.toml` or `setup.py` exists in the current\nworking directory (to confirm we're at the repo root). If found, run `pip install -e .`. If not\nfound, ask the user for the correct repo root path.\n\nAfter installing, re-run the diagnostic via the shell tool to confirm everything works. Only proceed when all imports succeed.\n\n### 0b. Check MCP Availability\n\n> **MCP configuration is handled centrally by `SKILL.md` Phase 0 (via `kernelgen-mcp-setup.md`).**\n> By the time execution reaches this sub-skill, MCP should already be configured and ready.\n> If MCP tools are unavailable or calls fail at any point, stop and tell the user:\n> \"MCP 服务不可用，请检查配置或访问 https://kernelgen.flagos.io/mcp 获取 Token 后重新配置。\"\n\n**If the MCP server IS configured**, proceed to Step 1.\n\n## Step 1: Understand the Operator Request\n\nParse the user's description to determine:\n- `kernel_name`: operator name (e.g., `relu`, `softmax`, `gelu`, `layer_norm`).\n Prefer names exactly as defined in `torch.ops.aten` (e.g., `relu`, `silu`, `neg`),\n since registration uses the aten name.\n- `torch_call`: the actual Python call to invoke this op on a tensor. This is NOT always\n `torch.`. Determine it as follows:\n 1. If `torch.` exists (`hasattr(torch, '')` is True) → use `torch.`\n 2. If not, check `torch.nn.functional.` (e.g., `torch.nn.functional.layer_norm`)\n 3. If not, check `torch.special.` or `torch.linalg.`\n 4. As a last resort, use `torch.ops.aten..default`\n Do NOT use `torch._C._nn` — it is an internal API that can break across PyTorch versions.\n **Mandatory extra arguments**: If the chosen `torch_call` requires mandatory arguments beyond the\n input tensor(s), you MUST include them with sensible defaults. Common cases:\n - `torch.nn.functional.softmax(x, dim=-1)`\n - `torch.nn.functional.log_softmax(x, dim=-1)`\n - `torch.nn.functional.layer_norm(x, normalized_shape=x.shape[-1:])`\n - `torch.nn.functional.dropout(x, p=0.5, training=False)`\n - `torch.nn.functional.normalize(x, dim=-1)`\n When in doubt, check the existing tests in the repo or the PyTorch docs for required arguments.\n Omitting mandatory arguments (e.g., `torch.nn.functional.softmax(x)` without `dim`) will cause\n runtime errors.\n Store `` (including any mandatory extra arguments) — it will be used in tests and\n benchmarks wherever the reference op is needed.\n- `func_desc`: what the operator does\n- `func_type`: one of the following categories (aligned with FlagGems test structure).\n **Auto-infer from the operator signature when possible:**\n - If 1 tensor input, no dim arg → `unary_pointwise`\n - If operator has >=2 tensor inputs, no dim arg → `binary_pointwise`\n - If has a `dim` argument → `reduction`\n - If name contains `norm`, `softmax` (also matches `log_softmax`), `rmsnorm`, or `batchnorm` → `normalization`\n - If name is `mm`, `matmul`, `bmm`, `addmm` → `blas`\n - Otherwise → `other`\n Confirm the inferred `func_type` with the user if ambiguous. Categories:\n - `unary_pointwise` — single-input elementwise ops (relu, sigmoid, abs, etc.)\n - `binary_pointwise` — two-input elementwise ops (add, mul, etc.)\n - `reduction` — ops that reduce dimensions (sum, mean, max, etc.)\n - `normalization` — norm ops (layer_norm, batch_norm, group_norm, etc.)\n - `blas` — matrix operations (matmul, mm, bmm, etc.)\n - `other` — everything else (indexing, special, etc.)\n- `arg_names`, `arg_type`, `arg_descs`, `output_arg_desc`: parameter information\n- `target_device`: determine target device from user context:\n - If user mentions NVIDIA/英伟达/nv or doesn't specify → `\"nvidia\"` (default)\n - If user mentions 华为/昇腾/Ascend → `\"huawei\"`\n - If user mentions 海光/Hygon/DCU → `\"haiguang\"`\n - If user mentions 摩尔/Moore/摩尔线程 → `\"moore\"`\n - If user mentions 天数/Tianshu/天数智芯 → `\"tianshu\"`\n - If user mentions 寒武纪/Cambricon → `\"cambricon\"`\n\nIf the operator name is ambiguous, first check this common alias table before asking the user:\n\n| User says | Correct `kernel_name` (torch.ops.aten) |\n|---|---|\n| swish | `silu` |\n| negative / negate | `neg` |\n| power | `pow` |\n| square | `pow` (exp=2) |\n| cube | `pow` (exp=3) |\n| absolute | `abs` |\n| multiply | `mul` |\n| divide | `div` |\n| clip | `clamp` |\n| hard_swish | `hardswish` |\n| hard_sigmoid | `hardsigmoid` |\n| logarithm / ln | `log` |\n| exponential | `exp` |\n| hyperbolic_tangent | `tanh` |\n| square_root | `sqrt` |\n| cube_root | `cbrt` |\n| inverse_sqrt | `rsqrt` |\n| reciprocal / inverse | `reciprocal` |\n| minimum | `min` |\n| maximum | `max` |\n| floor_divide | `floor_divide` |\n| modulo / mod | `remainder` |\n\nIf the alias is not in this table, ask the user to clarify.\n\nAfter determining `func_type`, derive the `` value used for test/benchmark file paths:\n\n| func_type | category (for file paths) |\n|---|---|\n| `unary_pointwise` | `unary_pointwise` |\n| `binary_pointwise` | `binary_pointwise` |\n| `reduction` | `reduction` |\n| `normalization` | `norm` |\n| `blas` | `blas` |\n| `other` | `special` |\n\nStore this `` value — it determines which test and benchmark files to modify.\n\n## Step 1.5: Determine Placement (NEW — CRITICAL DECISION)\n\nBased on the user's request and the operator analysis, determine ``:\n\n### Decision Rules\n\n1. **Backend placement** (`backend`): If `target_device` is NOT `\"nvidia\"` (i.e., non-NVIDIA chip), the operator goes into the backend directory:\n - Operator: `src/flag_gems/runtime/backend/_/ops/.py`\n - Tests: Follow the existing backend vendor test pattern (read existing tests in that vendor's directory first)\n\n2. **Core placement** (`core`): If the user explicitly asks to **modify/replace an existing core operator** in `src/flag_gems/ops/`, or the user specifically says \"add to src/core\":\n - Operator: `src/flag_gems/ops/.py`\n - Tests: `tests/test__ops.py` (append to existing file)\n - Benchmarks: `benchmark/test__perf.py` (append to existing file)\n\n3. **Experimental placement** (`experimental`) — **DEFAULT for new operators**:\n - New operators should go here unless the user explicitly requests core placement\n - Operator: `src/flag_gems/experimental_ops/.py`\n - Tests: `experimental_tests/_test.py` (standalone file)\n - Benchmarks: Included in the same test file\n\n### Ask the user if ambiguous\n\nIf the placement is not obvious from context, ask:\n- \"This is a new operator. I recommend placing it in `experimental_ops/` (experimental). Would you prefer to place it in `src/flag_gems/ops/` (core) instead?\"\n- If the user mentions a non-NVIDIA chip, confirm: \"I'll place this in the `_` backend directory. Is that correct?\"\n\nStore `` as one of: `core`, `experimental`, `backend`.\n\n## Step 2: Check Whether the Operator Already Exists\n\nBefore calling the MCP generator, **thoroughly search** the codebase for existing implementations:\n\n1. **Core ops**: Use the Glob tool to check for `src/flag_gems/ops/.py`\n2. **Experimental ops**: Use the Glob tool to check for `src/flag_gems/experimental_ops/.py`\n3. **Backend ops**: If `` is `backend`, check `src/flag_gems/runtime/backend/_/ops/.py`\n4. **Registration**: Use the Grep tool to search `src/flag_gems/ops/__init__.py` and `src/flag_gems/__init__.py` for the op name\n5. **Aten registration**: Use the Grep tool to search for `torch.ops.aten.` in the codebase (some registrations use the aten op reference instead of string names)\n6. **Tests**: Use the Grep tool to search `tests/test_*_ops.py` and `experimental_tests/_test.py` for the op name\n7. **Benchmarks**: Use the Grep tool to search `benchmark/test_*_perf.py` for the op name in `forward_operations`\n\n### If the operator already exists\n\nPresent findings to the user and ask them to choose one of the following:\n\n**Option A — Skip generation**: The operator already exists; do nothing.\n\n**Option B — Replace existing**: Overwrite the current implementation with MCP-generated code\n(adapted to FlagGems conventions). The file stays in its **current location** (core stays in core, experimental stays in experimental).\n\n**Option C — Create experimental variant (side-by-side)**: Generate the operator under the same\nname in `experimental_ops/` so it coexists with the original. This will:\n - Create `src/flag_gems/experimental_ops/.py`\n - Create a standalone test in `experimental_tests/_test.py`\n - Include a perf benchmark in that same test file\n - Do **NOT** register it in `_FULL_CONFIG` (it won't override the aten dispatch)\n\nOnly proceed to Step 3 after the user has made a choice.\n\n### If the operator does NOT exist\n\nProceed directly to Step 3. Use `` determined in Step 1.5.\n\n## Step 3: Research Context (flagos_wiki)\n\nBefore calling the MCP generator, use the Read tool to gather reference materials that improve generation quality:\n\n1. **Read similar operator code** from the codebase. Choose based on ``:\n - **Core placement** — read from `src/flag_gems/ops/`:\n - Unary pointwise → read `src/flag_gems/ops/abs.py` or `src/flag_gems/ops/sigmoid.py`\n - Binary pointwise → read `src/flag_gems/ops/add.py`\n - Reduction → read `src/flag_gems/ops/sum.py` or `src/flag_gems/ops/mean.py`\n - Normalization → read `src/flag_gems/ops/layer_norm.py`\n - BLAS → read `src/flag_gems/ops/mm.py`\n - **Experimental placement** — read from `src/flag_gems/experimental_ops/`:\n - Read `src/flag_gems/experimental_ops/abs.py` or a similar experimental op\n - **Backend placement** — read existing ops in the target vendor directory\n\n2. **Read the test pattern** from the matching test location:\n - **Core**: read `tests/test__ops.py` — note the import style, utility functions, parametrize patterns\n - **Experimental**: read an existing `experimental_tests/_test.py` — note the direct import style, `to_reference()`, `GenericBenchmark` usage\n - **Backend**: read existing tests in the backend vendor directory\n\n3. **Note the decorators used** in similar operators:\n - Core pointwise ops use `@pointwise_dynamic(promotion_methods=...)` + `@triton.jit`\n - Experimental ops use `@triton.jit` only (raw pointer-based)\n - Some ops may use other decorators like `@triton.autotune` — note these\n\n4. **If replacing an existing op** (Option B), read the current implementation so the new version\n can be compared / improved upon.\n\n5. Collect all findings and **summarize into concise notes** (not full file contents) to pass as\n the `flagos_wiki` parameter. For example:\n - `\"abs.py uses pointwise_dynamic with DEFAULT promotion, returns tl.abs(x)\"`\n - `\"sigmoid.py uses INT_TO_FLOAT promotion, returns 1/(1+tl.exp(-x))\"`\n - `\"test pattern: @pytest.mark.xxx, POINTWISE_SHAPES, FLOAT_DTYPES, gems_assert_close\"`\n - `\"experimental test style: from flag_gems.experimental_ops.abs import abs as gems_abs, uses torch.ops.aten.abs for reference\"`\n\n If the `mcp__kernelgen-mcp__generate_kernel` tool supports a `flagos_wiki` (or similarly named\n `context`/`references`) parameter, pass the collected notes. If the MCP call fails with an\n error containing `unexpected argument`, `unknown field`, `invalid schema`, or similar, retry\n the call without the `flagos_wiki` field.\n\n## Step 4: Call kernelgen-mcp\n\nInvoke `mcp__kernelgen-mcp__generate_kernel` with the parameters gathered above, including the\n`flagos_wiki` list for reference context.\n\nIf `` is `backend` and target device is not nvidia, pass the `device` parameter to the\nMCP call (e.g., `device=\"huawei\"` for Ascend).\n\n**Set the iteration counter**: `iteration_count = 1`. This tracks total MCP calls for the final report.\n\nThe MCP returns four code blocks:\n- `torch_code` — PyTorch reference implementation\n- `triton_code` — Triton kernel implementation\n- `test_func_code` — accuracy test code\n- `benchmark_func_code` — performance benchmark code\n\n## Step 4.5: MCP Report Review\n\nIf the MCP response includes test results (accuracy test reports, benchmark/speedup reports),\npresent them to the user **in full** before proceeding to local code adaptation:\n\n```\n=== MCP Generation Test Report ===\n\nAccuracy Test Results:\n\n\nPerformance Benchmark Results:\n\n```\n\nThen ask the user:\n\n> The MCP service has returned the above test reports from its generation environment.\n> Would you like to:\n>\n> **A — Skip local testing**: Trust the MCP reports and proceed directly to code placement\n> and summary. No local tests will be run.\n>\n> **B — Run local tests**: Run accuracy and performance tests on your local machine as well\n> to verify the results in your specific environment.\n\nWait for the user's choice before proceeding:\n- If **A (skip local testing)**: After placing code in Step 5, skip Steps 6 and 7 (local\n accuracy tests and benchmarks) and proceed directly to Step 8 (Summary). Use the MCP\n report numbers for the summary.\n- If **B (run local tests)**: Proceed normally through Steps 5 → 6 → 7 → 8. Before\n running local tests, perform the **Chip Compatibility Check** described in Step 5.5c.\n\n**If the MCP response does NOT include test reports** (only code blocks), proceed normally\nto Step 5 — local testing will be required.\n\n## Step 5: Adapt and Place Code into the FlagGems Project\n\nThis is the most critical step. The generated code must be transformed to match FlagGems conventions.\nThe exact placement and transformation depends on ``.\n\n---\n\n### Placement: Core (`src/flag_gems/ops/`)\n\n#### 5a. Operator Implementation → `src/flag_gems/ops/.py`\n\n**Transformation rules for pointwise ops**: The MCP generator will almost always output raw\npointer-based Triton code. For core pointwise ops, you MUST **always** rewrite it into the\n`pointwise_dynamic` elementwise style:\n- The `@triton.jit` kernel function receives **scalar elements** (not pointers)\n- `pointwise_dynamic` handles all pointer arithmetic, tiling, and masking automatically\n- The kernel only contains the elementwise math logic (e.g., `return tl.maximum(x, 0)`)\n- **Remove ALL** `tl.load()`, `tl.store()`, pointer arithmetic (`ptr + offsets`), mask logic,\n and BLOCK_SIZE parameters — these are incompatible with `pointwise_dynamic`\n- Only keep the pure math expression that transforms input element(s) to output element(s)\n\nExamples of correct `pointwise_dynamic` kernels (keep ONLY the scalar `tl.*` math):\n```python\n# relu: return tl.maximum(x, 0)\n# exp: return tl.exp(x)\n# silu: return x * tl.sigmoid(x)\n# add: return x + y\n```\n\nDo NOT keep `offsets`, `mask`, `n_elements`, `BLOCK_SIZE`, or any tiling logic.\n\n**CRITICAL**: The MCP Triton code MUST NOT be copied directly for core ops. Always rewrite it\nfollowing FlagGems `pointwise_dynamic` conventions. Never trust the raw MCP code — it will\nalmost always be pointer-style which is incompatible with `pointwise_dynamic`.\n\nUse the Write tool (new file) or Edit tool (replacing existing) to create the operator file.\n\n**Unary pointwise template:**\n```python\nimport logging\n\nimport triton\nimport triton.language as tl\n\nfrom flag_gems.utils import pointwise_dynamic\n\nlogger = logging.getLogger(__name__)\n\n\n@pointwise_dynamic(promotion_methods=)\n@triton.jit\ndef _forward(x):\n return ...\n\n\ndef (A):\n logger.debug(\"GEMS FORWARD\")\n return _forward(A)\n```\n\n**Binary pointwise template:**\n```python\nimport logging\n\nimport triton\nimport triton.language as tl\n\nfrom flag_gems.utils import pointwise_dynamic\n\nlogger = logging.getLogger(__name__)\n\n\n@pointwise_dynamic(promotion_methods=)\n@triton.jit\ndef _forward(x, y):\n return ...\n\n\ndef (A, B):\n logger.debug(\"GEMS FORWARD\")\n return _forward(A, B)\n```\n\nChoose the correct template based on `func_type`. For non-pointwise ops (reduction, BLAS,\nnormalization), do NOT use either template — instead read a similar existing op and follow\nits pattern.\n\n**For in-place variants**: Read an existing in-place operator (e.g., `relu.py`) to find the\nexact call pattern used in this repo, then follow it. Do NOT assume the in-place parameter\nname — it may be `out0=A`, `out=A`, or another convention. Always verify by reading an existing\nin-place operator first.\n\nKey conventions:\n- Use `pointwise_dynamic` decorator from `flag_gems.utils` for pointwise ops\n- Use `triton.jit` decorator\n- The kernel function takes raw tensor elements (not pointers) — `pointwise_dynamic` handles the boilerplate\n- **Unary ops**: kernel has one parameter `(x)`, e.g. `def relu_forward(x)`\n- **Binary ops**: kernel has two parameters `(x, y)`, e.g. `def add_forward(x, y)`\n- Promotion methods — **always read a similar existing operator first** and copy its promotion\n method. If no similar operator exists, use these guidelines as fallback:\n - `INT_TO_FLOAT`: ops that always produce float output (exp, log, sigmoid, tanh, sqrt, etc.)\n - `COMPLEX_TO_FLOAT`: ops that reduce complex to real (abs)\n - `DEFAULT`: ops that preserve input dtype (relu, neg, add, mul, clamp, etc.)\n - **Template placeholder** `` must expand to the full list:\n - Unary ops: `[(0, \"DEFAULT\")]` or `[(0, \"INT_TO_FLOAT\")]`\n - Binary ops: `[(0, \"DEFAULT\"), (1, \"DEFAULT\")]` or `[(0, \"INT_TO_FLOAT\"), (1, \"INT_TO_FLOAT\")]`\n - **When in doubt, prefer reading the repo** — promotion behavior is subtle and repo-specific.\n- The wrapper function takes `A` as the input tensor parameter name (NOT `self`)\n- Include `logger.debug(\"GEMS FORWARD\")` in the wrapper\n- For non-pointwise ops (reduction, BLAS, normalization, etc.), follow the specific pattern of similar existing ops — do NOT force the `pointwise_dynamic` pattern on them\n\n#### 5b. Register the Operator\n\nUse the Edit tool to make these changes:\n\n1. **`src/flag_gems/ops/__init__.py`**: Add import and `__all__` entry (in alphabetical order):\n ```python\n from flag_gems.ops. import , _\n ```\n\n2. **`src/flag_gems/__init__.py`**: Add to `_FULL_CONFIG` tuple (in alphabetical order).\n **IMPORTANT**: First read the existing `_FULL_CONFIG` entries to match the exact registration\n format used by this repo. Different versions may use different styles:\n - String style: `(\"relu\", relu)`\n - Aten op style: `(torch.ops.aten.relu.default, relu)`\n Copy the pattern from existing entries. Do NOT guess the format.\n Insert the new entry in **alphabetical order** — do not change the order of existing entries.\n\n#### 5c. Accuracy Test → `tests/test__ops.py`\n\nDo NOT create a new test file. Use the Edit tool to add the test function to the appropriate existing test file:\n- Unary pointwise: `tests/test_unary_pointwise_ops.py`\n- Binary pointwise: `tests/test_binary_pointwise_ops.py`\n- Reduction: `tests/test_reduction_ops.py`\n- BLAS: `tests/test_blas_ops.py`\n- Norm: `tests/test_norm_ops.py`\n- Special: `tests/test_special_ops.py`\n\n**IMPORTANT — Match existing style**: Before writing the test, **read at least 2 existing test\nfunctions** in the target file. Ensure you match:\n- Import style (from `.accuracy_utils import ...`)\n- The `flag_gems.use_gems()` context manager pattern\n- Assertion function (`gems_assert_close` or `gems_assert_equal`)\n- `to_reference()` usage pattern\n- Parametrize decorators (`POINTWISE_SHAPES`, `FLOAT_DTYPES`, etc.)\n\nFollow the existing test pattern:\n```python\n@pytest.mark.\n@pytest.mark.parametrize(\"shape\", POINTWISE_SHAPES)\n@pytest.mark.parametrize(\"dtype\", FLOAT_DTYPES)\ndef test_accuracy_(shape, dtype):\n inp = torch.randn(shape, dtype=dtype, device=flag_gems.device)\n ref_inp = to_reference(inp)\n\n ref_out = (ref_inp)\n with flag_gems.use_gems():\n res_out = torch.(inp)\n\n gems_assert_close(res_out, ref_out, dtype)\n```\n\nKey conventions:\n- Use `flag_gems.device` not `\"cuda\"`\n- Use `to_reference()` to create reference tensors\n- Use `flag_gems.use_gems()` context manager for the FlagGems call\n- Use `gems_assert_close` or `gems_assert_equal` for comparison\n- Mark with `@pytest.mark.`\n- If the marker `` is not already registered, find where markers are defined\n (check `pytest.ini`, `pyproject.toml`, or `setup.cfg`) and add it.\n- Use `POINTWISE_SHAPES`, `FLOAT_DTYPES` etc. from `accuracy_utils`\n\n#### 5d. Performance Benchmark → `benchmark/test__perf.py`\n\nDo NOT create a new benchmark file. Read the target benchmark file first, then use the Edit tool\nto append the new tuple into the existing `forward_operations` list (find the list definition and\ninsert in alphabetical order). Follow the exact tuple format already used in the file.\n\n---\n\n### Placement: Experimental (`src/flag_gems/experimental_ops/`)\n\nThis is the **default for new operators**.\n\n#### 5a. Operator Implementation → `src/flag_gems/experimental_ops/.py`\n\nUse the Write tool to create the file. Use the **raw Triton pointer-based style** (standard for experimental ops).\nKeep the MCP-generated Triton code mostly as-is but ensure it:\n- Is self-contained (no `flag_gems.utils` imports needed, no `pointwise_dynamic`)\n- Has proper `@triton.jit` kernel with pointer args, mask, BLOCK_SIZE\n- **Has a Python wrapper function** that computes the grid and launches the kernel\n- Exports the wrapper function named `` (and optionally `_out`)\n\n**Ensure all kernel meta parameters (BLOCK_SIZE, etc.) are passed** — either via\n`@triton.autotune` or as explicit keyword arguments.\n\nExample structure (reference `src/flag_gems/experimental_ops/abs.py`):\n```python\nimport torch\nimport triton\nimport triton.language as tl\n\n\n@triton.jit\ndef __kernel(in_ptr, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):\n pid = tl.program_id(axis=0)\n block_start = pid * BLOCK_SIZE\n offsets = block_start + tl.arange(0, BLOCK_SIZE)\n mask = offsets < n_elements\n x = tl.load(in_ptr + offsets, mask=mask)\n y = ... # elementwise computation\n tl.store(out_ptr + offsets, y, mask=mask)\n\n\ndef (x: torch.Tensor) -> torch.Tensor:\n output = torch.empty_like(x)\n n_elements = x.numel()\n grid = lambda meta: (triton.cdiv(n_elements, meta[\"BLOCK_SIZE\"]),)\n __kernel[grid](x, output, n_elements, BLOCK_SIZE=1024)\n return output\n```\n\n#### 5b. Registration (limited)\n\nFor experimental ops, do **NOT** modify `src/flag_gems/__init__.py` or `_FULL_CONFIG`.\nOptionally add an entry to `src/flag_gems/experimental_ops/__init__.py` if needed.\n\n#### 5c. Standalone Test + Benchmark → `experimental_tests/_test.py`\n\nUse the Write tool to create a single self-contained test file.\n\n**IMPORTANT — Match experimental_tests/ style**: Before writing the test, **read 1-2 existing\ntest files** in `experimental_tests/` (e.g., `abs_test.py`, `sigmoid_test.py`). Match their exact style:\n- Direct import: `from flag_gems.experimental_ops. import as gems_`\n- Reference call: `torch.ops.aten.(ref_x)` (not `torch.`)\n- `to_reference()` with `TO_CPU` support\n- `GenericBenchmark` from `benchmark.performance_utils` for perf tests\n- `triton.testing.do_bench` for inline benchmarks\n\nFollow this pattern (adapted from existing experimental tests):\n```python\n# operator test\n\nimport os\nimport sys\n\nimport pytest\nimport torch\n\nimport flag_gems\nfrom flag_gems.experimental_ops. import as gems_\n\n# Add parent directory to path to import flag_gems\nsys.path.insert(0, os.path.join(os.path.dirname(__file__), \"..\"))\ntry:\n from tests.accuracy_utils import TO_CPU, gems_assert_close # noqa: E402\nexcept ImportError:\n # Fallback values when running outside pytest\n TO_CPU = False # fallback\n\n def gems_assert_close(res, ref, dtype, **kwargs):\n # Simple fallback comparison\n torch.testing.assert_close(res, ref, **kwargs)\n\n\nsys.path.insert(0, os.path.join(os.path.dirname(__file__), \"..\"))\nfrom benchmark.performance_utils import GenericBenchmark # noqa: E402\n\n\ndef to_reference(inp, upcast=False):\n if inp is None:\n return None\n if TO_CPU:\n ref_inp = inp.to(\"cpu\")\n else:\n ref_inp = inp.clone()\n if upcast:\n if ref_inp.is_complex():\n ref_inp = ref_inp.to(torch.complex128)\n else:\n ref_inp = ref_inp.to(torch.float64)\n return ref_inp\n\n\n@pytest.mark.\n@pytest.mark.parametrize(\"shape\", [(2, 3), (128, 256), (512, 512)])\n@pytest.mark.parametrize(\"dtype\", [torch.float32, torch.float16, torch.bfloat16])\ndef test__tensor(shape, dtype):\n x = torch.randn(shape, dtype=dtype, device=flag_gems.device)\n ref_x = to_reference(x)\n ref_out = torch.ops.aten.(ref_x)\n with flag_gems.use_gems():\n act_out = gems_(x)\n gems_assert_close(act_out, ref_out, dtype=dtype)\n\n\n@pytest.mark.\ndef test_perf_aten_():\n def _input_fn(shape, dtype, device):\n inp = torch.randn(shape, dtype=dtype, device=flag_gems.device)\n yield inp,\n\n bench = GenericBenchmark(\n input_fn=_input_fn,\n op_name=\"\",\n torch_op=torch.ops.aten.,\n dtypes=[torch.float32, torch.float16, torch.bfloat16],\n )\n\n return bench.run()\n```\n\n---\n\n### Placement: Backend (`src/flag_gems/runtime/backend/_/`)\n\n#### 5a. Operator Implementation\n\nRead existing operators in `src/flag_gems/runtime/backend/_/ops/` to understand the\nexact patterns used by this vendor. Follow the same style.\n\n#### 5b. Registration\n\nFollow the vendor's existing registration pattern (check `__init__.py` in the vendor directory).\n\n#### 5c. Tests\n\nFollow the vendor's existing test patterns. Read existing test files in the vendor's test\ndirectory before writing new ones.\n\n---\n\n## Step 5.5: Pre-test Validation\n\nBefore running the full test suite, perform two quick checks:\n\n### 5.5a. Lint / Format Check\n\nIf the repo uses linting tools (ruff, black, isort, flake8), run them on the changed files to\ncatch style issues early. Use the shell tool:\n\n```bash\n# Adjust path based on \npython -m ruff check --fix\npython -m ruff format \n```\n\nIf `ruff` is not installed, check for other formatters in `pyproject.toml` or `setup.cfg` and\nuse those. If no linter is configured, skip this step.\n\n### 5.5b. Triton Compile Smoke Test\n\nTriton kernel compile errors only surface on first invocation (JIT compilation).\n\n**For Core placement:**\n```bash\npython -c \"\nimport torch, flag_gems\nfor dtype in [torch.float32, torch.float16, torch.bfloat16]:\n for shape in [(1,), (4,), (128,), (4, 128), (0,)]:\n inp = torch.empty(shape, dtype=dtype, device=flag_gems.device) if shape == (0,) else torch.randn(shape, dtype=dtype, device=flag_gems.device)\n with flag_gems.use_gems():\n out = # e.g. torch.relu(inp)\n print(f'Smoke test passed: {dtype}, shape={shape}')\n\"\n```\n\n**For Experimental placement:**\n```bash\npython -c \"\nimport torch, flag_gems\nfrom flag_gems.experimental_ops. import \nfor dtype in [torch.float32, torch.float16, torch.bfloat16]:\n for shape in [(1,), (4,), (128,), (4, 128), (0,)]:\n inp = torch.empty(shape, dtype=dtype, device=flag_gems.device) if shape == (0,) else torch.randn(shape, dtype=dtype, device=flag_gems.device)\n out = (inp)\n print(f'Smoke test passed: {dtype}, shape={shape}')\n\"\n```\n\nThe `(0,)` shape tests the empty-tensor edge case (numel==0), which often causes kernel crashes.\n\nIf this fails, read the error and fix the kernel before proceeding to Step 6.\n\n### 5.5c. Chip Compatibility Check (before local testing)\n\nBefore running local tests, verify that the local hardware matches the operator's\n`target_device` (determined in Step 1). Use the Bash tool:\n\n```bash\npython - <<'PY'\nimport subprocess\ntry:\n import torch\n if torch.cuda.is_available():\n for i in range(torch.cuda.device_count()):\n print(f\"nvidia: {torch.cuda.get_device_name(i)}\")\n else:\n print(\"No CUDA device available\")\nexcept Exception:\n print(\"torch not available\")\ntry:\n r = subprocess.run([\"npu-smi\", \"info\"], capture_output=True, text=True, timeout=5)\n if r.returncode == 0 and \"NPU\" in r.stdout:\n print(\"huawei: Ascend NPU detected\")\nexcept Exception:\n pass\ntry:\n r = subprocess.run([\"rocm-smi\"], capture_output=True, text=True, timeout=5)\n if r.returncode == 0:\n kind = \"haiguang\" if \"hygon\" in r.stdout.lower() else \"amd\"\n print(f\"{kind}: ROCm device detected\")\nexcept Exception:\n pass\nPY\n```\n\nCompare detected hardware with `target_device`:\n\n| Target Device | Required Local Hardware |\n|---|---|\n| `nvidia` (default) | NVIDIA GPU |\n| `huawei` | Huawei Ascend NPU (`npu-smi`) |\n| `haiguang` | Hygon DCU (`rocm-smi` + Hygon) |\n| `moore` | Moore Threads GPU |\n| `tianshu` | Tianshu GPU |\n| `cambricon` | Cambricon MLU |\n\n**If the local hardware does NOT match the target device**, warn the user:\n\n```\n⚠️ Hardware mismatch detected:\n - Operator target device: \n - Local hardware: \n\nRunning tests for a operator on hardware will likely\nfail or produce incorrect results. Please choose:\n\n A — Skip local testing: Do not run local tests; rely on MCP test reports (if available).\n B — Proceed anyway: Run local tests despite the mismatch (results may be unreliable).\n```\n\nIf the user chooses A, skip Steps 6 and 7 and proceed to Step 8 (Summary).\nIf the user chooses B, proceed with local tests but note the mismatch in the final report.\n\n## Step 6: Run Accuracy Tests\n\nRun the accuracy tests for the newly added operator.\nAssume the current working directory is the repository root.\n\n**For Core placement:**\n```bash\npython -m pytest tests/test__ops.py -m -v\n```\n\n**For Experimental placement:**\n```bash\npython -m pytest experimental_tests/_test.py -v -k \"test_\"\n```\n\n**For Backend placement:**\n```bash\n# Use the vendor's test running convention (read existing CI scripts or Makefile first)\npython -m pytest -m -v\n```\n\nReport the results to the user. If tests fail, follow the **error classification and retry\nprotocol** below. This protocol strictly limits self-fix attempts to prevent infinite loops.\n\n### Error Classification\n\nFirst, classify the error into one of two categories:\n\n**Category A — Compilation / Import errors** (model may attempt 1 self-fix):\n- `ImportError`, `ModuleNotFoundError` — wrong import path or missing module\n- `SyntaxError` — Python syntax issue\n- `TritonCompilationError`, `CompilationError` — Triton kernel compile failure\n- `NameError` — undefined variable or function\n- `TypeError` in function call signature — wrong number/type of arguments\n- `AttributeError` — wrong attribute access on a module or object\n\n**Category B — Algorithm / Numerical accuracy errors** (do NOT self-fix):\n- `AssertionError` from `torch.testing.assert_close` or `gems_assert_close` — numerical mismatch\n- Wrong output values (results don't match reference)\n- Shape mismatch in output tensors\n- NaN or Inf in output\n- Any error that indicates the kernel logic is incorrect\n\n### Retry Protocol (strictly follow this order)\n\n**Step 6a. (Category A only) Self-fix — maximum 1 attempt:**\nIf the error is Category A (compilation/import), you may attempt exactly ONE fix:\n1. Use the Read tool to examine the error traceback carefully.\n2. Apply a targeted fix using the Edit tool (e.g., fix import path, fix syntax, fix argument name).\n3. Use the shell tool to re-run the tests.\n4. If the test **passes** → proceed to Step 7.\n5. If the test **still fails** → proceed to Step 6b. Do NOT attempt a second self-fix.\n\n**If the error is Category B (algorithm/accuracy), skip Step 6a entirely** — go directly\nto Step 6b. Do NOT attempt to fix algorithm logic yourself, as this typically leads to\nan endless fix-retry loop without converging. The MCP service has better optimization\ncapabilities for these issues.\n\n**Step 6b. MCP re-generation — pass error context to generate_kernel:**\nRe-call `mcp__kernelgen-mcp__generate_kernel` with the **same parameters as Step 4**,\nbut add the error information to `flagos_wiki` as additional hints.\n**Increment**: `iteration_count += 1`.\nKeep `flagos_wiki` concise — maximum 10 items total. If retrying multiple times, replace\nearlier error entries rather than appending, to avoid bloating the prompt.\nReplace the kernel code with the new MCP output, re-run tests.\n- If tests **pass** → proceed to Step 7.\n- If tests **still fail** → proceed to Step 6c.\n\n**Step 6c. MCP optimization — pass error context to optimize_kernel:**\nTry `mcp__kernelgen-mcp__optimize_kernel` with the current kernel code and the\n`check_result` parameter containing the error traceback.\n**Increment**: `iteration_count += 1`.\nReplace the kernel code with the optimized output, re-run tests.\n- If tests **pass** → proceed to Step 7.\n- If tests **still fail** → proceed to Step 6d.\n\n**Step 6d. Stop and report:**\nDo not keep retrying. Report the failure to the user with:\n- The specific test failures and error messages\n- Your analysis of what might be wrong\n- Suggestion to try with different `func_type` or additional `flagos_wiki` hints\n\n## Step 7: Run Performance Benchmark\n\nRun the performance benchmark.\n\n**For Core placement:**\n```bash\npython -m pytest benchmark/test__perf.py -m -v\n```\n\n**For Experimental placement:**\n```bash\npython -m pytest experimental_tests/_test.py -v -k \"perf\"\n```\n\nLook for lines in the output containing keywords like `speedup`, `latency`, `gems`, `torch`, or\ntiming values. Use regex patterns to extract numbers.\n\n**Do not just copy the raw table** — always compute and report actual speedup ratios.\n\nIf a speedup metric is printed directly, extract and report it. Otherwise compute:\n`speedup = torch_latency / gems_latency` (>1.0 means gems is faster).\n\n**Note**: Triton kernels are JIT-compiled on first invocation, which adds significant overhead.\nIf benchmark results look unreasonably slow on the first config, check whether the benchmark\nframework includes a warmup phase. If not, the first data point may be an outlier — **if the\nfirst timing sample is >5x larger than the median of remaining samples, exclude it** from the\naverage speedup calculation and note \"first sample excluded (Triton JIT compile overhead)\" in\nthe report.\n\n## Step 8: Summary\n\nProvide a clear summary to the user with **exact numbers** extracted from test and benchmark\noutput:\n\n```\n=== KernelGen Operator Generation Report ===\n\nOperator Name: \nPlacement: Core / Experimental / Backend (_)\nGeneration Mode: New / Replace Existing / Experimental Variant\nOperator Type: \nTarget Device: \n\nFile Changes:\n - [New/Modified] \n - [New/Modified] \n - [Modified] (if applicable)\n - [Modified] src/flag_gems/ops/__init__.py (if core)\n - [Modified] src/flag_gems/__init__.py (if core)\n\nDecorator Style: pointwise_dynamic / raw triton.jit / other\nTest Style: use_gems() + torch.() / direct import from experimental_ops\n\nAccuracy Tests: passed, failed (total test cases)\n Pass Rate: %\n Failed Cases: \n\nPerformance Benchmark:\n Avg Speedup: x vs PyTorch reference\n Best Speedup: x (shape=, dtype=)\n Worst Speedup: x (shape=, dtype=)\n (If benchmark was not run or failed, write \"Incomplete\" and explain why)\n\nMCP Iterations: (write 1 if first generation passed)\n\nPerformance Analysis:\n 5.0x → \"Extreme fusion success, typical for fused multi-op kernels\"\n - Speedup > 3.0x → \"Excellent compute-bound optimization, significant kernel fusion benefit\"\n - Speedup > 2.0x → \"Compute-bound optimization effective, significant speedup achieved\"\n - Speedup 1.2x ~ 2.0x → \"Moderate speedup, kernel likely balanced between compute and memory\"\n - Speedup 0.5x ~ 1.2x → \"Kernel likely memory-bound, limited optimization headroom\"\n - Speedup < 0.5x → \"Kernel slower than PyTorch reference — likely suboptimal memory access pattern\">\n\nIssues and Fixes: (if any)\n```\n\n**After presenting the summary, check the speedup:**\n\n- **If average speedup < 0.5x** (kernel slower than PyTorch), proactively warn the user.\n- **If average speedup is 0.5x ~ 1.2x**, suggest optimization.\n- **If average speedup > 1.2x**, no action needed — report the result normally.\n\n## Step 9: Post-Completion Actions\n\nAfter the summary is presented (whether based on MCP reports or local test results), perform\nthe following two checks.\n\n### 9a. Chat Tool Notification\n\nIf the user's task was received via a chat tool (e.g., Feishu/飞书, Discord, Slack, Teams,\nDingTalk/钉钉, WeChat Work/企业微信, Telegram, or any similar messaging platform), or if\nthe user mentions sending results to a client or team, ask:\n\n> The operator generation is complete. Would you like me to prepare a message to send the\n> generated files or summary to your client/team via ****?\n\nIf the user confirms:\n1. Prepare a concise summary message including: operator name, file paths, accuracy pass\n rate, speedup metrics, and any issues.\n2. If the chat tool has a CLI or API integration available in the environment, use it to\n send the message directly.\n3. If no CLI/API is available, format the message as copy-paste ready text for the user.\n4. If files need to be sent, list the exact file paths the user should attach.\n\n### 9b. Git Repository PR Submission\n\nCheck whether the current project is managed by a git-based hosting service:\n\n```bash\ngit remote -v 2>/dev/null | head -5\n```\n\nDetect the hosting platform from the remote URL:\n- `github.com` → GitHub\n- `gitlab.com` (or self-hosted GitLab) → GitLab\n- `gitee.com` → Gitee\n- `bitbucket.org` → Bitbucket\n\n**If a git hosting platform is detected**, ask the user:\n\n> This project is hosted on ****. Would you like me to automatically create a\n> Pull Request with the generated operator code?\n\n**If the user confirms**, follow the standard PR creation workflow:\n\n1. Create a new branch: `kernelgen/`\n2. Stage all changed/new files related to this operator generation (never use `git add -A`)\n3. Create a commit with a descriptive message\n4. Push the branch to the remote with `-u` flag\n5. Create the PR using the platform's CLI tool:\n\n **For GitHub** (`gh`):\n ```bash\n gh pr create --title \"Add operator via KernelGen\" --body \"$(cat <<'EOF'\n ## Summary\n Add `` operator () generated via KernelGen MCP service.\n\n ## Changes\n - [New/Modified] `` — Triton kernel implementation\n - [New/Modified] `` — Accuracy tests\n - [New/Modified] `` — Performance benchmarks\n - [Modified] `` — Operator registration (if applicable)\n\n ## Test Results\n ### Accuracy\n - **Pass rate**: / (%)\n - **Failed cases**: \n\n ### Performance\n - **Avg speedup**: x vs PyTorch reference\n - **Best**: x | **Worst**: x\n\n ## Generation Details\n - **MCP iterations**: \n - **Target device**: \n - **Operator type**: \n - **Placement**: \n\n ---\n Generated by KernelGen MCP skill\n\n EOF\n )\"\n ```\n\n **For Gitee / GitLab**: Provide the PR/MR creation URL and prepared description.\n\n6. Return the PR URL to the user.\n\n**If PR creation fails**, report the error, provide the branch name, and suggest manual PR\ncreation with the prepared description.\n\n**If the user declines**, do nothing — changes remain as local files.\n\n## Important Notes\n\n- **`kernel_name` must match the `torch` API name** (e.g., `silu` not `swish`, `neg` not `negate`).\n The accuracy test calls `torch.(...)`, so a mismatch will cause test failures.\n- **Never overwrite existing operators** without explicit user permission (Step 2 choice)\n- **Placement determines decorator style**:\n - Core ops → `pointwise_dynamic` for pointwise, follow existing patterns for others\n - Experimental ops → raw `@triton.jit` pointer-based style\n - Backend ops → follow vendor-specific patterns\n- **Placement determines test style**:\n - Core (`tests/`) → `flag_gems.use_gems()` + `torch.()`\n - Experimental (`experimental_tests/`) → direct `from flag_gems.experimental_ops. import `\n - Backend → vendor-specific test patterns\n- **Default placement for new operators is experimental** — only use core if user explicitly requests it\n- **Follow alphabetical ordering** when adding entries to `__init__.py` and `__all__`\n- **Match the existing code style** exactly — read existing files first and copy the pattern\n (imports, logging, naming conventions, registration format, test patterns)\n- **Speedup formula**: `speedup = torch_latency / gems_latency` (values > 1.0 mean gems is faster)\n- **Always use built-in tools** (shell, Read, Write, Edit, Grep, Glob) instead of\n outputting commands or code snippets for the user to execute manually\n\nFile v1.1.0:kernelgen-generate-for-vllm.md\n\n# KernelGen Skill — Generate vLLM Operators via MCP\n\nYou are an expert at generating GPU kernel operators for the vLLM project using the `kernelgen-mcp` MCP service.\n\n> **⚠️ MCP Prerequisite**: MCP configuration is handled by `SKILL.md` Phase 0. If MCP tools fail during execution,\n> stop and tell the user to visit https://kernelgen.flagos.io/mcp 获取 Token 后重新配置。\n\n## Execution Rules\n\nFollow these rules strictly to avoid getting lost in the multi-step workflow:\n\n1. **Follow the steps in order**: Step 0 → Step 1 → Step 2 → ... → Step 8. Never jump ahead.\n2. **Never skip Step 2** (operator existence check) — always verify before generating.\n3. **Do not generate or write any code before Step 4** (MCP call). Steps 0–3 are preparation only.\n4. **Always confirm with the user before destructive actions**: replacing files, installing packages,\n overwriting operators.\n5. **If any step fails, stop and report the error** to the user. Do not silently skip or retry\n with different parameters.\n6. **Report progress to the user when entering each step**. Print a short status line like\n `Entering Step 2 — checking whether the operator already exists...` so the user knows\n the workflow is progressing and not stuck.\n7. **Never fabricate repository files or paths**. If a required file cannot be found using the\n Glob or Grep tools, report it to the user instead of guessing. Do not assume a file exists\n at a path without verifying it first.\n8. **CRITICAL — MCP is mandatory**: ALL operator code generation MUST go through the\n `mcp__kernelgen-mcp__generate_kernel` MCP tool. NEVER generate Triton kernels, PyTorch\n wrappers, or operator implementations yourself — even if the operator seems simple (e.g.,\n relu, abs). If MCP is not configured, not reachable, or fails after all retries in Step 4,\n STOP and report the issue to the user. Do NOT fall back to writing kernel code manually.\n The MCP service produces optimized, tested code that manual writing cannot match.\n\n## Step 0: Pre-flight — Environment & MCP Check\n\n### 0a. Check Python Environment\n\nUse the Bash tool to run a diagnostic that checks each package independently (so one failure\ndoes not prevent checking the rest):\n\n```bash\npython - <<'PY'\nimport importlib\n\ndef check(pkg):\n try:\n m = importlib.import_module(pkg)\n print(pkg, getattr(m, \"__version__\", \"installed\"))\n return True\n except Exception:\n print(pkg, \"NOT installed\")\n return False\n\ncheck(\"torch\")\ntriton_ok = check(\"triton\")\n\ntry:\n import torch\n print(\"cuda available\", torch.cuda.is_available())\n if torch.cuda.is_available():\n print(\"torch.cuda version:\", torch.version.cuda)\n n = torch.cuda.device_count()\n print(\"cuda devices:\", n)\n for i in range(n):\n print(f\" cuda:{i}\", torch.cuda.get_device_name(i))\nexcept Exception:\n pass\n\nif triton_ok:\n try:\n import triton.runtime\n print(\"triton runtime OK\")\n except Exception as e:\n print(\"triton runtime issue:\", e)\n\ncheck(\"vllm\")\ncheck(\"pytest\")\ncheck(\"numpy\")\n\ntry:\n import subprocess\n result = subprocess.run([\"nvidia-smi\"], capture_output=True, text=True, timeout=5)\n if result.returncode == 0:\n for line in result.stdout.splitlines()[:4]:\n print(line)\n else:\n print(\"nvidia-smi: failed\")\nexcept Exception:\n print(\"nvidia-smi: not found\")\nPY\n```\n\n**If any import fails**, stop and ask the user before installing anything. Present the missing\npackages and let the user confirm:\n\n| Missing package | Suggested install command | Notes |\n|---|---|---|\n| `torch` | `pip install torch --index-url https://download.pytorch.org/whl/cu121` | User MUST pick the correct CUDA variant (cu118/cu121/cu124). Ask which CUDA version they have. |\n| `triton` | `pip install triton` | |\n| `vllm` | `pip install -e .` (from repo root) | Confirm the current directory is the vLLM repo root first. |\n| `pytest` | `pip install pytest` | |\n| `numpy` | `pip install numpy` | |\n\n**Important**: Never install `torch` without asking the user which CUDA version they need.\nA wrong CUDA version will cause silent failures.\n\nIf `torch` is installed but `torch.cuda.is_available() == False`, warn the user that GPU\ntests will fail and ask how to proceed before continuing.\n\nAfter the user confirms and packages are installed, re-run the diagnostic. Only proceed when\nall imports succeed.\n\n### 0b. Check MCP Availability\n\n> **MCP configuration is handled centrally by `SKILL.md` Phase 0 (via `kernelgen-mcp-setup.md`).**\n> By the time execution reaches this sub-skill, MCP should already be configured and ready.\n> If MCP tools are unavailable or calls fail at any point, stop and tell the user:\n> \"MCP 服务不可用，请检查配置或访问 https://kernelgen.flagos.io/mcp 获取 Token 后重新配置。\"\n\n**If the MCP server IS configured**, proceed to Step 1.\n\n## Step 1: Understand the Operator Request\n\nParse the user's description to determine:\n- `kernel_name`: operator name (e.g., `rms_norm`, `silu_and_mul`, `rotary_embedding`, `paged_attention`)\n- `func_desc`: what the operator does\n- `func_type`: one of `activation`, `norm`, `attention`, `quantization`, `moe`, `sampling`, `other`.\n **Auto-infer from the operator signature when possible:**\n - If name contains `relu`, `gelu`, `silu`, `swish`, `tanh`, `sigmoid` → `activation`\n - If name contains `norm`, `rmsnorm`, `layernorm`, `batchnorm` → `norm`\n - If name contains `attention`, `flash`, `paged` → `attention`\n - If name contains `quant`, `awq`, `gptq`, `squeezellm` → `quantization`\n - If name contains `moe`, `expert` → `moe`\n - If name contains `sample`, `top_k`, `top_p` → `sampling`\n - If operator has >=2 tensor inputs → `binary_pointwise`\n - Otherwise → `other`\n Confirm the inferred `func_type` with the user if ambiguous.\n- `arg_names`, `arg_type`, `arg_descs`, `output_arg_desc`: parameter information\n\n**Normalize `kernel_name` to snake_case** before using it anywhere:\n- `RMSNorm` → `rms_norm`\n- `SiluAndMul` → `silu_and_mul`\n- `RMS-Norm` → `rms_norm`\n- `layerNorm` → `layer_norm`\n\nThis ensures consistent file naming and codebase search in subsequent steps.\n\nCheck this common alias table before asking the user:\n\n| User says | Correct `kernel_name` |\n|---|---|\n| swish | `silu` |\n| negative / negate | `neg` |\n| power | `pow` |\n| absolute | `abs` |\n| multiply | `mul` |\n| divide | `div` |\n| clip | `clamp` |\n| square | `pow` (exp=2) |\n| square_root | `sqrt` |\n| cube | `pow` (exp=3) |\n| cube_root | `cbrt` |\n| hard_swish | `hardswish` |\n| hard_sigmoid | `hardsigmoid` |\n\nIf the alias is not in this table, ask the user to clarify.\n\n**Keep `kernel_name` concise** — ideally 30 characters or fewer. If the normalized name is\ntoo long (e.g., `fused_swiglu_layernorm_kernel`), suggest a shorter alias to the user (e.g.,\n`fused_swiglu_ln`). Long names create unwieldy file paths like\n`tests/kernels/test_fused_swiglu_layernorm_kernel.py`.\n\n**If argument information is missing or the description is vague** (e.g., \"generate a fused swiglu kernel\"):\n1. Use the Grep tool to search for the operator name across the repository, specifically in:\n - `csrc/` — CUDA kernel implementations (function signatures)\n - `vllm/model_executor/layers/` — Python layer wrappers (forward method signatures)\n - `vllm/kernels/` — existing Triton kernels\n - `vllm/_custom_ops.py` — custom op registrations (full function signatures)\n2. From the search results, infer the argument names, types, and descriptions.\n3. Present the inferred signature to the user for confirmation before proceeding.\n4. **Fallback**: If no matches are found in the codebase, use a generic tensor signature as\n the starting point and ask the user to confirm or refine:\n ```\n arg_names: [\"x\"]\n arg_type: [\"torch.Tensor\"]\n arg_descs: [\"Input tensor, typically shaped (num_tokens, hidden_size) for most vLLM kernels, or (num_heads, seq_len, head_dim) for attention/rope kernels\"]\n output_arg_desc: \"Output tensor of same shape as input\"\n ```\n Note: vLLM kernels almost always use `(num_tokens, hidden_size)` as the input shape\n convention — NOT `(batch, seq_len, hidden)`. Use 2D shapes by default.\n5. If the operator is truly ambiguous even after presenting the fallback, stop and ask\n the user to provide the full signature.\n\n## Step 2: Check Whether the Operator Already Exists\n\nBefore calling the MCP generator, **thoroughly search** the codebase using the Glob and Grep tools:\n\n1. **Triton kernels**: Use `Glob(\"vllm/kernels/**/**\")` to find Triton-based kernel files.\n2. **Triton kernels (alt)**: Use `Glob(\"vllm/triton_kernels/**/**\")` — some kernels like rms_norm, silu_and_mul live here.\n3. **CUDA kernels**: Use `Glob(\"csrc/**/**\")` to find CUDA kernel files.\n4. **Attention backends**: Use `Glob(\"vllm/attention/**/**\")` — many attention/flash kernels live here.\n5. **Custom ops**: First use `Glob(\"vllm/**/*custom_op*\")` to locate the actual custom op file\n (commonly `vllm/_custom_ops.py` or `vllm/model_executor/custom_op.py`), then use Grep on\n the found file(s) to check for the op name.\n6. **Model layers**: Use `Grep(pattern=\"\", path=\"vllm/model_executor/layers/\")` for layer impls.\n7. **Tests**: Use `Glob(\"tests/kernels/**\")` to find test files.\n8. **Benchmarks**: Use `Glob(\"benchmarks/kernels/**\")` to find benchmark files.\n\nAlso search for common aliases and naming variants:\n- Underscore vs concatenated: `rms_norm` vs `rmsnorm`, `silu_mul` vs `silu_and_mul`\n- CamelCase vs snake_case: `RMSNorm` vs `rms_norm`, `SiluAndMul` vs `silu_and_mul`\n- All-lowercase: `rmsnorm`, `layernorm`, `rotaryembedding`\n- **Partial name components**: for compound names like `silu_and_mul`, also search for `silu`\n and `mul` individually — the kernel may live inside a file with a different name (e.g.,\n `activation.py` contains `silu_and_mul`).\n\n### If the operator already exists\n\nPresent the findings to the user and ask them to choose one of the following. Stop and wait\nfor the user to respond before continuing:\n\n> The operator `` already exists in the codebase:\n> - Implementation: ``\n> - Tests: ``\n> - Benchmarks: ``\n>\n> Please choose how to proceed:\n>\n> **A — Skip generation**: The operator already exists; do nothing.\n>\n> **B — Replace existing**: Overwrite the current implementation with MCP-generated code.\n>\n> **C — Create a custom variant (side-by-side)**: Generate as `_v2` alongside\n> the original, without overriding the existing dispatch.\n\nOnly proceed to Step 3 after the user has made a choice.\n\n### If the operator does NOT exist\n\nProceed directly to Step 3.\n\n## Step 3: Research Context (flagos_wiki)\n\nBefore calling the MCP generator, use the Read tool to gather reference materials. This improves\ngeneration quality significantly.\n\n1. **Read similar operator code** from the codebase:\n - Activation op → read `csrc/activation_kernels.cu` or `vllm/model_executor/layers/activation.py`\n - Norm op → read `csrc/layernorm_kernels.cu` or `vllm/model_executor/layers/layernorm.py`\n - Attention op → read files under `csrc/attention/` or `vllm/attention/`\n - Quantization op → read files under `csrc/quantization/`\n - MoE op → read files under `csrc/moe/` or `vllm/model_executor/layers/fused_moe/`\n\n2. **Read the test pattern** from the matching test file to understand shapes, dtypes, assertions.\n\n3. **Read the benchmark pattern** from the matching benchmark file to understand timing methodology.\n\n4. **If replacing an existing op** (Option B), read the current implementation so the new version\n can be compared / improved upon.\n\n **If reference files cannot be found or the Read tool fails**, do not stop. Proceed with\n general algorithm hints based on your knowledge of the operator type instead. An empty or\n minimal `flagos_wiki` is acceptable — it is better to continue than to block the workflow.\n\n5. Collect all findings into the `flagos_wiki` parameter as a `List[str]`. Rules:\n - Each string should be a concise reference snippet, **under 200 characters** (not tokens).\n - **Maximum 10 items** — avoid token explosion.\n - **Prefer algorithm insights over file path descriptions** — the MCP generator benefits\n more from understanding *how* an algorithm works than *where* a file lives.\n - Avoid redundant items — do not repeat the same concept in different words (e.g., \"block\n reduction\", \"shared memory reduction\", \"warp reduction\" are all the same idea).\n - **Always include a kernel pattern hint** as the first item — this helps the MCP choose\n the right code structure. The first item **MUST** start with the literal prefix\n `\"Kernel pattern: \"` followed by one of: `elementwise`, `row-wise`, `reduction`,\n `matmul-like`, `fused-multi-op`. Example: `\"Kernel pattern: row-wise reduction over hidden_size\"`.\n Do not use alternative prefixes like `\"Pattern:\"` or `\"Type:\"` — the MCP parser expects\n exactly `\"Kernel pattern: \"`.\n - **Always include a memory layout hint**: `\"Input tensors are contiguous row-major layout\"`\n — this reduces Triton miscompile risk for strided inputs.\n\n```python\nflagos_wiki = [\n # GOOD: algorithm insights\n \"Layernorm uses Welford online algorithm for numerically stable variance computation\",\n \"Block reduction via shared memory with warp-level shuffle for final reduction step\",\n \"Vectorized loads using float4 for memory bandwidth optimization\",\n \"RMSNorm computes 1/sqrt(mean(x^2) + eps) * x, no mean subtraction unlike LayerNorm\",\n # OK: structural hints when relevant\n \"Test parametrizes over NUM_TOKENS=[1,17,86,1234,3045] and HIDDEN_SIZES=[16,48,128,1562,4096]\",\n \"Benchmark uses triton.testing.do_bench_cudagraph with quantiles=[0.5, 0.2, 0.8]\",\n # BAD (avoid): bare file paths with no insight\n # \"File: csrc/layernorm_kernels.cu\"\n]\n```\n\n## Step 4: Call kernelgen-mcp\n\nInvoke `mcp__kernelgen-mcp__generate_kernel` with the parameters gathered above:\n\n```\nkernel_name: \"\"\nfunc_desc: \"\"\nfunc_type: \"\"\narg_names: []\narg_type: []\narg_descs: []\noutput_arg_desc: \"\"\nflagos_wiki: []\n```\n\nThe MCP returns code blocks which may include:\n- `torch_code` — PyTorch reference implementation\n- `triton_code` — Triton kernel implementation\n- `test_func_code` — accuracy test code\n- `benchmark_func_code` — performance benchmark code\n\n**If `torch_code` is missing from the MCP response**, construct a `ref_impl` yourself using\nstandard PyTorch operations based on `func_desc`. For example, for `rms_norm`:\n```python\ndef ref_impl(x, weight, eps=1e-6):\n return x * torch.rsqrt(x.pow(2).mean(-1, keepdim=True) + eps) * weight\n```\nThis is needed so that accuracy tests have a reference to compare against.\n\n## Step 4.5: MCP Report Review\n\nIf the MCP response includes test results (accuracy test reports, benchmark/speedup reports),\npresent them to the user **in full** before proceeding to local code adaptation:\n\n```\n=== MCP Generation Test Report ===\n\nAccuracy Test Results:\n\n\nPerformance Benchmark Results:\n\n```\n\nThen ask the user:\n\n> The MCP service has returned the above test reports from its generation environment.\n> Would you like to:\n>\n> **A — Skip local testing**: Trust the MCP reports and proceed directly to code placement\n> and summary. No local tests will be run.\n>\n> **B — Run local tests**: Run accuracy and performance tests on your local machine as well\n> to verify the results in your specific environment.\n\nWait for the user's choice before proceeding:\n- If **A (skip local testing)**: After placing code in Step 5, skip Steps 6 and 7 (local\n accuracy tests and benchmarks) and proceed directly to Step 8 (Summary). Use the MCP\n report numbers for the summary.\n- If **B (run local tests)**: Proceed normally through Steps 5 → 6 → 7 → 8. Before\n running local tests, perform the **Chip Compatibility Check** described in Step 5.5.\n\n**If the MCP response does NOT include test reports** (only code blocks), proceed normally\nto Step 5 — local testing will be required.\n\n## Step 5: Adapt and Place Code into the vLLM Project\n\nThis is the most critical step. The generated code must be transformed to match vLLM conventions.\nThe exact placement depends on the operator type and the user's choice in Step 2.\n\n---\n\n### Path 1: New operator / Replace existing (Option B or new op)\n\n#### 5a. Determine the Correct Location\n\nBased on operator type, place the code in the appropriate location:\n\n| Operator Type | Triton Kernel Location | Layer Wrapper Location | CUDA Alternative |\n|---|---|---|---|\n| Activation | `vllm/kernels/` or inline in layer | `vllm/model_executor/layers/activation.py` | `csrc/activation_kernels.cu` |\n| Normalization | `vllm/kernels/` or inline in layer | `vllm/model_executor/layers/layernorm.py` | `csrc/layernorm_kernels.cu` |\n| Attention | `vllm/attention/backends/` | `vllm/attention/layer.py` | `csrc/attention/` |\n| Quantization | `vllm/kernels/` | `vllm/model_executor/layers/quantization/` | `csrc/quantization/` |\n| MoE | `vllm/kernels/` | `vllm/model_executor/layers/fused_moe/` | `csrc/moe/` |\n| Position Encoding | `vllm/kernels/` | `vllm/model_executor/layers/rotary_embedding.py` | `csrc/pos_encoding_kernels.cu` |\n| Cache | `vllm/kernels/` | N/A | `csrc/cache_kernels.cu` |\n| Sampling | `vllm/kernels/` | `vllm/model_executor/layers/sampler.py` | `csrc/sampler.cu` |\n| Other | `vllm/kernels/.py` | As needed | As needed |\n\n#### 5b. Transform the Triton Kernel Code\n\nAdapt the MCP-generated Triton code to follow vLLM conventions:\n\n```python\n# SPDX-License-Identifier: Apache-2.0\n# SPDX-FileCopyrightText: Copyright contributors to the vLLM project\n\nimport torch\nimport triton\nimport triton.language as tl\n\nfrom vllm.logger import init_logger\n\nlogger = init_logger(__name__)\n\n\n@triton.autotune(\n configs=[\n triton.Config({\"BLOCK_SIZE\": 64}, num_warps=2),\n triton.Config({\"BLOCK_SIZE\": 128}, num_warps=4),\n triton.Config({\"BLOCK_SIZE\": 256}, num_warps=4),\n triton.Config({\"BLOCK_SIZE\": 512}, num_warps=8),\n ],\n # Choose autotune key based on kernel pattern:\n # elementwise → key=[\"n_elements\"]\n # row-wise (layernorm, softmax) → key=[\"hidden_size\"] or [\"n_cols\"]\n # IMPORTANT: the key MUST match an existing non-pointer, non-constexpr\n # kernel argument name. If the key does not exist in the kernel signature,\n # Triton will raise a KeyError at compile time.\n key=[\"n_elements\"], # ← replace with the actual problem-size arg name\n)\n@triton.jit\ndef __kernel(\n # pointer args\n input_ptr,\n output_ptr,\n # dimensions\n n_elements,\n # meta parameters\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 x = tl.load(input_ptr + offsets, mask=mask)\n # kernel logic here\n output = x # placeholder\n tl.store(output_ptr + offsets, output, mask=mask)\n\n\ndef (x: torch.Tensor) -> torch.Tensor:\n \"\"\"Python wrapper for the Triton kernel.\"\"\"\n assert x.is_cuda, \"Input tensor must be on CUDA device\"\n assert x.is_contiguous(), \"Input tensor must be contiguous\"\n output = torch.empty_like(x)\n n_elements = x.numel()\n\n def grid(meta):\n return (triton.cdiv(n_elements, meta[\"BLOCK_SIZE\"]),)\n\n __kernel[grid](\n x, output, n_elements,\n )\n return output\n```\n\n**Important**: The example above is an **elementwise** kernel pattern. For **row-wise** kernels\n(layernorm, softmax, rmsnorm), the grid and indexing are different:\n\n```python\ndef (x: torch.Tensor, ...) -> torch.Tensor:\n assert x.is_cuda and x.is_contiguous()\n output = torch.empty_like(x)\n num_rows, hidden_size = x.shape\n\n # Row-wise: one program per row, each processes hidden_size elements\n grid = (num_rows,)\n\n __kernel[grid](\n x, output, hidden_size,\n # BLOCK_SIZE is autotuned via @triton.autotune\n )\n return output\n```\n\n**Choose the correct pattern based on kernel type:**\n- `activation` (elementwise): `grid = (triton.cdiv(n_elements, BLOCK_SIZE),)` — one program per block of elements\n- `norm` / `softmax` (row-wise): `grid = (num_rows,)` — one program per row\n- `attention` / `moe`: typically custom 2D/3D grids — follow MCP output\n- Always include `SPDX-License-Identifier: Apache-2.0` and copyright header\n- Use `vllm.logger.init_logger(__name__)` for logging\n- Kernel function names start with `_` and end with `_kernel`\n- Public wrapper function has a clean Python signature; use `x` (not `input`) as the tensor\n parameter name to match vLLM convention and avoid shadowing Python's built-in `input()`\n- Use `torch.empty_like` or explicit allocation for outputs\n- Use `triton.cdiv` for grid computation\n- Use `tl.constexpr` for meta parameters like `BLOCK_SIZE`\n- **Use `@triton.autotune`** with power-of-two block sizes (64, 128, 256, 512) instead of\n hardcoding a single BLOCK_SIZE. Avoid 1024 by default (register spill risk). Omit autotune\n only for very simple kernels where a fixed block size is clearly optimal.\n- **Choose the correct autotune `key`**: use the main problem-size dimension that the launch grid\n depends on. For elementwise kernels use `key=[\"n_elements\"]`, for row-wise kernels (e.g.,\n layernorm, softmax) use `key=[\"hidden_size\"]` or `key=[\"n_cols\"]`. The key must be a\n non-pointer, non-constexpr parameter of the kernel.\n\n#### 5c. Register the Operator (if applicable)\n\nIf the kernel needs to be callable via `vllm._custom_ops`:\n\n1. Use the Edit tool to add a Python wrapper function in `vllm/_custom_ops.py` (or the project's\n actual custom op registration file if the path differs — verify with Glob first) that calls\n the Triton kernel.\n2. If wrapping as a layer, use the Edit tool to add or modify the appropriate file under\n `vllm/model_executor/layers/`.\n\nFor Triton kernels that don't need custom op registration (standalone usage), just ensure\nthe module is importable from `vllm.kernels.`.\n\n#### 5d. Accuracy Test → `tests/kernels/test_.py`\n\nUse the Write tool to create a test file following vLLM test conventions:\n\n```python\n# SPDX-License-Identifier: Apache-2.0\n# SPDX-FileCopyrightText: Copyright contributors to the vLLM project\n\nimport pytest\nimport torch\n\n\nDTYPES = [torch.float16, torch.bfloat16, torch.float32]\nNUM_TOKENS = [1, 17, 128, 1024]\nHIDDEN_SIZES = [128, 512, 1024, 4096]\nSEEDS = [0]\nif not torch.cuda.is_available():\n pytest.skip(\"CUDA not available\", allow_module_level=True)\n\nCUDA_DEVICES = [f\"cuda:{i}\" for i in range(torch.cuda.device_count())]\n\n# Tolerance per dtype — fp16/bf16 need looser tolerances to avoid flaky CI\nTOLERANCES = {\n torch.float32: {\"rtol\": 1e-5, \"atol\": 1e-5},\n torch.float16: {\"rtol\": 1e-2, \"atol\": 1e-2},\n torch.bfloat16: {\"rtol\": 2e-2, \"atol\": 2e-2},\n}\n\n\ndef ref_impl(x: torch.Tensor) -> torch.Tensor:\n \"\"\"Reference PyTorch implementation.\n\n This function uses pure PyTorch ops to compute the expected output.\n It is used as ground truth for accuracy comparison.\n \"\"\"\n # Use the torch_code from MCP as reference\n ...\n\n\n@pytest.mark.parametrize(\"num_tokens\", NUM_TOKENS)\n@pytest.mark.parametrize(\"hidden_size\", HIDDEN_SIZES)\n@pytest.mark.parametrize(\"dtype\", DTYPES)\n@pytest.mark.parametrize(\"seed\", SEEDS)\n@pytest.mark.parametrize(\"device\", CUDA_DEVICES)\n@torch.inference_mode()\ndef test_(\n num_tokens: int,\n hidden_size: int,\n dtype: torch.dtype,\n seed: int,\n device: str,\n) -> None:\n torch.random.manual_seed(seed)\n if torch.cuda.is_available():\n torch.cuda.manual_seed(seed)\n x = torch.randn(num_tokens, hidden_size, dtype=dtype, device=device)\n\n ref_out = ref_impl(x)\n\n from vllm.kernels. import as kernel_fn\n act_out = kernel_fn(x)\n\n tol = TOLERANCES[dtype]\n torch.testing.assert_close(act_out, ref_out, rtol=tol[\"rtol\"], atol=tol[\"atol\"])\n```\n\nKey conventions:\n- Always include the SPDX license header\n- Use `@torch.inference_mode()` decorator\n- Parametrize with `DTYPES`, `NUM_TOKENS`, `HIDDEN_SIZES`, `SEEDS`, `CUDA_DEVICES`\n- Use `torch.testing.assert_close` for comparison\n- **Use per-dtype tolerances**: fp32 can use 1e-5, but fp16/bf16 need 1e-2 to avoid false failures\n- Provide `ref_impl` using **pure PyTorch ops** (from MCP's `torch_code`) — not `torch.` which\n may not exist for custom fused kernels\n- Use `torch.random.manual_seed` for reproducibility; guard `torch.cuda.manual_seed` with\n `if torch.cuda.is_available()` to avoid crash on CPU-only environments\n\n#### 5e. Performance Benchmark → `benchmarks/kernels/benchmark_.py`\n\nUse the Write tool to create a benchmark file following vLLM benchmark conventions:\n\n```python\n# SPDX-License-Identifier: Apache-2.0\n# SPDX-FileCopyrightText: Copyright contributors to the vLLM project\n\nimport itertools\n\nimport torch\n\nif not torch.cuda.is_available():\n raise RuntimeError(\"CUDA is required to run this benchmark\")\n\nimport triton\nimport triton.testing\n\nfrom vllm.utils.torch_utils import set_random_seed\n\nbatch_size_range = [1, 16]\nseq_len_range = [1, 64, 1024]\nhidden_size_range = [1024, 4096]\n# 2 * 3 * 2 = 12 configs — keeps benchmark under 2 minutes\nconfigs = list(itertools.product(batch_size_range, seq_len_range, hidden_size_range))\n\n\ndef ref_impl(x: torch.Tensor) -> torch.Tensor:\n \"\"\"Reference PyTorch implementation (same as in the test file).\n\n Used as the baseline for performance comparison.\n \"\"\"\n # Paste the same ref_impl from the test file\n ...\n\n\ndef benchmark_(\n batch_size: int,\n seq_len: int,\n hidden_size: int,\n provider: str,\n dtype: torch.dtype,\n):\n device = torch.device(\"cuda\")\n num_tokens = batch_size * seq_len\n set_random_seed(42)\n\n x = torch.randn(num_tokens, hidden_size, dtype=dtype, device=device).contiguous()\n\n from vllm.kernels. import as kernel_fn\n\n # Warmup — compile Triton kernel and stabilize GPU clocks before benchmarking\n for _ in range(5):\n kernel_fn(x)\n torch.cuda.synchronize()\n\n if provider == \"triton\":\n fn = lambda: kernel_fn(x)\n elif provider == \"torch\":\n fn = lambda: ref_impl(x)\n else:\n raise ValueError(f\"Unknown provider: {provider}\")\n\n ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(\n fn, quantiles=[0.5, 0.2, 0.8]\n )\n return ms, min_ms, max_ms\n\n\n@triton.testing.perf_report(\n triton.testing.Benchmark(\n x_names=[\"batch_size\", \"seq_len\", \"hidden_size\"],\n x_vals=configs,\n line_arg=\"provider\",\n line_vals=[\"triton\", \"torch\"],\n line_names=[\"Triton Kernel\", \"PyTorch Reference\"],\n styles=[(\"blue\", \"-\"), (\"green\", \"-\")],\n ylabel=\"Time (ms)\",\n plot_name=\"-benchmark\",\n # float16 is the default because Triton kernels are typically optimized for fp16 inference\n args={\"dtype\": torch.float16},\n )\n)\ndef bench_(batch_size, seq_len, hidden_size, provider, dtype):\n return benchmark_(batch_size, seq_len, hidden_size, provider, dtype)\n\n\nif __name__ == \"__main__\":\n bench_.run(print_data=True)\n```\n\nKey conventions:\n- Use `import triton; import triton.testing` directly (not `vllm.triton_utils`) for benchmarks,\n as `vllm.triton_utils` may not export `triton.testing` in all versions\n- Use `triton.testing.do_bench_cudagraph` for accurate GPU timing\n- Use `triton.testing.perf_report` and `triton.testing.Benchmark` for structured output\n- **Always use `ref_impl` for the torch baseline**, never `torch.` — many custom/fused\n kernels have no single torch equivalent\n- Parametrize across batch sizes, sequence lengths, and hidden sizes\n- Use `set_random_seed` for reproducibility\n\n---\n\n### Path 2: Custom variant side-by-side (Option C)\n\n#### 5a. Operator Implementation → `vllm/kernels/_v2.py`\n\nUse the Write tool to create the file. Use the **raw Triton pointer-based style** for the kernel.\nKeep the MCP-generated Triton code mostly as-is but ensure it:\n- Has the SPDX license header\n- Is self-contained (minimal vLLM imports)\n- Has proper `@triton.jit` kernel with pointer args, mask, and autotuned BLOCK_SIZE\n- Exports a clean Python wrapper function named `_v2`\n\n#### 5b. Registration (limited)\n\nDo **NOT** register the variant in `vllm/_custom_ops.py` — it must NOT override the existing\noperator. The variant should be importable directly from `vllm.kernels._v2`.\n\n#### 5c. Standalone Test → `tests/kernels/test__v2.py`\n\nUse the Write tool to create a self-contained test file:\n\n```python\n# SPDX-License-Identifier: Apache-2.0\n# SPDX-FileCopyrightText: Copyright contributors to the vLLM project\n\nimport pytest\nimport torch\n\nfrom vllm.kernels._v2 import _v2 as kernel_fn\n\n\nDTYPES = [torch.float16, torch.bfloat16, torch.float32]\nSHAPES = [(1, 128), (16, 512), (128, 1024), (1024, 4096)]\nSEEDS = [0]\n\nTOLERANCES = {\n torch.float32: {\"rtol\": 1e-5, \"atol\": 1e-5},\n torch.float16: {\"rtol\": 1e-2, \"atol\": 1e-2},\n torch.bfloat16: {\"rtol\": 2e-2, \"atol\": 2e-2},\n}\n\n\ndef ref_impl(x: torch.Tensor) -> torch.Tensor:\n \"\"\"Reference PyTorch implementation.\"\"\"\n ...\n\n\n@pytest.mark.parametrize(\"shape\", SHAPES)\n@pytest.mark.parametrize(\"dtype\", DTYPES)\n@pytest.mark.parametrize(\"seed\", SEEDS)\n@torch.inference_mode()\ndef test__v2_accuracy(shape, dtype, seed):\n torch.random.manual_seed(seed)\n x = torch.randn(shape, dtype=dtype, device=\"cuda\")\n ref_out = ref_impl(x)\n act_out = kernel_fn(x)\n tol = TOLERANCES[dtype]\n torch.testing.assert_close(act_out, ref_out, rtol=tol[\"rtol\"], atol=tol[\"atol\"])\n```\n\n#### 5d. Standalone Benchmark → `benchmarks/kernels/benchmark__v2.py`\n\nUse the Write tool to create a self-contained benchmark file following the same pattern as\nPath 1 Section 5e, but importing from `vllm.kernels._v2` and using `ref_impl`\nas the torch baseline.\n\n---\n\n## Step 6: Run Accuracy Tests\n\n### Step 5.5: Chip Compatibility Check (before local testing)\n\nBefore running any local tests, verify that the local hardware matches the target device.\nvLLM operators typically target NVIDIA GPUs, but verify this explicitly:\n\n```bash\npython - <<'PY'\nimport subprocess\ntry:\n import torch\n if torch.cuda.is_available():\n for i in range(torch.cuda.device_count()):\n print(f\"nvidia: {torch.cuda.get_device_name(i)}\")\n else:\n print(\"No CUDA device available\")\nexcept Exception:\n print(\"torch not available\")\ntry:\n r = subprocess.run([\"npu-smi\", \"info\"], capture_output=True, text=True, timeout=5)\n if r.returncode == 0 and \"NPU\" in r.stdout:\n print(\"huawei: Ascend NPU detected\")\nexcept Exception:\n pass\nPY\n```\n\n**If the detected hardware does NOT match the operator's target device** (e.g., operator\ntargets `huawei` but local machine has NVIDIA GPU, or vice versa), warn the user:\n\n```\n⚠️ Hardware mismatch detected:\n - Operator target device: \n - Local hardware: \n\nRunning tests for a operator on hardware will likely\nfail or produce incorrect results. Please choose:\n\n A — Skip local testing: Do not run local tests; rely on MCP test reports (if available).\n B — Proceed anyway: Run local tests despite the mismatch (results may be unreliable).\n```\n\nIf the user chooses A, skip Steps 6 and 7 and proceed to Step 8 (Summary).\nIf the user chooses B, proceed with local tests but note the mismatch in the final report.\n\n---\n\nUse the Bash tool to run the accuracy tests for the newly added operator.\n\nFirst, verify the working directory is the vLLM repo root:\n```bash\npwd\n```\nIf the output is NOT the vLLM repo root (should contain `vllm/`, `csrc/`, `tests/`), stop\nand ask the user for the correct path.\n\nThen run with `PYTHONPATH=.` to ensure vllm is importable even if not pip-installed.\n\n**First, do a quick smoke test** to catch Triton compilation errors before running the full\npytest suite (which produces messy output on compilation failures):\n\n```bash\nPYTHONPATH=. python -c \"\nfrom vllm.kernels. import as fn\nimport torch\nx = torch.randn(4, 128, dtype=torch.float16, device='cuda')\n# If kernel requires extra arguments (e.g. weight, bias, eps for norm kernels),\n# create matching dummy tensors:\n# weight = torch.ones(128, dtype=torch.float16, device='cuda')\n# out = fn(x, weight, eps=1e-6)\nout = fn(x)\nprint('Smoke test passed, output shape:', out.shape)\n\"\n```\n\nIf this fails with `triton.compiler.CompilationError` or `ImportError`, fix the kernel code\nfirst — do not proceed to pytest. **Adapt the smoke test arguments to match the kernel's\nactual signature** — if the kernel requires `weight`, `bias`, `eps`, or other parameters,\ncreate appropriate dummy tensors or scalar values. Use `inspect.signature` to infer missing\narguments (e.g., `weight`, `bias`, `eps` for `layer_norm`) before constructing the smoke test.\n\n**If smoke test passes, run the full test suite:**\n\n**For Path 1 (new / replace):**\n```bash\nPYTHONPATH=. python -m pytest tests/kernels/test_.py -v\n```\n\n**For Path 2 (custom variant):**\n```bash\nPYTHONPATH=. python -m pytest tests/kernels/test__v2.py -v\n```\n\nReport the results to the user. If tests fail, follow the **error classification and retry\nprotocol** below. This protocol strictly limits self-fix attempts to prevent infinite loops.\n\n### Error Classification\n\nFirst, classify the error into one of two categories:\n\n**Category A — Compilation / Import errors** (model may attempt 1 self-fix):\n- `ImportError`, `ModuleNotFoundError` — wrong import path or missing module\n- `SyntaxError` — Python syntax issue\n- `TritonCompilationError`, `CompilationError` — Triton kernel compile failure\n- `NameError` — undefined variable or function\n- `TypeError` in function call signature — wrong number/type of arguments\n- `AttributeError` — wrong attribute access on a module or object\n\n**Category B — Algorithm / Numerical accuracy errors** (do NOT self-fix):\n- `AssertionError` from `torch.testing.assert_close` — numerical mismatch\n- Wrong output values (results don't match reference)\n- Shape mismatch in output tensors\n- NaN or Inf in output\n- Any error that indicates the kernel logic is incorrect\n\n### Retry Protocol (strictly follow this order)\n\n**Step 6a. (Category A only) Self-fix — maximum 1 attempt:**\nIf the error is Category A (compilation/import), you may attempt exactly ONE fix:\n1. Read the error traceback carefully.\n2. Apply a targeted fix using the Edit tool (e.g., fix import path, fix syntax, fix argument name).\n3. Re-run the tests using the Bash tool.\n4. If the test **passes** → proceed to Step 7.\n5. If the test **still fails** → proceed to Step 6b. Do NOT attempt a second self-fix.\n\n**If the error is Category B (algorithm/accuracy), skip Step 6a entirely** — go directly\nto Step 6b. Do NOT attempt to fix algorithm logic yourself, as this typically leads to\nan endless fix-retry loop without converging. The MCP service has better optimization\ncapabilities for these issues.\n\n**Step 6b. MCP re-generation — pass error context to generate_kernel:**\nRe-call `mcp__kernelgen-mcp__generate_kernel` with the **same parameters as Step 4**,\nbut add the error information to `flagos_wiki` as additional hints:\n```python\nflagos_wiki = [\n # ... original flagos_wiki items from Step 3 ...\n \"Previous generation failed accuracy test: \",\n \"Error was: \",\n \"Fix hint: \"\n]\n```\nReplace the kernel code with the new MCP output, re-run tests.\n- If tests **pass** → proceed to Step 7.\n- If tests **still fail** → proceed to Step 6c.\n\n**Step 6c. MCP optimization — pass error context to optimize_kernel:**\nTry `mcp__kernelgen-mcp__optimize_kernel` with the current kernel code and the\n`check_result` parameter containing the error traceback. This endpoint can fix\nmemory access patterns, index calculations, and numerical issues.\nReplace the kernel code with the optimized output, re-run tests.\n- If tests **pass** → proceed to Step 7.\n- If tests **still fail** → proceed to Step 6d.\n\n**Step 6d. Stop and report:**\nDo not keep retrying. Report the failure to the user with:\n- The specific test failures and error messages\n- Your analysis of what might be wrong\n- Suggestion to try with different `func_type` or additional `flagos_wiki` hints\n\n## Step 7: Run Performance Benchmark\n\nUse the Bash tool to run the benchmark with `PYTHONPATH=.`:\n\n**For Path 1 (new / replace):**\n```bash\nPYTHONPATH=. python benchmarks/kernels/benchmark_.py\n```\n\n**For Path 2 (custom variant):**\n```bash\nPYTHONPATH=. python benchmarks/kernels/benchmark__v2.py\n```\n\nReport the speedup results to the user.\n\n## Step 8: Summary\n\nProvide a clear summary to the user with **exact numbers** extracted from test and benchmark\noutput:\n\n```\n=== KernelGen Operator Generation Report ===\n\nOperator Name: \nGeneration Mode: New / Replace Existing / Custom Variant (v2)\n\nFile Changes:\n - [New/Modified] vllm/kernels/.py (Triton kernel)\n - [Modified] vllm/_custom_ops.py (op registration, if applicable)\n - [Modified] vllm/model_executor/layers/.py (layer wrapper, if applicable)\n - [New/Modified] tests/kernels/test_.py (accuracy test)\n - [New/Modified] benchmarks/kernels/benchmark_.py (benchmark)\n\nAccuracy Tests: passed, failed (total test cases)\n Pass Rate: %\n Failed Cases: \n\nPerformance Benchmark:\n Avg Speedup: x vs PyTorch reference\n Best Speedup: x (batch_size=, seq_len=, hidden_size=)\n Worst Speedup: x (batch_size=, seq_len=, hidden_size=)\n (If benchmark was not run or failed, write \"Incomplete\" and explain why)\n\nMCP Iterations: <1, 2, or 3> (write 1 if first generation passed)\n\nPerformance Analysis:\n 5.0x → \"Extreme fusion success, typical for fused multi-op kernels (e.g. fused SwiGLU, fused GeLU+bias)\"\n - Speedup > 3.0x → \"Excellent compute-bound optimization, significant kernel fusion benefit\"\n - Speedup > 2.0x → \"Compute-bound optimization effective, significant speedup achieved\"\n - Speedup 1.2x ~ 2.0x → \"Moderate speedup, kernel likely balanced between compute and memory\"\n - Speedup 0.5x ~ 1.2x → \"Kernel likely memory-bound, limited optimization headroom\"\n - Speedup < 0.5x → \"Kernel slower than PyTorch reference — likely suboptimal memory access pattern\">\n\nIssues and Fixes: (if any)\n```\n\n**After presenting the summary, check the speedup:**\n\n- **If average speedup < 0.5x** (kernel slower than PyTorch), proactively warn the user:\n ```\n ⚠️ The generated Triton kernel is slower than the PyTorch reference implementation.\n This is typically caused by suboptimal memory access patterns or improper block size configuration.\n Would you like to use /kernelgen_optimizer to optimize this kernel's performance?\n ```\n\n- **If average speedup is 0.5x ~ 1.2x**, ask the user:\n ```\n Current speedup is low (x), there may be room for optimization.\n Would you like to use /kernelgen_optimizer to try improving performance?\n ```\n\n- **If average speedup > 1.2x**, no action needed — report the result normally.\n\n**How to extract the numbers:**\n- **Pass/fail counts**: Parse pytest output for the line matching `X passed` or `X passed, Y failed`.\n Use regex pattern `(\\d+) passed` and `(\\d+) failed` to extract.\n- **Speedup**: The benchmark outputs a table with columns for each provider. For each config row:\n 1. Find the `Triton Kernel` time (ms) and `PyTorch Reference` time (ms)\n 2. Calculate `speedup = torch_time / triton_time`\n 3. Compute min, max, and average across all config rows\n - If the output format is not a table, look for lines containing `provider=triton` and\n `provider=torch` with time values, and compute the ratio.\n - Use these regex patterns to extract timing values:\n - `(\\d+(?:\\.\\d+)?)\\s*(us|ms|s)` — number with unit\n - `(?i)mean:\\s*(\\d+(?:\\.\\d+)?)\\s*(us|ms|s)` — mean timing (pytest-benchmark format)\n - `(?i)avg:\\s*(\\d+(?:\\.\\d+)?)\\s*(us|ms|s)` — average timing\n - `(?i)median:\\s*(\\d+(?:\\.\\d+)?)\\s*(us|ms|s)` — median timing value\n - `(?i)Triton.*?(\\d+(?:\\.\\d+)?)\\s*(us|ms|s)` — Triton kernel time\n - `(?i)(torch|pytorch).*?(\\d+(?:\\.\\d+)?)\\s*(us|ms|s)` — PyTorch reference time\n - **Do not just copy the raw table** — always compute and report the actual speedup ratios.\n\n## Step 9: Post-Completion Actions\n\nAfter the summary is presented (whether based on MCP reports or local test results), perform\nthe following two checks.\n\n### 9a. Chat Tool Notification\n\nIf the user's task was received via a chat tool (e.g., Feishu/飞书, Discord, Slack, Teams,\nDingTalk/钉钉, WeChat Work/企业微信, Telegram, or any similar messaging platform), or if\nthe user mentions sending results to a client or team, ask:\n\n> The operator generation is complete. Would you like me to prepare a message to send the\n> generated files or summary to your client/team via ****?\n\nIf the user confirms:\n1. Prepare a concise summary message including: operator name, file paths, accuracy pass\n rate, speedup metrics, and any issues.\n2. If the chat tool has a CLI or API integration available in the environment, use it to\n send the message directly.\n3. If no CLI/API is available, format the message as copy-paste ready text for the user.\n4. If files need to be sent, list the exact file paths the user should attach.\n\n### 9b. Git Repository PR Submission\n\nCheck whether the current project is managed by a git-based hosting service:\n\n```bash\ngit remote -v 2>/dev/null | head -5\n```\n\nDetect the hosting platform from the remote URL:\n- `github.com` → GitHub\n- `gitlab.com` (or self-hosted GitLab) → GitLab\n- `gitee.com` → Gitee\n- `bitbucket.org` → Bitbucket\n\n**If a git hosting platform is detected**, ask the user:\n\n> This project is hosted on ****. Would you like me to automatically create a\n> Pull Request with the generated operator code?\n\n**If the user confirms**, follow the standard PR creation workflow:\n\n1. Create a new branch: `kernelgen/`\n2. Stage all changed/new files related to this operator generation (never use `git add -A`)\n3. Create a commit with a descriptive message\n4. Push the branch to the remote with `-u` flag\n5. Create the PR using the platform's CLI tool:\n\n **For GitHub** (`gh`):\n ```bash\n gh pr create --title \"Add operator via KernelGen\" --body \"$(cat <<'EOF'\n ## Summary\n Add `` operator () generated via KernelGen MCP service.\n\n ## Changes\n - [New/Modified] `` — Triton kernel implementation\n - [New/Modified] `` — Accuracy tests\n - [New/Modified] `` — Performance benchmarks\n - [Modified] `` — Operator registration (if applicable)\n\n ## Test Results\n ### Accuracy\n - **Pass rate**: / (%)\n - **Failed cases**: \n\n ### Performance\n - **Avg speedup**: x vs PyTorch reference\n - **Best**: x | **Worst**: x\n\n ## Generation Details\n - **MCP iterations**: \n - **Target device**: \n - **Operator type**: \n\n ---\n Generated by KernelGen MCP skill\n\n EOF\n )\"\n ```\n\n **For Gitee / GitLab**: Provide the PR/MR creation URL and prepared description.\n\n6. Return the PR URL to the user.\n\n**If PR creation fails**, report the error, provide the branch name, and suggest manual PR\ncreation with the prepared description.\n\n**If the user declines**, do nothing — changes remain as local files.\n\n## Important Notes\n\n- **Never overwrite existing operators** without explicit user permission (Step 2 choice).\n- **Always ask the user before installing packages** — never run `pip install` without confirmation.\n- **Follow vLLM code style** exactly: SPDX headers, logging via `init_logger`, type hints.\n- **Use explicit tools for all file operations**: Read tool to read files, **Write tool to create\n new files** (only when the file does not exist), **Edit tool to modify existing files** (never\n use Write to overwrite an existing file — use Edit to apply targeted changes), Glob tool to\n find files, Grep tool to search content, Bash tool to run commands.\n- **Use `@triton.autotune`** with multiple power-of-two BLOCK_SIZE configs (64/128/256/512)\n instead of hardcoding a single value. Include 64 for small hidden sizes (e.g., 96, 192).\n Avoid 1024 by default as it risks register spill on complex kernels — only add it if the\n kernel is very simple (e.g., elementwise with no shared state).\n- **Use per-dtype tolerances** in tests: fp32 → rtol/atol=1e-5, fp16 → 1e-2, bf16 → 2e-2.\n For complex kernels (layernorm, softmax, attention) consider even looser bf16 tolerances (3e-2)\n if CI is flaky.\n- **Always use `ref_impl` for benchmarks**, not `torch.` — many custom/fused kernels have\n no single torch equivalent function.\n- **Match the existing code patterns** when adding to existing files (e.g., `_custom_ops.py`, layer files).\n- If accuracy tests fail, try to fix the operator. If benchmark is slow, consider calling\n `mcp__kernelgen-mcp__optimize_kernel` to optimize the kernel.\n- When creating a custom variant, always make it clear that it does NOT override existing ops\n to avoid conflicts.\n- **Always use `@torch.inference_mode()`** in tests for consistency with the codebase.\n- **Use `import triton; import triton.testing`** directly in benchmarks. Do not use\n `vllm.triton_utils.triton` as it may not export `triton.testing` in all versions.\n- If the generated kernel is a CUDA kernel (not Triton), place it in `csrc/` and add bindings\n in `csrc/torch_bindings.cpp` and `csrc/ops.h` — but prefer Triton when possible.\n- Check if the kernel needs to be registered in `vllm/model_executor/custom_op.py` if it's\n an activation or layer-level op that should participate in the custom op dispatch system.\n- **Never modify unrelated files** in the repository. Do not reformat imports, fix lint, or\n make stylistic changes to files that are not directly part of the operator being generated.\n Keep the diff minimal and focused.\n\nFile v1.1.0:kernelgen-generate.md\n\n# KernelGen Skill — Generate GPU Operators via MCP (General Purpose)\n\nYou are an expert at generating GPU kernel operators using the `kernelgen-mcp` MCP service and\nintegrating them into **any** Python/Triton repository. Unlike the specialized FlagGems or vLLM\nskills, this skill dynamically discovers the target repository's structure and conventions before\nplacing code.\n\n**Tool usage**: This skill relies on the following capabilities:\n- **Bash**: execute shell commands (python, pip, pytest, etc.)\n- **Read**: read files from the codebase\n- **Write / Edit**: create or modify files\n- **Grep**: search file contents\n- **Glob**: find files by pattern\n- **MCP tools**: `mcp__kernelgen-mcp__generate_kernel` and `mcp__kernelgen-mcp__optimize_kernel`\n\n> **⚠️ MCP Prerequisite**: MCP configuration is handled by `SKILL.md` Phase 0. If MCP tools fail during execution,\n> stop and tell the user to visit https://kernelgen.flagos.io/mcp 获取 Token 后重新配置。\n\n## Execution Rules\n\nFollow these rules strictly to avoid getting lost in the multi-step workflow:\n\n1. **Follow the steps in order**: Step 0 → Step 1 → ... → Step 9. Never jump ahead.\n2. **Never skip Step 2** (repo structure discovery) — always analyze before generating.\n3. **Never skip Step 3** (operator existence check) — always verify before generating.\n4. **Do not generate or write any code before Step 5** (MCP call). Steps 0–4 are preparation only.\n5. **Always confirm with the user before destructive actions**: replacing files, installing packages,\n overwriting operators.\n6. **Report progress to the user when entering each step**. Print a short status line like\n `Entering Step 2 — Analyzing repository structure...` so the user knows the workflow is progressing.\n7. **Never fabricate repository files or paths**. If a required file cannot be found using the\n Glob or Grep tools, report it to the user instead of guessing.\n8. **Minimize file reads during discovery and search.** Always use Grep or Glob to locate\n targets first. Only Read a file when you have a specific reason (e.g., learning conventions\n from a reference operator, checking registration format). Never speculatively read files\n \"just in case\".\n9. **Never scan the entire repository.** In Step 3 and beyond, only search inside directories\n discovered in Step 2. Never run Grep on `.` or the repo root — always scope searches to\n the specific discovered source, test, or benchmark directories. Unscoped searches cause\n token explosion and match vendored code, docs, and build artifacts.\n10. **CRITICAL — MCP is mandatory**: ALL operator code generation MUST go through the\n `mcp__kernelgen-mcp__generate_kernel` MCP tool. NEVER generate Triton kernels, PyTorch\n wrappers, or operator implementations yourself — even if the operator seems simple (e.g.,\n relu, abs). If MCP is not configured, not reachable, or fails after all retries in Step 5,\n STOP and report the issue to the user. Do NOT fall back to writing kernel code manually.\n The MCP service produces optimized, tested code that manual writing cannot match.\n\n---\n\n## Step 0: Pre-flight — Environment & MCP Check\n\n### 0a. Check Python Environment\n\nUse the Bash tool to run a diagnostic that checks each package independently:\n\n```bash\npython - <<'PY'\nimport importlib, sys\n\ndef check(pkg, alias=None):\n try:\n m = importlib.import_module(pkg)\n ver = getattr(m, \"__version__\", \"installed\")\n print(f\"{alias or pkg}: {ver}\")\n return True\n except Exception:\n print(f\"{alias or pkg}: NOT installed\")\n return False\n\ncheck(\"torch\")\ncheck(\"triton\")\n\ntry:\n import torch\n print(f\"cuda_version: {torch.version.cuda}\")\n print(f\"cuda_available: {torch.cuda.is_available()}\")\n if torch.cuda.is_available():\n for i in range(torch.cuda.device_count()):\n print(f\" cuda:{i} {torch.cuda.get_device_name(i)}\")\nexcept Exception:\n pass\n\ncheck(\"pytest\")\ncheck(\"numpy\")\nPY\n```\n\n**If any import fails**, stop and ask the user before installing anything. Present the missing\npackages and let the user confirm:\n\n| Missing package | Suggested install command | Notes |\n|---|---|---|\n| `torch` | Ask the user | User MUST pick the correct CUDA variant. Never auto-install torch. |\n| `triton` | `pip install triton` | |\n| `pytest` | `pip install pytest` | |\n| `numpy` | `pip install numpy` | |\n\n**Important**: Never install `torch` without asking the user which CUDA version they need.\n\nIf `torch` is installed but `torch.cuda.is_available() == False`, warn the user that GPU tests\nwill fail and ask how to proceed.\n\nAfter the user confirms and packages are installed, re-run the diagnostic. Only proceed when\ntorch and triton imports succeed.\n\n### 0b. Check MCP Availability\n\n> **MCP configuration is handled centrally by `SKILL.md` Phase 0 (via `kernelgen-mcp-setup.md`).**\n> By the time execution reaches this sub-skill, MCP should already be configured and ready.\n> If MCP tools are unavailable or calls fail at any point, stop and tell the user:\n> \"MCP 服务不可用，请检查配置或访问 https://kernelgen.flagos.io/mcp 获取 Token 后重新配置。\"\n\n**If the MCP server IS configured**, proceed to Step 1.\n\n---\n\n## Step 1: Understand the Operator Request\n\nParse the user's description to determine:\n- `kernel_name`: operator name in **snake_case** (e.g., `rms_norm`, `relu`, `softmax`)\n- `func_desc`: what the operator does\n- `func_type`: one of `unary_pointwise`, `binary_pointwise`, `reduction`, `normalization`,\n `attention`, `activation`, `quantization`, `moe`, `blas`, `sampling`, `other`.\n **Auto-infer from the operator signature when possible:**\n - If 1 tensor input, no dim arg → `unary_pointwise`\n - If operator has >=2 tensor inputs → `binary_pointwise`\n - If has a `dim` argument → `reduction`\n - If name contains `norm`, `softmax`, `rmsnorm`, `layernorm`, `batchnorm` → `normalization`\n - If name contains `relu`, `gelu`, `silu`, `swish`, `tanh`, `sigmoid` → `activation`\n - If name is `mm`, `matmul`, `bmm`, `addmm` → `blas`\n - If name contains `attention`, `flash`, `paged` → `attention`\n - If name contains `quant`, `awq`, `gptq` → `quantization`\n - If name contains `moe`, `expert` → `moe`\n - Otherwise → `other`\n Confirm the inferred `func_type` with the user if ambiguous.\n- `arg_names`, `arg_type`, `arg_descs`, `output_arg_desc`: parameter information\n\n**Normalize `kernel_name` to snake_case** before using it anywhere:\n- `RMSNorm` → `rms_norm`\n- `SiluAndMul` → `silu_and_mul`\n- `LayerNorm` → `layer_norm`\n\nCheck this common alias table before asking the user:\n\n| User says | Correct `kernel_name` |\n|---|---|\n| swish | `silu` |\n| negative / negate | `neg` |\n| power | `pow` |\n| absolute | `abs` |\n| multiply | `mul` |\n| divide | `div` |\n| clip | `clamp` |\n| hard_swish | `hardswish` |\n| hard_sigmoid | `hardsigmoid` |\n| logarithm / ln | `log` |\n| exponential | `exp` |\n| hyperbolic_tangent | `tanh` |\n| square | `pow (exp=2)` |\n| square_root | `sqrt` |\n| cube | `pow (exp=3)` |\n| cube_root | `cbrt` |\n| inverse_sqrt | `rsqrt` |\n| reciprocal / inverse | `reciprocal` |\n| minimum | `min` |\n| maximum | `max` |\n| floor_divide | `floor_divide` |\n| modulo / mod | `remainder` |\n\n### Complexity Gate\n\nBefore proceeding, check the `func_type`. If it is one of the following **complex** types,\nwarn the user and ask for explicit confirmation before continuing:\n\n- `attention`\n- `moe`\n- `quantization`\n- `sampling`\n\n```\nNote: Operators of type are relatively complex. KernelGen's auto-generation success rate is approximately 60%,\nand manual tuning may be required. Do you want to proceed with generation?\n```\n\nFor simpler types (`unary_pointwise`, `binary_pointwise`, `reduction`, `normalization`,\n`activation`, `blas`), proceed without warning — these have >90% success rate.\n\n### Argument Inference\n\n**If argument information is missing or the description is vague**:\n\n1. Use the Grep tool to search for the operator name **only inside directories discovered in\n Step 2** (source dirs, test dirs, kernel dirs) to infer argument signatures from existing\n code. **Never grep the repo root or unscoped paths** — this causes token explosion in\n large repos.\n2. Present the inferred signature to the user for confirmation before proceeding.\n3. **Fallback**: If no matches are found, use a default signature based on `func_type`:\n\n | func_type | Default args |\n |---|---|\n | `unary_pointwise`, `activation` | `arg_names: [\"x\"]`, `arg_type: [\"torch.Tensor\"]` |\n | `binary_pointwise` | `arg_names: [\"x\", \"y\"]`, `arg_type: [\"torch.Tensor\", \"torch.Tensor\"]` |\n | `normalization` | `arg_names: [\"x\", \"weight\"]`, `arg_type: [\"torch.Tensor\", \"torch.Tensor\"]` |\n | `reduction` | `arg_names: [\"x\"]`, `arg_type: [\"torch.Tensor\"]` |\n | `attention` | `arg_names: [\"q\", \"k\", \"v\"]`, `arg_type: [\"torch.Tensor\", \"torch.Tensor\", \"torch.Tensor\"]` |\n | `blas` | `arg_names: [\"a\", \"b\"]`, `arg_type: [\"torch.Tensor\", \"torch.Tensor\"]` |\n | other | `arg_names: [\"x\"]`, `arg_type: [\"torch.Tensor\"]` |\n\n Present the default to the user and ask them to confirm or refine. **Do not send defaults\n directly to MCP without user confirmation** — a wrong signature causes schema errors.\n\n---\n\n## Step 2: Discover Repository Structure\n\n**This step is critical.** Unlike the FlagGems/vLLM-specific skills, we do NOT assume any fixed\ndirectory layout. Instead, we dynamically discover the project structure.\n\n### 2a. Identify the Project Type\n\nUse the Glob and Read tools to check for project identity:\n\n```\nGlob: pyproject.toml, setup.py, setup.cfg\n```\n\nRead `pyproject.toml` (or `setup.py` / `setup.cfg`) to determine:\n- **Project name** (e.g., `flag_gems`, `vllm`, `lmdeploy`, etc.)\n- **Source root** (e.g., `src/`, `.`, `lib/`)\n- **Package name** — the importable top-level module\n\n**Known-project shortcut**: If the project name matches a known specialized skill, inform\nthe user and suggest switching:\n\n| Project name | Specialized skill | Suggestion |\n|---|---|---|\n| `flag_gems` / `FlagGems` | `/kernelgen_for_flaggems` | \"Detected FlagGems repository. It is recommended to use the specialized skill `/kernelgen_for_flaggems` for a higher success rate. Would you like to switch?\" |\n| `vllm` | `/kernelgen_for_vllm` | \"Detected vLLM repository. It is recommended to use the specialized skill `/kernelgen_for_vllm` for a higher success rate. Would you like to switch?\" |\n\nIf the user agrees, stop and tell them to invoke the specialized skill. If the user declines\nor the project is not in the table, continue with the general skill.\n\n### 2b. Map Key Directories\n\nUse the Glob tool with **shallow patterns** to avoid token explosion in large repos. Do NOT\nuse deep recursive `**` patterns that could match thousands of files.\n\n**Important**: If any single Glob returns more than 20 matches, **stop inspecting that\npattern** — just record the directory prefix (e.g., `src/pkg/ops/`) and move on. Do NOT\nread or list all matched files; only note the directory structure.\n\n```\nGlob: src/*/ops/*.py (source ops — one level under src/)\nGlob: src/*/kernels/*.py (source kernels)\nGlob: */ops/*.py (top-level package ops)\nGlob: */kernels/*.py (top-level package kernels)\nGlob: tests/*.py (top-level tests)\nGlob: tests/kernels/*.py (kernel-specific tests)\nGlob: benchmark*/*.py (benchmarks)\nGlob: benchmark*/kernels/*.py (kernel benchmarks)\n```\n\nIf these shallow patterns return no results for a category, try ONE level deeper:\n```\nGlob: src/*/*/ops/*.py\nGlob: src/*/*/kernels/*.py\n```\n\nIf **still** no operator directory is found, run a generic fallback scan:\n```\nGlob: src/*/*kernel*.py\nGlob: src/*/*op*.py\n```\nThis catches non-standard directory names like `triton_kernels/`, `gpu_ops/`, `cuda_ops/`.\n\n**Stop expanding once you find matches.** Do not keep globbing deeper.\n\n**Early-stop rule**: Once you have discovered both an **operator source directory** and a\n**test directory**, stop all further glob exploration. Do not continue searching for additional\ndirectories — two anchors are sufficient to infer the rest of the layout.\n\n**Disambiguation rule**: If multiple operator directories are discovered (e.g., `src/foo/ops/`\nand `src/bar/ops/`), **prefer the one that belongs to the main package name** detected in\nStep 2a (`pyproject.toml`). If still ambiguous, ask the user which directory to use.\n\nRecord the discovered layout in this table (fill in dynamically):\n\n| Purpose | Discovered Path |\n|---------|----------------|\n| **Operator source code** | _(e.g., `src/pkg/ops/`, `pkg/kernels/`, etc.)_ |\n| **Tests (accuracy)** | _(e.g., `tests/`, `tests/kernels/`, etc.)_ |\n| **Benchmarks** | _(e.g., `benchmark/`, `benchmarks/kernels/`, etc.)_ |\n| **Init/registry files** | _(e.g., `src/pkg/ops/__init__.py`, etc.)_ |\n| **Custom op registration** | _(e.g., `pkg/_custom_ops.py`, etc.)_ |\n\n### 2c. Study Existing Operators as Reference\n\nFind an existing operator implementation similar to the requested one and read it to learn:\n\n1. **Import style** — what modules are imported, decorator patterns\n2. **Kernel style** — raw Triton pointer-based, or higher-level wrappers like `pointwise_dynamic`\n3. **Function naming** — snake_case, prefix/suffix conventions (e.g., `_kernel`, `_forward`)\n4. **Wrapper pattern** — how the Python wrapper calls the Triton kernel\n5. **License headers** — whether files start with SPDX or other license headers\n6. **Logging** — which logging approach is used (e.g., `vllm.logger.init_logger`, stdlib `logging`)\n7. **Autotune usage** — whether the repo uses `@triton.autotune` and which BLOCK_SIZE configs.\n If autotune is used, also note `num_warps` and `num_stages` values from existing configs —\n these will be replicated for the new kernel's autotune configs in Step 6.\n\nAlso read the matching test and benchmark files to learn:\n- **Test pattern**: parametrize style, shapes, dtypes, assertion utilities\n- **Benchmark pattern**: timing methodology, reporting format\n\n**Store all discovered conventions** — they will be applied in Step 6.\n\n---\n\n## Step 3: Check Whether the Operator Already Exists\n\nBefore calling the MCP generator, **thoroughly search** the codebase for existing implementations\nusing the Glob and Grep tools.\n\n### 3a. Generate Search Variants\n\nFrom `kernel_name`, generate **all naming variants** to search:\n- snake_case: `layer_norm`\n- concatenated lowercase: `layernorm`\n- CamelCase: `LayerNorm`\n- Partial components (for compound names): only if `kernel_name` contains `_and_` or `_with_`,\n split and search components **longer than 4 characters**. Example:\n - `silu_and_mul` → search `silu` (5 chars, OK), skip `mul` (3 chars, too short)\n - `add_relu` → do NOT split (no `_and_` / `_with_`), search as whole name only\n Short fragments like `mul`, `add`, `div`, `neg` match too many unrelated symbols.\n\n### 3b. Search\n\nFor **each** naming variant, search:\n1. **Source files**: Glob for `**.py` in discovered source directories (shallow)\n2. **Function and class definitions**: Grep for these patterns in discovered source directories:\n - `def ` — direct function definition\n - `_kernel` — kernel function naming convention\n - `_forward` — forward function naming convention\n - `_backward` — backward function naming convention\n - `class ` — class-based operator (e.g., `class LayerNorm`,\n `class SoftmaxKernel`, `class RMSNormOp`)\n - `register_op(\"\")` / `OP_REGISTRY[\"\"]` / `aten_map[\"\"]` —\n dispatch-style registration where the implementation function name may differ from\n the operator name\n This catches operators that live inside files with different names (e.g., `softmax_kernel`\n inside `normalization.py`, or `silu_and_mul` inside `activation.py`).\n3. **Init/registry**: Grep for `` in all `__init__.py` files and any custom op\n registration files found in Step 2\n4. **Tests**: Grep for `` in discovered test directories\n5. **Benchmarks**: Grep for `` in discovered benchmark directories\n\n### If the operator already exists\n\nPresent findings to the user and ask them to choose. **Stop and wait** for the user's explicit\nchoice before continuing:\n\n> Detected that the `` operator already exists in the repository:\n> - Implementation: ``\n> - Tests: ``\n> - Benchmarks: ``\n>\n> Please choose how to proceed:\n>\n> **A — Cancel generation**: Keep the existing operator unchanged.\n>\n> **B — Replace existing**: Replace the current implementation with a new version generated by KernelGen.\n>\n> **C — Add custom variant side-by-side**: Keep the existing operator and place the new version as `_v2` alongside it, so both coexist.\n\nOnly proceed to Step 4 after the user has made a choice.\n\n### If the operator does NOT exist\n\nProceed directly to Step 4.\n\n---\n\n## Step 4: Research Context (flagos_wiki)\n\nBefore calling the MCP generator, use the Read tool to gather reference materials that improve\ngeneration quality:\n\n1. **Read similar operator code** from the codebase — pick the most similar existing operator\n to the one being generated (same `func_type`).\n\n2. **Read the test pattern** from the matching test file to understand shapes, dtypes, assertions.\n\n3. **Read the benchmark pattern** from the matching benchmark file to understand timing methodology.\n\n4. **If replacing an existing op** (Option B), read the current implementation so the new version\n can be compared / improved upon.\n\n **If reference files cannot be found**, do not stop. Proceed with general algorithm hints\n based on your knowledge of the operator type. An empty or minimal `flagos_wiki` is acceptable.\n\n5. Collect all findings into the `flagos_wiki` parameter as a `List[str]`. Rules:\n - Each string should be a concise reference snippet, **under 200 characters**.\n - **Maximum 10 items** — avoid token explosion.\n - **Prefer algorithm insights over file path descriptions**.\n - **Always include a kernel pattern hint** as the first item (e.g.,\n `\"Kernel pattern: elementwise\"`, `\"Kernel pattern: row-wise reduction over hidden_size\"`).\n - **Always include `\"Use Triton for GPU parallelization\"` as the second item** — this\n anchors the generator toward Triton output and reduces fallback failures.\n - **Always include a memory layout hint** (e.g.,\n `\"Input tensors are contiguous row-major layout\"`).\n - **Include a memory access pattern hint if inferable** from the `func_type` (e.g.,\n `\"Memory access: coalesced contiguous loads\"` for pointwise/elementwise,\n `\"Memory access: row-wise strided reads over hidden_size\"` for normalization/reduction).\n This significantly improves generated kernel quality.\n - Include relevant conventions discovered in Step 2 (e.g.,\n `\"Repo uses pointwise_dynamic wrapper for elementwise ops\"`,\n `\"Repo uses @triton.autotune with BLOCK_SIZE configs [64,128,256,512]\"`).\n\n---\n\n## Step 5: Call kernelgen-mcp\n\nInvoke `mcp__kernelgen-mcp__generate_kernel` with the parameters gathered above:\n\n```\nkernel_name: \"\"\nfunc_desc: \"\"\nfunc_type: \"\"\narg_names: []\narg_type: []\narg_descs: []\noutput_arg_desc: \"\"\nflagos_wiki: []\n```\n\nThe MCP returns code blocks which may include:\n- `torch_code` — PyTorch reference implementation\n- `triton_code` — Triton kernel implementation\n- `test_func_code` — accuracy test code\n- `benchmark_func_code` — performance benchmark code\n\n**If `torch_code` is missing**, construct a `ref_impl` yourself using standard PyTorch ops\nbased on `func_desc`.\n\n### MCP Failure Retry Strategy\n\nIf the MCP call fails, apply retries in this order:\n\n1. **First retry**: Remove `flagos_wiki` parameter and call again (handles `unexpected argument`\n / `unknown field` / `invalid schema` errors).\n2. **Second retry**: Use minimal parameters only (`kernel_name`, `func_desc`, `func_type`,\n `arg_names`, `arg_type`) — drop `arg_descs`, `output_arg_desc`, and `flagos_wiki`.\n3. **Third retry**: Change `func_type` to `\"other\"` and keep only minimal parameters. This\n handles `func_type` mismatch errors where the MCP backend does not recognize the original\n `func_type` value (e.g., `activation` vs `unary_pointwise`).\n4. **If still failing**: Stop and report the exact error to the user. Do NOT keep retrying.\n\n---\n\n## Step 5.5: MCP Report Review\n\nIf the MCP response includes test results (accuracy test reports, benchmark/speedup reports),\npresent them to the user **in full** before proceeding to local code adaptation:\n\n```\n=== MCP Generation Test Report ===\n\nAccuracy Test Results:\n\n\nPerformance Benchmark Results:\n\n```\n\nThen ask the user:\n\n> The MCP service has returned the above test reports from its generation environment.\n> Would you like to:\n>\n> **A — Skip local testing**: Trust the MCP reports and proceed directly to code placement\n> and summary. No local tests will be run.\n>\n> **B — Run local tests**: Run accuracy and performance tests on your local machine as well\n> to verify the results in your specific environment.\n\nWait for the user's choice before proceeding:\n- If **A (skip local testing)**: After placing code in Step 6, skip Steps 7 and 8 (local\n accuracy tests and benchmarks) and proceed directly to Step 9 (Summary). Use the MCP\n report numbers for the summary.\n- If **B (run local tests)**: Proceed normally through Steps 6 → 7 → 8 → 9. Before\n running local tests, perform the **Chip Compatibility Check** described in Step 6.5d.\n\n**If the MCP response does NOT include test reports** (only code blocks), proceed normally\nto Step 6 — local testing will be required.\n\n---\n\n## Step 6: Adapt and Place Code into the Repository\n\nThis is the most critical step. The generated code must be **transformed to match the local\nrepository's conventions** discovered in Step 2.\n\n### 6a. Transform the Kernel Code\n\nApply the conventions discovered in Step 2c. **The decision to rewrite depends entirely on\nwhat the repo actually uses** — never assume:\n\n**Only rewrite to wrapper style if BOTH conditions are met**:\n1. The repo explicitly uses wrapper utilities (e.g., FlagGems `pointwise_dynamic`).\n2. The reference operator's kernel function signature receives **scalar elements** (not pointers).\n\nIf either condition is unclear, **keep the MCP-generated raw Triton pointer-based code**.\nMisidentifying wrapper style produces completely broken kernels.\n\nWhen rewriting to wrapper style:\n- The `@triton.jit` kernel should receive **scalar elements** (not pointers).\n- **Remove ALL** `tl.load()`, `tl.store()`, pointer arithmetic, mask logic, and BLOCK_SIZE\n parameters — the wrapper handles these automatically.\n- Only keep the pure math expression.\n\n**Otherwise, keep the MCP-generated raw Triton pointer-based code.** This is the safer default\nfor most repositories. Apply these adaptations:\n- **Match the repo's autotune convention exactly**:\n - If existing kernels use `@triton.autotune` → add it, using the same BLOCK_SIZE configs\n found in existing kernels.\n - If existing kernels do **NOT** use `@triton.autotune` (e.g., torchinductor, tinygrad, some\n inference engines explicitly avoid it) → do **NOT** introduce it. Use a fixed BLOCK_SIZE\n matching the repo's pattern instead.\n - If no existing kernels are found for reference, use a default autotune set:\n ```python\n @triton.autotune(\n configs=[\n triton.Config({\"BLOCK_SIZE\": 64}, num_warps=4),\n triton.Config({\"BLOCK_SIZE\": 128}, num_warps=4),\n triton.Config({\"BLOCK_SIZE\": 256}, num_warps=8),\n triton.Config({\"BLOCK_SIZE\": 512}, num_warps=8),\n ],\n key=[\"n_elements\"], # replace with actual problem-size arg name\n )\n ```\n- Ensure proper `tl.constexpr` for meta parameters.\n- Use `triton.cdiv` for grid computation.\n\n**Common adaptations for all repos:**\n- Add the repo's standard license header (if any, discovered in Step 2c).\n- Use the repo's logging approach (e.g., `vllm.logger.init_logger`, `logging.getLogger`).\n- Follow the repo's naming conventions for kernel functions (e.g., `_kernel` suffix, `_forward`\n suffix).\n- Match import style exactly.\n\n### 6b. Determine File Placement\n\nPlace files according to the layout\n\nArchive v1.0.0: 9 files, 68655 bytes\n\nFiles: kernelgen-for-flaggems.md (50321b), kernelgen-for-vllm.md (49198b), kernelgen-general.md (51634b), kernelgen-submit-feedback.md (11593b), LICENSE.txt (11262b), README_zh.md (4192b), README.md (5495b), SKILL.md (5442b), _meta.json (135b)","readmeExcerpt":"Skill: Kernelgen Flagos Owner: wbavon Summary: Unified GPU kernel operator generation and optimization skill. Automatically detects the target repository type (FlagGems, vLLM, or general Python/Triton) an... Tags: latest:1.1.0 Version history: v1.1.0 | 2026-05-19T07:04:17.149Z | user Sync from flagos-ai/skills latest v1.0.0 | 2026-03-23T04:33:20.901Z | auto - Initial release of kernelgen-flagos as a unified GPU kerne","codeSnippets":[],"executableExamples":[{"language":"text","snippet":"Glob: **/skills/kernelgen-flagos/kernelgen-mcp-setup.md"},{"language":"text","snippet":"Glob: pyproject.toml\nGlob: setup.py\nGlob: setup.cfg"},{"language":"text","snippet":"Glob: **/skills/kernelgen-flagos/kernelgen-generate.md"},{"language":"bash","snippet":"# === Generation ===\n# Generate a kernel operator (auto-detects repo type)\n/kernelgen-flagos relu\n\n# Generate with explicit function type\n/kernelgen-flagos rms_norm --func-type normalization\n\n# === Optimization ===\n# Optimize an existing Triton kernel (auto-detects repo type)\n# Just say \"optimize the relu kernel\" or \"improve kernel performance\"\n# The skill will automatically dispatch to the right optimization sub-skill\n\n# The skill will automatically:\n# - Detect if you're in a FlagGems repo → use FlagGems-specific workflow\n# - Detect if you're in a vLLM repo → use vLLM-specific workflow\n# - Otherwise → use the general-purpose workflow"},{"language":"text","snippet":"or"},{"language":"text","snippet":"**Note**: You may need to restart your agent to let the settings take effect. Check this in the documentation of the relevant AI agent.\n\n3. Verify KernelGen Operator Development MCP Toolkit connection, send a prompt:"}],"parameters":null,"dependencies":[],"permissions":[],"extractedFiles":[{"path":"SKILL.md","content":"---\nname: kernelgen-flagos\ndescription: >\n Unified GPU kernel operator generation and optimization skill. Automatically detects the target\n repository type (FlagGems, vLLM, or general Python/Triton) and dispatches to the appropriate\n specialized sub-skill. Includes operator generation, MCP-based iterative optimization, and\n feedback submission sub-skills. Use this skill when the user wants to generate or optimize a\n GPU kernel operator, create a Triton kernel, or says things like \"generate an operator\",\n \"create a kernel for X\", \"optimize triton kernel\", or \"/kernelgen-flagos\".\nargument-hint: \" [--func-type ]\"\nuser-invokable: true\ncompatibility: \"Python 3.8+, PyTorch with CUDA, Triton\"\nmetadata:\n version: \"1.0.0\"\n author: flagos-ai\n category: gpu-kernel-generation\n tags: [kernelgen, triton, gpu, mcp, operator-generation, operator-optimization, flaggems, vllm, feedback]\nallowed-tools:\n - Bash\n - Bash(gh:*)\n - Bash(python:*)\n - Bash(python3:*)\n - Bash(command:*)\n - Read\n - Write\n - Edit\n - Glob\n - Grep\n - AskUserQuestion\n---\n\n# kernelgen-flagos — Unified GPU Operator Generation Skill\n\nThis is a **unified entry point** that bundles generation and optimization sub-skills into one:\n\n| Sub-skill file | Purpose |\n|---|---|\n| **Generation** | |\n| `kernelgen-generate.md` | Generate GPU kernels for **any** Python/Triton repository |\n| `kernelgen-generate-for-flaggems.md` | Specialized generation for **FlagGems** repositories |\n| `kernelgen-generate-for-vllm.md` | Specialized generation for **vLLM** repositories |\n| **Optimization** | |\n| `kernelgen-optimize.md` | Optimize existing Triton kernels via MCP iterative optimization (general purpose) |\n| `kernelgen-optimize-for-flaggems.md` | Optimize Triton operators and integrate into **FlagGems** (3 modes: built-in/external/experimental) |\n| `kernelgen-optimize-for-vllm.md` | Optimize Triton operators and integrate into **vLLM** (with CustomOp registration) |\n| **Platform Specialization** | |\n| `kernelgen-specialize.md` | Specialize Triton operators to target platforms (e.g., GPU → Ascend NPU) via MCP `specialize_kernel` |\n| `kernelgen-specialize-for-flaggems.md` | Platform specialization + **FlagGems** integration (4 modes: vendor-ops/vendor-fused/override-builtin/experimental) |\n| **MCP Configuration** | |\n| `kernelgen-mcp-setup.md` | Check and auto-configure the `kernelgen-server` MCP service (URL built-in, user only provides Token) |\n| **Feedback** | |\n| `kernelgen-submit-feedback.md` | Submit bug reports and feedback via GitHub or email |\n\nAll sub-skill files are located in the **same directory** as this `SKILL.md` file.\n\n---\n\n## Routing Protocol — Follow This BEFORE Doing Anything Else\n\n### Phase 0: MCP Configuration Check\n\nBefore anything else, ensure the `kernelgen-server` MCP service is configured and ready.\n\nUse the Glob tool to find `kernelgen-mcp-setup.md` in this skill's directory:\n\n```\nGlob: **/skills/kernelgen-flagos/kernelgen-mcp-setup.md\n```\n\nThen use "},{"path":"README.md","content":"# kernelgen-flagos: Unified GPU Kernel Generation Skills\n\n[中文版](README_zh.md)\n\n## Overview\n\nThe KernelGen `kernelgen-flagos` skill is a unified AI programming skill launched by FlagOS. By leveraging the `kernelgen-mcp` KernelGen Operator Development MCP Toolkit service, it is compatible with various domestic and international chip platforms to generate high-precision GPU kernel operators. It automatically detects the type of the target code repository and dispatches it to the corresponding dedicated workflow.\n\n## Quick Start\n\nThis section introduces how to quickly start a task, for example, creating, optimizing, specialize Kernels for a specific platform, or submitting feedback. You can use either prompt or the command line to execute these tasks. In this section, we only teach you how to use prompts. For how to use command lines or manual configurations, see [KernelGen Skills User Guide](https://docs.flagos.io/projects/kernelgen/en/latest/skills_user_guide/skills-user-guide.html#).\n\nPerform all of the following steps except for obtaining the KernelGen Token in your agent client, for example, OpenClaw, VSCode with activated Github Copilot, or Claude Code.\n\n### Prerequisites\n\n* Claude Code version 2.1 and later\n\n* OpenClaw version 2026.3.2 and later\n\n* VSCode with Github Copilot activated\n\n### Steps\n\n1. Obtain the **KernelGen Token** from: [https://kernelgen.flagos.io/mcp](https://kernelgen.flagos.io/mcp).\n\n2. In your agent client, send a prompt to connect to the KernelGen Operator Development MCP Toolkit, for example:\n\n ```{code-block} shell\n Based on the Claude Code configuration documentation: https://code.claude.com/docs/en/mcp, connect to the MCP. The MCP URL is https://kernelgen.flagos.io/sse, and the token is . Configure this in the claude.json file.\n ```\n\n or\n\n ```{code-block} shell\n Based on the VSCode documentation: https://code.visualstudio.com/docs/copilot/customization/mcp-servers, configure the kernelgen MCP. The MCP URL is https://kernelgen.flagos.io/sse, and the token is . Configure this in the mcp.json file.\n ```\n\n **Note**: You may need to restart your agent to let the settings take effect. Check this in the documentation of the relevant AI agent.\n\n3. Verify KernelGen Operator Development MCP Toolkit connection, send a prompt:\n\n ```{code-block} shell\n Please verify the kernelgen mcp connection is successful.\n ```\n\n4. Send a prompt to setup the `kernelgen-flagos` unified skill, including all sub-skills, for example:\n\n * **Copilot**:\n \n ```{code-block} shell\n Setup kernelgen-flagos skills in Copilot from .\n ```\n\n * **Other agent clients**:\n\n ```{code-block} shell\n Setup kernelgen-flagos skills from .\n ```\n\n5. Verify skills are successfully installed，send a prompt\n\n ```{code-block} shell\n Please verify if th"},{"path":"_meta.json","content":"{\n \"ownerId\": \"kn7c1yzknmhkmxasrtdbxa3fhd81z80k\",\n \"slug\": \"kernelgen-flagos\",\n \"version\": \"1.1.0\",\n \"publishedAt\": 1779174257149\n}"},{"path":"kernelgen-generate-for-flaggems.md","content":"# KernelGen Skill — Generate FlagGems Operators via MCP\n\nYou are an expert at generating GPU kernel operators for the FlagGems project using the `kernelgen-mcp` MCP service.\n\n**Tool usage**: This skill relies on the following capabilities:\n- **Shell**: execute shell commands (python, pip, pytest, etc.) via the Bash/shell tool\n- **Read**: read files from the codebase\n- **Write / Edit**: create or modify files\n- **Grep**: search file contents\n- **Glob**: find files by pattern\n- **MCP tools**: `mcp__kernelgen-mcp__generate_kernel` and `mcp__kernelgen-mcp__optimize_kernel`\n\n> **⚠️ MCP Prerequisite**: MCP configuration is handled by `SKILL.md` Phase 0. If MCP tools fail during execution,\n> stop and tell the user to visit https://kernelgen.flagos.io/mcp 获取 Token 后重新配置。\n\nWhen you need user input or clarification, ask a question directly and wait for their reply.\nAlways use the appropriate built-in tool rather than outputting commands for the user to run manually.\n\n## Execution Rules\n\n1. **Follow the steps in order**. Never jump ahead.\n2. **Never skip Step 2** (operator existence check) — always verify before generating.\n3. **Do not generate or write any code before Step 4** (MCP call). Steps 0–3 are preparation only.\n4. **Always confirm with the user before destructive actions**.\n5. **Report progress to the user when entering each step**.\n6. **Never fabricate repository files or paths**.\n7. **CRITICAL — MCP is mandatory**: ALL operator code generation MUST go through the\n `mcp__kernelgen-mcp__generate_kernel` MCP tool. NEVER generate Triton kernels, PyTorch\n wrappers, or operator implementations yourself — even if the operator seems simple (e.g.,\n relu, abs). If MCP is not configured, not reachable, or fails after all retries in Step 4,\n STOP and report the issue to the user. Do NOT fall back to writing kernel code manually.\n The MCP service produces optimized, tested code that manual writing cannot match.\n\n---\n\n## FlagGems Project Layout Reference\n\nThe FlagGems project has three operator locations and corresponding test/benchmark locations:\n\n### Operator Locations\n| Location | Path | Description |\n|---|---|---|\n| **Core (src)** | `src/flag_gems/ops/.py` | Mature, production-quality operators. Uses `pointwise_dynamic` decorators for pointwise ops. Registered in `_FULL_CONFIG` for aten dispatch. |\n| **Experimental** | `src/flag_gems/experimental_ops/.py` | New/experimental operators. Uses raw Triton pointer-based style (self-contained, no `pointwise_dynamic`). NOT registered in `_FULL_CONFIG`. |\n| **Backend** | `src/flag_gems/runtime/backend/_/ops/.py` | Non-NVIDIA chip-specific operators. Vendor names: `_ascend` (Huawei), `_cambricon`, `_hygon` (DCU), `_mthreads` (Moore), `_iluvatar` (Tianshu), `_metax`, `_amd`, `_kunlunxin`, `_arm`, `_aipu`. |\n\n### Test Locations\n| Operator Location | Accuracy Tests | Performance Tests |\n|---|---|---|\n| **Core (src)** | `tests/test__ops.py` (append to e"},{"path":"kernelgen-generate-for-vllm.md","content":"# KernelGen Skill — Generate vLLM Operators via MCP\n\nYou are an expert at generating GPU kernel operators for the vLLM project using the `kernelgen-mcp` MCP service.\n\n> **⚠️ MCP Prerequisite**: MCP configuration is handled by `SKILL.md` Phase 0. If MCP tools fail during execution,\n> stop and tell the user to visit https://kernelgen.flagos.io/mcp 获取 Token 后重新配置。\n\n## Execution Rules\n\nFollow these rules strictly to avoid getting lost in the multi-step workflow:\n\n1. **Follow the steps in order**: Step 0 → Step 1 → Step 2 → ... → Step 8. Never jump ahead.\n2. **Never skip Step 2** (operator existence check) — always verify before generating.\n3. **Do not generate or write any code before Step 4** (MCP call). Steps 0–3 are preparation only.\n4. **Always confirm with the user before destructive actions**: replacing files, installing packages,\n overwriting operators.\n5. **If any step fails, stop and report the error** to the user. Do not silently skip or retry\n with different parameters.\n6. **Report progress to the user when entering each step**. Print a short status line like\n `Entering Step 2 — checking whether the operator already exists...` so the user knows\n the workflow is progressing and not stuck.\n7. **Never fabricate repository files or paths**. If a required file cannot be found using the\n Glob or Grep tools, report it to the user instead of guessing. Do not assume a file exists\n at a path without verifying it first.\n8. **CRITICAL — MCP is mandatory**: ALL operator code generation MUST go through the\n `mcp__kernelgen-mcp__generate_kernel` MCP tool. NEVER generate Triton kernels, PyTorch\n wrappers, or operator implementations yourself — even if the operator seems simple (e.g.,\n relu, abs). If MCP is not configured, not reachable, or fails after all retries in Step 4,\n STOP and report the issue to the user. Do NOT fall back to writing kernel code manually.\n The MCP service produces optimized, tested code that manual writing cannot match.\n\n## Step 0: Pre-flight — Environment & MCP Check\n\n### 0a. Check Python Environment\n\nUse the Bash tool to run a diagnostic that checks each package independently (so one failure\ndoes not prevent checking the rest):\n\n```bash\npython - <<'PY'\nimport importlib\n\ndef check(pkg):\n try:\n m = importlib.import_module(pkg)\n print(pkg, getattr(m, \"__version__\", \"installed\"))\n return True\n except Exception:\n print(pkg, \"NOT installed\")\n return False\n\ncheck(\"torch\")\ntriton_ok = check(\"triton\")\n\ntry:\n import torch\n print(\"cuda available\", torch.cuda.is_available())\n if torch.cuda.is_available():\n print(\"torch.cuda version:\", torch.version.cuda)\n n = torch.cuda.device_count()\n print(\"cuda devices:\", n)\n for i in range(n):\n print(f\" cuda:{i}\", torch.cuda.get_device_name(i))\nexcept Exception:\n pass\n\nif triton_ok:\n try:\n import triton.runtime\n print(\"triton runtime OK\")\n except Exception as e:\n p"}],"languages":[],"docsSourceLabel":"CLAWHUB","editorialOverview":null,"editorialQuality":{"score":100,"threshold":65,"status":"thin","wordCount":1192,"uniquenessScore":43,"reasons":["uniqueness-below-45"]}},"media":{"evidence":{"source":"no-media","verified":false,"confidence":"low","updatedAt":"2026-10-11T23:08:28.846Z","emptyReason":"No screenshots, media assets, or demo links are available."},"primaryImageUrl":null,"mediaAssetCount":0,"assets":[],"demoUrl":null},"ownerResources":{"evidence":{"source":"unclaimed","verified":false,"confidence":"low","updatedAt":"2026-10-11T23:08:28.846Z","emptyReason":"This page has not been claimed by the agent owner."},"hasCustomPage":false,"customPageUpdatedAt":null,"customLinks":[],"structuredLinks":{"docsUrl":null,"demoUrl":null,"supportUrl":null,"pricingUrl":null,"statusUrl":null},"customPage":null},"relatedAgents":{"evidence":{"source":"protocol-neighbors","verified":false,"confidence":"medium","updatedAt":"2026-10-15T03:21:35.702Z","emptyReason":null},"items":[{"id":"8ebccd8e-3863-4187-8355-c3f14e1f9edf","entityType":"agent","canonicalPath":"/agent/iofficeai-aionui","slug":"iofficeai-aionui","name":"AionUi","description":"Free, local, open-source 24/7 Cowork app and OpenClaw for Gemini CLI, Claude Code, Codex, OpenCode, Qwen Code, Goose CLI, Auggie, and more | 🌟 Star if you like it!","url":"https://github.com/iOfficeAI/AionUi","homepage":"https://www.aionui.com","source":"GITHUB_REPOS","protocols":["MCP","OPENCLAW"],"capabilities":[],"safetyScore":100,"overallRank":70,"updatedAt":"2026-10-09T19:11:12.944Z","createdAt":"2026-02-25T03:38:16.584Z","downloads":null},{"id":"b917f68a-ebff-438e-84f8-3f4b2494c0bc","entityType":"agent","canonicalPath":"/agent/activepieces-activepieces","slug":"activepieces-activepieces","name":"activepieces","description":"AI Agents & MCPs & AI Workflow Automation • (~400 MCP servers for AI agents) • AI Automation / AI Agent with MCPs • AI Workflows & AI Agents • MCPs for AI Agents","url":"https://github.com/activepieces/activepieces","homepage":"https://www.activepieces.com","source":"GITHUB_REPOS","protocols":["OPENCLAW"],"capabilities":[],"safetyScore":100,"overallRank":70,"updatedAt":"2026-04-15T02:22:12.426Z","createdAt":"2026-02-25T03:38:12.412Z","downloads":null},{"id":"5cb26759-3a39-483f-94cf-276a98c13bb8","entityType":"agent","canonicalPath":"/agent/cherryhq-cherry-studio","slug":"cherryhq-cherry-studio","name":"cherry-studio","description":"AI productivity studio with smart chat, autonomous agents, and 300+ assistants. Unified access to frontier LLMs","url":"https://github.com/CherryHQ/cherry-studio","homepage":"https://cherry-ai.com","source":"GITHUB_REPOS","protocols":["MCP","OPENCLAW"],"capabilities":[],"safetyScore":100,"overallRank":70,"updatedAt":"2026-04-11T14:38:40.986Z","createdAt":"2026-02-25T03:38:19.379Z","downloads":null},{"id":"6f6582d0-5d76-4f0f-b81d-86520247950b","entityType":"agent","canonicalPath":"/agent/copilotkit-copilotkit","slug":"copilotkit-copilotkit","name":"CopilotKit","description":"The Frontend for Agents & Generative UI. React + Angular","url":"https://github.com/CopilotKit/CopilotKit","homepage":"https://docs.copilotkit.ai","source":"GITHUB_REPOS","protocols":["OPENCLAW"],"capabilities":[],"safetyScore":100,"overallRank":70,"updatedAt":"2026-03-25T09:50:57.846Z","createdAt":"2026-02-25T03:39:14.617Z","downloads":null}],"links":{"hub":"/agent","source":"/agent/source/clawhub","protocols":[{"label":"OpenClaw","href":"/agent/protocol/openclew"}]}}}