AI-COMPASS

AI Compute Performance Analysis & Statistics Suite

A unified RDNA1–4 AI compute performance toolset for Windows that traces every kernel dispatch, measures GPU utilization in real time, reads hardware performance counters, and generates actionable optimization reports — from model loading through prompt processing to token generation.

What It Does

AI-COMPASS instruments your AI inference workload end-to-end and answers three questions:

Where is the GPU spending its time? — per-kernel breakdown (MMQ, MMVQ, Attention, MoE, RoPE, Norm, etc.) across prompt processing and token generation phases.
What is limiting performance? — occupancy bottlenecks, low GPU utilization, specific kernels dominating the critical path.
Did my optimization help? — before/after comparison with per-category regression detection.

Example Output

Summary: 4,956 kernels | 68.88 ms total | GPU busy: 47.7%

Phases:
  Prompt Processing: 990 kernels, 41.21 ms (59.8%)
  Token Generation:  3966 kernels, 27.67 ms (40.2%)

Category Breakdown:
  MMVQ         848 kernels   27.3 ms   39.6%   occ:27.6%
  Vector      1332 kernels   20.2 ms   29.3%   occ:58.8%
  Attention   2472 kernels   12.8 ms   18.5%   occ:51.7%
  Norm          96 kernels    4.8 ms    6.9%   occ:100%
  Quantize     192 kernels    3.9 ms    5.6%   occ:95.8%

Bottlenecks:
  MMVQ dominates at 39.6% with low occupancy (27.6%)
  MMVQ has low occupancy — likely occupancy-bound

Optimization Targets:
  [MMVQ] Check Split-K heuristic and small_k path for RDNA4.
  [Attention] Consider flash attention tuning.

Reports are generated as both JSON (for scripting) and standalone HTML (for sharing).

How It Helps

Problem	AI-COMPASS Solution
"Why is inference slow?"	Per-kernel timing breakdown shows exact cost
"Is the GPU fully utilized?"	ADLX real-time GPU metrics (util %, clocks, power, temp, VRAM)
"Which kernel should I optimize?"	Sorted by time %, bounded by occupancy analysis
"Did my patch help?"	`--compare` mode diffs before vs after runs — measured +7.6% TG on Gemma-4 12B
"Where is the bottleneck?"	Automatic bottleneck detection per category
"Can I visualize this?"	RCV integration for GPU trace visualization

Session Results (Jul 2026)

All benchmarks on RX 9070 XT (gfx1201, 16GB VRAM, Wave32) with Ryzen 9 5900XT / 96GB DDR4.

RDNA4 Kernel Optimizations Committed

Change	File	Gemma-4 12B TG	Qwen3.5-9B NVFP4 TG
Baseline	—	58.3 t/s	103.0 t/s
nwarps=4 for K-quant	mmvq.cu	58.5 t/s	—
Thread mapping fix	mmvq.cu	62.7 t/s (+7.6%)	—
nwarps=2 for FP4	mmvq.cu	—	128.1 t/s (+24.4%)
CPU AVX2+FMA enabled	CMake	+23% on CPU TG	—
Best combined	—	62.7 t/s	128.1 t/s

Quantization Format Comparison

Model	Format	Size	PP256	TG128
Gemma-4 12B	Q4_K_M	6.86GB	3,648 t/s	62.7 t/s
Gemma-4 12B	IQ4_XS	6.22GB	4,510 t/s	68.4 t/s
Qwen3.5-9B	NVFP4	4.94GB	4,810 t/s	103.0 t/s
Qwen3-8B	Q4_K_M	4.68GB	5,854 t/s	99.1 t/s
Qwen3.5-4B	NVFP4	2.36GB	6,405 t/s	149.7 t/s
Laguna XS 33B MoE	IQ4_XS	16.84GB	698 t/s	69.0 t/s
Laguna XS 33B MoE	Q4_K_M	18.88GB	557 t/s	51.6 t/s
GLM-4.7-Flash-APEX (30B MoE)	Q6_K	17.88GB	728 t/s	41.5 t/s
Gemma-4-31B-Fable-5-Distill	Q4_K_M	17.39GB	262 t/s	6.3 t/s

ROCmFP4 Support (New)

ROCmFP4 (AMD-native FP4 format) is now buildable in the HIP backend:

Add GGML_TYPE_Q4_0_ROCMFP4 (type 100) and Q4_0_ROCMFP4_FAST (101)
Full vec_dot, dequantize, MMVQ dispatch, nwarps=2 tuning
Status: Builds clean, model loading crashes at runtime (HIP copy path needs completion)

