feat: 支持 GLM-4.7-Flash 推理

.agents/skills/add-sparse-method/SKILL.md

@@ -28,7 +28,7 @@ Follow this placement order.
 5. Keep `src/sparsevllm/layers/attention.py` method-agnostic. It may call generic hooks, but should not grow method-specific branches unless adding a new reusable hook.
 6. Put cross-layer observation, attention-score collection, or scheduler-facing sparse orchestration in `src/sparsevllm/engine/sparse_controller.py`.
 7. Use `src/sparsevllm/utils/` only for truly generic helpers shared by multiple methods. Do not place an entire method implementation there.
-8. Add custom kernels under `src/sparsevllm/triton_kernel/` or another explicit runtime module, then call them through the method's cache manager or shared decode path.
+8. Add custom kernels under `src/sparsevllm/kernels/triton/` or another explicit runtime module, then call them through the method's cache manager or shared decode path.
 
 ## Decision Rules
 

.agents/skills/add-sparse-method/references/file-map.md

@@ -48,7 +48,7 @@ Do not bury a full method implementation in `attention.py`.
 
 ## Add Kernel Code Only When Needed
 
-Touch `src/sparsevllm/triton_kernel/` or another explicit kernel module when:
+Touch `src/sparsevllm/kernels/triton/` or another explicit kernel module when:
 
 - the existing decode or prefill kernels are the bottleneck
 - the method requires a new layout-aware fused operator

.agents/skills/optimize-sparsevllm-kernel/SKILL.md

