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 ("