Skip to content

[MIOpen] Feature-engineered 2D AI heuristics (Tunanet + KTN candidate selection) for gfx942 & gfx950 - #7109

Merged
amd-bartgips merged 68 commits into
developfrom
heuristic_2d_model_update
Jun 19, 2026
Merged

amd-bartgips merged 68 commits into
developfrom
heuristic_2d_model_update

Conversation

@amd-ahyttine

@amd-ahyttine amd-ahyttine commented May 6, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

MIOpen's 2D AI heuristics previously fed raw convolution and kernel-config parameters directly into the fdeep models. Training showed that explicit, stateless feature engineering improves both solver selection (Tunanet) and kernel-config ranking (KTN candidate selection) for grouped implicit-GEMM convolutions. This PR ports that feature engineering into the C++ inference path so training and inference agree, and extends the engineered path from gfx950 to gfx942 (MI300).

Beyond keeping training and inference in agreement, the refreshed models are also required to support the new large-tensor grouped CK xdlops kernels — those that #7663 (now merged into develop) makes applicable for large-stride/large-tensor problems. MIOpen marks those solvers applicable, but the heuristics can only select a kernel that exists in the model's vocabulary, and the previous models predate these kernels and never rank them. The updated Tunanet and KTN models add the new large-tensor kernel variants to their output space (and the KTN candidate decode handles their variable kernel-config layout), so the heuristics can actually rank and select them rather than falling back.

Risk Assessment

Moderate, and well contained. The changes are confined to MIOpen's 2D AI heuristic selection path (Tunanet solver selection and KTN candidate ranking) for grouped implicit-GEMM convolutions, gated to gfx942/gfx950 and 2D; other
architectures and 3D use unchanged paths (3D candidate selection now explicitly takes a raw-feature passthrough). The AI heuristic is an immediate-mode/predictive selector whose output is a choice of solver/kernel-config: it is wrapped in a
try/catch that falls back to the existing heuristic on any failure, and the chosen kernel still goes through MIOpen's normal applicability and (in Find) verification flow — so a mispredicted or failed inference degrades to the fallback rather than producing incorrect results. Unsupported precisions and kernel configs the model cannot encode now consistently throw and are caught at the same fallback boundary (rather than feeding a degraded all-zero/NaN feature vector into the model), and a nightly test drives every grouped-CK kernel instance through the encoder to assert it degrades gracefully instead of crashing. No public/user-facing API, ABI, or header changes; a few internal (MIOPEN_INTERNALS_EXPORT) helper signatures gained parameters, and all call sites are updated. The metadata schema gains optional fields (conditional_layouts, the INT8 precision class); the decode treats them as optional, so a model without them follows the prior static path. The main risk surface — the variable-length kernel-config decode and the per-model feature widths — is covered by new unit tests. Bundled .tn.model weights/metadata are updated; any quality regression would be in selection quality (performance), not correctness, and is bounded by the fallback.

Technical Details

Tunanet (solver selection), 2D:

  • ExtractTunaNetND2dFeatures() builds the engineered 2D feature vector (one-hot layout/precision/direction encodings, raw conv parameters, and derived complexity/GEMM/utilization features) consumed by the Tunanet ND model.
  • The precision one-hot includes INT8, so the engineered vector is 47 elements; MetadataND::EncodePrecision maps miopenInt8 accordingly.
  • Enabled for 2D on both gfx950 and gfx942; other archs/dimensionalities keep the existing path.

KTN candidate selection (kernel-config ranking), 2D:

  • EngineerCandidateSelectionInputFeatures() (input encoder) and EngineerCandidateSelectionKernelConfigFeatures() (kernel-config encoder) produce the engineered vectors (one-hot categoricals + raw numerics + derived interaction features). Widths are derived from the model metadata (e.g. the precision one-hot width, which differs per solver once INT8 is present), rather than hardcoded.
  • 3D grouped solvers retain the raw (non-engineered) feature path; the engineered path is selected per-model from the spatial_dim constant in the metadata, so 3D models continue to work unchanged.

Variable-layout CK kernel decode (metadata-driven):

  • Some WRW CK kernels have a variable type-string layout: NumGroupsToMerge is emitted only when > 1, SplitK is appended after the closing > as a +<int> suffix, and the Wavelet kernel packs its first token as "{load}l+{math}m"
    (split into two features).
  • EncodeKernelParams() decodes these from a declarative conditional_layouts section in the model metadata (kinds present_if_gt_one / appended_suffix / inline_after_optionals, and a named packed-token codec), instead of hardcoded kernel-name/index logic.
    Kernels without such an entry use the static kernel_str_mapping path.

Robustness / fixes:

  • Metadata decodings/outputs parsing tolerates numeric (not only string) values.
  • Updated gfx942/gfx950 2D Tunanet and per-solver KTN model + metadata files.

Test Plan

Automated (this PR, in CI):

  • Build MIOpen with AI heuristics enabled on gfx942 and gfx950.
  • Engineered-encoder unit tests: Tunanet 2D (47-feature vector) and KTN candidate-selection input (43 for Bwd/Wrw, 44 for Fwd) + kernel-config encoders — golden and hardcoded tests pinning assembly order, metadata-driven one-hot widths/indices, and num_cu.
  • Variable-layout CK decode (conditional_layouts: present_if_gt_one / appended_suffix / inline_after_optionals + the packed Wavelet codec) covered by crafted unit tests (these instances are not enumerated in this build).
  • Encoder robustness gate (nightly/weekly): every grouped-CK kernel instance of all six solvers is driven through the encoder and must not throw and must produce a well-formed row (a graceful NaN sentinel for configs the model has not seen) — guards the variable-layout decode-crash class that motivated this work.
  • Graceful fallback: an unsupported precision or an unencodable kernel config throws → is caught → falls back to the non-AI heuristic (both paths unit-tested).
  • Direction-aware channel/spatial swap for Bwd/Wrw.
  • clang-format and clang-tidy clean on the changed files.

End-to-end & performance:

  • MIOpenDriver on gfx942/gfx950: the new large-tensor grouped CK xdlops kernels enabled by [MIOpen] Enable grouped CK xdlops solvers for large-stride tensors (ROCM-23997) #7663 are present and are selected by the AI heuristics.
  • Performance vs the develop baseline (gfx950, 120-shape 2D sample, 40 per direction): no regression (see Test Result).
  • Large-shape behaviour: on a separate large-problem workload, the heuristic selects the new large-tensor kernels (no naive fallback).
  • Regression: confirm non-gfx942/gfx950 and 3D conv paths still use the legacy raw-feature extraction.

Test Result