@@ -0,0 +1,160 @@
+---
+name: optimize-sparsevllm-kernel
+description: Optimize and integrate Sparse-vLLM GPU kernels across Triton, TileLang, CUDA/CuTe, and external SGL kernels. Use when Codex needs to identify an LLM inference hotspot, choose a kernel implementation path, write or fuse a kernel, tune an existing kernel, build correctness or microbenchmark coverage, analyze Nsight Compute results, integrate a provider, or validate kernel and end-to-end performance.
+---
+
+# Optimize Sparse-vLLM Kernel
+
+Use one evidence loop from the serving workload to the kernel and back:
+
+```text
+matched LLM workload
+  -> hotspot evidence
+  -> implementation choice
+  -> correctness oracle
+  -> stable microbenchmark
+  -> targeted tuning
+  -> Nsight Compute
+  -> operator/provider integration
+  -> matched end-to-end validation
+```
+
+Do not assume that a kernel rewrite is useful before locating its contribution
+to the requested workload. If the user already specifies a kernel, proceed but
+state whether end-to-end hotspot evidence exists.
+
+## Load the Relevant Guidance
+
+Read [benchmark-protocol.md](references/benchmark-protocol.md) before running
+any performance experiment.
+
+Read exactly one implementation guide first, then load another only when a
+measured comparison requires it:
+
+- Triton: [triton.md](references/triton.md)
+- TileLang: [tilelang.md](references/tilelang.md)
+- CUDA, CuTe, CUTLASS, or an external compiled kernel:
+  [cuda-cute.md](references/cuda-cute.md)
+
+Read [nsight-playbook.md](references/nsight-playbook.md) before collecting or
+interpreting Nsight Compute data. Read
+[operator-integration.md](references/operator-integration.md) before changing
+provider selection, dependencies, workspaces, layouts, model call sites, or
+production dispatch.
+
+Use [reference-sources.md](references/reference-sources.md) only when an
+upstream example is needed. Pin the exact source revision and inspect its
+license before adapting code.
+
+When available, use companion skills for their focused expertise:
+
+- `llm-torch-profiler-analysis` for SGLang/vLLM/TRT-LLM trace analysis.
+- `kernel-triton-writing` for Triton implementation details.
+- `add-jit-kernel` for JIT CUDA integration without a large C++ project.
+- `add-sgl-kernel` for CUTLASS or complex AOT integration.
+- `kernel-cute-writing` for CuTe-specific implementation.
+- `perf-nsight-compute-analysis` for deep Nsight Compute interpretation.
+- `debug-cuda-crash` for illegal access, misalignment, or graph-capture faults.
+
+Do not block when a companion skill is unavailable; follow this skill's local
+references and report the missing capability.
+
+## Execute the Workflow
+
+### 1. Freeze the Scope
+
+Inspect Git status before editing and preserve unrelated tracked and untracked
+work. Define the operator, phase, workload, hardware, dtype, shapes, layouts,
+parallel topology, CUDA Graph mode, and comparison baseline. Check all devices
+and select an idle permitted GPU before starting a GPU task.
+
+### 2. Locate the Cost
+
+Profile a representative end-to-end workload. Separate prefill, decode,
+sampling, communication, host overhead, graph replay, and compilation. Rank
+hotspots by total contribution rather than kernel latency alone. Record fusion
+and overlap opportunities, but treat them as hypotheses until measured.
+
+### 3. Choose the Implementation Path
+
+Prefer the smallest path that can express the required computation:
+
+- Keep or improve Triton for broadly applicable repository-owned kernels.
+- Use TileLang when explicit tiling, pipelining, shared-memory layouts,
+  tensor-core scheduling, TMA, or warp specialization materially helps.
+- Prefer JIT CUDA when custom CUDA is needed without CUTLASS or a large AOT
+  project.
+- Use compiled SGL/CUTLASS/CuTe integration only when the required primitives,
+  layouts, or performance cannot be reached cleanly with a JIT path.
+
+Do not replace a mature provider solely because another DSL looks promising.
+
+### 4. Establish Correctness
+
+Create an independent Torch or mathematically direct oracle. Specify input and
+output shapes, dtypes, strides, aliases, mutation, padding, empty cases,
+numerical tolerances, and reduction semantics. Cover real model shapes,
+boundary shapes, non-contiguous inputs when supported, repeated execution, and
+CUDA Graph capture/replay when claimed.
+
+Reject a candidate immediately when correctness fails. Never tune against a
+known-wrong implementation.
+
+### 5. Build the Microbenchmark
+
+Benchmark the actual callable boundary used by serving, including required
+workspace initialization, output reset, synchronization, or materialization.
+Separate compile and cold-start cost from steady-state latency. Use shapes
+derived from the workload, fixed inputs and seeds, sufficient warmup, raw
+samples, and identical conditions for baseline and candidate.
+
+### 6. Tune with Bounded Hypotheses
+
+Change one explained dimension or run a declared finite matrix. Record every
+candidate, including failures. Typical dimensions include tile shape, program
+grid, threads or warps, pipeline stages, split count, vector width,
+shared-memory layout, swizzle, async copy/TMA, fusion boundary, register use,
+and workspace layout.
+
+Tune offline. Do not benchmark or search configurations in the serving hot
+path. Bind the selected configuration deterministically from validated shape
+and device facts.
+
+### 7. Profile Finalists
+
+Run Nsight Compute only after correctness and stable timing identify a small
+set of finalists. Compare the same shape and call boundary. Use measured
+occupancy, achieved bandwidth, tensor-core utilization, register or local
+memory use, shared-memory behavior, scheduler issue rate, and warp stalls to
+support the next hypothesis. Do not infer a bottleneck from occupancy alone.
+
+### 8. Integrate through Operators
+
+Keep kernels under the repository-owned kernel packages and keep dependency
+checks, support predicates, workspace ownership, launch selection, and
+fallback policy under `src/sparsevllm/operators/`. Resolve and bind a provider
+before the forward hot path. Route unsupported configurations before launch;
+never catch a runtime kernel failure and silently switch providers.
+
+### 9. Return to End-to-End Measurement
+
+Repeat the original workload with the same model, request trace, concurrency,
+context/output lengths, TP/EP topology, cache state, graph mode, and metric
+window. Report kernel latency improvement separately from end-to-end latency,
+throughput, TTFT, or TPOT. A faster microbenchmark is not an end-to-end win.
+
+## Acceptance Gates
+
+Require all applicable gates before calling the work complete:
+
+1. Independent correctness equivalence passes.
+2. Provider selection and rejection paths are tested.
+3. Claimed devices, dtypes, shapes, and graph modes are exercised on hardware.
+4. Microbenchmark raw samples and summary are saved.
+5. Nsight evidence exists for hardware-level bottleneck claims.
+6. Matched end-to-end validation supports serving-level claims.
+7. Commands, Git state, environment, selected provider, and artifacts are
+   recorded.
+
+Mark any unrun gate explicitly. Preserve the best verified implementation and
+the baseline; do not leave an unverified candidate as the production default.

