fix(MIOpen): correct deterministic SetNextValue polarity for grouped Bwd/Wrw CK solvers

[JonathanLichtnerAMD]

JIRA ID: ALMIOPEN-2359

Motivation

miopenFindConvolutionBackwardDataAlgorithm() (and the analogous BackwardWeights find)
could hang indefinitely on MI300X whenever three conditions coincided: TF32 data type,
the MIOPEN_CONVOLUTION_ATTRIB_DETERMINISTIC attribute set on the convolution, and an
exhaustive find with an empty perf-db cache. Root cause: an inverted return-value
polarity bug in the is_deterministic branch of SetNextValue() for
PerformanceConfigHipImplicitGemmGroupBwdXdlops and
PerformanceConfigHipImplicitGemmGroupWrwXdlops.

NextLinear() returns true on wraparound (search exhausted) and false on a normal
advance. The old code did if(!NextLinear(...)) return false;, i.e. it treated a
normal advance (!NextLinear == true) as "exhausted" and a wraparound
(NextLinear == true, so !NextLinear == false) as "keep going" — exactly
backwards. The practical fallout:

• With exactly one valid CK kernel candidate, the index space never wraps, so the
  branch never returned false and the enumeration spun forever (the observed hang).
• With more than one valid kernel, the very first call incorrectly reported "exhausted"
  and returned false, so the search silently visited zero performance configs
  (under-tuning: the deterministic path never actually explored the kernel list).

The unaffected sibling solvers, PerformanceConfigHipImplicitGemmGroupFwdXdlops and
PerformanceConfigHipImplicitGemm3DGroupBwdXdlops, already implement the equivalent
deterministic-mode advance correctly using a direct boundary check instead of
NextLinear()'s return value, which is why only the grouped Bwd/Wrw Xdlops solvers
were affected.

Technical Details

• projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp and
  .../conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp: replaced the
  if(!NextLinear(0, valid_kernels.size() - 1, index)) idiom in the is_deterministic
  branch of SetNextValue() with the same explicit boundary-check idiom already used
  correctly by the Fwd and 3D Bwd siblings:

  if((index + 1) < static_cast<int>(valid_kernels.size()))
  {
      ++index;
      split_k   = 1;
      kernel_id = valid_kernels[index] + "+1";
      return true;
  }
  return false; // All kernels exhausted

  This makes the deterministic branch self-contained (no dependence on NextLinear's
  wraparound semantics) and consistent with the rest of the codebase. The
  non-deterministic branch (which still uses NextTwoPower/NextCKSplitkValue) is
  untouched.
• projects/miopen/src/include/miopen/conv/solvers.hpp: added MIOPEN_INTERNALS_EXPORT
  to both affected SetNextValue() declarations, matching the annotation already
  present on the 3D Bwd sibling's SetNextValue(). This is what lets the new gtest
  binary call SetNextValue() directly from a host-only test without linking against
  the CK shared library or requiring a GPU.
• projects/miopen/test/gtest/perf_config_group_deterministic_set_next_value.cpp
  (new file): host-only regression test with 6 cases covering, for both Bwd and Wrw:
  • the hang regression (one valid kernel — must return false immediately, not loop),
  • the under-tuning regression (three valid kernels — must visit index 1 then index 2
    before exhausting),
  • a non-deterministic-branch smoke guard (asserts the untouched branch still visits
    more than one index/split_k combination and terminates within a bounded iteration
    count).
    The tests populate valid_kernels directly on the performance config before calling
    SetNextValue(), bypassing the lazy CK-library-backed candidate lookup, so they run
    entirely on the host with no GPU.

Risk Assessment

Risk Level: 🟡 Medium

The code change itself is minimal (two near-identical 5-line branch rewrites plus two
export annotations), and the affected code path is gated behind the opt-in
MIOPEN_CONVOLUTION_ATTRIB_DETERMINISTIC attribute, so most callers are unaffected.
It is rated Medium rather than Low because it touches core convolution solver
performance-config search logic (conv_hip_implicit_gemm_grouped_{bwd,wrw}_xdlops.cpp)
— exactly the kind of file whose subtle polarity/boundary bugs are hard to catch by
inspection (as this bug itself demonstrates) and which is exercised by every grouped
Bwd/Wrw Xdlops find/tune call.