HIP Shared Memory Fix (New)

Fixed shared memory corruption in HIP backend MMA kernel (mul_mat_q) when physical batch size > 16:

Changed smpbo from sharedMemPerBlock (~48KB) to sharedMemPerBlockOptin (~64KB+) for HIP builds
Implemented CUDA_SET_SHARED_MEMORY_LIMIT for HIP using hipFuncSetAttribute to raise kernel shared memory limit
Added HIP_CHECK macro for HIP error checking
Status: Fixed, verified with --ubatch-size 17 test (previously corrupted, now correct)

Unified Multi-Backend Toolkit (New)

Created unified PowerShell toolkit for building, testing, and validating AI inference across all backends:

toolkit-unified.ps1 — Single entry point for HIP, Vulkan, DirectX 12, and WinML
validate-system.ps1 — System validation for all backends
validate-model.ps1 — Cross-backend model validation and output comparison
toolkit-hip-vulkan.ps1 — HIP + Vulkan specific toolkit
toolkit-dx.ps1 — DirectX 12 specific toolkit with DRED and PIX support
Status: Complete, committed to AI-compass-DX branch

DirectX 12 Backend Findings (2026-07-28)

Metric	Value
GPU	AMD RX 9070 XT (gfx1201, 16GB VRAM, Wave32)
Driver	AMD 26.10 RC26 required for `test-backend-ops -o MUL_MAT` stability
WaveMMATier	Tier 10 (WaveMMA supported)
Cooperative Matrix	✅ Available (LinAlg tier 1.0)
UFM Integration	✅ Vulkan UFM (`UFM_FA=1`, `UFM_MULMAT=1`) works in parallel to DX12

Performance Comparison (acrux-500m Q6_K)

Backend	pp128	tg64
Vulkan	20,438 t/s	603 t/s
DX12	4,058 t/s	272 t/s

Performance Issue Found

Q6_K quant is NOT in the DXLA allowed quant types (dx12_select_gemm_path). Falls back to DX12_GEMM_STANDARD instead of DX12_GEMM_DXLA_WAVE. Fix needed: add Q6_K to the allowed quant list and add corresponding DXLA shader.

DX12 Optimal Switches (RX 9070 XT)

Switch	Value	Notes
`-ngl 99`	Full GPU offload	Default 40 may under-utilize
`--ubatch-size 32`	Batch size 32	Higher throughput than default 1-8
`--threads 8`	8 threads	Matches RDNA4 CU count (32) with threading

DX12 Tests Verified

Test	Result
`test_dx12_gemm`	✅ Pass
`test_dx12_e2e`	✅ Pass
`test_dx12_layer`	✅ Pass
`test_dx12_ops`	✅ Pass
`test_dx12_stability`	✅ Pass

Large Model CPU Inference

Model	Size	CPU TG	Bottleneck
Laguna 117B Q4_K	68GB	4.3 t/s	DDR4 BW (~50 GB/s)
Laguna XS 33B Q4_K	18.9GB	13.0 t/s	DDR4 BW (~50 GB/s)
MiniCPM 1B Q8_0	1.1GB	38.7 t/s	CPU compute

Quick Start

Prerequisites

Windows 10/11 with AMD Radeon RX 5000+ GPU (RDNA1–4)
ROCm 7.13+ (for HIP runtime)
Visual Studio 2022 + CMake + Ninja
Vulkan SDK 1.4+ (for GPUPerfAPI counter access)

Build

git clone https://github.com/Maxritz/AMD-AI-COMPASS.git
cd AI-COMPASS
mkdir build && cd build
cmake .. -G Ninja -DCMAKE_BUILD_TYPE=Release
cmake --build .

For the llama.cpp ROCm build (RDNA4 / RX 9070 XT):

scripts/llama_build_cpu_gpu.ps1

For RDNA2 (RX 6700 XT / ROCm 7.3):

scripts/llama_build_rdna2.ps1

Profile a Model

# Quick analysis from an existing HIP trace
python tools/analyze.py hip_trace.csv --output report/

# Full end-to-end benchmark (RDNA4 default)
python tools/run_benchmark.py --model path/to/model.gguf --pp 256 --tg 128

# RDNA2-specific analysis
python tools/analyze.py trace.csv --arch rdna2 --bench-pp 262 --bench-tg 6.3
python tools/run_benchmark.py --model model.gguf --arch rdna2 --pp 256 --tg 128

Unified Multi-Backend Testing Toolkit

