MIOpen exhaustive benchmarking fails on gfx1201 (RX 9070) with workspace allocation returning null (ptr: 0 size: 0)

[Xarianne]

MIOpen exhaustive benchmarking fails on gfx1201 (RX 9070) with workspace allocation returning null (ptr: 0 size: 0)

Environment

• GPU: AMD Radeon RX 9070 (gfx1201)
• CPU: Ryzen 7 9800X3D
• Motherboard: Gigabyte B650M DS3H
• OS: Fedora Kinoite (Linux 7.0.14-201.fc44.x86_64)
• ROCm: 7.2
• PyTorch: 2.12.1+rocm7.2
• MIOpen: bundled with ROCm 7.2

Context

This issue was discovered while running Voicebox (an open-source TTS application, https://github.com/jamiepine/voicebox) in a Docker container with ROCm PyTorch. The inference workload involves Qwen3-TTS (Qwen CustomVoice 1.7B) model, which uses 1D convolution and matrix multiplication operations that trigger MIOpen's kernel selection.

Problem

MIOpen's default exhaustive benchmarking mode (MIOPEN_FIND_MODE=DEFAULT) fails on gfx1201. During inference, MIOpen tries multiple kernel variants for matrix multiplication operations (GemmFwdRest) but every variant fails to allocate workspace memory, returning a null pointer with size 0:

MIOpen(HIP): Warning [IsEnoughWorkspace] [GetSolutionsFallback WTI] Solver <GemmFwdRest>, workspace required: 79994880, provided ptr: 0 size: 0
MIOpen(HIP): Warning [IsEnoughWorkspace] [EvaluateInvokers] Solver <GemmFwdRest>, workspace required: 79994880, provided ptr: 0 size: 0
MIOpen(HIP): Warning [IsEnoughWorkspace] [GetSolutionsFallback WTI] Solver <GemmFwdRest>, workspace required: 239984640, provided ptr: 0 size: 0
MIOpen(HIP): Warning [IsEnoughWorkspace] [EvaluateInvokers] Solver <GemmFwdRest>, workspace required: 239984640, provided ptr: 0 size: 0
MIOpen(HIP): Warning [IsEnoughWorkspace] [GetSolutionsFallback WTI] Solver <GemmFwdRest>, workspace required: 516096000, provided ptr: 0 size: 0
MIOpen(HIP): Warning [IsEnoughWorkspace] [EvaluateInvokers] Solver <GemmFwdRest>, workspace required: 516096000, provided ptr: 0 size: 0

The workspace size requested grows with each generation (80MB, 240MB, 516MB, 774MB), suggesting different tensor shapes trigger new failed allocations each time.

Symptoms

1. System stuttering during inference, MIOpen thrashes through failing kernels on every generation, stalling the GPU
2. Repeated warnings on every generation, the cache is being written (gfx1201 kernel database files exist) but MIOpen still re-benchmarks every call instead of using cached results
3. Without explicit cache path environment variables (MIOPEN_USER_DB_PATH, MIOPEN_CUSTOM_CACHE_DIR), the error escalates to RuntimeError: miopenStatusUnknownError

Workaround

Setting MIOPEN_FIND_MODE=FAST (equivalent to MIOPEN_FIND_MODE=1) resolves the issue completely. Fast mode uses heuristics to select kernels instead of exhaustive benchmarking, avoiding the failing workspace allocation path.

Without the cache path environment variables, miopenStatusUnknownError still occurs even with MIOPEN_FIND_MODE=FAST. Both fixes are needed together.

Context

• The GPU is natively supported by ROCm 7.2 (no HSA_OVERRIDE_GFX_VERSION needed). GPU detection and basic CUDA operations work correctly.
• This appears to be the same IsEnoughWorkspace / GemmFwdRest warning pattern as Excessive warning messages on workspace provided and required (IsEnoughWorkspace) after #2947 #2981, which was closed and reportedly fixed by PR Fix excessive IsEnoughWorkspace warning messages after #2947 #2986. The fix does not appear to cover gfx1201 / RDNA4.
• Possibly related: Workspace is zero bytes for TTS workloads #3650 (workspace zero bytes on gfx1100 / RDNA3, resolved by tuning performance database), same class of symptom on a different architecture

Questions

1. Is exhaustive benchmarking expected to work on gfx1201 with ROCm 7.2, or is this a known gap?
2. Could the ptr: 0 size: 0 workspace allocation be related to IOMMU behavior on AM5 platforms?
3. Are there plans to include gfx1201 performance data in the default MIOpen kernel database so exhaustive mode has heuristics to fall back on?