feat(conv): large-tensor grouped CK xdlops conv + split_k workspace backport

[JonathanLichtnerAMD]

JIRA ID: ROCM-27526

Motivation

MIOpen's grouped Composable Kernel (CK) xdlops convolution solvers
(ConvHipImplicitGemm{,3D}Group{Fwd,Bwd,Wrw}Xdlops) hard-rejected any problem
whose tensor dimensions or element strides did not fit in int32
(AllTensorsDimsFitIntoInt()). Large NHWC 2D conv shapes that cross the
INT_MAX element-count boundary therefore fell through to slower fallbacks or
had no applicable CK path at all. This is the ROCM-27526 large-tensor
enablement for the ROCm 7.0.2 backport branch, building on the ROCM-23997
large-tensor work.

Separately, the grouped bwd / wrw / 3D-wrw solvers previously sized the CK
workspace via GetCKAlphaBetaWorkspace (the alpha/beta backward-weights
workspace query), which does not account for CK's split_k reduction workspace.
Once a tuned split_k > 1 configuration was selected, the CK reduction
therefore ran against an under-allocated workspace, silently returning wrong
results (ALMIOPEN-2305). This PR backports the workspace portion of PR #1426 /
ALMIOPEN-706 to fix that, and adds test coverage for both the large-tensor and
the split_k-workspace behaviors.

Technical Details

This PR has three parts:

Large-tensor support for grouped CK xdlops conv.

• New src/include/miopen/solver/ck_grouped_conv_narrow.hpp: shared helpers for
  the six inline CKArgs structs. Defines ToCKIndexArray (int64→int32
  narrowing that asserts the value fits — guards the >INT_MAX contract), and
  the NarrowedCKArrays{2,3}D bundles plus MakeNarrowedCKArrays. The bundles
  are stored as mutable members of the owning CKArgs because CK's
  MakeArgumentPointer captures references into the length/stride arrays; a
  function-local temporary would be a stack-use-after-scope.
• Widened every CKArgs dim member (G,N,K,C,Hi,Wi,...) to int64_t and the
  length/stride arrays to std::array in all six solver
  .cpp files, so the NCHW packed-stride products (e.g. Hi*Wi*G*C) are
  computed in 64-bit before they can overflow.
• Added the large-tensor detection to implicitgemm_ck_util.hpp:
  RequiresLargeTensorCKInstance (true when !AllTensorsDimsFitIntoInt()) and
  IsLargeTensorCKInstance (matches "Large_Tensor" in the CK instance's type
  string). FillValidKernelsIDs, IsCKArgsSupported, and IsCKApplicable now
  filter to large-tensor instances for out-of-range shapes, and
  FillValidKernelsIDs may now legitimately return an empty list (asserts
  relaxed accordingly). PerformanceConfig::Init no longer indexes an empty
  valid_kernels.
