[hipDNN] Forward MIOpen's tuning policy through the MIOpen provider (suggestion: a miopen.tuning_policy knob) #14372

Problem

MIOpen's hipDNN shim forwards conv and batchnorm calls to the MIOpen provider, which runs MIOpen on a handle of its own. The tuning policy set on the caller's handle (miopenSetTuningPolicy) never reaches that handle. The provider can express only two policies, through global.benchmarking:

DbUpdate, SearchDbUpdate and DbClean can't be expressed. So today:

• Batchnorm (spatial training and backward, the tunable solvers) serves None and Search and declines the other three (feat(miopen): forward batchnorm entry points to hipDNN behind MIOPEN_ENABLE_HIPDNN_WRAPPER #14362).
• Convolution ignores the caller's policy entirely and always runs under None, with no error.

MIOPEN_FIND_ENFORCE isn't affected because it is process-wide.

Suggested approach

This is a suggestion; the knob API is the hipDNN team's call. One option is a provider engine knob miopen.tuning_policy (int, range 0–5):

• 0 (default): follow global.benchmarking, so current behaviour is unchanged.
• 1–5: run under that miopenTuningPolicy_t. These values equal FindEnforceAction.

Precedence could be: explicit miopen.tuning_policy > HIPDNN_FORCE_BENCHMARKING > global.benchmarking.

Things to consider:

• Advertise the knob on every engine with tunable solvers: the batchnorm engines as well as the conv engines.
• create_execution_plan_ext drops knobs an engine doesn't advertise, so callers must check EngineConfigInfo::knobs before relying on it.
• Conv workspace sizing has to cover the Search policies (2/3/4).
• DbUpdate and DbClean write to the user's perf-db during plan creation, so plans built under different policies shouldn't share a cache entry.
• The name follows RFC 0004 §4.5 (<plugin-namespace>.<knob>).

Alternatives

• Widen benchmarking to an int. This means one knob and no precedence rule. But benchmarking is a global setting that every provider reads, and values like DbClean only mean something to MIOpen. That's why a miopen.* knob is suggested above.
• Decline instead of forwarding. Conv could decline policies it can't serve, as batchnorm already does. That's cheaper, and may be enough if no caller uses policies 3–5. It would still fix the silent mismatch in conv.

Happy to go with whichever the team prefers.

Follow-up in the MIOpen shim

• Batchnorm: pass the policy through PlanKey instead of declining. Only the two shared cores (training and backward) change. Inference and per-activation solvers aren't tunable.
• Convolution: read the handle's policy into PlanKey, or decline unsupported policies.

Out of scope

The per-descriptor conv find mode (miopenSetConvolutionFindMode) is a separate, conv-only gap. Batchnorm has no find mode: it picks the first applicable solver and doesn't use the find-db. The MIOPEN_FIND_MODE env var is process-wide, so the provider already sees it.

Related

• [hipDNN] Defaulting benchmarking=false affects the find mode of regular convolution interfaces. #14096: MIOPEN_FIND_MODE under the provider (closed, not reproduced).
• ALMIOPEN-2847: honor or decline the conv settings hipDNN can't carry. The conv side of this issue fits there as a sixth setting.
• ALMIOPEN-2176: env-var translation under forwarding (covers MIOPEN_FIND_ENFORCE).