TLDR: No regression vs develop; correct large-tensor kernel selection; ~2–3% over prior release on near-roofline shapes

  • AI/heuristics gtest suites: 242 passed, 0 failed (candidate-selection 86, ND heuristics 31, kernel-tuning-utils 82, kernel-tuning-net 20, tuning-policy 6, kernel-tuning-mode logs 17; 23 skipped, all architecture/feature-gated — e.g. Wavelet/large-tensor instances not enumerated in this build and 3D variants of 2D-only tests).
  • Encoder robustness gate green across all six grouped solvers. Coverage note: the 3D-Wrw model currently encodes 44/108 enumerated CK instances (the MultipleD_Xdl_CShuffle family awaits a retrain); these degrade gracefully (ranked last), not a crash.
  • End-to-end, gfx950 (MI355), implicit-GEMM forced: the grouped-CK xdlops solvers run for the covered shapes and KTN candidate selection ranks and selects the new large-tensor configs — including CShuffle / CShuffleV3 large-tensor variants with SplitK and BlkGemmPipelineScheduler/Version tokens — with zero decode/encoding failures, and every KTN-selected kernel Verifies OK against the CPU reference. (One Wrw shape failed verification on an unrelated ConvAsmImplicitGemmGTCDynamicWrwXdlops path — not the AI heuristics or the CK group solver.)
  • Regression, gfx942/gfx950: 3D convolutions invoke the ND model but use the raw 3D feature path (no engineered-2D extraction) and Verify OK; the engineered path is confirmed active only for 2D. Architecture gating confirms the engineered ND path is enabled for gfx942/gfx950 only (gfx90a/gfx908 report NOT SUPPORTED → legacy path), for both 2D and 3D.
  • Inference-path validation: the C++ engineered features match the feature definitions the models were trained with, and the deployed (frugally-deep) models reproduce their source models' outputs exactly across the smoke-test shapes. Any remaining Tunanet solver-selection quality is therefore a model-training matter, not the inference code.
  • Performance vs develop (gfx950, 120-shape 2D sample, 40 per direction): no regression. Solver selection is identical to develop on all shapes; the updated models change the KTN-selected config on 34/120 shapes with net-neutral impact (median Δ −0.05%, 29/34 within ±3%). Larger run-to-run swings were traced to multi-GPU benchmarking contention on identical kernels, not heuristic differences. (This general-shape sample is already well served by existing solvers, so it does not isolate the new large-tensor kernels — see below.)
  • Selection quality vs full tuning (gfx950): the heuristic selects an implicit-GEMM solver for 86 of the 88 shapes where iGEMM is the exhaustively-tuned optimum (the 2 misses are sub-millisecond shapes), confirming it correctly chooses iGEMM over naive. The remaining gap to the tuned optimum is the model-tuning target for follow-up.
  • Large-problem workload (independent testing, gfx950/MI355): on a separate large-problem set, the updated heuristic selects the new large-tensor CK kernels for the large shapes (never falling back to naive), and improves end-to-end runtime by ~2–3% over the prior release on the non-large shapes — modest because those shapes are already near roofline.

Submission Checklist

@amd-ahyttine amd-ahyttine changed the title Heuristic 2d model update DRAFT NOT TO BE MERGED Heuristic 2d model update DRAFT May 15, 2026
root and others added 19 commits May 18, 2026 12:04
…OCM-23997)

Allow the six grouped CK xdlops solvers to handle convolutions whose
tensor element-strides exceed the int32 range:

  - ConvHipImplicitGemmGroup{Fwd,Bwd,Wrw}Xdlops       (2D)
  - ConvHipImplicitGemm3DGroup{Fwd,Bwd,Wrw}Xdlops     (3D)

Changes:

* Remove the AllTensorsDimsFitIntoInt() applicability guard from all six
  solvers.

* Widen CKArgs dimension members from int to int64_t and length/stride
  arrays from std::array<ck::index_t,N> to std::array<ck::long_index_t,N>
  across the shared CRTP base (CKArgsSplitK, used by 2D Bwd/Wrw), the 2D
  Fwd CKArgs, and all three 3D CKArgs templates (Fwd/Bwd/Wrw). The NCHW
  stride builder is the prime overflow site for shapes like
  (1,96,4736,4736) or (1,96,1024,1024,1024) — operands are now int64 so
  the contiguous stride no longer wraps before assignment. CKConvDims in
  the common helpers is widened in lock step. Argument construction now
  binds to CK's long_index_t MakeArgumentPointer overload.