AI-COMPASS includes a unified PowerShell toolkit for building, testing, and validating AI inference across HIP (ROCm), Vulkan, DirectX 12, and WinML backends.

# Validate system for all backends (GPU, drivers, SDKs, build tools)
.\scripts\validate-system.ps1 -Detailed

# Build all available backends
.\scripts\toolkit-unified.ps1 -Backend all -Action build

# Run coherence test on all backends (detects garbled output)
.\scripts\toolkit-unified.ps1 -Backend all -Action coherence -ModelPath "model.gguf" -Prompt "Hi"

# Full test suite across all backends
.\scripts\toolkit-unified.ps1 -Backend all -Action test -ModelPath "model.gguf"

# Debug mode with DRED (DirectX 12 crash debugging)
.\scripts\toolkit-unified.ps1 -Backend dx12 -Action debug -ModelPath "model.gguf" -EnableDRED

# GPU capture for frame analysis
.\scripts\toolkit-unified.ps1 -Backend all -Action capture -ModelPath "model.gguf"

# Cross-backend model validation (compares outputs across all backends)
.\scripts\validate-model.ps1 -ModelPath "model.gguf" -Prompt "Explain quantum computing"

Backend-Specific Toolkit Scripts

Script	Backend	Purpose
`scripts/toolkit-unified.ps1`	All (HIP, Vulkan, DX12, WinML)	Unified entry point for build/test/debug/capture
`scripts/toolkit-hip-vulkan.ps1`	HIP, Vulkan	HIP + Vulkan build, test, coherence, capture
`scripts/toolkit-dx.ps1`	DirectX 12	DX12 build, test, DRED, PIX capture
`scripts/validate-system.ps1`	All	System validation (GPU, SDKs, build tools)
`scripts/validate-model.ps1`	All	Cross-backend model validation and output comparison
`scripts/run-dx-ai.ps1`	DirectX 12	Generic DirectX AI application runner
`scripts/validate-dx-model.ps1`	DirectX 12	Model validation for DX12 backend

Toolkit Actions

Action	Description
`build`	Build the specified backend(s)
`test`	Run coherence tests (detect garbled output)
`debug`	Debug mode with DRED, debug layers, validation
`validate`	System validation (GPU, SDKs, tools)
`capture`	GPU frame capture (PIX, RenderDoc)
`profile`	Performance profiling
`coherence`	Cross-backend output comparison

Key Features

HIP/ROCm: Full ROCm 7.13 support with rocwmma v2+, RDNA4 optimizations, shared memory fixes
Vulkan: Vulkan SDK integration with GLSL compute shaders, RenderDoc capture
DirectX 12: DirectML, D3D12 debug layers, DRED crash debugging, PIX capture
WinML: Windows ML via DirectML, ONNX model support
Cross-backend coherence: Detects garbled output by comparing results across backends
Performance testing: Batch size testing (1, 16, 17, 32, 64) to identify corruption thresholds

Compare Before/After Optimization

# Run baseline
python tools/run_benchmark.py --model model.gguf --pp 256 --tg 128 -o baseline/

# Apply optimization, rebuild, run again
python tools/run_benchmark.py --model model.gguf --pp 256 --tg 128 -o optimized/ \
  --compare baseline/hip_trace.csv

Visualize with RCV

build/rocprof-compute-viewer.exe output_dir/

Components

Component	Purpose	RDNA4	RDNA3	RDNA2	RDNA1	Source
HIP Tracer	Kernel dispatch hooking via MinHook	✅	✅	✅	✅	Custom
GPUPerfAPI	GPU SQ/SPI/TCP/GL2C hardware counters	✅	✅	✅	✅	AMD GPUOpen
ADLX	Real-time GPU metrics (util %, clocks, power, temp, VRAM)	✅	✅	✅	✅	AMD ADLX SDK
rocprof-compute-viewer	GPU trace visualization (SQTT viewer)	✅	✅	✅	✅	ROCm
rocprofv3	Compute profiling CLI	✅	✅	✅	✅	ROCm
ROCmFP4	Native AMD FP4 format (HIP backend, WIP)	⚠	⚠	❌	❌	Ciru ROCmFPX
RDNA2 Optimizer	Wave64 + L2 + occupancy tuning recommendations	✅	✅	✅	❌	Custom plugin
GEAK	Agent-based workload optimization	✅	⚠	⚠	❌	Custom port
Hyperloom	Performance tuning framework	✅	⚠	⚠	❌	Custom port
AQLProfile	Low-level SQ/SX/TA/TD counter access	✅	✅	✅	⚠	ROCm
Magpie	Trace analysis	✅	✅	✅	⚠	Custom port
intellikit	AI toolkit	✅	✅	✅	⚠	Custom port
Unified Toolkit	PowerShell build/test/validate for HIP, Vulkan, DX12, WinML	✅	✅	✅	✅	Custom
Vulkan Backend	Vulkan compute shaders (143+ GLSL, 19,579 lines C++)	✅	✅	✅	✅	llama.cpp fork
DirectX 12 Backend	DirectML, D3D12, PIX capture, DRED debugging	✅	✅	✅	✅	llama.cpp fork
WinML Backend	Windows ML via DirectML, ONNX support	✅	✅	✅	✅	Custom integration