• MakeArgPtr dispatches per instance: large-tensor instances bind CK's int64
  long_index_t overload with the un-narrowed members; all other instances go
  through GetNarrowedArrays() and bind the int32 overload. Forward,
  backward-data, and weight (2D + 3D) are all enabled, dispatching to CK's
  grouped-conv Large_Tensor int64 instances provided by the container CK
  (composable_kernel PR feat(grouped-conv): port large-tensor / global load-store support (#9258) composable_kernel#3755).
  The six solvers dropped their !AllTensorsDimsFitIntoInt() early-return in
  IsApplicable.
• convolutionocl.cpp hardening: both PrepareInvoker and CompileSolution
  now check that a coarsely-applicable solver actually produces a usable
  solution / is applicable to the concrete problem, throwing MIOPEN_THROW
  instead of dereferencing an unset invoker_factory (previously undefined
  behavior / crash for a large-tensor shape with no matching CK instance).
• CHANGELOG.md entry; a unit_TensorDescriptor.cpp case added to pin the
  exactly-INT_MAX vs one-past boundary (strict >).

Large 2D grouped conv coverage.

• unit_conv_solver_ConvHipImplicitGemmGroupXdlops_Large2D.cpp: on-device
  execute+verify (fwd/bwd/wrw) over large NHWC 2D grouped conv shapes spanning
  three regimes — sub-INT_MAX element counts; tensors above 2 GB whose element
  count still fits int32; and tensors whose element count exceeds INT_MAX
  (which require the CK large-tensor instances). All geometries are 3×3 with
  SAME padding, exercising the native NHWC path. Shapes are anonymized.
• conv_api_solution_count_large2d.cpp: fast applicability/compile probe over
  the same shapes.

CK split_k reduction workspace backport (PR #1426 / ALMIOPEN-706).

• implicitgemm_ck_util.hpp: added NextCKSplitkValue (which now includes CK's
  split_k = -1 autodeduce value, CkSplitkAutoDeduce) and
  GetCKSplitkMaxWorkspaceSize, which sweeps split_k per instance and queries CK
  for the real workspace size. For >INT_MAX problems this sweep is filtered to
  large-tensor instances (mirroring FillValidKernelsIDs / IsCKArgsSupported /
  IsCKApplicable), so it never drives a non-large instance through the int32
  narrowing path (ToCKIndexArray) — which otherwise asserts in a debug build at
  GetWorkspaceSize for a >INT_MAX-stride grouped-wrw shape.
  GetWorkspaceSizeLayoutTransformConv now accepts
  an optional CK workspace size and uses it in the NCHW/NHWC totals (falling
  back to GetCKAlphaBetaWorkspace only when CK reports 0); the invoker
  factories bind a correctly-sized workspace pointer (_ck_buff_des sized from
  GetCKSplitkWorkspaceSize) instead of the alpha/beta-only estimate.
• Added GetCKMaxWorkspaceSize to the grouped bwd, grouped wrw, and 3D grouped
  wrw solvers and wired it into their GetWorkspaceSize (declarations added in
  conv/solvers.hpp).
• Removed the generic-search assertion in generic_search.hpp that workspace
  size must be independent of the PerformanceConfig — with split_k it
  legitimately varies.
• Fixes silently-wrong grouped bwd/wrw results when a tuned split_k > 1 config
  needed a reduction workspace (ALMIOPEN-2305).

Risk Assessment

Roughly 3000 lines across 25 files, but ~1900 are new self-contained test
files. Behavioral source changes are concentrated in the six grouped CK xdlops
solvers plus two shared CK utility headers: the large-tensor path (fwd/bwd/wrw)
only activates for shapes that already fail AllTensorsDimsFitIntoInt() —
sub-INT_MAX shapes keep the identical int32 code path via the narrowing
helpers. The workspace backport is the highest-blast-radius piece because it
changes reported workspace sizes and invoker workspace binding for all grouped
bwd/wrw problems (not just large tensors), and it removes a generic-search
invariant.

Risk Level: 🟡 Medium

Impacted Components

• Grouped CK xdlops conv solvers (2D + 3D, fwd/bwd/wrw):
  conv_hip_implicit_gemm_{,3d_}grouped_{fwd,bwd,wrw}_xdlops.cpp.
• Shared CK utilities: implicitgemm_ck_util.hpp (workspace + large-tensor
  dispatch), new ck_grouped_conv_narrow.hpp.
• conv/solvers.hpp (new GetCKMaxWorkspaceSize declarations).
• generic_search.hpp (removed workspace-invariant assertion — affects the
  generic tuning loop broadly, not just these solvers).
• convolutionocl.cpp explicit-compile / prepare-invoker paths.
• Tests + CHANGELOG.

Potential Side Effects

• Grouped bwd/wrw problems now report a larger GetWorkspaceSize (includes the
  CK reduction workspace); callers that pre-size workspace buffers will see the
  increase. This is the intended fix.
• Widening CKArgs members to int64 is arithmetically transparent for
  in-range shapes but is a struct-layout change across the six solvers.
• Removing the generic-search workspace-invariant assertion means a genuine
  future workspace-sizing bug would no longer be caught by that check.
• The convolutionocl.cpp guards convert some previously-UB/crash paths into
  clean miopenStatus* throws — a behavior change for callers that hit an
  inapplicable/unsatisfiable solver by explicit ID.

Mitigation Steps

• Sub-INT_MAX shapes are unchanged: they take the narrowing path and bind the
  same int32 CK overload as before; ToCKIndexArray asserts if the contract is
  ever bypassed.
• Workspace backport validated by an explicit before/after A/B (see Test
  Result).
• Heavyweight numerical tests (*LargeStride*, *Large2D*) are excluded from
  the default gtest filter so pre-checkin cost is unaffected; run explicitly.

Test Plan

• API applicability sweeps bracketing the INT_MAX element-stride boundary:
  conv_api_solution_count_{2d,3d}_large_stride.cpp and
  conv_api_solution_count_large2d.cpp (shared scaffolding in
  conv_api_solution_count_large_stride_common.hpp), restricted to a CI arch
  allowlist (gfx90A / gfx94X / gfx950).
• Numerical execute+verify unit tests for the large-stride fwd/bwd/wrw path
  (unit_conv_solver_*_LargeStride.cpp) and the large 2D NHWC path
  (unit_conv_solver_ConvHipImplicitGemmGroupXdlops_Large2D.cpp) across the
  three element-count regimes.
• unit_TensorDescriptor.cpp boundary cases at exactly INT_MAX and one past.
• Explicit A/B on the grouped bwd numerical tests to confirm the split_k
  workspace backport.

Test Result

• Forward, sub-INT_MAX, and the >2 GB-but-element-count-fits-int32 cases pass
  on gfx942-class hardware.
• The split_k workspace backport was confirmed by a before/after A/B: on
  int32-safe shapes, workspace-requiring split_k > 1 CK configurations are now
  selected and produce correct results, where before they returned silently
  wrong results (ALMIOPEN-2305).
• Additional shapes have been manually tested in combination with composable_kernel PR https://github.com/ROCm/ composable_kernel/pull/3755).
• One high-batch wrw case overflows the fp16 numerical reference — a
  test-harness precision artifact of the reference computation, not a solver
  bug.

[SreecharanGundaboluAMD]

LGTM