* Widen the GetCKSplitkMaxWorkspaceSize accumulator from \`auto\` (deduced
  as int) to std::size_t so the max-reduction over per-instance workspace
  sizes cannot itself overflow when large-tensor instances report
  multi-GB workspaces.

* Add a CK large-tensor instance filter
  (RequiresLargeTensorCKInstance / IsLargeTensorCKInstance) applied at
  every CK-instance enumeration site:
    - FillValidKernelsIDs         (GetSolution path)
    - IsCKApplicable              (IsApplicable path)
    - IsCKArgsSupported           (Check*ArgSupported* / HeuristicInit path)
    - IsCKSplitKSupported         (per-dtype; Generic dispatch inherits)
    - GetCKSplitkMaxWorkspaceSize (Bwd/Wrw workspace sizing)
  Without the filter, most non-large-tensor CK impls silently narrow the
  int64 overload back to int32 and would corrupt kernel arguments for
  out-of-range strides; in debug builds the ToCKIndexArray assert in
  ck_grouped_conv_impl_helpers.hpp catches the narrow. When the
  predicate fires and no large-tensor instance is registered, the solver
  is correctly marked inapplicable.

Tests:

* Add gtest unit tests (conv_api_solution_count_{2d,3d}_large_stride)
  that invoke each of the six solvers directly via
  miopenConvolution{Forward,BackwardData,BackwardWeights}CompileSolution
  to confirm CK is selected past the 2^31 element-stride boundary.
  Shapes with no matching CK kernel instance are listed in IsKnownFailing
  helpers and skipped. The IsKnownFailing* lists were hand-tuned for
  gfx942 CK tile coverage, so these sweeps are gated to the CI-covered
  arches (gfx90A | gfx94X | gfx950 | gfx115X) via IsTestSupportedByDevice
  in the shared RunCompile{Fwd,BwdData,Wrw} helpers; non-allowlisted
  arches SKIP rather than emit stale FAILED lines.

* Add kernel-launching numerical Fwd tests on shapes whose
  element-strides exceed INT_MAX (2D: 1x96x4736x4736; 3D:
  1x96x512x512x88) via UnitTestConvSolverGroupXDlops, so the CPU
  reference compare catches int32 wrap-around inside the CK kernel
  itself. These are excluded from standard/pre-checkin via
  test_categories.yaml (~4-5+ GB each) and run in comprehensive/full;
  skipped on memory-constrained systems. Gated to gfx90A | gfx94X (CI
  covered and manually qualified for the large-stride kernel-launch
  path) via UnitTestConvSolverParams::supported_devs.

* Add four single-position TestCheckDimsFitIntoInt boundary cases
  (INT_MAX and INT_MAX+1 in one length and one stride slot) so the
  strict-\`>\` semantics is greppable and survives future refactors.

Add a Resolved-issues bullet to projects/miopen/CHANGELOG.md.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…eserve filter contract

The int32 narrowed length/stride arrays handed to CK BWD/WRW
MakeArgumentPointer must outlive the returned Argument (CK captures
references). Store them as a member of CKArgs instead of as locals in
MakeArgPtr, which previously caused stack-use-after-scope under ASAN.

Populate the bundle lazily via a const method on the (now `mutable`)
member, not in the constructor. CKArgs is constructed unconditionally
in FillValidKernelsIDs before the RequiresLargeTensorCKInstance filter
runs; narrowing in the constructor would trip ToCKIndexArray's
debug-only overflow assert on >INT_MAX shapes even when no kernel is
ultimately selected. Lazy population means narrowing only runs from
MakeArgPtr, which is only reached for kernels that survived the filter
-- restoring the original safety contract.

Applies to 2D BWD/WRW (CKArgsSplitK base) and 3D BWD/WRW CKArgs.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The IsKnownFailing* lists in conv_api_solution_count_{2d,3d}_large_stride
are hand-tuned for gfx942 CK tile coverage. A gfx1151 CI run produced 26
failures across both sub- and >INT_MAX shapes (including the smallest 3D
FP32 shape), because the CK *Xdlops solvers target CDNA MFMA and have not
been characterized on RDNA 3.5. Remove gfx115X from IsArchInCiAllowlist
so these tests SKIP on gfx115X instead of FAILing, matching the header's
documented design intent. Re-add once gfx115X CK coverage is characterized
and the IsKnownFailing* lists grow per-arch entries.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…s gfx90a/gfx942/gfx950

CK upstream has registered large-tensor xdlops kernels for the >INT_MAX
element-stride paths on gfx90a/gfx942/gfx950, so the per-shape
IsKnownFailing* predicates in the 2D/3D API sweeps are no longer needed
and the unit tests can be widened to include gfx950 alongside gfx90A/gfx94X.

- Drop IsKnownFailing{Fwd,BwdData,Wrw}{2D,3D} predicates and MatchesDhw
  helper; simplify Run* callsites accordingly.
- Drop the KnownFailingFn template parameter and per-shape GTEST_SKIP
  block from RunCompile{Fwd,BwdData,Wrw} in the shared header.
- Add Gpu::gfx950 to supportedDevices in the two
  ConvHipImplicitGemm{,3D}GroupFwdXdlops_LargeStride unit tests.
- Refresh stale comments that referenced the removed predicates.

Verified on gfx942: 606/606 tests pass across all four binaries with
zero SKIPPED and zero FAILED.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds end-to-end kernel-launch tests for the grouped CK xdlops BackwardData
and BackwardWeights solvers on shapes whose element-stride exceeds INT_MAX,
mirroring the existing Forward LargeStride coverage. These complement the
API-level applicability sweeps by catching int32 wraparound that can occur
inside the CK kernel even after MIOpen's host-side widening. Excluded from
the per-PR test category via test_categories.yaml.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…nt64 MakeArgumentPointer

PR 7663 widened CK grouped-conv CKArgs length/stride arrays to
ck::long_index_t so the NCHW stride builder doesn't silently overflow
on tensors with stride > INT_MAX. Commit 23059ec patched 2D/3D
BWD/WRW to narrow back to int32 in MakeArgPtr via a mutable
NarrowedCKArrays member. 2D/3D Fwd were missed, regressing the
StaticFDBSync gtest on gfx942 (1862 "Invalid perf config found" errors
in ConvHipImplicitGemm{,3D}GroupFwdXdlops on sub-INT_MAX shapes
because the int64 MakeArgumentPointer overload was being bound for
non-large-tensor CK instances that were only registered for int32).

Unlike BWD/WRW, large-tensor Fwd registrations DO exist, so Fwd
cannot unconditionally narrow. Dispatch at runtime on
miopen::solver::IsLargeTensorCKInstance(conv_ptr): large-tensor
instances bind the long_index_t overload directly; non-large-tensor
instances go through the new NarrowedArrays() bridge and bind the
int32 overload. The IsCKArgsSupported filter only admits non-large-
tensor instances for sub-INT_MAX problems, so ToCKIndexArray's
overflow assert cannot trip on the narrow path.

Applies to 2D Fwd (single MakeArgPtr) and 3D Fwd (Bilinear/Scale/
Default sub-methods). For Bilinear's two braced output-tensor slots,
the int32 branch uses {a.out_l}/{a.out_s} so the brace-wrapped array
carries int32, not int64.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
@amd-ahyttine amd-ahyttine changed the title Heuristic 2d model update DRAFT Update MIOpen gfx950 2D AI heuristics with feature-engineered Tunanet and KTN models DRAFT Jun 9, 2026
root and others added 4 commits June 9, 2026 07:55
…ve tier

The LargeStride conv solver unit tests are slow on the {1,96,4736,4736} (2D)
and {1,96,512,512,88} (3D) shapes. Keep only the BFP16 instances in the
comprehensive (nightly) tier and leave FP16/FP32 in full, since all three
dtypes share the same dispatch code path.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Fixes pre-commit trailing-whitespace hook failure in CI.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
amd-bartgips and others added 7 commits June 16, 2026 10:14
…+ add loud model-load test

A heuristic model that frugally-deep cannot load throws a non-miopen std::exception
(e.g. an incompatibly-exported model -> "tensor_shape_variable needs to be an array").
The AI fallback boundaries only caught miopen::Exception, and the TunaNet ND model is
loaded lazily in Forward (outside GetNDModel's try/catch), so such a failure escaped the
fallback and hard-failed the convolution ("Error getting workspace size, status = 7")
instead of degrading to the non-AI heuristic.

Production robustness:
- PredictSolver now wraps the model evaluation and returns empty (-> WTI fallback) on any
  std::exception; the predictor's contract is "a prediction or nothing, never throws".
- The candidate-selection boundary (ai_conv_nd_kernel_tuning_utils.hpp) catches
  std::exception, not just miopen::Exception.

So a bad/incompatible model degrades gracefully rather than breaking conv. To make sure
this does NOT silently mask a broken bundled model in CI, add a dedicated, loud test:
- BundledFdeepModelsLoad loads every shipped fdeep .tn.model (all except the plain-JSON
  *_metadata.tn.model) via fdeep::load_model and fails if any cannot be parsed.

This keeps the current (incompatibly-exported) models RED in CI until they are re-exported
in a frugally-deep 0.15.30-compatible (Keras 2) format, while production stays robust.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…red features

The derived feature block (common::EngineeredConvFeatures) computed the implicit-GEMM
(M, N, K) dimensions with a single, mislabeled formula: N_gemm and K_gemm were assigned
the opposite way from the feature definitions the models were trained with (even for
Forward), and the direction was not taken into account at all -- a conv lowers to a
different GEMM for Forward / BackwardData / BackwardWeights. As a result 5 of the 18
derived features (log N_gemm, log K_gemm, M/N, M/K, N/K) were wrong on every problem, and
backward/wrw were further scrambled. The model therefore received out-of-distribution
features and made poor solver/kernel selections (e.g. falling back to naive).

Make EngineeredConvFeatures direction-aware: add a ConvDirection parameter and select the
(M, N, K) assignment per direction to match the trained feature engineering exactly.
Callers pass dimensions in the forward (driver) convention and the actual direction:
- ExtractTunaNetND2dFeatures maps problem.GetDirection().
- EngineerCandidateSelectionInputFeatures maps its direction feature (0=Fwd,1=Bwd,2=Wrw).

The existing golden tests used C_in == C_out (Forward), which hid the N/K swap. Rewrite the
CPU golden to use C_in != C_out across all three directions so the GEMM assignment is fully
pinned. Verified against the training feature pipeline: the rebuilt EngineeredConvFeatures
now reproduces its derived vectors exactly for Fwd/Bwd/Wrw.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
… derived features

The kernel-config encoder's derived block clamped MXdlPerWave and NXdlPerWave with the
missing->1 "safe" helper before using them in xdl_ops_per_block / xdl_coverage_m /
xdl_coverage_n. The trained feature definitions use the raw values there (only the other
GEMM-tile parameters are clamped), so a missing (sentinel) MXdlPerWave/NXdlPerWave produced
a different value (a sign flip in the xdl_coverage features). Use the raw value for those two
params to match. Latent in practice (these are core xdlops params, normally present) but a
real divergence from the trained features.

Verified with a C++<->Python kernel-config parity check across the gfx942/gfx950 Fwd/Bwd/Wrw
solvers, including the missing-parameter edge case.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Retrained MI355 2D solver-selection TunaNet with soft-label cross-entropy
at lr=5e-5. Test ACC ~85.2%, AVG_REL_DT ~1.03 (epoch 385 checkpoint),
improving over the previous gfx950 2D model.
@amd-bartgips
amd-bartgips marked this pull request as ready for review June 17, 2026 09:02
@amd-bartgips
amd-bartgips requested a review from a team as a code owner June 17, 2026 09:02
@amd-bartgips

Copy link
Copy Markdown
Contributor

I got confirmation from @jdcampbe that his testing showed improvment wrt the ROCm 7.2 release, and specifically that the tunanet picked igemm solvers (instead of Naive) for the large shapes that now have access to new CK kernel configs

Small refinement of the currently committed 2D solver-selection models for
gfx942 (MI300) and gfx950 (MI355). The previously committed models were used as
a starting point and briefly fine-tuned with a training objective that directly
targets predicted runtime rather than top-pick classification accuracy. This is
a minor update, not a full retrain.

Quality is measured by the relative slowdown of the solver the model picks
versus the fastest available solver per problem, averaged over a held-out test
set and the training set (1.0 = always picks the optimal solver; lower is
better):

  gfx950 (MI355): test 1.104 -> 1.077, train 1.049 -> 1.027
  gfx942 (MI300): test 1.025 -> 1.021, train 1.026 -> 1.025

Both models pick faster solvers on average than the previously committed ones,
on both the test and training sets. The models' raw top-1 selection accuracy is
slightly lower by design, since the objective optimizes for actual runtime of
the chosen solver, which is what matters at inference time.

@jdcampbe jdcampbe left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM. I've run tests on the branch for MI355x to positive results. MI300x evaluations in progress. The branch is on a stale develop so be wary merging it in.

@amd-bartgips

amd-bartgips commented Jun 18, 2026 •

Copy link
Copy Markdown
Contributor

I see the CI pipeline fails:

[2026-06-17T13:37:32.532Z] 21: [ RUN      ] Smoke/GPU_HipGraphExistTest_FP32.HipGraphExist/2
[2026-06-17T13:37:32.532Z] 21: Executing: /var/jenkins/workspace/_libraries-folder_MIOpen_PR-7109/projects/miopen/build/bin/MIOpenDriver bnorm -F 2 -n 8 -c 64 -H 8 -W 8 -m 1 -r 1 -i 1 -V 1 -t 1 --use_hip_graph 1
[2026-06-17T13:37:32.532Z] 21: Command return code: 134
[2026-06-17T13:37:32.532Z] 21: Driver output:
[2026-06-17T13:37:32.532Z] 21: MIOpenDriver bnorm -F 2 -n 8 -c 64 -H 8 -W 8 -m 1 -r 1 -i 1 -V 1 -t 1 --use_hip_graph 1
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :443 : 2815562398309 us: [pid:986592 tid: 0xfbfc9] Initalizing runtime stack, Enumerated GPU agents = 1
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :178 : 2815562398600 us: [pid:986592 tid: 0xfbfc9] Numa selects cpu agent[0]=0x59d177f5c1a0(fine=0x59d177f5c400,coarse=0x59d177f59330) for gpu agent=0x59d177f5cf00 CPU<->GPU XGMI=0
[2026-06-17T13:37:32.532Z] 21: :3:comgrctx.cpp             :112 : 2815562398642 us: [pid:986592 tid: 0xfbfc9] Loaded COMGR library version 3.0.
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :1643: 2815562398730 us: [pid:986592 tid: 0xfbfc9] addressableNumVGPRs=512, totalNumVGPRs=512, vGPRAllocGranule=8, availableRegistersPerCU_=131072
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :1654: 2815562398740 us: [pid:986592 tid: 0xfbfc9] imageSupport=0
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :1701: 2815562398742 us: [pid:986592 tid: 0xfbfc9] Gfx Major/Minor/Stepping: 9/4/2, Device ID: 0x74b9
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :1702: 2815562398742 us: [pid:986592 tid: 0xfbfc9] Using dev kernel arg wa = 2
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :1704: 2815562398743 us: [pid:986592 tid: 0xfbfc9] HMM support: 1, XNACK: 0, Direct host access: 0
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :1706: 2815562398744 us: [pid:986592 tid: 0xfbfc9] Max SDMA Read Mask: 0xffff, Max SDMA Write Mask: 0xffff
[2026-06-17T13:37:32.532Z] 21: :4:rocdevice.cpp            :2099: 2815562399351 us: [pid:986592 tid: 0xfbfc9] Allocate hsa host memory 0x7eae93a00000, size 0x400000, numa_node = 0, mem_seg = 0
[2026-06-17T13:37:32.532Z] 21: :4:rocdevice.cpp            :2099: 2815562399589 us: [pid:986592 tid: 0xfbfc9] Allocate hsa host memory 0x7eaed5d13000, size 0x38, numa_node = 0, mem_seg = 1
[2026-06-17T13:37:32.532Z] 21: :3:runtime.cpp              :70  : 2815562399603 us: [pid:986592 tid: 0xfbfc9] ROCclr version: 8017ffed34
[2026-06-17T13:37:32.532Z] 21: :3:hip_context.cpp          :40  : 2815562399609 us: [pid:986592 tid: 0xfbfc9] HIP Version: 7.13.61040, Direct Dispatch: 1
[2026-06-17T13:37:32.532Z] 21: :3:os_posix.cpp             :967 : 2815562399613 us: [pid:986592 tid: 0xfbfc9] Runtime Library Path: /opt/rocm/lib/libamdhip64.so.7
[2026-06-17T13:37:32.532Z] 21: :3:hip_stream.cpp           :275 : 2815562400219 us: [pid:986592 tid: 0xfbfc9]  hipStreamCreate ( 0x7ffc988638e0 ) 
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :3037: 2815562400231 us: [pid:986592 tid: 0xfbfc9] Number of allocated hardware queues with low priority: 0, with normal priority: 0, with high priority: 0, maximum per priority is: 4
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :3111: 2815562424566 us: [pid:986592 tid: 0xfbfc9] Created SWq=0x7eae9a8ea000 to map on HWq=0x7eae93700000 with size 16384 with priority 1, cooperative: 0
[2026-06-17T13:37:32.532Z] 21: :3:rocdevice.cpp            :3211: 2815562424638 us: [pid:986592 tid: 0xfbfc9] acquireQueue refCount: 0x7eae93700000 (1) 
[2026-06-17T13:37:32.532Z] 21: :4:rocdevice.cpp            :2283: 2815562424665 us: [pid:986592 tid: 0xfbfc9] Allocate hsa device memory 0x7eae93200000, size 0x400000, hsa_mem_flags 0x4h
[2026-06-17T13:37:32.532Z] 21: LLVM ERROR: out of memory
[2026-06-17T13:37:32.532Z] 21: Allocation failed
[2026-06-17T13:37:32.532Z] 21: Aborted
[2026-06-17T13:37:32.532Z] 21: 
[2026-06-17T13:37:32.532Z] 21: /var/jenkins/workspace/_libraries-folder_MIOpen_PR-7109/projects/miopen/test/gtest/hip_graph_exist.cpp:182: Failure
[2026-06-17T13:37:32.532Z] 21: Failed
[2026-06-17T13:37:32.532Z] 21: MIOpenDriver failed with return code 134
[2026-06-17T13:37:32.532Z] 21: 
[2026-06-17T13:37:32.532Z] 21: [  FAILED  ] Smoke/GPU_HipGraphExistTest_FP32.HipGraphExist/2, where GetParam() = bnorm_bnorm_hip_graph (160 ms)

Which seems to be unrelated to the current PR. I'll merge in the current develop again and hope it'll pass this time.

@amd-bartgips

Copy link
Copy Markdown
Contributor

It seems I lost write access, so I asked @johannes-graner to update the branch for me :). thanks!
I requested renewed write access here: https://github.com/orgs/ROCm/teams/rocm-core-write
Let's hope I don't have to fix any bugs and get access again soon.

@amd-bartgips
amd-bartgips enabled auto-merge (squash) June 19, 2026 08:49
@amd-bartgips
amd-bartgips merged commit c24ad41 into develop Jun 19, 2026
59 checks passed
@amd-bartgips
amd-bartgips deleted the heuristic_2d_model_update branch June 19, 2026 21:43
JonathanLichtnerAMD added a commit that referenced this pull request Jul 7, 2026
JIRA ID: ROCM-27477

## Motivation

`GetFdeepModel` and `GetModel` in
`projects/miopen/src/conv/heuristics/ai_heuristics.cpp` cache
loaded AI heuristic models in `static std::map` containers with no
synchronization. When
convolutions run concurrently from multiple threads on the same device
(e.g. per-thread MIOpen
handles), two callers can race on the same cache key: both find the key
missing, both build a
model, and both insert. In `GetFdeepModel` the second insert
(`models[key] = std::move(model)`)
move-assigns over the first thread's `unique_ptr<fdeep::model>`,
destroying the model object while
the first thread is still using the reference it was handed — a
use-after-free that segfaults deep
inside `fdeep`'s prediction code.

This was surfaced by the PyTorch test
`test_cuda.py::TestCuda::test_cudnn_multiple_threads_same_device`
on ROCm 7.14.0-RC1 / gfx942 (MI300 series), with an observed crash rate
of roughly 3-17% per
triggering run. The exposing commit (#7109, which enabled the 2D AI
heuristic default path for
gfx942/gfx950) predates the last known-good nightly, but the underlying
missing-lock defect itself
is older and not a regression introduced within the RC window. The race
is warm-up-bounded: it can
only occur during first concurrent population of a given cache key,
since entries are never
overwritten afterward.

A runtime workaround exists without rebuilding:
`MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS_AI_HEUR=0`
disables the AI heuristic path
entirely, but that is not a substitute for fixing the underlying race.

## Technical Details

The fix mirrors the already-correct sibling `GetCandidateSelectionModel`
in
`ai_candidate_selection.cpp`, which guards its structurally identical
cache with a
`static std::mutex` + `std::lock_guard`.

- `GetModel(arch, solver)`: added a `static std::mutex models_mutex` and
moved the
`models.find(solver)` lookup inside a `std::lock_guard` that also covers
the insert-on-miss
path, so find+insert is atomic with respect to other callers. The
function now returns the
cached `std::shared_ptr<Model>` by value while still holding the lock,
which is safe since
`shared_ptr` copies are cheap and the returned value keeps the model
alive independently of the
  map.
- `GetFdeepModel(path, key)`: added a `static std::mutex models_mutex`
guarding find+insert.
Critically, the insertion path was changed from `models[key] =
std::move(model); return ref;`
(which risks a move-assign clobbering an existing entry on a race, and
previously took the
returned reference *before* the map write) to `it = models.emplace(key,
std::move(model)).first;`
followed by `return *(it->second);`. Because entries are only ever
inserted once and never
reassigned, a reference into the map obtained after the lock is released
remains valid for the
lifetime of the process — the lock only needs to serialize the
population itself, not every
  subsequent read.

No public API, call signatures, or heuristic selection logic changed;
this is purely a
thread-safety fix around existing caching behavior.

## Risk Assessment

**Risk Level: 🟢 Low**

The change is a single-file, 27-line diff that adds standard
mutex-guarded caching to two
functions, following the exact pattern already used correctly elsewhere
in the same file. It does
not alter model selection, loading logic, or any public interface.

### Impacted Components

- `projects/miopen/src/conv/heuristics/ai_heuristics.cpp`: `GetModel`
(used by the general AI
heuristic model cache) and `candidate_selection::GetFdeepModel` (used by
the AI candidate
selection path). Any convolution solver selection path that goes through
the AI heuristics
(currently defaulted on for gfx942/gfx950 2D convolutions) is indirectly
affected.

### Potential Side Effects

- Serializing model cache population introduces a mutex acquisition on
every call to `GetModel`/
`GetFdeepModel`, not only on cache misses. Since these are in-process,
uncontended locks after
warm-up (steady-state calls hit the fast `find` path while holding a
brief, uncontended lock),
the performance impact is expected to be negligible relative to the cost
of running the model
itself, but reviewers should watch for any contention under very high
thread-count concurrent
  convolution workloads.
- No behavioral change to which model is loaded or which heuristic
decision is made — only the
  concurrency safety of caching is affected.

### Mitigation Steps

- The fix follows an existing, proven pattern
(`GetCandidateSelectionModel`) already in production
  in the same file, minimizing the chance of a novel defect.
- The runtime env var workaround
(`MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS_AI_HEUR=0`)
remains available as an immediate mitigation if an unexpected issue
surfaces post-merge, without
  requiring a revert or rebuild.
- Revert plan: the diff is small and self-contained to one file, making
a straightforward revert
  possible if needed.

## Test Plan

This will be validated manually on gfx942 hardware using the PyTorch
repro test
(`test_cuda.py::TestCuda::test_cudnn_multiple_threads_same_device`) that
originally surfaced the
crash, running it repeatedly to build confidence that the segfault no
longer reproduces.

In addition, unit tests have been added that reproduce the race
condition. These have been run manually, and will also run as part of
the standard MIOpen test category.

## Test Result

- [x] manual validation on gfx942 via the PyTorch repro
- [x] CI passes

---------

Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Jul 15, 2026
JIRA ID: ALMIOPEN-2312

## Motivation

`projects/miopen/CHANGELOG.md`'s "MIOpen 3.5.2 for ROCm 7.14.0" section
only had 4 entries, but 68 commits touched `projects/miopen` between the
`therock-7.13` tag and `release/therock-7.14`. Auditing that commit
range (via `git log --oneline
therock-7.13..remotes/origin/release/therock-7.14 -- projects/miopen`)
found 18 additional user-facing changes — new gfx950/gfx1250
architecture support, SystemDB refreshes, AI-heuristic improvements,
several correctness/crash bug fixes, and OpenCL solver removals — that
were never documented. This PR backfills those entries so the 7.14.0
changelog accurately reflects what shipped.

## Technical Details

Added 24 lines to the existing "MIOpen 3.5.2 for ROCm 7.14.0" section of
`projects/miopen/CHANGELOG.md`, merging new bullets into the existing
`Added`/`Changed`/`Resolved Issues` sub-sections (no new entries were
duplicated or reworded) and adding a new `Removed` sub-section for
solver removals. No existing content was altered.

Each new bullet was traced to a specific commit + PR:

| Entry | Commit | PR |
|---|---|---|
| Added gfx950 (MI350X/MI355X) SystemDB | `60d8e1e5175` | #8081 |
| Added initial gfx1250 architecture support | `d3f057b2f22` | #7587 |
| Added JSON-formatted performance logs (`MIOPEN_PERFORMANCE_LOGS`) |
`ffbde871260` | #5736 |
| Refreshed gfx942 SystemDB | `faf754a9ef6` | #8174 |
| Refined 2D AI heuristics (Tunanet/KTN) for gfx942 & gfx950 |
`c24ad415884` | #7109 |
| Fixed thread-safety issue in AI-heuristic model caches | `cd957402309`
| #9149 (cherry-pick of #9131) |
| Fixed CK grouped-conv solvers not always registered in host |
`c971d54d7ae` | #9105 |
| Fixed legacy CK buffer resource descriptor (gfx1250) | `1dee011b8b0` |
#8758 |
| Removed stale gfx942 SystemDB naive-solver entries | `f42668f6316` |
#7028 |
| Fixed softmax NaN when `beta=0` | `fb8157c8512` | #8426 |
| Fixed BN backward spatial variant 1 vectorized memory access bug |
`1351a7e8b10` | #6765 |
| Fixed use-after-free of `kernel_id` in `InitInvokerFactoryNCHW` |
`394ebffb50d` | #7922 |
| Fixed SoftMarginLoss forward reduce-kernel kernarg corruption |
`1e65a31570f` | #8065, #8135 |
| Fixed build with `MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF` | `e6288c3ca61`
| #6800 |
| Fixed install RPATH for MIOpenDriver / CK backend libs | `ef8a4cf2f93`
| #7485 |
| Fixed dangling lambda causing log-level 6 segfault | `190f9791737` |
#6953 |
| Enabled Windows long-path support for `addkernels` | `1f9d69724cf` |
#7387 |
| Removed OCL solver `ConvOclDirectFwd1x1` | `126c720cadd` | #6491 |
| Removed OCL solver `conv_ocl_dir2Dfwdgen` | `2d499ddf709` | #6633 |
| Removed OCL solvers `ConvOclBwdWrW1x1`, `ConvOclBwdWrW2` |
`b5d43a33dc6` | #6859 |
| Removed OCL solvers `OpTensorLeadingOnes`, `Op5dTensorGeneric` |
`8642c83b7de` | #7347 |

The remaining ~40 commits in the 68-commit range were deliberately
excluded (tests/CI/build/docs/internal refactors/1:1 OCL→HIP swaps with
no behavior change, per each commit's own stated scope) and are not
represented here.

## Risk Assessment

**Risk Level: 🟢 Low**

This is a documentation-only change to a single Markdown file. No code,
build configuration, tests, or public API surface is touched.

### Impacted Components

- `projects/miopen/CHANGELOG.md` only.

### Potential Side Effects

None expected. This does not change any runtime behavior; it only
documents changes that already shipped on `release/therock-7.14`.

### Mitigation Steps

None required beyond normal review — a reviewer familiar with the 7.14
release can spot-check a few entries against the linked commit hashes/PR
numbers for accuracy.

## Test Plan

Not applicable (no code changes). Verification was done by
cross-referencing every added bullet against a specific commit hash and
PR number pulled from `git log`, and by diffing the edited section
against `git show
remotes/origin/release/therock-7.14:projects/miopen/CHANGELOG.md` to
confirm no existing entry was duplicated or altered.

## Test Result

All 21 new bullets map 1:1 to a real commit/PR in the
`therock-7.13..release/therock-7.14` range for `projects/miopen`. No
pre-existing changelog text was modified — `git diff` shows a pure
addition (24 insertions, 0 deletions). Two borderline commits
(`a67834413e8` "Standalone build", `7c7cd5774a9` "enable warnings about
deprecated API/features") were deliberately left out pending author
judgment on whether they warrant an entry.
pschang-phy pushed a commit that referenced this pull request Jul 17, 2026
JIRA ID: ALMIOPEN-2312

## Motivation

`projects/miopen/CHANGELOG.md`'s "MIOpen 3.5.2 for ROCm 7.14.0" section
only had 4 entries, but 68 commits touched `projects/miopen` between the
`therock-7.13` tag and `release/therock-7.14`. Auditing that commit
range (via `git log --oneline
therock-7.13..remotes/origin/release/therock-7.14 -- projects/miopen`)
found 18 additional user-facing changes — new gfx950/gfx1250
architecture support, SystemDB refreshes, AI-heuristic improvements,
several correctness/crash bug fixes, and OpenCL solver removals — that
were never documented. This PR backfills those entries so the 7.14.0
changelog accurately reflects what shipped.

## Technical Details

Added 24 lines to the existing "MIOpen 3.5.2 for ROCm 7.14.0" section of
`projects/miopen/CHANGELOG.md`, merging new bullets into the existing
`Added`/`Changed`/`Resolved Issues` sub-sections (no new entries were
duplicated or reworded) and adding a new `Removed` sub-section for
solver removals. No existing content was altered.

Each new bullet was traced to a specific commit + PR:

| Entry | Commit | PR |
|---|---|---|
| Added gfx950 (MI350X/MI355X) SystemDB | `60d8e1e5175` | #8081 |
| Added initial gfx1250 architecture support | `d3f057b2f22` | #7587 |
| Added JSON-formatted performance logs (`MIOPEN_PERFORMANCE_LOGS`) |
`ffbde871260` | #5736 |
| Refreshed gfx942 SystemDB | `faf754a9ef6` | #8174 |
| Refined 2D AI heuristics (Tunanet/KTN) for gfx942 & gfx950 |
`c24ad415884` | #7109 |
| Fixed thread-safety issue in AI-heuristic model caches | `cd957402309`
| #9149 (cherry-pick of #9131) |
| Fixed CK grouped-conv solvers not always registered in host |
`c971d54d7ae` | #9105 |
| Fixed legacy CK buffer resource descriptor (gfx1250) | `1dee011b8b0` |
#8758 |
| Removed stale gfx942 SystemDB naive-solver entries | `f42668f6316` |
#7028 |
| Fixed softmax NaN when `beta=0` | `fb8157c8512` | #8426 |
| Fixed BN backward spatial variant 1 vectorized memory access bug |
`1351a7e8b10` | #6765 |
| Fixed use-after-free of `kernel_id` in `InitInvokerFactoryNCHW` |
`394ebffb50d` | #7922 |
| Fixed SoftMarginLoss forward reduce-kernel kernarg corruption |
`1e65a31570f` | #8065, #8135 |
| Fixed build with `MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF` | `e6288c3ca61`
| #6800 |
| Fixed install RPATH for MIOpenDriver / CK backend libs | `ef8a4cf2f93`
| #7485 |
| Fixed dangling lambda causing log-level 6 segfault | `190f9791737` |
#6953 |
| Enabled Windows long-path support for `addkernels` | `1f9d69724cf` |
#7387 |
| Removed OCL solver `ConvOclDirectFwd1x1` | `126c720cadd` | #6491 |
| Removed OCL solver `conv_ocl_dir2Dfwdgen` | `2d499ddf709` | #6633 |
| Removed OCL solvers `ConvOclBwdWrW1x1`, `ConvOclBwdWrW2` |
`b5d43a33dc6` | #6859 |
| Removed OCL solvers `OpTensorLeadingOnes`, `Op5dTensorGeneric` |
`8642c83b7de` | #7347 |

The remaining ~40 commits in the 68-commit range were deliberately
excluded (tests/CI/build/docs/internal refactors/1:1 OCL→HIP swaps with
no behavior change, per each commit's own stated scope) and are not
represented here.

## Risk Assessment

**Risk Level: 🟢 Low**

This is a documentation-only change to a single Markdown file. No code,
build configuration, tests, or public API surface is touched.

### Impacted Components

- `projects/miopen/CHANGELOG.md` only.

### Potential Side Effects

None expected. This does not change any runtime behavior; it only
documents changes that already shipped on `release/therock-7.14`.

### Mitigation Steps

None required beyond normal review — a reviewer familiar with the 7.14
release can spot-check a few entries against the linked commit hashes/PR
numbers for accuracy.

## Test Plan

Not applicable (no code changes). Verification was done by
cross-referencing every added bullet against a specific commit hash and
PR number pulled from `git log`, and by diffing the edited section
against `git show
remotes/origin/release/therock-7.14:projects/miopen/CHANGELOG.md` to
confirm no existing entry was duplicated or altered.

## Test Result

All 21 new bullets map 1:1 to a real commit/PR in the
`therock-7.13..release/therock-7.14` range for `projects/miopen`. No
pre-existing changelog text was modified — `git diff` shows a pure
addition (24 insertions, 0 deletions). Two borderline commits
(`a67834413e8` "Standalone build", `7c7cd5774a9` "enable warnings about
deprecated API/features") were deliberately left out pending author
judgment on whether they warrant an entry.
amd-ahyttine added a commit that referenced this pull request Aug 28, 2026
## Motivation

<!-- Explain the purpose of this PR and the goals it aims to achieve.
-->

Corresponding 2D PR: #7109

JIRA ID: https://amd-hub.atlassian.net/browse/ALMIOPEN-2265

The 2D PR (#7109) ported MIOpen's AI heuristics (TunaNet solver
selection + KTN candidate
selection) onto an explicit feature-engineered path and extended it from
gfx950 to gfx942, but
it deliberately left **3D grouped implicit-GEMM convolutions on the
legacy raw-feature
passthrough**: `ExtractTunaNetND2dFeatures` was added for 2D only, and
the 3D models had not yet
been retrained for the engineered path.

This PR delivers the 3D counterpart — retrained, feature-engineered 3D
TunaNet and KTN models
for **gfx942 (MI300)** and **gfx950 (MI355)**, together with the **C++
inference changes** that
build the engineered 3D feature vector. As a result the 3D heuristic
path now uses the same
feature engineering the models are trained with, and the models can rank
the new large-tensor
grouped CK xdlops kernels (the V3 / `Large_Tensor` / `MergedGroups`
variants) instead of falling
back.

The **2D model assets are left unchanged from develop**; this PR ships
only the retrained 3D
TunaNet/KTN models plus the shared C++ generalization.

## Technical Details

<!-- Explain the changes along with any relevant GitHub links. -->

This PR has **two parts**: retrained 3D model assets, and the companion
C++ change that
generalizes the engineered feature path from 2D-only to N-D.

### C++ inference path (companion to #7109)

- `ExtractTunaNetND2dFeatures` is generalized to
**`ExtractTunaNetNDFeatures(problem,
spatial_dim)`** (default `spatial_dim = 0` → inferred from
`problem.Is3d()`), building the
engineered TunaNet vector for both 2D and 3D. The engineered path is now
selected for any model
  whose metadata carries the `spatial_dim` constant.
- Candidate-selection input/kernel-config encoding gains a 3D branch:
`IsThreeDimensional` reads
the `spatial_dim` metadata constant, and
`EngineerCandidateSelectionInputFeatures` folds depth
into the volume terms. The derived-feature math is shared with TunaNet
via
  `common::EngineeredConvFeatures`.
- `metadata_nd.cpp`: adds `spatial_dim` handling and the `INT8`
precision class.
- Robustness is unchanged: an unsupported precision or unencodable
config throws → is caught →
  falls back to the legacy heuristic.

This resolves the cross-component dependency that the earlier revision
of this PR flagged as
"confirm before merge": the engineered 3D vector is now produced by the
C++ path in this same PR,
so the 3D models are fed a correctly-shaped vector rather than falling
back on every inference.

### TunaNet (solver selection), 3D

- Input feature width grows from the raw count to the engineered 3D
vector: `batch_input_shape`
**24 → 51** on both gfx942 and gfx950 (the metadata `num_inputs` stays
at 24 = raw conv-param
count, mirroring 2D, where the raw count and the wider engineered vector
differ).
- Output solver vocabulary: **gfx950 9 → 12**; gfx942 stays at 12. Both
now expose the three
`ConvDirectNaiveConv{Fwd,Bwd,Wrw}` solvers as active entries (gfx942
drops the `_DEACTIVATE`
suffixes; gfx950 adds them at indices 0–2). Hidden layers
(64/128/256/64…) are unchanged —
  these are full retrains, not a weights-only tweak.
- Metadata: the `precision` one-hot gains **INT8** (3 → 4 classes);
refreshed to a single new
  tuning session, `miopen_version` 3.5.2.

### KTN candidate selection (kernel-config ranking), 3D

- Encoder networks retrained: input-encoder input widths grow (Fwd 23 →
48, Bwd/Wrw 23 → 47) and
the final embedding `units` grows **16 → 64** across all six encoders;
kernel-config-encoder
input widths grow correspondingly (e.g. Fwd 18/16 → 40/38, Bwd 10/13 →
28, Wrw 21/22 → 45).
- Metadata schema gains the declarative variable-layout decode fields
introduced for 2D in
#7109 — `conditional_layouts` (Fwd/Bwd/Wrw) and `split_k_values`
(Bwd/Wrw) — plus, for Wrw, a
`base_param_count` / `conditional_params` / `packed_params` block
(NumGroupsToMerge,
BlkGemmPipelineScheduler/Version, SplitK, and the packed
`TileLoadMathThreadGroupSize` codec).
- Output kernel vocabulary expands with the new large-tensor CK
variants:
- **Fwd** 3 → 5 kernels (`..._Xdl_CShuffle_V3_Large_Tensor`,
`..._V3_MergedGroups`; the mapping
also adds `_V3_DirectLoad` / `_V3_Large_Tensor` / `_V3_MergedGroups`).
- **Bwd** adds
`DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffleV3_Large_Tensor`; output
params
grow 5 → 10 (`SplitK`, `MRepeat`, `NRepeat`,
`CShuffle{M,N}RepeatPerShuffle`),
    `num_output_params` 16 → 21.
- **Wrw** adds
`DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle_Large_Tensor` with the
new
    conditional-/packed-param schema.
- `precision` gains INT8 in the KTN metadata too; `num_cu` / `arch` are
unchanged.

### Tests

- New golden test
`GPU_ConvNDAIHeuristics_FP32.ExtractTunaNet3dFeaturesGolden` pins the 3D
engineered TunaNet vector (layout / precision / direction one-hots + raw
block +
`EngineeredConvFeatures`) for a representative
`ConvHipImplicitGemm3DGroupFwdXdlops` shape, plus
  additional ND metadata/model coverage.
- The ND heuristics gtest is renamed `conv_ai_nd_heuristics.cpp` →
`conv_ai_nd_heuristics_test.cpp`
so the PR-bot Unit-Test policy recognizes it. Sources are collected via
`file(GLOB TEST_SOURCES *.cpp)`, so no CMake change is required, and
MIOpen's own GTest-naming
  hook only validates in-file suite names.
- The shared group-conv backward tolerance in
`test/gtest/group_conv.hpp` is raised **0.4% → 0.5%**.
The retrained 3D models expand the Wrw kernel vocabulary to include the
CK V3 large-tensor
variants (`CShuffleV3_Large_Tensor`, `TwoStage…_Large_Tensor`,
`…DirectLoad_MergedGroups`) that
develop's smaller vocab could not select. For a few bf16
backward-weights shapes the heuristic now
picks one of these — a fast but slightly less accurate kernel (~0.44%
RMS). This is an accepted
speed/accuracy trade for the new large-tensor variants (the AI heuristic
ranks on timing, not
numerical error), **not** a correctness regression; the chosen kernel
still passes verification
  within the widened bf16 tolerance.

## Risk Assessment

The change centers on the 3D grouped implicit-GEMM AI heuristic path on
gfx942/gfx950; the 2D
model assets are unchanged from develop, and other archs and non-grouped
solvers are untouched.
Note the shared `EngineerKernelConfigFeaturesImpl` missing-feature
convention (missing derived
input → `-1` sentinel) applies to the 2D KTN inference path as well, so
the 2D models see a
slightly different derived-feature encoding than develop; this affects
only *selection quality*
and is bounded by the fallback below. The AI heuristic is a predictive
*selector* wrapped in try/catch that falls
back to the legacy heuristic on any failure, and the chosen kernel still
goes through normal
applicability and (in Find) verification checks — so a misprediction or
an unencodable config
degrades to the fallback rather than producing wrong results. Any
regression would be in
*selection quality* (performance), bounded by the fallback, not in
correctness.

## Test Plan

3D performance comparison on **gfx950** over 1,000 random 3D
grouped-conv shapes:

- **Baseline:** exhaustive tuning (`MIOPEN_FIND_ENFORCE=3`).
- **Experimental:** fast-find with KTN heuristics
(`MIOPEN_FIND_MODE=2`).
- Kernel time recorded over 25 iterations per shape.
- Comparison restricted to the **989 shapes** where both runs produced
timings.

## Test Result

Relative to exhaustive tuning 3D gfx950, the heuristic has 90% cases
within 0.17ms, 99% cases within 2.29 ms of the running time of the
optimal kernel. 10 grouped INT8 forward shapes failed to produce
experimental timings due to a verification gap (missing naive reference
kernel), tracked in a separate bug report. 1 shape timed out in
exhaustive and heuristic run. This was over 1000 random shapes.

Relative to previous 3D heuristic on gfx950, the new heuristic is on
par, with some gains especially for bigger more time consuming shapes:
<img width="474" height="349" alt="image"
src="https://github.com/user-attachments/assets/f4d0f968-ec16-4a60-8b70-cd62839f9f45"
/>
Note that this is in addition to supporting new kernels and precisions.

For 2D — where the models are unchanged from develop but the shared C++
now applies the new missing-value (`-1` sentinel) convention to the 2D
KTN inference path — the measured performance is essentially identical
to before (New/Old = 1.0093). This **proves that the earlier concern
about the missingness mechanism is not a problem in practice**: feeding
the develop-trained 2D models the new derived-feature encoding does not
materially change 2D selection quality.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

---------

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants