diff --git a/Makefile b/Makefile index 3783387c..4fa267bf 100755 --- a/Makefile +++ b/Makefile @@ -7,7 +7,7 @@ SHELL := /bin/bash .PHONY: help docker-shell docker-check-agents docker-smoke docker-run docker-parallel-run docker-setup-flydsl \ - check-docker-runner check-evaluator check-held-out check-visualization \ + check-docker-runner check-evaluator check-gpu-architecture check-held-out check-visualization \ visualization-build visualization-serve visualization-run \ sync-perf-helpers check-perf-helpers materialize-perf-workspace \ materialize-perf-task cleanup-works install-cursor-agent vllm @@ -23,12 +23,13 @@ help: @echo "make docker-smoke - Verify Docker Python, ROCm tools, imports, and GPU access" @echo "make docker-run CONFIG=example_configs/quickstart_claude_mi300.yaml RUN_ARGS=\"--run-suffix test\" - Run an experiment in Docker" @echo "make docker-parallel-run CONFIG=example_configs/benchmark_cursor_mi355x.yaml GPU_IDS=0,1 - Run an experiment across one worker container per GPU" - @echo " Default CONFIG is the MI300/MI300X Claude quickstart" + @echo " Default CONFIG is the generic MI300-series Claude quickstart" @echo " On other GPUs, pass a matching CONFIG explicitly" @echo " Images: gfx942->mi30x, gfx950->mi35x; override with AKA_DOCKER_IMAGE=..." @echo "make docker-setup-flydsl - Install FlyDSL when absent (for flydsl2flydsl, torch2flydsl, and triton2flydsl)" @echo "make check-docker-runner - Check Docker runner syntax and runtime-specific arguments" @echo "make check-evaluator - Run centralized evaluator unit tests" + @echo "make check-gpu-architecture - Check GPU model mappings and composed architecture prompts" @echo "make check-held-out - Run held-out module unit tests" @echo "make visualization-run - Build and serve the local comparison dashboard" @echo "make check-visualization - Run visualization module unit tests" @@ -80,6 +81,9 @@ check-docker-runner: check-evaluator: @python3 -m unittest discover -s tests -p 'test_evaluator_*.py' +check-gpu-architecture: + @python3 -m unittest discover -s tests -p 'test_gpu_architecture.py' + check-held-out: @python3 -m unittest discover -s tests -p 'test_held_out.py' diff --git a/README.md b/README.md index b57b934c..c0644fcc 100755 --- a/README.md +++ b/README.md @@ -187,17 +187,22 @@ installation. The npm path requires Node.js 22+ and npm. See the [official Claude Code setup guide](https://code.claude.com/docs/en/installation) for the current alternatives. -The repository provides three ready-to-use run configurations: +The repository provides generic and model-specific quickstarts plus a longer +MI355X benchmark configuration: | Configuration | Purpose | | --- | --- | -| `example_configs/quickstart_claude_mi300.yaml` | One Claude Code GELU task on MI300/MI300X (`gfx942`); use this for a first run on MI300-series hardware. | +| `example_configs/quickstart_claude_mi300.yaml` | Generic MI300-series (`gfx942`) Claude Code GELU task; retains backward compatibility without assuming a physical SKU. | +| `example_configs/quickstart_claude_mi300x.yaml` | Model-specific MI300X (`gfx942`) Claude Code GELU task. | +| `example_configs/quickstart_claude_mi325x.yaml` | Model-specific MI325X (`gfx942`) Claude Code GELU task. | +| `example_configs/quickstart_claude_mi300a.yaml` | Model-specific MI300A (`gfx942`) Claude Code GELU task. | | `example_configs/quickstart_claude_mi355x.yaml` | One Claude Code GELU task on MI355X (`gfx950`); use this for a first run on MI355X. | | `example_configs/benchmark_cursor_mi355x.yaml` | Curated 60-task Cursor Agent benchmark on MI355X; use this for a longer benchmark only after installing and authenticating Cursor Agent. | -Running `make docker-run` without `CONFIG` uses the MI300/MI300X Claude -quickstart. On another GPU, pass the matching configuration explicitly; for -example, use `CONFIG=example_configs/quickstart_claude_mi355x.yaml` on MI355X. +Running `make docker-run` without `CONFIG` uses the generic MI300-series Claude +quickstart. Pass the exact model configuration to give the optimization agent +the correct XCD, CU, and HBM limits. For example, use +`CONFIG=example_configs/quickstart_claude_mi325x.yaml` on MI325X. FlyDSL tasks require FlyDSL in the container. The pinned image may already provide it; otherwise run: diff --git a/agents/mini_swe_triton/launch_agent.py b/agents/mini_swe_triton/launch_agent.py index ae87a6e5..7796af31 100644 --- a/agents/mini_swe_triton/launch_agent.py +++ b/agents/mini_swe_triton/launch_agent.py @@ -23,6 +23,7 @@ import yaml from agents import register_agent +from src.prompt_builder import _load_cheatsheet def _read_stream(stream, lines: list, prefix: str, log_func): @@ -122,6 +123,16 @@ def launch_agent(eval_config: dict[str, Any], task_config_dir: str, workspace: s num_parallel = agent_config.get("agent", {}).get("num_parallel", 2) model = agent_config.get("agent", {}).get("model", "claude-opus-4-6") step_limit = agent_config.get("agent", {}).get("step_limit", 100) + target_gpu_model = str(eval_config.get("target_gpu_model") or "MI300") + project_root = Path(__file__).resolve().parents[2] + architecture_context, gfx_arch = _load_cheatsheet( + task_config.get("task_type", "triton2triton"), + target_gpu_model, + project_root, + task_config, + logger, + include_knowledge=False, + ) # PYTHONPATH for mini-swe-agent modules. GEAK_SRC must point to the # absolute path of the GEAK source tree (the directory containing @@ -149,6 +160,7 @@ def launch_agent(eval_config: dict[str, Any], task_config_dir: str, workspace: s logger.info(f" num_parallel: {num_parallel}") logger.info(f" model: {model}") logger.info(f" step_limit: {step_limit}") + logger.info(f" target_gpu: {target_gpu_model} ({gfx_arch or 'unknown architecture'})") logger.info("=" * 60) all_output: list[str] = [] @@ -161,7 +173,11 @@ def launch_agent(eval_config: dict[str, Any], task_config_dir: str, workspace: s else: kernel_snippet = kernel_code - task_prompt = f"""Optimize this Triton GPU kernel for maximum performance on AMD MI300X (gfx942/gfx950). + target_description = target_gpu_model + if gfx_arch: + target_description += f" ({gfx_arch})" + + task_prompt = f"""Optimize this Triton GPU kernel for maximum performance on AMD {target_description}. The kernel is at: {kernel_path.name} The test harness is at: {harness_path.name} @@ -176,7 +192,9 @@ def launch_agent(eval_config: dict[str, Any], task_config_dir: str, workspace: s - Correctness must pass after your changes - Focus on real kernel-body optimizations (block sizes, memory access patterns, vectorization, loop unrolling, warp-level primitives) -- Target: AMD MI300X with gfx942/gfx950 architecture, 304 CUs, HBM3 + +Target hardware context: +{architecture_context or f"No model profile is available for {target_description}; inspect the visible GPU before tuning."} Current kernel code: ```python diff --git a/docs/examples/examples.md b/docs/examples/examples.md index ea3b93e6..fc154676 100644 --- a/docs/examples/examples.md +++ b/docs/examples/examples.md @@ -17,11 +17,11 @@ multi-GPU examples use `make docker-parallel-run`. Run the single-task Claude Code quickstart for the physical GPU. -1. Select the MI300/MI300X example (use - `example_configs/quickstart_claude_mi355x.yaml` on MI355X): +1. Select the example matching the physical GPU (use the adjacent MI325X, + MI300A, or MI355X quickstart for those models): ```bash - CONFIG_PATH=example_configs/quickstart_claude_mi300.yaml + CONFIG_PATH=example_configs/quickstart_claude_mi300x.yaml ``` 2. Run: diff --git a/docs/how-to/run-evaluation.md b/docs/how-to/run-evaluation.md index ddf0f36f..6875e2f2 100644 --- a/docs/how-to/run-evaluation.md +++ b/docs/how-to/run-evaluation.md @@ -15,18 +15,23 @@ resume, and inspect a run. ## Choose or create a run configuration A run configuration selects the agent, tasks, and target GPU. The repository -ships three examples: +ships generic and model-specific quickstarts plus a longer benchmark example: | Configuration | Purpose | | --- | --- | -| `example_configs/quickstart_claude_mi300.yaml` | One Claude Code GELU task on MI300/MI300X (`gfx942`). | +| `example_configs/quickstart_claude_mi300.yaml` | Generic MI300-series (`gfx942`) fallback. | +| `example_configs/quickstart_claude_mi300x.yaml` | One Claude Code GELU task on MI300X (`gfx942`). | +| `example_configs/quickstart_claude_mi325x.yaml` | One Claude Code GELU task on MI325X (`gfx942`). | +| `example_configs/quickstart_claude_mi300a.yaml` | One Claude Code GELU task on MI300A (`gfx942`). | | `example_configs/quickstart_claude_mi355x.yaml` | One Claude Code GELU task on MI355X (`gfx950`). | | `example_configs/benchmark_cursor_mi355x.yaml` | Curated 60-task Cursor Agent benchmark on MI355X; use only after installing and authenticating Cursor Agent. | For a first run, select the quickstart that matches the physical GPU: ```bash -CONFIG_PATH=example_configs/quickstart_claude_mi300.yaml +CONFIG_PATH=example_configs/quickstart_claude_mi300x.yaml +# MI325X: example_configs/quickstart_claude_mi325x.yaml +# MI300A: example_configs/quickstart_claude_mi300a.yaml # For MI355X instead: # CONFIG_PATH=example_configs/quickstart_claude_mi355x.yaml ``` diff --git a/docs/install/install.md b/docs/install/install.md index 4d459b0f..3265175a 100644 --- a/docs/install/install.md +++ b/docs/install/install.md @@ -93,24 +93,31 @@ their own runtime dependencies. Review the corresponding directory under ## Choose an example configuration Choose the configuration that matches the physical GPU and installed agent. -The two quickstart configurations each run one GELU task; the benchmark -configuration is a longer 60-task Cursor Agent run. +Each quickstart configuration runs one GELU task; the benchmark configuration +is a longer 60-task Cursor Agent run. | Configuration | Purpose | | --- | --- | -| `example_configs/quickstart_claude_mi300.yaml` | First Claude Code run on MI300/MI300X (`gfx942`). | +| `example_configs/quickstart_claude_mi300.yaml` | Generic MI300-series (`gfx942`) fallback. | +| `example_configs/quickstart_claude_mi300x.yaml` | First Claude Code run on MI300X (`gfx942`). | +| `example_configs/quickstart_claude_mi325x.yaml` | First Claude Code run on MI325X (`gfx942`). | +| `example_configs/quickstart_claude_mi300a.yaml` | First Claude Code run on MI300A (`gfx942`). | | `example_configs/quickstart_claude_mi355x.yaml` | First Claude Code run on MI355X (`gfx950`). | | `example_configs/benchmark_cursor_mi355x.yaml` | Curated 60-task Cursor Agent benchmark on MI355X; requires an installed and authenticated Cursor Agent CLI. | -The default `make docker-run` configuration is the MI300/MI300X quickstart. -On another GPU, always pass the matching `CONFIG`; for MI355X, use -`example_configs/quickstart_claude_mi355x.yaml`. +The default `make docker-run` configuration is the generic MI300-series +quickstart. Prefer the matching model-specific `CONFIG` so the optimization +agent receives accurate hardware limits. Select one quickstart configuration in the current shell: ```bash -# MI300/MI300X: -CONFIG_PATH=example_configs/quickstart_claude_mi300.yaml +# MI300X: +CONFIG_PATH=example_configs/quickstart_claude_mi300x.yaml + +# For MI325X or MI300A instead: +# CONFIG_PATH=example_configs/quickstart_claude_mi325x.yaml +# CONFIG_PATH=example_configs/quickstart_claude_mi300a.yaml # For MI355X instead: # CONFIG_PATH=example_configs/quickstart_claude_mi355x.yaml diff --git a/docs/reference/api-reference.md b/docs/reference/api-reference.md index ce076500..c811cafb 100644 --- a/docs/reference/api-reference.md +++ b/docs/reference/api-reference.md @@ -19,7 +19,7 @@ A run configuration defines a single experiment. Start from a file under | --- | --- | --- | | `agent.template` | string | Agent to run. One of the [supported agents](../how-to/agents.md#supported-agents). | | `tasks` | list of strings | Task selectors relative to `tasks/`. Use `all` for every task, a category prefix for a group, or a full path for a single task. | -| `target_gpu_model` | string | Target GPU model, for example `MI300` or `MI355X`. Used to select the Docker image architecture, set `PYTORCH_ROCM_ARCH`, and name the workspace. | +| `target_gpu_model` | string | Target GPU model: `MI300X`, `MI325X`, `MI300A`, `MI355X`, or the generic `MI300` compatibility target. Used to select the Docker image architecture, compose model-aware optimization context, set `PYTORCH_ROCM_ARCH`, and name the workspace. | | `log_directory` | string | Directory for run logs. | | `workspace_directory_prefix` | string | Prefix for the workspace directory. The full name is `__`. | @@ -56,7 +56,7 @@ The in-container `main.py` entrypoint accepts these flags: | Flag | Description | | --- | --- | -| `--config_name ` | Config file to load (default `example_configs/quickstart_claude_mi300.yaml` for MI300/MI300X). Pass a matching config explicitly on another GPU | +| `--config_name ` | Config file to load (default `example_configs/quickstart_claude_mi300.yaml` for the generic MI300-series target). Pass a model-specific config for exact hardware context | | `--run-suffix ` | Suffix appended to the run directory name (letters, numbers, `.`, `_`, `-` only). Useful for labeling A/B runs | | `--resume-run ` | Resume a specific run directory, skipping completed tasks | | `--resume-latest` | Resume the most recent run in the workspace | diff --git a/docs/reference/compatibility-matrix.md b/docs/reference/compatibility-matrix.md index 780d0ee2..019f950c 100644 --- a/docs/reference/compatibility-matrix.md +++ b/docs/reference/compatibility-matrix.md @@ -9,12 +9,23 @@ myst: ## Hardware -The following hardware configurations are supported and tested. +The following GPU models have explicit runtime and optimization-prompt +configuration. Model-to-architecture and Docker routing are covered by automated +tests. Hardware task validation is listed separately so that configured support +is not mistaken for results from unavailable physical GPUs. -| Component | Supported | Notes | -| --- | --- | --- | -| GPU architecture | AMD Instinct™ MI300 series | `target_gpu_model: MI300` | -| GPU architecture | AMD Instinct™ MI355X | `target_gpu_model: MI355X` | +| GPU model | Configuration value | LLVM target | Validation status | +| --- | --- | --- | --- | +| AMD Instinct™ MI300X | `target_gpu_model: MI300X` | `gfx942` | Supported by the existing MI300-series runtime path. | +| AMD Instinct™ MI325X | `target_gpu_model: MI325X` | `gfx942` | Configuration tested; physical GPU task validation pending. | +| AMD Instinct™ MI300A | `target_gpu_model: MI300A` | `gfx942` | Configuration tested; physical GPU task validation pending. | +| Generic AMD Instinct™ MI300 series | `target_gpu_model: MI300` | `gfx942` | Backward-compatible fallback; runtime detection is required for SKU-specific limits. | +| AMD Instinct™ MI355X | `target_gpu_model: MI355X` | `gfx950` | Supported and tested. | + +MI300X, MI325X, and MI300A share the `gfx942` compiler target but load separate +hardware profiles. The shared CDNA 3 guidance contains ISA-level constraints; +each profile supplies its model-specific compute topology, HBM limits, and +partitioning guidance. ## Software diff --git a/docs/reference/release-notes.md b/docs/reference/release-notes.md index 477523ca..00b9dfdc 100644 --- a/docs/reference/release-notes.md +++ b/docs/reference/release-notes.md @@ -27,8 +27,10 @@ experimentation platform and RL-ready task environment for GPU kernel agents. - **Docker-first execution**: select pinned ROCm/SGLang images from the target GPU architecture. - **Example run configurations**: provide one-task Claude Code quickstarts for - MI300/MI300X and MI355X plus a curated 60-task Cursor benchmark for MI355X - under `example_configs/`. + the generic MI300-series target, MI300X, MI325X, MI300A, and MI355X plus a + curated 60-task Cursor benchmark for MI355X under `example_configs/`. +- **Model-aware MI300 profiles**: compose shared CDNA 3 (`gfx942`) guidance with + separate MI300X, MI325X, and MI300A topology and memory profiles. - **Centralized outcomes**: independently measure compilation, correctness, and GPU performance, then compute a configurable score. - **Visualization module**: build and serve the comparison dashboard through diff --git a/example_configs/quickstart_claude_mi300.yaml b/example_configs/quickstart_claude_mi300.yaml index 1784c0ba..1b00d96b 100644 --- a/example_configs/quickstart_claude_mi300.yaml +++ b/example_configs/quickstart_claude_mi300.yaml @@ -1,4 +1,5 @@ -# Minimal first-run example for Claude Code on MI300/MI300X (gfx942). +# Generic MI300-series Claude Code quickstart (gfx942). +# Prefer the model-specific MI300X, MI325X, or MI300A config for accurate tuning context. agent: template: claude_code diff --git a/example_configs/quickstart_claude_mi300a.yaml b/example_configs/quickstart_claude_mi300a.yaml new file mode 100644 index 00000000..b2d932f0 --- /dev/null +++ b/example_configs/quickstart_claude_mi300a.yaml @@ -0,0 +1,10 @@ +# Minimal first-run example for Claude Code on MI300A (gfx942). +agent: + template: claude_code + +tasks: + - hip2hip/gpumode/GELU + +target_gpu_model: MI300A +log_directory: logs +workspace_directory_prefix: workspace diff --git a/example_configs/quickstart_claude_mi300x.yaml b/example_configs/quickstart_claude_mi300x.yaml new file mode 100644 index 00000000..d8458d0f --- /dev/null +++ b/example_configs/quickstart_claude_mi300x.yaml @@ -0,0 +1,10 @@ +# Minimal first-run example for Claude Code on MI300X (gfx942). +agent: + template: claude_code + +tasks: + - hip2hip/gpumode/GELU + +target_gpu_model: MI300X +log_directory: logs +workspace_directory_prefix: workspace diff --git a/example_configs/quickstart_claude_mi325x.yaml b/example_configs/quickstart_claude_mi325x.yaml new file mode 100644 index 00000000..385830b0 --- /dev/null +++ b/example_configs/quickstart_claude_mi325x.yaml @@ -0,0 +1,10 @@ +# Minimal first-run example for Claude Code on MI325X (gfx942). +agent: + template: claude_code + +tasks: + - hip2hip/gpumode/GELU + +target_gpu_model: MI325X +log_directory: logs +workspace_directory_prefix: workspace diff --git a/main.py b/main.py index d1bf1ee6..7f0a3faf 100755 --- a/main.py +++ b/main.py @@ -38,8 +38,8 @@ default="example_configs/quickstart_claude_mi300.yaml", help=( "run configuration for AgentKernelArena (default: " - "example_configs/quickstart_claude_mi300.yaml for MI300/MI300X). " - "Select a matching config explicitly when using another GPU." + "example_configs/quickstart_claude_mi300.yaml for the generic MI300-series target). " + "Select a model-specific config for exact hardware tuning context." ), ) parser.add_argument( diff --git a/src/prompt_builder.py b/src/prompt_builder.py index bc799194..e3d8c1aa 100755 --- a/src/prompt_builder.py +++ b/src/prompt_builder.py @@ -79,17 +79,19 @@ def task_result_format(task_result_template: str) -> str: def _load_cheatsheet(task_type_name: str, target_gpu_model: str, project_root: Path, - task_config: dict, logger: logging.Logger) -> tuple[str, str | None]: + task_config: dict, logger: logging.Logger, *, + include_knowledge: bool = True) -> tuple[str, str | None]: """ Load the combined cheatsheet prompt and resolve the target gfx arch string. The cheatsheet config (default_cheatsheet.yaml) has two independent sections: - - architecture: maps GPU model name → hardware spec document + gfx_arch string + - architecture: maps GPU model name → hardware spec document(s) + gfx_arch string - knowledge: maps target language → language best-practice guide (GPU-agnostic) For task_type ``repository``, the task config must set ``repository_language`` to a key under ``knowledge`` (e.g. hip, triton). Add new stacks by extending that map and the - referenced markdown file. + referenced markdown file. Set ``include_knowledge=False`` for integrations that need + only the target architecture and model profile. Returns: (cheatsheet_text, gfx_arch) where gfx_arch may be None if not found. @@ -116,16 +118,27 @@ def _load_cheatsheet(task_type_name: str, target_gpu_model: str, project_root: P ) if arch_entry: gfx_arch = arch_entry.get('gfx_arch') - arch_file = arch_entry.get('file') - if arch_file: + arch_files = arch_entry.get('files') + if arch_files is None: + legacy_arch_file = arch_entry.get('file') + arch_files = [legacy_arch_file] if legacy_arch_file else [] + elif isinstance(arch_files, str): + arch_files = [arch_files] + + for arch_file in arch_files: arch_path = project_root / arch_file parts.append(arch_path.read_text()) logger.info(f"Loaded architecture context for '{target_gpu_model}': {arch_path}") - else: - logger.warning(f"Architecture entry for '{target_gpu_model}' has no 'file' key") + if not arch_files: + logger.warning( + f"Architecture entry for '{target_gpu_model}' has no 'file' or 'files' key" + ) else: logger.warning(f"No architecture entry for GPU '{target_gpu_model}' in default_cheatsheet.yaml") + if not include_knowledge: + return "\n\n---\n\n".join(parts) if parts else "", gfx_arch + # --- Knowledge section --- # L3 repository tasks: language is configured per task (repository_language → key in default_cheatsheet.yaml). if task_type_name == 'repository': diff --git a/src/prompts/cheatsheet/CDNA3_gfx942_architecture.md b/src/prompts/cheatsheet/CDNA3_gfx942_architecture.md new file mode 100644 index 00000000..05e5d701 --- /dev/null +++ b/src/prompts/cheatsheet/CDNA3_gfx942_architecture.md @@ -0,0 +1,66 @@ +# AMD CDNA 3 (`gfx942`) Kernel Optimization Context + +Use this shared architecture guidance together with the target GPU's model +profile. The model profile is authoritative for XCD count, active compute units, +HBM capacity and bandwidth, and partition topology. + +## Execution model + +- CDNA 3 executes Wave64 wavefronts. Cross-lane operations must preserve the + 64-lane execution model. +- A workgroup can contain multiple wavefronts, up to 1024 work-items. +- An XCD is the compute chiplet. Each XCD contains 40 physical CUs, with 38 + active CUs at the aggregate device level and two disabled for yield. +- Each XCD has 4 MiB of L2 cache. Never infer aggregate L2 capacity without + consulting the model profile and the visible accelerator partition. + +## Memory hierarchy and locality + +- Each CU has 64 KiB of LDS and a 32 KiB L1 vector cache. Keep LDS accesses + bank-friendly and use padding when a tile layout would otherwise conflict. +- The MI300-series packages described by the model profiles have 256 MiB of + last-level Infinity Cache. It is shared package infrastructure, not a single + low-latency cache local to every XCD. +- Coalesce adjacent lanes onto aligned contiguous memory regions. Avoid random + traffic that repeatedly crosses XCD or memory-partition boundaries. +- Size tiles against the resources of the *visible logical device*. A CPX, QPX, + or DPX partition exposes fewer XCDs and less HBM than the full SPX device. +- Do not hardcode HBM capacity or peak bandwidth from another `gfx942` model. + +## Matrix and scalar data types + +- CDNA 3 matrix cores support FP64, FP32, TF32, FP16, BF16, FP8, and INT8 + operations. Use an MFMA shape that is valid for the selected input and + accumulator types. +- Align LDS loading, register tiling, and output layout with the selected MFMA + instruction. Inspect generated code when performance depends on a particular + lowering. +- CDNA 3 uses the FNUZ FP8 variants. Do not assume the OCP FP8 or MXFP formats + available on later CDNA generations. + +## Runtime topology checks + +Before topology-sensitive tuning, inspect the environment rather than assuming +the full physical package is visible: + +1. Confirm `gcnArchName` is `gfx942` and inspect the device name, CU count, and + visible memory through PyTorch/HIP device properties. +2. Use `amd-smi` or `rocminfo` when available to identify accelerator and memory + partition modes. +3. Benchmark and compare kernels under the same partition mode. A launch tuned + for an eight-XCD SPX device may be inappropriate for a one-XCD CPX device. + +## Kernel-generation constraints + +1. Preserve synchronization semantics. Replace a workgroup barrier with + wave-level synchronization only when all communicating work-items are proven + to remain in the same wavefront. +2. Bound VGPR and LDS use based on measured occupancy; avoid register spilling + and do not assume that maximum occupancy always gives maximum throughput. +3. Fuse memory-bound operations when it reduces measured global-memory traffic, + but retain numerically equivalent behavior and validate every supported + shape. +4. Prefer runtime-derived CU and memory counts over constants embedded in launch + heuristics. + +Reference: [AMD Instinct MI300 series microarchitecture](https://rocm.docs.amd.com/en/latest/conceptual/gpu-arch/mi300.html) diff --git a/src/prompts/cheatsheet/MI300A_profile.md b/src/prompts/cheatsheet/MI300A_profile.md new file mode 100644 index 00000000..504b8e90 --- /dev/null +++ b/src/prompts/cheatsheet/MI300A_profile.md @@ -0,0 +1,30 @@ +# AMD Instinct MI300A Model Profile + +## Physical resources + +- Architecture target: CDNA 3, `gfx942` +- Compute topology: 6 XCDs, 228 active CUs (38 active CUs per XCD) +- Aggregate L2: 24 MiB (4 MiB per XCD) +- Last-level Infinity Cache: 256 MiB +- Memory: 128 GB HBM3 shared by the integrated CPU and GPU, with approximately + 5.3 TB/s peak theoretical bandwidth +- Package topology: three Zen 4 CPU chiplets (24 CPU cores) and six CDNA 3 XCDs + +## APU and partition-aware tuning + +MI300A is an APU, not a discrete MI300X-class accelerator. CPU and GPU traffic +can share HBM capacity and bandwidth, so leave allocation headroom and avoid +assuming that the GPU owns all 128 GB. Unified physical memory also does not make +every placement or access path equally fast; use the framework's device tensors +and measure migrations and CPU/GPU contention when they are part of a task. + +Accelerator partitioning can expose fewer than six XCDs, including one-XCD CPX +logical devices. Detect the visible CU count, memory capacity, accelerator +partition, and memory partition before choosing persistent grids or whole-device +tile counts. Never reuse MI300X/MI325X constants of eight XCDs or 304 CUs. + +References: + +- [AMD CDNA 3 and MI300A overview](https://www.amd.com/en/technologies/cdna.html) +- [ROCm accelerator hardware specifications](https://rocm.docs.amd.com/en/latest/reference/gpu-arch-specs.html) +- [AMD SMI GPU partitioning](https://rocm.docs.amd.com/projects/amdsmi/en/latest/conceptual/partition.html) diff --git a/src/prompts/cheatsheet/MI300X_architecture.md b/src/prompts/cheatsheet/MI300X_architecture.md deleted file mode 100644 index e8639f08..00000000 --- a/src/prompts/cheatsheet/MI300X_architecture.md +++ /dev/null @@ -1,52 +0,0 @@ -# AMD MI300X (CDNA 3) Kernel Optimization Context & Directives - -## 1. Role & Objective -You are an expert AMD GPU Kernel Engineer. Your objective is to generate, optimize, and debug HIP/ROCm C++ and assembly kernels for the AMD Instinct™ MI300X (CDNA 3 architecture). Your optimizations must strictly adhere to the execution models, memory hierarchies, and hardware limits detailed below. - -## 2. Execution Model & Compute Topology - -AMD's execution model has specific terminologies and constraints. The MI300X is NOT a single monolithic die; it is a complex multi-chiplet module. -* **Wavefront:** CDNA uses **Wave64** (64 work-items per wavefront). When using LDS permute semantics or cross-lane operations, explicitly assume `thread_id` ranges from 0 to 63. -* **Workgroup:** Composed of multiple Wave64s. Maximum workgroup size is 1024. -* **XCD (Accelerated Compute Die):** The compute engines. The MI300X contains **8 XCDs**. -* **Compute Unit (CU):** The MI300X features **304 active CUs**. (Each of the 8 XCDs physically has 40 CUs, with 38 active and 2 disabled for yield). -* **IOD (I/O Die):** 4 IODs handle the Infinity Cache, HBM3 interfaces, and Infinity Fabric links. - -## 3. Memory Hierarchy & Locality Rules -Memory locality is the primary bottleneck on MI300X. You must design kernels to keep data accesses localized to the correct XCD and its corresponding HBM stack. - -### 3.1 Memory Specifications -* **LDS (Local Data Share):** **64 KB per CU**. - * *Constraint:* LDS allocation is performed at a 512-byte granularity. - * *Rule:* You must pad arrays in LDS to avoid severe bank conflicts. When writing low-level ISA/assembly, use the `M0` register for bounds clamping to prevent out-of-bounds LDS access. -* **L1 Cache:** 32 KB per CU. -* **L2 Cache:** 4 MB per XCD (16-way, 16 channels). - * *Rule:* L2 is the critical point for coalescing local traffic and maintaining intra-XCD coherency. Attempt to size your workgroup working sets to fit within this 4 MB boundary per XCD. -* **L3 / Infinity Cache (LLC):** 256 MB located on the IODs. - * *Architecture:* It acts as a memory-side cache and does NOT participate in coherency evictions, using a snoop filter instead. It resolves cross-XCD coherent requests. Peak bandwidth is ~17.2 TB/s. -* **HBM3 (Global Memory):** 192 GB capacity across 8 stacks, **5.3 TB/s peak bandwidth** (8192-bit bus). - -### 3.2 Memory Optimization Directives -1. **Target the L2/LLC:** Do not treat MI300X as having a unified L2 and HBM. Force hot working sets to stay in the 4MB L2 or 256MB Infinity Cache. Bandwidth drops precipitously if your kernel causes random accesses across IODs to HBM. -2. **Coalesced Access:** Global memory accesses must be coalesced. Ensure adjacent work-items in a Wave64 access contiguous 256-byte aligned memory segments. - -## 4. Matrix Cores & MFMA Instructions -For AI workloads, utilize the Matrix Cores via **MFMA (Matrix Fused Multiply-Add)** instructions. - -* **Target Data Types:** MI300X natively supports FP64, FP32, **TF32**, FP16, BF16, FP8, and INT8. - * *Note:* CDNA 3 **natively supports TF32** in hardware. You may use TF32 MFMA instructions. It DOES NOT support FP6 or FP4. -* **Tile Alignment:** MFMA instructions operate on specific matrix tile dimensions (e.g., 32x32x8). Strictly align your LDS block loading to match these wave-level matrix shapes. - -## 5. Execution Environment & NUMA Partitioning (Critical for Tuning) - -The MI300X exposes its multi-die nature through Compute and Memory partitions. As an expert agent, you must design and tune kernels with these partitions in mind: -* **SPX (Single Partition):** The system sees one giant GPU. Workgroups are distributed round-robin across all 8 XCDs. You cannot control placement. -* **CPX (Core Partition):** The GPU is split into 8 logical partitions (1 per XCD). Workgroups are pinned to a specific XCD. -* **NPS4 (Memory Partition):** HBM is divided into 4 NUMA quadrants to enforce locality. - -* **Directive for Kernel Tuning:** When generating code for microbenchmarks or extreme tuning, assume the environment is set to **CPX + NPS4**. Optimize the kernel to achieve maximum throughput on a *single XCD* (Phase A: Single XCD Locality) before scaling it up to rely on the Infinity Cache in SPX mode (Phase B: Full GPU Scaling). - -## 6. Strict Kernel Generation Constraints -1. **Never use `__syncthreads()` unnecessarily:** Replace with wave-level synchronization (`__builtin_amdgcn_s_barrier()`). -2. **Register Allocation & Spilling:** MI300X performance collapses if registers spill to memory. Keep VGPR usage tightly bounded to allow at least 2-4 waves per SIMD. Use `__launch_bounds__`. -3. **Kernel Fusion:** With 5.3 TB/s HBM bandwidth but incredibly high compute throughput, memory-bound operations (RoPE, SiLU, Softmax, RMSNorm) must be fused into compute-bound kernels (like GEMM) whenever possible to prevent unnecessary trips to HBM. \ No newline at end of file diff --git a/src/prompts/cheatsheet/MI300X_profile.md b/src/prompts/cheatsheet/MI300X_profile.md new file mode 100644 index 00000000..c371d5a4 --- /dev/null +++ b/src/prompts/cheatsheet/MI300X_profile.md @@ -0,0 +1,25 @@ +# AMD Instinct MI300X Model Profile + +## Physical resources + +- Architecture target: CDNA 3, `gfx942` +- Compute topology: 8 XCDs, 304 active CUs (38 active CUs per XCD) +- Aggregate L2: 32 MiB (4 MiB per XCD) +- Last-level Infinity Cache: 256 MiB +- Memory: 192 GB HBM3, 8192-bit interface, 5.3 TB/s peak theoretical bandwidth + +## Partition-aware tuning + +The full SPX device exposes eight XCDs and 192 GB. Accelerator partitioning can +instead expose DPX (four XCDs, 96 GB), QPX (two XCDs, 48 GB), or CPX (one XCD, +24 GB) logical devices. Memory partitioning is independent. Detect the visible +configuration and do not assume CPX, SPX, or a particular NPS mode. + +Tune grid sizing and persistent-work scheduling against the visible CU count. +Use 192 GB only as a physical-SPX upper bound; leave headroom for the runtime and +other allocations. + +References: + +- [AMD Instinct MI300X accelerator specifications](https://www.amd.com/en/products/accelerators/instinct/mi300/mi300x.html) +- [ROCm MI300-series workload optimization](https://rocm.docs.amd.com/en/latest/how-to/rocm-for-ai/inference-optimization/workload.html) diff --git a/src/prompts/cheatsheet/MI300_series_profile.md b/src/prompts/cheatsheet/MI300_series_profile.md new file mode 100644 index 00000000..29d9b38a --- /dev/null +++ b/src/prompts/cheatsheet/MI300_series_profile.md @@ -0,0 +1,21 @@ +# Generic AMD Instinct MI300-Series Profile + +`target_gpu_model: MI300` is the backward-compatible generic `gfx942` target. +It does not identify a physical SKU. Do not assume one model's topology or HBM +limits until the visible device has been inspected. + +| Model | XCDs | Active CUs | HBM | Peak HBM bandwidth | +| --- | ---: | ---: | ---: | ---: | +| MI300X | 8 | 304 | 192 GB HBM3 | 5.3 TB/s | +| MI325X | 8 | 304 | 256 GB HBM3E | 6.0 TB/s | +| MI300A | 6 | 228 | 128 GB shared HBM3 | approximately 5.3 TB/s | + +Identify the model from `torch.cuda.get_device_name()`, HIP device properties, +or `amd-smi static`, then apply the matching resource limits. For reproducible +model-specific tuning, prefer `target_gpu_model: MI300X`, `MI325X`, or `MI300A` +so that AgentKernelArena loads the exact profile. + +References: + +- [ROCm accelerator hardware specifications](https://rocm.docs.amd.com/en/latest/reference/gpu-arch-specs.html) +- [ROCm MI300-series workload optimization](https://rocm.docs.amd.com/en/latest/how-to/rocm-for-ai/inference-optimization/workload.html) diff --git a/src/prompts/cheatsheet/MI325X_profile.md b/src/prompts/cheatsheet/MI325X_profile.md new file mode 100644 index 00000000..186b7d68 --- /dev/null +++ b/src/prompts/cheatsheet/MI325X_profile.md @@ -0,0 +1,25 @@ +# AMD Instinct MI325X Model Profile + +## Physical resources + +- Architecture target: CDNA 3, `gfx942` +- Compute topology: 8 XCDs, 304 active CUs (38 active CUs per XCD) +- Aggregate L2: 32 MiB (4 MiB per XCD) +- Last-level Infinity Cache: 256 MiB +- Memory: 256 GB HBM3E, 6.0 TB/s peak theoretical bandwidth + +## Partition-aware tuning + +The full SPX device exposes eight XCDs and 256 GB. Accelerator partitioning can +instead expose DPX (four XCDs, 128 GB), QPX (two XCDs, 64 GB), or CPX (one XCD, +32 GB) logical devices. Memory partitioning is independent. Detect the visible +configuration and do not carry over MI300X's smaller per-partition memory limits. + +MI300X and MI325X have the same active-CU topology, but their memory technology, +capacity, and peak bandwidth differ. Re-benchmark memory-bound launch choices on +MI325X rather than treating MI300X results as interchangeable. + +References: + +- [AMD Instinct MI325X accelerator overview](https://www.amd.com/en/products/accelerators/instinct/mi300/mi325x.html) +- [ROCm MI300-series workload optimization](https://rocm.docs.amd.com/en/latest/how-to/rocm-for-ai/inference-optimization/workload.html) diff --git a/src/prompts/cheatsheet/default_cheatsheet.yaml b/src/prompts/cheatsheet/default_cheatsheet.yaml index c7963f80..405db9d7 100755 --- a/src/prompts/cheatsheet/default_cheatsheet.yaml +++ b/src/prompts/cheatsheet/default_cheatsheet.yaml @@ -2,20 +2,37 @@ # Cheatsheet Configuration # # Two independent dimensions: -# architecture: GPU-model → hardware spec document + gfx arch string +# architecture: GPU-model → hardware spec document(s) + gfx arch string # knowledge: target language → language best-practice guide (GPU-agnostic) # # prompt_builder.py combines BOTH for every task: -# final cheatsheet = architecture[gpu_model].file + knowledge[language] +# final cheatsheet = architecture[gpu_model].files + knowledge[language] +# +# ``file`` remains supported for architecture entries that only need one +# document. ``files`` composes shared architecture guidance with a model profile. # ============================================================ architecture: MI300: gfx_arch: gfx942 - file: src/prompts/cheatsheet/MI300X_architecture.md + files: + - src/prompts/cheatsheet/CDNA3_gfx942_architecture.md + - src/prompts/cheatsheet/MI300_series_profile.md MI300X: gfx_arch: gfx942 - file: src/prompts/cheatsheet/MI300X_architecture.md + files: + - src/prompts/cheatsheet/CDNA3_gfx942_architecture.md + - src/prompts/cheatsheet/MI300X_profile.md + MI325X: + gfx_arch: gfx942 + files: + - src/prompts/cheatsheet/CDNA3_gfx942_architecture.md + - src/prompts/cheatsheet/MI325X_profile.md + MI300A: + gfx_arch: gfx942 + files: + - src/prompts/cheatsheet/CDNA3_gfx942_architecture.md + - src/prompts/cheatsheet/MI300A_profile.md MI355X: gfx_arch: gfx950 file: src/prompts/cheatsheet/MI355X_architecture.md diff --git a/src/scripts/docker_benchmark.sh b/src/scripts/docker_benchmark.sh index fe601cb3..fca0d8f7 100755 --- a/src/scripts/docker_benchmark.sh +++ b/src/scripts/docker_benchmark.sh @@ -31,8 +31,8 @@ Usage: src/scripts/docker_benchmark.sh smoke Default run config: - example_configs/quickstart_claude_mi300.yaml (MI300/MI300X). - On another GPU, pass --config_name with a matching run configuration. + example_configs/quickstart_claude_mi300.yaml (generic MI300-series target). + Pass --config_name with a model-specific run configuration for exact tuning context. Environment overrides: GPU_IDS Comma/space separated GPU indices for parallel-run. diff --git a/tests/test_docker_benchmark.sh b/tests/test_docker_benchmark.sh index d68e2ffd..17ef70e0 100755 --- a/tests/test_docker_benchmark.sh +++ b/tests/test_docker_benchmark.sh @@ -156,7 +156,21 @@ assert_has "$NATIVE_CLAUDE_HOME/.claude:$NATIVE_CLAUDE_HOME/.claude" "${args[@]} assert_has "$NATIVE_CLAUDE_HOME/.claude.json:$NATIVE_CLAUDE_HOME/.claude.json" "${args[@]}" assert_has "claude_code" "${args[@]}" -# Omitting --config_name uses the one-task MI300/MI300X Claude quickstart. +# Every explicit MI300-series model resolves to the shared gfx942 runtime. +for gpu_model in MI300X MI325X MI300A; do + model_config="$TEST_HOME/${gpu_model,,}-config.yaml" + printf 'agent:\n template: claude_code\ntarget_gpu_model: %s\n' "$gpu_model" > "$model_config" + mapfile -t args < <( + env \ + HOME="$NATIVE_CLAUDE_HOME" \ + bash "$RUNNER" preflight --config_name "$model_config" 2>/dev/null + ) + assert_has "lmsysorg/sglang:v0.5.12-rocm720-mi30x" "${args[@]}" + assert_has "AGENT_KERNEL_ARENA_GPU_ARCH=gfx942" "${args[@]}" + assert_has "PYTORCH_ROCM_ARCH=gfx942" "${args[@]}" +done + +# Omitting --config_name uses the one-task generic MI300-series Claude quickstart. mapfile -t args < <( env \ HOME="$NATIVE_CLAUDE_HOME" \ diff --git a/tests/test_gpu_architecture.py b/tests/test_gpu_architecture.py new file mode 100644 index 00000000..933dab38 --- /dev/null +++ b/tests/test_gpu_architecture.py @@ -0,0 +1,106 @@ +import logging +import tempfile +import unittest +from pathlib import Path + +import yaml + +from src.preprocessing import _resolve_gfx_arch +from src.prompt_builder import _load_cheatsheet + + +PROJECT_ROOT = Path(__file__).resolve().parents[1] + + +class GPUArchitectureConfigTests(unittest.TestCase): + def test_all_mi300_models_resolve_to_gfx942(self) -> None: + for model in ("MI300", "MI300X", "MI325X", "MI300A"): + with self.subTest(model=model): + self.assertEqual(_resolve_gfx_arch(model), "gfx942") + + def test_exact_mi300_models_load_shared_and_model_specific_context(self) -> None: + expected_headings = { + "MI300X": "# AMD Instinct MI300X Model Profile", + "MI325X": "# AMD Instinct MI325X Model Profile", + "MI300A": "# AMD Instinct MI300A Model Profile", + } + + for model, profile_heading in expected_headings.items(): + with self.subTest(model=model): + prompt, gfx_arch = _load_cheatsheet( + "hip2hip", model, PROJECT_ROOT, {}, logging.getLogger(__name__) + ) + + self.assertEqual(gfx_arch, "gfx942") + self.assertIn("# AMD CDNA 3 (`gfx942`) Kernel Optimization Context", prompt) + self.assertIn(profile_heading, prompt) + self.assertLess( + prompt.index("# AMD CDNA 3 (`gfx942`) Kernel Optimization Context"), + prompt.index(profile_heading), + ) + + def test_generic_mi300_target_loads_non_sku_specific_profile(self) -> None: + prompt, gfx_arch = _load_cheatsheet( + "triton2triton", "MI300", PROJECT_ROOT, {}, logging.getLogger(__name__) + ) + + self.assertEqual(gfx_arch, "gfx942") + self.assertIn("# Generic AMD Instinct MI300-Series Profile", prompt) + self.assertNotIn("# AMD Instinct MI300X Model Profile", prompt) + + def test_architecture_only_context_omits_language_knowledge(self) -> None: + prompt, gfx_arch = _load_cheatsheet( + "triton2triton", + "MI300A", + PROJECT_ROOT, + {}, + logging.getLogger(__name__), + include_knowledge=False, + ) + + self.assertEqual(gfx_arch, "gfx942") + self.assertIn("# AMD CDNA 3 (`gfx942`) Kernel Optimization Context", prompt) + self.assertIn("# AMD Instinct MI300A Model Profile", prompt) + self.assertNotIn("# Triton Kernel Best Practices", prompt) + + def test_legacy_file_and_composed_files_are_both_supported(self) -> None: + with tempfile.TemporaryDirectory() as temporary_directory: + project_root = Path(temporary_directory) + config_dir = project_root / "src" / "prompts" / "cheatsheet" + config_dir.mkdir(parents=True) + (project_root / "legacy.md").write_text("legacy architecture") + (project_root / "shared.md").write_text("shared architecture") + (project_root / "profile.md").write_text("model profile") + (project_root / "hip.md").write_text("hip knowledge") + config = { + "architecture": { + "LEGACY": {"gfx_arch": "gfx900", "file": "legacy.md"}, + "COMPOSED": { + "gfx_arch": "gfx942", + "files": ["shared.md", "profile.md"], + }, + }, + "knowledge": {"hip": "hip.md"}, + } + (config_dir / "default_cheatsheet.yaml").write_text( + yaml.safe_dump(config), encoding="utf-8" + ) + + legacy_prompt, legacy_arch = _load_cheatsheet( + "hip2hip", "LEGACY", project_root, {}, logging.getLogger(__name__) + ) + composed_prompt, composed_arch = _load_cheatsheet( + "hip2hip", "COMPOSED", project_root, {}, logging.getLogger(__name__) + ) + + self.assertEqual(legacy_arch, "gfx900") + self.assertEqual(legacy_prompt, "legacy architecture\n\n---\n\nhip knowledge") + self.assertEqual(composed_arch, "gfx942") + self.assertEqual( + composed_prompt, + "shared architecture\n\n---\n\nmodel profile\n\n---\n\nhip knowledge", + ) + + +if __name__ == "__main__": + unittest.main()