fix(conv): gate ASM-GTC fwd NHWC solver off int32-overflow tensors

ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC (the ASM-GTC forward solver
from the unmaintained MISA project) indexes global tensor memory with
32-bit element indices and does not implement large-tensor support.

Its IsApplicable() relied on AllTensorsDimsFitIntoInt(), which only
checks that each individual length/stride fits in int32 -- it never
bounds the flattened element count. A tensor whose total element count
exceeds INT_MAX therefore overflows the kernel's 32-bit indexing and is
silently computed incorrectly (observed forward verification error
~0.37 on gfx942 and gfx950).

Gate the solver not-applicable when any tensor's element count exceeds
INT_MAX, so a large-tensor-capable solver (e.g. grouped CK xdlops) is
selected instead. This is a no-op for every int32-safe shape, so
existing workloads are unaffected.

Adds a CPU device-applicability unit test that shares geometry across a
batch-size boundary so the element count crosses INT_MAX, asserting the
solver stays applicable just below the limit and is gated off just above.

Related to ROCM-27526.

Author: yassinsolim

projects/miopen/CHANGELOG.md

@@ -19,6 +19,7 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 ### Resolved Issues
 * [RNN] Fix RNN workspace tensor descriptor int overflow
 * [Conv] Enabled grouped Composable Kernel (CK) xdlops fwd, bwd, and wrw convolution (2D and 3D) for tensors whose strides exceed the int32 range.
+* [Conv] Gated the unmaintained ASM-GTC forward NHWC solver (`ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC`) off tensors whose element count exceeds the int32 range, preventing it from being selected and silently returning incorrect results on large convolutions.
 * [Conv] Fixed `miopenStatusInternalError` thrown by Find on depthwise NHWC grouped convolutions under `MIOPEN_FIND_MODE=NORMAL`.
 * Fixed a thread-safety issue where concurrent access to the AI-heuristic model caches was not guarded by a mutex.
 * [Conv] Fixed Composable Kernel (CK) grouped-convolution solvers not always being registered in host code.

projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp

@@ -923,6 +923,22 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsApplicable(
     if(!problem.AllTensorsDimsFitIntoInt())
         return false;
 
+    // This solver originates in the (now unmaintained) MISA project and addresses
+    // global tensor memory with 32-bit element indexing. AllTensorsDimsFitIntoInt()
+    // above only validates that each individual length/stride fits in int32; it does
+    // NOT bound the flattened element count. A tensor whose total element count
+    // exceeds INT_MAX therefore overflows the kernel's 32-bit indexing and silently
+    // returns wrong results. MISA does not implement large-tensor support, so gate the
+    // solver off such shapes; a large-tensor-capable solver is selected instead.
+    {
+        constexpr std::size_t max_int32 =
+            static_cast<std::size_t>(std::numeric_limits<int>::max());
+        if(problem.GetIn().GetElementSize() > max_int32 ||
+           problem.GetOut().GetElementSize() > max_int32 ||
+           problem.GetWeights().GetElementSize() > max_int32)
+            return false;
+    }
+
     if(!problem.IsFp32() && !problem.IsFp16() &&
        !(problem.IsBfp16() &&
          (device_name == "gfx90a" || device_name == "gfx942" || StartsWith(device_name, "gfx95"))))

projects/miopen/test/gtest/unit_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC_LargeTensor.cpp

@@ -0,0 +1,97 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2026 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#include "unit_conv_solver.hpp"
+
+// ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC (the MISA ASM-GTC forward solver)
+// indexes global tensor memory with 32-bit element indices and does not implement
+// large-tensor support. IsApplicable() therefore gates the solver off any problem
+// whose flattened element count exceeds INT_MAX, so it cannot be selected for a
+// shape it would silently compute incorrectly.
+//
+// The two cases below share geometry and differ only in batch size N, so the
+// element count crosses the INT_MAX boundary and isolates the gate:
+//   C*H*W per sample = 1024*162*92 = 15,261,696 elements
+//     N=140 -> 2,136,637,440 <= INT_MAX (2,147,483,647): int32-safe, applicable
+//     N=141 -> 2,151,899,136 >  INT_MAX               : gated off, not applicable
+
+namespace {
+
+auto GetInRangeConvCase()
+{
+    using miopen::unit_tests::ConvTestCase;
+    return ConvTestCase{{miopenHalf, miopenTensorNHWC, {140, 1024, 162, 92}},
+                        {miopenHalf, miopenTensorNHWC, {1024, 1024, 3, 3}},
+                        miopenHalf,
+                        {{1, 1}, {1, 1}, {1, 1}}};
+}
+
+auto GetOverInt32ConvCase()
+{
+    using miopen::unit_tests::ConvTestCase;
+    return ConvTestCase{{miopenHalf, miopenTensorNHWC, {141, 1024, 162, 92}},
+                        {miopenHalf, miopenTensorNHWC, {1024, 1024, 3, 3}},
+                        miopenHalf,
+                        {{1, 1}, {1, 1}, {1, 1}}};
+}
+
+// int32-safe shape: the solver keeps its usual device applicability (regression guard).
+auto GetInRangeParams()
+{
+    auto p = miopen::unit_tests::UnitTestConvSolverParams(Gpu::gfx908 | Gpu::gfx90A | Gpu::gfx94X |
+                                                          Gpu::gfx950);
+    p.CheckXnackDisabled();
+    return p;
+}
+
+// element count > INT_MAX: the solver must be applicable on no device (gated off).
+auto GetGatedParams()
+{
+    auto p = miopen::unit_tests::UnitTestConvSolverParams(Gpu::None);
+    p.CheckXnackDisabled();
+    return p;
+}
+
+} // namespace
+
+using CPU_UnitTestConvSolverAsmGTCFwdNHWCLargeTensorDevApplicability_NONE =
+    CPU_UnitTestConvSolverDevApplicabilityFwd_NONE;
+
+TEST_P(CPU_UnitTestConvSolverAsmGTCFwdNHWCLargeTensorDevApplicability_NONE,
+       ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC)
+{
+    this->RunTest(miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC{});
+};
+
+INSTANTIATE_TEST_SUITE_P(Int32Safe,
+                         CPU_UnitTestConvSolverAsmGTCFwdNHWCLargeTensorDevApplicability_NONE,
+                         testing::Combine(testing::Values(GetInRangeParams()),
+                                          testing::Values(GetInRangeConvCase())));
+
+INSTANTIATE_TEST_SUITE_P(OverInt32ElemCount,
+                         CPU_UnitTestConvSolverAsmGTCFwdNHWCLargeTensorDevApplicability_NONE,
+                         testing::Combine(testing::Values(GetGatedParams()),
+                                          testing::Values(GetOverInt32ConvCase())));