# GPU Goblin — curated ROCm / MI300X optimization rules. # # Each entry is validated against agent.schemas.Rule at module import. # Rules are hand-curated from ROCm docs + AMD blog posts. NO LLM-generated rules. # This is the moat — every recommendation in the final report is grounded here. # # Ordering note: file order is irrelevant; query_rocm_kb.py ranks by cosine # similarity between query symptom and rule.symptom embeddings. # --------------------------------------------------------------------------- # precision # --------------------------------------------------------------------------- - id: precision.bf16_over_fp16_on_mi300x category: precision targets_bucket: precision_path symptom: >- fp16 used on MI300X / CDNA3 matrix cores. Loss can NaN on long fine-tuning runs and there is no throughput benefit over bf16 — the reflexive NVIDIA default carries forward without thinking. detect: precision: fp16 transform: precision: bf16 expected_recovery_fraction: 0.85 expected_impact: >- MI300X CDNA3 matrix cores execute bf16 at the same throughput as fp16 with strictly better numerical stability. Switch removes NaN risk in long runs; no loss-scaling required. rocm_version_min: "6.0" citation: "ROCm MI300X Optimization Guide §3.2 — bf16 vs fp16" - id: precision.fp32_default_wastes_matrix_cores category: precision targets_bucket: precision_path symptom: >- Training in fp32 on MI300X. The CDNA3 matrix cores deliver ~2x throughput in bf16 vs fp32 and the LoRA fine-tuning loss surface tolerates bf16 cleanly — fp32 is simply leaving silicon idle. detect: precision: fp32 transform: precision: bf16 expected_recovery_fraction: 0.6 expected_impact: >- Switching fp32 → bf16 nets roughly 2x matmul throughput on MI300X and halves activation memory, which usually unlocks a larger batch size as a follow-on win. rocm_version_min: "6.0" citation: "AMD MI300X Tuning Guide — Matrix-core throughput by precision" # --------------------------------------------------------------------------- # attention # --------------------------------------------------------------------------- - id: attention.flash_rocm_over_eager category: attention targets_bucket: kernel_shape symptom: >- Naive (eager) attention on MI300X — no flash kernel loaded, attention materializes the full O(seq_len^2) matrix. Most MI300X waste at long context lives here. detect: attention_impl: eager transform: attention_impl: flash_rocm expected_recovery_fraction: 0.7 expected_impact: >- Use the ROCm-validated flash-attention kernel (Optimum-AMD ships a tested build). Eliminates the O(seq_len^2) attention memory blow-up; typically 2-3x faster on MI300X for seq_len >= 1024. rocm_version_min: "6.0" citation: "AMD Optimum-AMD docs — Flash Attention 2 validated on MI300" - id: attention.sdpa_over_eager category: attention targets_bucket: kernel_shape symptom: >- Eager attention used when a drop-in faster path exists. PyTorch SDPA on ROCm dispatches to a fused kernel without requiring the flash-attn fork install — the lowest-friction attention upgrade available. detect: attention_impl: eager transform: attention_impl: sdpa expected_recovery_fraction: 0.45 expected_impact: >- PyTorch's torch.nn.functional.scaled_dot_product_attention on ROCm is a fused, memory-efficient backend with no extra dependency. Slower than flash-attn ROCm but a one-line config change with no install pain. rocm_version_min: "6.0" citation: "PyTorch SDPA backend selection — ROCm support notes" # --------------------------------------------------------------------------- # memory # --------------------------------------------------------------------------- - id: memory.batch_too_small_for_192gb category: memory targets_bucket: memory_headroom symptom: >- HBM peak is far below MI300X's 192 GB ceiling — the user copied a batch size from an A100 80GB tutorial and is leaving more than half the HBM idle every step. detect: batch_size: {"$lt": 8} transform: batch_size: 12 expected_recovery_fraction: 0.55 expected_impact: >- Raising batch size to fill more of the 192 GB HBM increases tokens/sec proportionally for compute-bound LoRA workloads. Adjust grad-accum down to keep effective batch unchanged if needed. rocm_version_min: "6.0" citation: "AMD Instinct MI300X — 192 GB HBM3 fine-tuning playbook" - id: memory.gradient_checkpointing_for_long_seq category: memory targets_bucket: memory_headroom symptom: >- Long sequences (seq_len >= 4096) without gradient checkpointing are spending HBM on activations that could be recomputed instead — leaving no headroom to grow the batch. detect: seq_len: {"$gte": 4096} gradient_checkpointing: false transform: gradient_checkpointing: true expected_recovery_fraction: 0.4 expected_impact: >- Enables freeing activation memory to enlarge batch size at long context; typical net win 1.3-1.7x throughput at seq_len 4k-8k on MI300X despite the recompute tax. rocm_version_min: "6.0" citation: "HF Transformers gradient_checkpointing + AMD long-context guide" # --------------------------------------------------------------------------- # data # --------------------------------------------------------------------------- - id: data.dataloader_workers_zero category: data targets_bucket: data_wait symptom: >- DataLoader num_workers=0 — the GPU stalls while the main process tokenizes the next batch. Classic MI300X starvation pattern at high-throughput LoRA fine-tuning. detect: dataloader_workers: 0 transform: dataloader_workers: 8 expected_recovery_fraction: 0.8 expected_impact: >- Spinning up 4-8 worker processes hides tokenization and host transfer behind GPU compute. On MI300X with a fast NVMe dataset, this often nets 2x tokens/sec on its own. rocm_version_min: "6.0" citation: "PyTorch DataLoader best practices + MI300X fine-tuning notes" - id: data.pin_memory_false category: data targets_bucket: data_wait symptom: >- DataLoader pin_memory=False on a CUDA/ROCm host. Each batch copy from pageable host memory blocks the launch queue instead of overlapping with compute. detect: dataloader_pin_memory: false transform: dataloader_pin_memory: true expected_recovery_fraction: 0.5 expected_impact: >- pin_memory=True allows asynchronous host-to-device transfer that overlaps with the previous step's compute. Roughly free win, no other config changes needed. rocm_version_min: "6.0" citation: "PyTorch CUDA/ROCm DataLoader pin_memory documentation" - id: data.prefetch_factor_default category: data targets_bucket: data_wait symptom: >- DataLoader prefetch_factor left at the default of 2 while MI300X is fast enough to drain the queue and stall before the next batch is ready. Symptom: GPU util oscillates 100% → 30% in a sawtooth. detect: dataloader_prefetch_factor: {"$lte": 2} transform: dataloader_prefetch_factor: 4 expected_recovery_fraction: 0.3 expected_impact: >- Raising prefetch_factor to 4 keeps two extra batches buffered per worker, smoothing out the GPU-util sawtooth. Costs a few hundred MB of host RAM per worker; trivial on a typical MI300X box. rocm_version_min: "6.0" citation: "PyTorch DataLoader prefetch_factor — long-step jitter notes" - id: data.persistent_workers_false category: data targets_bucket: data_wait symptom: >- DataLoader persistent_workers=False — workers are torn down and respawned every epoch. On large datasets the warmup tax bleeds time between epochs even when the per-epoch throughput is good. detect: dataloader_persistent_workers: false transform: dataloader_persistent_workers: true expected_recovery_fraction: 0.25 expected_impact: >- Keeps DataLoader workers alive across epochs so the dataset doesn't re-initialize. Removes a 5-30 second hitch at every epoch boundary on Qwen2.5 LoRA workloads. rocm_version_min: "6.0" citation: "PyTorch DataLoader persistent_workers documentation" # --------------------------------------------------------------------------- # compile # --------------------------------------------------------------------------- - id: compile.torch_compile_off category: compile targets_bucket: host_gap symptom: >- torch.compile is disabled but the model is on the eager-mode-friendly list (Qwen, Llama, Mistral, Gemma). Eager dispatch leaves host-gap bubbles between MI300X kernels that compile would erase. detect: torch_compile: false transform: torch_compile: true expected_recovery_fraction: 0.6 expected_impact: >- torch.compile fuses small ops and removes Python dispatch overhead; typical 1.1-1.3x throughput on Qwen-class fine-tuning on MI300X. Verify no graph breaks in your forward by running with TORCH_LOGS=graph_breaks. rocm_version_min: "6.0" citation: "PyTorch torch.compile + ROCm support matrix" # --------------------------------------------------------------------------- # env_vars # --------------------------------------------------------------------------- - id: env.nccl_min_nchannels category: env_vars targets_bucket: comm_excess symptom: >- NCCL_MIN_NCHANNELS not set on a collective-heavy training run. RCCL on MI300X benefits from a larger channel floor than the upstream NCCL default to saturate the XGMI mesh. detect: {} transform: env_vars.NCCL_MIN_NCHANNELS: "112" expected_recovery_fraction: 0.35 expected_impact: >- Setting NCCL_MIN_NCHANNELS=112 forces RCCL to use enough channels to saturate MI300X XGMI bandwidth on all-reduce. Helps anywhere RCCL/NCCL appears in the top kernels list. rocm_version_min: "6.0" citation: "AMD ROCm RCCL tuning guide — NCCL_MIN_NCHANNELS recommendation" - id: env.numa_auto_balancing_disable category: env_vars targets_bucket: host_gap symptom: >- Linux NUMA auto-balancing left enabled — the kernel migrates pinned pages between NUMA nodes during the run, causing variance and occasional multi-second stalls in benchmark numbers. detect: {} transform: env_vars.GOBLIN_HINT_NUMA_AUTO_BALANCING: "disable_via_sysctl" expected_recovery_fraction: 0.15 expected_impact: >- Disable with `echo 0 | sudo tee /proc/sys/kernel/numa_balancing` before benchmarking. Removes a major source of step-time variance on MI300X multi-socket boxes; required for reproducible before/after numbers. rocm_version_min: "6.0" citation: "AMD ROCm performance tuning — NUMA auto-balancing disable" - id: env.hsa_force_fine_grain_pcie category: env_vars targets_bucket: host_gap symptom: >- HSA_FORCE_FINE_GRAIN_PCIE not set. Fine-grained PCIe transfers reduce host→device latency for the small control packets that bracket every kernel launch. detect: {} transform: env_vars.HSA_FORCE_FINE_GRAIN_PCIE: "1" expected_recovery_fraction: 0.15 expected_impact: >- Enables fine-grained PCIe coherent traffic on MI300X. Small but consistent host-gap reduction; no downside on supported platforms. rocm_version_min: "6.0" citation: "AMD ROCm env var reference — HSA_FORCE_FINE_GRAIN_PCIE" # --------------------------------------------------------------------------- # kernels # --------------------------------------------------------------------------- - id: kernels.hipblaslt_hint_logging category: kernels targets_bucket: kernel_shape symptom: >- Stable model + batch + seq shapes with hipBLASLt heuristic logging disabled. Each step is solving a tile-selection problem the library already solved; an offline-tuned hint file would eliminate the search. detect: {} transform: env_vars.HIPBLASLT_TUNING_FILE: "/tmp/hipblaslt.tuning" env_vars.HIPBLASLT_LOG_LEVEL: "1" expected_recovery_fraction: 0.3 expected_impact: >- Run once with HIPBLASLT_LOG_LEVEL=1 to capture the best tile heuristic per GEMM shape, then use HIPBLASLT_TUNING_FILE to reuse it. Typical win: 5-15% on long-running fine-tunes with stable shapes. rocm_version_min: "6.0" citation: "AMD hipBLASLt tuning + offline-hint workflow" - id: kernels.miopen_find_mode_2 category: kernels targets_bucket: kernel_shape symptom: >- MIOPEN_FIND_MODE not set. MIOpen falls back to a slow exhaustive search on the first iteration of every new shape; with FIND_MODE=2 the cached fast solution is reused immediately. detect: {} transform: env_vars.MIOPEN_FIND_MODE: "2" expected_recovery_fraction: 0.25 expected_impact: >- MIOPEN_FIND_MODE=2 (FAST + DB lookup) eliminates the cold-start search pause on conv/attention layers. First-step time drops dramatically; steady-state perf is unchanged but warmup is cheaper. rocm_version_min: "6.0" citation: "AMD MIOpen documentation — MIOPEN_FIND_MODE values" # --------------------------------------------------------------------------- # optimizer (warning-only rule) # --------------------------------------------------------------------------- - id: optimizer.bitsandbytes_not_supported_warning category: optimizer targets_bucket: kernel_shape symptom: >- bitsandbytes 8-bit Adam variant configured. bitsandbytes is NOT officially supported on ROCm — it may import on a community build but correctness on MI300X is unverified, and AMD-validated 8-bit/quant paths exist instead. detect: optimizer: {"$regex": "(?i)(bitsandbytes|bnb_)"} transform: {} expected_recovery_fraction: 0.0 expected_impact: >- WARNING: bitsandbytes is not officially supported on ROCm. Recommended alternatives are Optimum-AMD-validated paths (GPTQ/AWQ for inference, 8-bit Adam via Apex on ROCm, or DeepSpeed ZeRO-Offload). Do not assume correctness on MI300X without your own numerical check. rocm_version_min: "6.0" citation: "AMD Optimum-AMD compatibility matrix — bitsandbytes status on ROCm" # --------------------------------------------------------------------------- # collectives # --------------------------------------------------------------------------- - id: collectives.one_process_per_gpu category: collectives targets_bucket: comm_excess symptom: >- Single launcher process driving multiple MI300X GPUs. ROCm serializes kernel launches across devices when one process owns multiple HIP streams — every collective becomes a bottleneck. detect: extras.launch_num_processes: 1 extras.world_size: {"$gt": 1} transform: extras.launch_num_processes: 8 expected_recovery_fraction: 0.6 expected_impact: >- Use `accelerate launch --num_processes=N` (or `torchrun --nproc_per_node=N`) so each MI300X has its own driver process. Removes the launch-queue serialization; most multi-GPU "ROCm is slow" reports trace back to this. rocm_version_min: "6.0" citation: "AMD ROCm distributed training — one-process-per-GPU best practice" # --------------------------------------------------------------------------- # topology # --------------------------------------------------------------------------- - id: topology.tp_within_xgmi_island category: topology targets_bucket: comm_excess symptom: >- Tensor parallelism enabled (TP > 1) without pinning ranks to an XGMI island. TP all-reduces hop across the slow inter-island link instead of the fast intra-island XGMI mesh. detect: extras.tensor_parallel_size: {"$gt": 1} transform: extras.tp_placement_hint: "xgmi_island" expected_recovery_fraction: 0.45 expected_impact: >- Restrict TP rank placement to a single 4-GPU XGMI island per group so all-reduces stay on intra-island links. Use ROCR_VISIBLE_DEVICES to pin; prefer TP over EP on a single MI300X node. rocm_version_min: "6.0" citation: "AMD MI300X system topology — XGMI islands and TP placement" # --------------------------------------------------------------------------- # Bonus high-value rules to reach the 20-25 target # --------------------------------------------------------------------------- - id: precision.fp8_inference_stretch category: precision targets_bucket: precision_path symptom: >- Workload is inference / batched eval on MI300X without FP8. CDNA3 has native FP8 matrix-core throughput that doubles peak vs bf16 — applicable once the model has been calibrated. detect: extras.workload_kind: "inference" transform: precision: fp8 expected_recovery_fraction: 0.4 expected_impact: >- Switch calibrated inference to FP8 (E4M3 weight, E5M2 grad if needed). Doubles matmul peak throughput on MI300X. Requires offline calibration; verify accuracy on your eval set first. rocm_version_min: "6.0" citation: "AMD MI300X inference guide — FP8 calibration workflow" - id: kernels.hipblaslt_log_level_hint category: kernels targets_bucket: kernel_shape symptom: >- Top kernels show a long tail of small unfused GEMMs. Without HIPBLASLT_LOG_MASK enabled the user can't see which heuristics were selected and won't notice when the library falls back to a slow path. detect: {} transform: env_vars.HIPBLASLT_LOG_MASK: "32" expected_recovery_fraction: 0.1 expected_impact: >- HIPBLASLT_LOG_MASK=32 enables heuristic-selection logging. Diagnostic only — does not change perf — but lets you confirm hipBLASLt picked the expected tile and switch to the offline tuning workflow when it didn't. rocm_version_min: "6.0" citation: "AMD hipBLASLt logging environment variables" - id: env.miopen_user_db_path category: env_vars targets_bucket: kernel_shape symptom: >- Multiple training runs share a host but MIOpen's per-user perf DB is on the default ~/.config path. First run on a new host eats the autotune cost every time even though a previous run already solved the same shapes. detect: {} transform: env_vars.MIOPEN_USER_DB_PATH: "/var/cache/miopen" env_vars.MIOPEN_CUSTOM_CACHE_DIR: "/var/cache/miopen" expected_recovery_fraction: 0.1 expected_impact: >- Point MIOPEN_USER_DB_PATH at a shared host directory so autotune results survive container restarts and are reused across runs. Eliminates first-step warm-up cost after the first build. rocm_version_min: "6.0" citation: "AMD MIOpen perf-db environment variables"