✅ Full support ⚠ Partial/WIP ❌ Not supported

GPU Architecture Support

Auto-detected via hipGetDeviceProperties with ADLX fallback:

Arch	GFX IP	Example GPU	Tested
RDNA1	gfx1010–1012	RX 5700 XT	❌
RDNA2	gfx1030–1035	RX 6700 XT	✅ separate system (ROCm 7.3, Win11, 5600X/48GB)
RDNA3	gfx1100–1103	RX 7900 XTX	❌
RDNA3.5	gfx1150–1151	RX 8800 XT	❌
RDNA4	gfx1200–1201	RX 9070 XT	✅ primary target

CLI Reference

aicompass <command> [options]

Commands:
  trace    <app> [args]   Capture HIP kernel trace + GPU metrics
  profile  <app> [args]   Full profile with all plugins
  analyze  <dir>          Post-process trace data
  list                    List installed plugins
  help                    Show this help

Options:
  -o, --output <dir>     Output directory
  -i, --interval <ms>    Metrics poll interval (default: 100)
  -v, --verbose          Verbose output
  --counters             Enable hardware counter collection

Python Analysis Tools

python tools/analyze.py <trace.csv> [options]

Options:
  -o, --output <dir>    Output directory for reports
  --compare <baseline>  Compare against a baseline trace
  --cu-count <n>        GPU CU count (default: from arch profile)
  --arch <arch>         GPU architecture: rdna1|rdna2|rdna3|rdna3_5|rdna4 (default: rdna4)
  --bench-pp <t/s>      Prompt processing throughput from llama-bench
  --bench-tg <t/s>      Token generation throughput from llama-bench
  --html                Generate HTML report (default: on)

Architecture

┌──────────────────────────────────────────────────────────┐
│                   AI-COMPASS CLI                         │
│        trace | profile | analyze | visualize              │
├──────────────────────────────────────────────────────────┤
│                                                          │
│  ┌──────────┐  ┌──────────┐  ┌──────────┐  ┌─────────┐ │
│  │HIP Tracer│  │GPUPerfAPI│  │ADLX      │  │Plugins  │ │
│  │(MinHook) │  │(VK/DX12) │  │(Metrics) │  │GEAK/etc │ │
│  └────┬─────┘  └────┬─────┘  └────┬─────┘  └────┬────┘ │
│       │             │             │             │       │
│       └─────────────┴─────────────┴─────────────┘       │
│                         │                                │
│                  ┌──────▼──────┐                         │
│                  │  Analysis   │                         │
│                  │  Engine     │                         │
│                  │  (analyze.py)│                        │
│                  └──────┬──────┘                         │
│                         │                                │
│            ┌────────────▼────────────┐                   │
│            │  Output Formats         │                   │
│            │  HTML | JSON | RCV      │                   │
│            └─────────────────────────┘                   │
└──────────────────────────────────────────────────────────┘

RDNA2 Optimization Guide

Key Differences from RDNA4

Property	RDNA2 (RX 6700 XT)	RDNA4 (RX 9070 XT)
GFX IP	gfx1031	gfx1201
CUs	40	32
Wave Size	64 (Wave64)	32 (Wave32)
L2 Cache	4 MB	12 MB
LDS/CU	64 KB	128 KB
Memory BW	384 GB/s (GDDR6)	~960 GB/s (GDDR6)
Max waves	2,560	1,024

Recommendations for RDNA2

