[MIOpen] Enable grouped CK xdlops solvers for large-stride tensors (ROCM-23997)

[JonathanLichtnerAMD]

Motivation

PyTorch started dispatching large convolutions to MIOpen as of pytorch 2.11/2.12, but the MIOpen support is limited and requires additional work on the MIOpen side to fully support the large convolutions. In particular, the six grouped CK xdlops solvers (ConvHipImplicitGemmGroup{Fwd,Bwd,Wrw}Xdlops and their 3D counterparts) are marked inapplicable for convolutions whose tensor element-strides exceed the int32 range, such as (1,96,4736,4736) or (1,96,1024,1024,1024). Without these solvers, the remaining applicable options are limited to naive kernels (slow) and explicit GEMM (which can require very large workspaces).

As per ROCM-23997, this PR removes the int32 element-stride applicability restriction on the grouped CK xdlops solvers so that MIOpen can serve these shapes. Requires the matching CK changes in #7734 to be merged first.

Technical Details

• Drop 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), 2D Fwd CKArgs, and all three 3D CKArgs templates. The NCHW stride builder is the prime overflow site — operands are now int64 so the contiguous stride no longer wraps before assignment. CKConvDims in the common helpers is widened in lock step, and 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, IsCKApplicable, IsCKArgsSupported, IsCKSplitKSupported, GetCKSplitkMaxWorkspaceSize). 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. When no large-tensor instance is registered for a given shape, the solver is correctly marked inapplicable.
• For CK Bwd/Wrw, the int32 narrowed length/stride arrays handed to MakeArgumentPointer are stored as a mutable member of CKArgs (2D CKArgsSplitK base and 3D CKArgs templates) rather than as locals in MakeArgPtr, because CK captures references into those arrays and the locals went out of scope before IsSupportedArgument ran. The narrowed bundle is populated lazily from MakeArgPtr rather than in the CKArgs constructor, so the debug-only overflow assert in ToCKIndexArray only runs for kernels that survived the large-tensor filter (CKArgs is constructed unconditionally in FillValidKernelsIDs before the filter).
• projects/miopen/CHANGELOG.md: add Resolved-issues bullet.

Risk Assessment

Medium-low. Changes are scoped to the six grouped CK xdlops solvers and their shared CK helpers; widening from int → int64_t / std::size_t is additive at every call site verified. The new large-tensor instance filter guarantees the solver is correctly marked inapplicable when no large-tensor CK instance is registered, so existing shapes continue to dispatch to the same kernels as before.

Test Plan

Automated Tests:

• API-level solution-count gtests (conv_api_solution_count_{2d,3d}_large_stride) 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. Gated to gfx90A | gfx94X | gfx950 via IsTestSupportedByDevice.
• Kernel-launching numerical tests for all six solvers on shapes whose element-strides exceed INT_MAX (2D: 1x96x4736x4736; 3D: 1x96x512x512x88), so the CPU reference compare catches int32 wrap-around inside the CK kernel itself. One unit_conv_solver_ConvHipImplicitGemm{,3D}Group{Fwd,Bwd,Wrw}Xdlops_LargeStride.cpp per solver. Excluded from standard/pre-checkin via test_categories.yaml (~4–5+ GB each) and run in comprehensive/full; gated to gfx90A | gfx94X | gfx950.
• 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.

Manual tests:

• Run ASAN against the new tests
• Run the UnitTestConvSolverGroupXDlops large-stride tests manually on gfx90a
• Run the UnitTestConvSolverGroupXDlops large-stride tests manually on gfx942
• Run the UnitTestConvSolverGroupXDlops large-stride tests manually on gfx950

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
• [ci] Per-category timeouts for test matrix; migrate MIOpen TheRock#5507 has been merged and rocm-libraries bumped (or the slow tests disabled)
• also need CK PR [CK] Grouped Convolution Global Load/Store instances #7734

[JonathanLichtnerAMD]

This PR can be be merged once the associated CK changes in PR #7734 are merged. (Note that the added tests will prevent this from being accidentally merged.)

[JonathanLichtnerAMD]

This might also need ROCm/TheRock#5507 or there might be a risk of test timeouts on nightlies and rock bumps.