Impacted Components

• miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupBwdXdlops::SetNextValue
  (grouped backward-data CK Xdlops solver)
• miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupWrwXdlops::SetNextValue
  (grouped backward-weights CK Xdlops solver)
• projects/miopen/src/include/miopen/conv/solvers.hpp (export visibility only, no
  behavioral change)
• New test target: perf_config_group_deterministic_set_next_value.cpp

Potential Side Effects

• Deterministic-mode Find/tuning calls for grouped Bwd/Wrw Xdlops convolutions will
  now actually enumerate all valid kernel candidates instead of hanging (single
  candidate) or silently enumerating zero candidates (multiple candidates). Any
  downstream code or test that was inadvertently relying on the previous
  zero-candidates-enumerated behavior (e.g. assuming the deterministic path always
  falls back to a default/heuristic config) could see a different kernel selected.
• No change to the non-deterministic branch, to IsValidValue()/IsValid(), or to
  any other solver.
• The added MIOPEN_INTERNALS_EXPORT annotations only affect symbol visibility for
  internal builds/tests; they do not change runtime behavior.

Mitigation Steps

• The fix reuses an idiom already proven correct and shipping in two sibling solvers
  (GroupFwdXdlops, 3DGroupBwdXdlops) rather than inventing new logic, minimizing
  the chance of a new class of bug.
• Added a dedicated, fast, host-only regression test suite that directly exercises the
  exact hang and under-tuning scenarios, so any future regression on this branch is
  caught in CI without needing GPU hardware or a TF32/MI300X repro environment.
• Revert plan: the change is a self-contained 2-file, ~10-line functional diff plus a
  new test file; reverting conv_hip_implicit_gemm_grouped_{bwd,wrw}_xdlops.cpp to the
  prior NextLinear-based branch (and dropping the new test) fully restores prior
  behavior if needed.

Test Plan

Added projects/miopen/test/gtest/perf_config_group_deterministic_set_next_value.cpp,
a new host-only gtest file with 6 test cases (no GPU or CK shared library required):

• CPU_PerfConfigGroupBwdDeterministic_NONE.HangRegressionSingleKernel
• CPU_PerfConfigGroupBwdDeterministic_NONE.UnderTuningThreeKernels
• CPU_PerfConfigGroupWrwDeterministic_NONE.HangRegressionSingleKernel
• CPU_PerfConfigGroupWrwDeterministic_NONE.UnderTuningThreeKernels
• CPU_PerfConfigGroupBwdNonDeterministic_NONE.VisitsMultipleIndicesAndSplitKs
• CPU_PerfConfigGroupWrwNonDeterministic_NONE.VisitsMultipleIndicesAndSplitKs

The hang-regression cases use a single valid kernel and assert SetNextValue() returns
false on the first call (with a bounded iteration cap of 10 as a safety net against a
true infinite loop during test execution). The under-tuning cases use three valid
kernels and assert the config advances through index 1 and index 2 (with correct
split_k/kernel_id) before exhausting on the fourth call. The non-deterministic
smoke test guards the untouched branch, asserting it still visits more than one index
and more than one split_k value within a 500-iteration cap.

Test Result

All 6 new tests were built and run locally and pass. The fix was additionally verified
to be meaningful (not a no-op) by temporarily reverting the SetNextValue() change and
re-running the suite: with the reverted (buggy) code, the corresponding tests fail fast
via the bounded-iteration-cap assertion (ASSERT_LT(iterations, 10) /
ASSERT_LT(iterations, 500)) rather than hanging the test process, confirming the test
suite correctly detects both the hang and under-tuning regressions. No GPU hardware or
TF32/MI300X environment was needed to build or run these tests, since they call
SetNextValue() directly on a host-constructed ProblemDescription with a
manually-populated valid_kernels list.

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

• python -m pip install pre-commit
• pre-commit install
• pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.