Improve the Claude engine for SYCL generation

src/xe_forge/claude/generator.py

@@ -1,8 +1,9 @@
 """Generate a Claude Code workspace for kernel optimization.
 
-Creates CLAUDE.md, config.yaml, .claude/commands/, .claude/agents/,
-and copies kernel files into the workspace. All text artifacts are
-rendered from Jinja templates under ``templates/``.
+Creates CLAUDE.md, config.yaml, .claude/settings.local.json,
+.claude/commands/, .claude/agents/, and copies kernel files into
+the workspace. All text artifacts are rendered from Jinja templates
+under ``templates/``.
 """
 
 from __future__ import annotations
@@ -70,10 +71,11 @@ def generate_workspace(
 
     agent_dir = workspace / ".claude" / "agents"
     agent_dir.mkdir(parents=True, exist_ok=True)
-    (agent_dir / "tool-runner.md").write_text(_render("tool-runner.md.j2"))
+    (agent_dir / "tool-runner.md").write_text(_render("tool-runner.md.j2", dsl=dsl))
+    (workspace / ".claude" / "settings.local.json").write_text(_render("settings.local.json.j2"))
 
     _write_kernel_files(workspace, kernel_name, kernel_code, reference_code, spec_path)
-    _symlink_knowledge_base(workspace)
+    _copy_knowledge_base(workspace, dsl, device)
 
     if config.engine.git_init:
         _git_init(workspace)
@@ -96,11 +98,10 @@ def _write_kernel_files(
         shutil.copy2(spec_path, tk_dir / f"{kernel_name}.yaml")
 
 
-def _symlink_knowledge_base(workspace: Path) -> None:
-    """Create a symlink to the installed knowledge_base directory."""
-    kb_link = workspace / "knowledge_base"
-    if kb_link.exists() or kb_link.is_symlink():
-        return
+def _copy_knowledge_base(workspace: Path, dsl: str, device: str) -> None:
+    """Copy common and DSL/device-specific knowledge base subdirectories."""
+    kb_root = workspace / "knowledge_base"
+    kb_root.mkdir(parents=True, exist_ok=True)
 
     import xe_forge
 
@@ -112,7 +113,23 @@ def _symlink_knowledge_base(workspace: Path) -> None:
     ]
     for candidate in candidates:
         if candidate.is_dir():
-            kb_link.symlink_to(candidate.resolve())
+            src_root = candidate.resolve()
+
+            common_src = src_root / "common"
+            dsl_device_src = src_root / dsl / device
+            common_dst = kb_root / "common"
+            dsl_device_dst = kb_root / dsl / device
+
+            if common_dst.is_symlink():
+                common_dst.unlink()
+            if common_src.is_dir():
+                shutil.copytree(common_src, common_dst, dirs_exist_ok=True)
+
+            dsl_device_dst.parent.mkdir(parents=True, exist_ok=True)
+            if dsl_device_dst.is_symlink():
+                dsl_device_dst.unlink()
+            if dsl_device_src.is_dir():
+                shutil.copytree(dsl_device_src, dsl_device_dst, dirs_exist_ok=True)
             return
 
 

src/xe_forge/claude/templates/CLAUDE.md.j2

@@ -33,12 +33,12 @@ All runtime behavior below is gated by these values; re-read `config.yaml` if an
 | **Benchmark** | `xe-forge-skill benchmark <baseline> <optimized> --spec <spec.yaml> [--baseline-us <cached>]` |
 | **Init trials** | `xe-forge-skill trial init <kernel_name> <baseline_file>` |
 | **Save trial** | `xe-forge-skill trial save <kernel_name> <file> [--parent <id>] [--strategy "..."]` |
-| **Record result** | `xe-forge-skill trial result <kernel_name> <trial_id> --correctness <pass\|fail> --speedup <float> --baseline-us <float> --triton-us <float>` |
+| **Record result** | `xe-forge-skill trial result <kernel_name> <trial_id> --correctness <pass\|fail> --speedup <float> --baseline-us <float> --custom-us <float>` |
 | **Check status** | `xe-forge-skill trial status <kernel_name>` |
 | **Best trial** | `xe-forge-skill trial best <kernel_name>` |
 | **Baseline time** | `xe-forge-skill trial baseline-us <kernel_name>` |
 | **Finalize** | `xe-forge-skill trial finalize <kernel_name> <output_file>` |
-| **Profile** (only when `vtune_enabled: true`) | `xe-forge-skill profile <triton_file> --spec <spec.yaml>` |
+| **Profile** (only when `vtune_enabled: true`) | `xe-forge-skill profile <{{ dsl }}_file> --spec <spec.yaml>` |
 
 ## WORKFLOW — Follow these steps in order
 
@@ -58,10 +58,10 @@ For each trial:
 2. **Validate** — `xe-forge-skill validate <file> --dsl {{ dsl }}` (fix until passing).
 3. **Save** — `xe-forge-skill trial save {{ kernel_name }} <file> --parent <id> --strategy "description"`.
 4. **Benchmark** (MANDATORY every trial):
-   - **Trial t0:** `xe-forge-skill benchmark test_kernels/{{ kernel_name }}.py <triton_file> --spec test_kernels/{{ kernel_name }}.yaml`
-   - **Trials t1+:** Get cached baseline via `xe-forge-skill trial baseline-us {{ kernel_name }}`, then `xe-forge-skill benchmark test_kernels/{{ kernel_name }}.py <triton_file> --spec test_kernels/{{ kernel_name }}.yaml --baseline-us <cached>`
-5. **Record** — `xe-forge-skill trial result {{ kernel_name }} <trial_id> --correctness <pass|fail> --speedup <float> --baseline-us <float> --triton-us <float>`
-6. **Profile** — if `vtune_enabled: true` in `config.yaml` and this is trial t1 or later, run `xe-forge-skill profile <triton_file> --spec <spec.yaml>` and use the output to guide the next trial. If `vtune_enabled: false`, skip this step.
+   - **Trial t0:** `xe-forge-skill benchmark test_kernels/{{ kernel_name }}.py <{{ dsl }}_file> --spec test_kernels/{{ kernel_name }}.yaml`
+   - **Trials t1+:** Get cached baseline via `xe-forge-skill trial baseline-us {{ kernel_name }}`, then `xe-forge-skill benchmark test_kernels/{{ kernel_name }}.py <{{ dsl }}_file> --spec test_kernels/{{ kernel_name }}.yaml --baseline-us <cached>`
+5. **Record** — `xe-forge-skill trial result {{ kernel_name }} <trial_id> --correctness <pass|fail> --speedup <float> --baseline-us <float> --{{ dsl }}-us <float>`
+6. **Profile** — if `vtune_enabled: true` in `config.yaml` and this is trial t1 or later, run `xe-forge-skill profile <{{ dsl }}_file> --spec <spec.yaml>` and use the output to guide the next trial. If `vtune_enabled: false`, skip this step.
 7. **Decide next action**:
    - Speedup > 5x -> stop, finalize
    - Speedup improved -> continue on this branch
@@ -76,7 +76,9 @@ xe-forge-skill trial finalize {{ kernel_name }} output/{{ kernel_name }}_optimiz
 
 ## CRITICAL CORRECTNESS CONSTRAINTS
 
+{% if dsl == "triton" %}
 - NO default values for `@triton.autotune` meta-parameters in kernel signature
+{% endif %}
 - Use 1D grid when applying tile swizzling (GROUP_SIZE_M)
 - `boundary_check` uses dimension indices (0, 1), not booleans
 - Cast batch indices to `int64` before stride multiplication

src/xe_forge/claude/templates/settings.local.json.j2

@@ -0,0 +1,12 @@
+{
+  "permissions": {
+    "defaultMode": "acceptEdits",
+    "allow": [
+      "Bash(xe-forge-skill analyze *)",
+      "Bash(xe-forge-skill trial *)",
+      "Bash(xe-forge-skill validate *)",
+      "Bash(xe-forge-skill benchmark *)",
+      "Bash(xe-forge-skill profile *)"
+    ]
+  }
+}

src/xe_forge/claude/templates/tool-runner.md.j2

@@ -31,7 +31,7 @@ Safe to parallelize: analyze, validate, trial (CPU-only).
 ## Output Rules
 
 For benchmark:
-- Extract ONLY: Correctness (PASSED/FAILED), Performance (baseline_us, triton_us, speedup), Errors.
+- Extract ONLY: Correctness (PASSED/FAILED), Performance (baseline_us, custom_us, speedup), Errors.
 - Do NOT include configuration header or decorative separators.
 
 For profile:

src/xe_forge/core/trial_manager.py

@@ -40,6 +40,8 @@ def _load_state(self, kernel_name: str) -> dict:
         for trial in state.get("trials", {}).values():
             if "pytorch_us" in trial and "baseline_us" not in trial:
                 trial["baseline_us"] = trial.pop("pytorch_us")
+            if "triton_us" in trial and "custom_us" not in trial:
+                trial["custom_us"] = trial.pop("triton_us")
         return state
 
     def _save_state(self, kernel_name: str, state: dict) -> None:
@@ -118,7 +120,7 @@ def save_trial(
             "correctness": None,
             "speedup": None,
             "baseline_us": None,
-            "triton_us": None,
+            "custom_us": None,
             "status": "saved",
         }
         self._save_state(kernel_name, state)
@@ -134,6 +136,7 @@ def record_result(
         correctness: str | None = None,
         speedup: float | None = None,
         baseline_us: float | None = None,
+        custom_us: float | None = None,
         triton_us: float | None = None,
     ) -> dict:
         """Record benchmark results for a trial. Returns the trial dict."""
@@ -153,8 +156,10 @@ def record_result(
             trial["speedup"] = speedup
         if baseline_us is not None:
             trial["baseline_us"] = baseline_us
+        if custom_us is not None:
+            trial["custom_us"] = custom_us
         if triton_us is not None:
-            trial["triton_us"] = triton_us
+            trial["custom_us"] = triton_us
 
         if baseline_us is not None and state.get("baseline_us") is None:
             state["baseline_us"] = [baseline_us]
@@ -222,8 +227,8 @@ def _render(tid: str, prefix: str = "", is_last: bool = True) -> None:
             icon = status_icon.get(trial["status"], "?")
             speedup_str = f"{trial['speedup']:.2f}x" if trial["speedup"] is not None else "---"
             runtime = ""
-            if trial.get("baseline_us") is not None and trial.get("triton_us") is not None:
-                runtime = f" (bl={trial['baseline_us']:.0f}us, tr={trial['triton_us']:.0f}us)"
+            if trial.get("baseline_us") is not None and trial.get("custom_us") is not None:
+                runtime = f" (bl={trial['baseline_us']:.0f}us, opt={trial['custom_us']:.0f}us)"
             best_marker = " <<<< BEST" if tid == state["best_trial"] else ""
             strategy_short = (trial["strategy"] or "")[:60]
             lines.append(

src/xe_forge/pipeline.py

@@ -443,7 +443,7 @@ def optimize(
                             correctness="pass" if stage_result.success else "fail",
                             speedup=speedup,
                             baseline_us=(val_orig_ms or 0) * 1000,
-                            triton_us=(current_ms or 0) * 1000,
+                            custom_us=(current_ms or 0) * 1000,
                         )
                         if stage_result.success:
                             last_trial_id = trial_id

src/xe_forge/skills/__init__.py

@@ -66,7 +66,7 @@ def main():
     t_result.add_argument("--correctness", choices=["pass", "fail"])
     t_result.add_argument("--speedup", type=float)
     t_result.add_argument("--baseline-us", type=float)
-    t_result.add_argument("--triton-us", type=float)
+    t_result.add_argument("--custom-us", "--triton-us", dest="custom_us", type=float)
     t_result.add_argument("--trials-dir", default="./trials")
 
     t_status = trial_sub.add_parser("status")

src/xe_forge/skills/benchmark.py

@@ -32,14 +32,14 @@ def run(args):
             init_args=init_args,
             input_dtypes=input_dtypes,
         )
-        if optimized_result.success:
+        if optimized_result.success and optimized_result.execution_time_ms is not None:
             baseline_ms = sum(baseline_us) / len(baseline_us) / 1000.0
             opt_ms = optimized_result.execution_time_ms
             speedup = baseline_ms / opt_ms if opt_ms > 0 else 0
             print(f"Correctness: {'PASSED' if optimized_result.success else 'FAILED'}")
             print(
                 f"Performance: baseline_us={baseline_ms * 1000:.2f}, "
-                f"triton_us={opt_ms * 1000:.2f}, speedup={speedup:.2f}x"
+                f"custom_us={opt_ms * 1000:.2f}, speedup={speedup:.2f}x"
             )
         else:
             print("Correctness: FAILED")
@@ -58,7 +58,7 @@ def run(args):
         if result.original_time_us and result.optimized_time_us:
             print(
                 f"Performance: baseline_us={result.original_time_us:.2f}, "
-                f"triton_us={result.optimized_time_us:.2f}, speedup={result.speedup:.2f}x"
+                f"custom_us={result.optimized_time_us:.2f}, speedup={result.speedup:.2f}x"
             )
         if result.feedback_message:
             print(f"Feedback: {result.feedback_message}")

src/xe_forge/skills/trial.py

@@ -30,7 +30,7 @@ def run(args):
                 correctness=args.correctness,
                 speedup=args.speedup,
                 baseline_us=args.baseline_us,
-                triton_us=args.triton_us,
+                custom_us=getattr(args, "custom_us", getattr(args, "triton_us", None)),
             )
             status_icon = {"completed": "+", "failed": "X", "partial": "~", "saved": "?"}
             icon = status_icon.get(trial["status"], "?")