From c374c262f57547b69fd3597d7111b113b7469bab Mon Sep 17 00:00:00 2001
From: Duncan Tourolle <duncan@tourolle.paris>
Date: Sun, 9 Aug 2026 10:22:07 +0200
Subject: [PATCH] perf(rocm): enable MIOpen exhaustive search and TunableOp

The ROCm execution provider was appended with only device_id set, leaving
MIOpen on its no-workspace convolution fallback (the "GemmFwdRest,
provided ptr: 0 size: 0" warnings) and TunableOp off. Enable
miopen_conv_exhaustive_search and tunable_op_enable/tuning so MIOpen picks
the fast conv kernels and the GEMMs autotune; both cache to the MIOpen
user DB (MIOPEN_USER_DB_PATH), so the tuning cost is paid once per shape.
This helps the 2D-conv models that actually run on the ROCm GPU (SCRFD,
ArcFace). Opt out with SAE_ROCM_NOTUNE=1 for a quick no-warmup run.
---
 src/backends/ort_provider.hpp | 15 ++++++++++++++-
 1 file changed, 14 insertions(+), 1 deletion(-)

diff --git a/src/backends/ort_provider.hpp b/src/backends/ort_provider.hpp
index 7f27c26..dd6d8b6 100644
--- a/src/backends/ort_provider.hpp
+++ b/src/backends/ort_provider.hpp
@@ -124,8 +124,21 @@ inline OrtProvider apply_ort_provider(Ort::SessionOptions& opts,
         try {
             OrtROCMProviderOptions rocm{};
             rocm.device_id = 0;
+            // Without these MIOpen runs convolutions on the no-workspace GEMM
+            // fallback (the "GemmFwdRest, provided ptr: 0 size: 0" warnings), which
+            // is the slow path — most visible on the conv-heavy TransNetV2 scene
+            // detector. Exhaustive search lets MIOpen pick the fast conv kernel,
+            // and TunableOp autotunes the GEMMs; both cache to the MIOpen user DB
+            // (MIOPEN_USER_DB_PATH), so the tuning cost is paid once per shape.
+            // Opt-out via SAE_ROCM_NOTUNE=1 for a quick no-warmup run.
+            const bool tune = std::getenv("SAE_ROCM_NOTUNE") == nullptr;
+            rocm.miopen_conv_exhaustive_search = tune ? 1 : 0;
+            rocm.tunable_op_enable             = tune;
+            rocm.tunable_op_tuning_enable      = tune;
             opts.AppendExecutionProvider_ROCM(rocm);
-            std::cerr << "[" << label << "] ROCm provider\n";
+            std::cerr << "[" << label << "] ROCm provider"
+                      << (tune ? " (MIOpen exhaustive + TunableOp)" : " (untuned)")
+                      << "\n";
             return OrtProvider::ROCm;
         } catch (const Ort::Exception& e) {
             std::cerr << "[" << label << "] ROCm unavailable ("