.agents/skills/optimize-sparsevllm-kernel/agents/openai.yaml

@@ -0,0 +1,4 @@
+interface:
+  display_name: "Optimize Sparse-vLLM Kernel"
+  short_description: "Optimize Sparse-vLLM GPU kernels end to end"
+  default_prompt: "Use $optimize-sparsevllm-kernel to find, implement, tune, profile, and integrate a Sparse-vLLM GPU kernel."

.agents/skills/optimize-sparsevllm-kernel/references/benchmark-protocol.md

@@ -0,0 +1,74 @@
+# Kernel Benchmark Protocol
+
+Apply this protocol to every performance claim.
+
+## Before Launch
+
+1. Inspect every GPU's utilization, memory, and compute-process ownership.
+2. Select an idle permitted device. If none is idle, wait or report instead of
+   sharing a busy device.
+3. Inspect Git status and preserve unrelated tracked and untracked work.
+4. Record the kernel callable boundary and the end-to-end workload that
+   produced the target shapes.
+5. Establish correctness before collecting performance samples.
+
+## Record the Case
+
+Save a manifest containing:
+
+- repository path, Git SHA, branch, and dirty status
+- exact command and interpreter
+- Torch, Triton, TileLang, CUDA, driver, and relevant kernel-package versions
+- GPU name, compute capability, selected device, clocks or power constraints
+  when controlled
+- input shapes, dtypes, strides, seed, and data-generation method
+- provider, launch configuration, graph mode, TP/EP topology, and cache state
+- warmup count, timed repetitions, timing method, and synchronization boundary
+- output paths and explicit run status
+
+Do not overwrite previous raw results. Use a new run directory or immutable
+case identifier.
+
+## Microbenchmark Semantics
+
+- Time the serving-relevant wrapper, not only an internal launch, unless the
+  result is explicitly labeled kernel-only.
+- Include required output reset, workspace preparation, conversion, or
+  materialization. Report optional components separately when decomposition is
+  useful.
+- Exclude compilation from steady-state latency after recording cold compile
+  time separately.
+- Warm every specialization and synchronize before and after the timed region.
+- Use CUDA events or another GPU-aware timer correctly; never time asynchronous
+  launches with host wall time alone.
+- Use identical inputs, streams, graph mode, and synchronization for baseline
+  and candidate. Interleave them when long runs may drift.
+- Save raw samples and report at least sample count and median. Add tail or
+  dispersion statistics when they affect the decision.
+- State whether caches are intentionally warm or cold. Do not mix regimes.
+- Repeat suspicious gains and reject results affected by competing processes,
+  throttling, compilation, or changing clocks.
+
+## End-to-End Semantics
+
+Match model/checkpoint, request trace, prompt and output lengths, batch and
+concurrency, TP/EP, cache state, graph/provider settings, decoding parameters,
+and metric window. Keep TTFT, TPOT/inter-token latency, request latency,
+throughput, and kernel time distinct. Label partial serving runs separately
+from completed end-to-end results.
+
+## Artifact Minimum
+
+Persist:
+
+```text
+run_manifest.json
+raw_samples.jsonl
+summary.json
+stdout.log
+stderr.log
+```
+
+Add profiler traces and `.ncu-rep` files when collected. Use structured status
+and error fields; do not treat a non-empty output or a successful process start
+as benchmark success.

.agents/skills/optimize-sparsevllm-kernel/references/cuda-cute.md

