dtourolle/scene-actor-extraction
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

perf(rocm): enable MIOpen exhaustive search and TunableOp

Browse Source 
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.
This commit is contained in:
dtourolle
2026-08-09 10:22:07 +02:00
parent 84513d3fa7
commit c374c262f5
1 changed files with 14 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/backends/ort_provider.hpp
+14 -1
View File
@@ -124,8 +124,21 @@ inline OrtProvider apply_ort_provider(Ort::SessionOptions& opts,
try { try {
OrtROCMProviderOptions rocm{}; OrtROCMProviderOptions rocm{};
rocm.device_id = 0; 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); opts.AppendExecutionProvider_ROCM(rocm);
std::cerr << "[" << label << "] ROCm provider\n"; std::cerr << "[" << label << "] ROCm provider"
<< (tune ? " (MIOpen exhaustive + TunableOp)" : " (untuned)")
<< "\n";
return OrtProvider::ROCm; return OrtProvider::ROCm;
} catch (const Ort::Exception& e) { } catch (const Ort::Exception& e) {
std::cerr << "[" << label << "] ROCm unavailable (" std::cerr << "[" << label << "] ROCm unavailable ("