Wave64 MMVQ thread mapping — change tid/16 to tid/32 in mmvq.cu for correct start-position calculation. Wave64 needs /32 since each wave has 64 threads.
Smaller L2 cache — 4 MB vs 12 MB means larger models (30B+) may be L2-bound. Use smaller K-tiles (MMQ_ITER_K=4) and prefer IQ4_XS over Q4_K_M for better cache utilization.
Higher occupancy headroom — 40 CUs × 4 SIMDs × 16 = 2,560 wave slots. Use larger block sizes (256+) to fully utilize available wave slots.
ROCm 7.3 compatibility — use -DAMDGPU_TARGETS=gfx1031 and verify hipGetDeviceProperties populates gcnArchName correctly. Set HSA_OVERRIDE_GFX_VERSION=10.3.0 if detection fails.
Build with: scripts/llama_build_rdna2.ps1 (sets gfx1031 target automatically).

Analysis Command

python tools/analyze.py trace.csv --arch rdna2 --cu-count 40
python tools/run_benchmark.py --model model.gguf --pp 256 --tg 128 --arch rdna2

The RDNA2 Optimizer plugin (list shows RDNA2_OPTIMIZER) auto-activates on gfx103x hardware.

RDNA4 Optimization Guide

Key Findings for RX 9070 XT

MMVQ is memory-bandwidth bound — K-quant types saturate VRAM bandwidth. TG improvement comes from reducing memory traffic (IQ4_XS, NVFP4) rather than increasing occupancy.
nwarps=2 (64 threads) is the sweet spot for FP4 — avoids load imbalance from idle threads on small K tiles while providing enough parallelism.
Thread mapping matters — fixing tid/16 to tid/8 in MMVQ eliminated 2x start-position collision within warps, gaining +7.6% TG.
IQ4_XS beats Q4_K_M — simpler dequantization + less data = +17% TG on Gemma-4 12B.
CPU AVX2 was OFF — the HIP-targeted build had all CPU SIMD disabled. Enabling AVX2+FMA gave +23% on CPU TG.
ROCmFP4 uses dp4a — not native WMMA yet on RDNA4. A true WMMA FP4 kernel would bypass dp4a entirely.

Credits & Acknowledgements

AI-COMPASS integrates and builds upon several open-source projects:

Core Dependencies

GPU Performance API (GPUPerfAPI) — AMD GPUOpen. MIT license. Hardware performance counter access library.
ADLX SDK — AMD GPUOpen. MIT license. AMD Display Library Next.
rocprof-compute-viewer — ROCm. MIT license. GPU trace visualization.
rocprofiler-sdk / rocprofv3 — ROCm. MIT license. ROCm profiler SDK.
ROCm — AMD. HIP runtime, ROCr, compiler toolchain.
ROCmFPX — Ciru / Charlie. Custom ROCmFP4 Vulkan runtime.
MinHook — Tsuda Kageyu. BSD 2-Clause. API hooking library.
Qt — The Qt Company. GPL/LGPL. UI framework for RCV.
Vulkan SDK — LunarG / Khronos. Vulkan API headers.

AI Toolkit Ports

GEAK-RDNA — Agent-based optimization framework, ported for RDNA4.
Hyperloom-RDNA — Performance tuning framework, ported for RDNA4.
Magpie — Trace analysis tool, ported for RDNA4.
intellikit — AI toolkit integration, ported for RDNA4.
Apex — Optimization toolkit, ported for RDNA4.
TraceLens — Trace analysis visualization, ported for RDNA4.

Our Work

All integration code — the HIP tracer DLL, GPUPerfAPI HIP adapter, ADLX metrics poller, plugin system, kernel classification engine, bottleneck detection, HTML report generator, MMVQ RDNA4 optimizations (nwarps tuning, thread mapping fix, FP4 support), and CLI framework — is original work developed for this project.

License

MIT — See LICENSE for details. Third-party components retain their original licenses.

Name		Name	Last commit message	Last commit date
Latest commit History 21 Commits
.github		.github
.graphify_scripts		.graphify_scripts
benchmark_runs		benchmark_runs
docs		docs
scripts		scripts
skills		skills
src		src
third_party		third_party
tools		tools
vendor		vendor
.gitignore		.gitignore
AI-COMPASS-v0.1.0.zip		AI-COMPASS-v0.1.0.zip
ARCHITECTURE.md		ARCHITECTURE.md
CMakeLists.txt		CMakeLists.txt
README.md		README.md
RESEARCH-TRACE-ANALYZE-OPTIMIZE.md		RESEARCH-TRACE-ANALYZE-OPTIMIZE.md
aicompass.py		aicompass.py
hip_tracer.log		hip_tracer.log
hip_tracer_func_dump.txt		hip_tracer_func_dump.txt
requirements.txt		requirements.txt

Folders and files

Latest commit

History

Repository files navigation