@@ -0,0 +1,39 @@
+# CUDA, CuTe, CUTLASS, and SGL Kernel Guide
+
+Use a compiled or JIT CUDA path only when its required primitives, layout
+control, or measured performance justify the added integration cost.
+
+## Choose the Smallest Integration
+
+- Prefer JIT CUDA when the kernel does not require CUTLASS or a large C++
+  project. Use `add-jit-kernel` when available.
+- Use CuTe when explicit NVIDIA tensor layouts and architecture-specific
+  primitives are central. Use `kernel-cute-writing` when available.
+- Use SGL/CUTLASS AOT integration for complex compiled projects or packaged
+  kernels. Use `add-sgl-kernel` when available.
+- Keep a verified Triton or Torch implementation as the correctness baseline
+  and, where supported, the portable provider.
+
+## Control the Contract
+
+Declare supported compute capabilities, CUDA/toolchain versions, dtypes,
+alignments, layouts, workspace, streams, graph behavior, and mutation. Keep
+build and package availability checks inside the provider and import compiled
+extensions lazily.
+
+Do not expose CUTLASS packing, reordered projections, descriptor formats, or
+workspace details to model classes. The selected provider owns physical
+layouts and preparation. Reject unsupported configurations during resolution
+or preparation; do not catch a launch failure and switch implementations.
+
+## Validate
+
+Compare the actual wrapper against an independent oracle over production and
+boundary shapes. Exercise the minimum declared dependency/toolchain and every
+claimed architecture on real hardware. Check sanitizer or crash diagnostics
+for indexing and lifetime changes, and validate CUDA Graph capture/replay when
+advertised.
+
+Measure compile/startup separately from steady-state execution. Include any
+descriptor construction, packing, workspace clearing, or output conversion
+that remains in the service call boundary.

.agents/skills/optimize-sparsevllm-kernel/references/nsight-playbook.md

@@ -0,0 +1,62 @@
+# Nsight Compute Playbook
+
+Use Nsight Compute after correctness and stable microbenchmarks narrow the
+candidate set. Nsight replay and metric collection can perturb execution, so
+do not use profiler duration as the production latency result.
+
+## Capture Deliberately
+
+1. Choose one representative shape and a callable that launches the target
+   kernel predictably.
+2. Warm JIT compilation before capture.
+3. Filter the target kernel or limit launches when possible.
+4. Start with focused sections, then collect a broader set only when needed.
+5. Save the command, console output, and `.ncu-rep` artifact.
+6. Capture baseline and candidate under the same software and hardware state.
+
+Useful section families include Speed of Light, Launch Statistics, Occupancy,
+Memory Workload Analysis, Scheduler Statistics, Warp State Statistics, and
+Source Counters. Confirm the exact section names supported by the installed
+`ncu` version before scripting them.
+
+## Interpret as a Chain of Evidence
+
+### Launch and Occupancy
+
+Check grid size, waves per SM, block/thread shape, shared memory, registers per
+thread, theoretical occupancy, and achieved occupancy. Low occupancy matters
+only when it limits latency hiding or parallelism; high occupancy does not
+prove efficiency.
+
+### Memory
+
+Compare achieved DRAM/L2/shared throughput, transaction efficiency, cache hit
+rates, sectors, and shared-bank conflicts. Relate bytes moved to the algorithm
+and wrapper, including intermediate tensors eliminated or introduced by
+fusion.
+
+### Compute
+
+Check tensor-core or arithmetic-pipe utilization, instruction mix, issue rate,
+and dependency stalls. Verify that the chosen tile and dtype actually reach
+the intended hardware path.
+
+### Registers and Stalls
+
+Inspect register count, local-memory traffic, spills, scoreboard/dependency
+stalls, barrier stalls, memory throttling, and not-selected warps. Connect a
+stall change to a concrete code or schedule change before acting on it.
+
+### Roofline
+
+Estimate arithmetic intensity using the measured callable boundary. Classify
+memory- versus compute-limited behavior only when achieved bandwidth/compute
+and the traffic model agree. Re-evaluate after fusion because the boundary and
+bytes moved have changed.
+
+## Close the Loop
+
+Use the profile to form one next hypothesis, benchmark the resulting change,
+and retain it only if stable latency improves without breaking correctness.
+Do not optimize a metric that does not move the microbenchmark or end-to-end
+result.