Skip to content
Merged
Show file tree
Hide file tree
Changes from 19 commits
Commits
Show all changes
23 commits
Select commit Hold shift + click to select a range
311d11e
fix(conv): gate ASM-GTC fwd NHWC solver off int32-overflow tensors
yassinsolim Jul 17, 2026
1eba5bc
fix(conv): prune ASM-GTC fwd NHWC from find-db entries over INT_MAX
yassinsolim Jul 17, 2026
6989dbe
fix(conv): gate ASM-GTC bwd NHWC solver off int32-overflow tensors
yassinsolim Jul 31, 2026
afcd08a
fix(conv): gate ASM-GTC wrw NHWC solver off int32-overflow tensors
yassinsolim Jul 31, 2026
48fd8c5
test(conv): applicability test for the ASM-GTC bwd NHWC INT_MAX gate
yassinsolim Jul 31, 2026
f9b183f
test(conv): applicability test for the ASM-GTC wrw NHWC INT_MAX gate
yassinsolim Jul 31, 2026
33a00f0
docs(changelog): cover bwd and wrw in the ASM-GTC INT_MAX gate entry
yassinsolim Jul 31, 2026
95825d1
test(conv): consolidate ASM-GTC large-tensor tests into one file
yassinsolim Jul 31, 2026
24f7ebb
test(conv): drop test_conv_asm_gtc_fwd_nhwc_large_tensor.cpp, superse…
yassinsolim Jul 31, 2026
c270ba5
test(conv): drop test_conv_asm_gtc_bwd_nhwc_large_tensor.cpp, superse…
yassinsolim Jul 31, 2026
77b9e33
test(conv): drop test_conv_asm_gtc_wrw_nhwc_large_tensor.cpp, superse…
yassinsolim Jul 31, 2026
4f198de
fix(conv): DBSync prune gfx90a68 (auto)
yassinsolim Aug 5, 2026
8393599
fix(conv): DBSync prune gfx90a68 (auto)
yassinsolim Aug 5, 2026
e1ff811
fix(conv): DBSync prune gfx90a6e (auto)
yassinsolim Aug 5, 2026
c376df7
fix(conv): compare bytes, not elements, in ASM-GTC NHWC overflow gates
yassinsolim Aug 19, 2026
03ef09d
fix(db): drop stale ASM-GTC NHWC find-db rows above the byte-based ov…
yassinsolim Aug 20, 2026
eb5a76b
Merge branch 'develop' into users/ysoliman/gate-asm-gtc-fwd-large-tensor
yassinsolim Aug 21, 2026
8e1eee3
Merge branch 'develop' into users/ysoliman/gate-asm-gtc-fwd-large-tensor
yassinsolim Aug 25, 2026
ce70fff
Merge remote-tracking branch 'refs/remotes/origin/develop-x' into mrg…
yassinsolim Aug 26, 2026
7488e45
refactor(miopen): share the ASM-GTC large-tensor guard across fwd/bwd…
yassinsolim Sep 2, 2026
bc3ffa3
Merge branch 'develop' into gate-asm-gtc-fwd-large-tensor
yassinsolim Sep 2, 2026
8d2dd81
Merge branch 'develop' into gate-asm-gtc-fwd-large-tensor
yassinsolim Sep 3, 2026
f37181b
docs(miopen): document the ASM-GTC find-db prune under Changed
yassinsolim Sep 4, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions projects/miopen/CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -41,6 +41,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 NHWC solvers (`ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC`, `ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC` and `ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC`) off tensors whose element count exceeds the int32 range, preventing them 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.
Expand Down
Binary file modified projects/miopen/src/kernels/gfx90878.HIP.fdb.txt.bz2
Binary file not shown.
Binary file modified projects/miopen/src/kernels/gfx90a68.HIP.fdb.txt.bz2
Binary file not shown.
Binary file modified projects/miopen/src/kernels/gfx90a6e.HIP.fdb.txt.bz2
Binary file not shown.
Binary file modified projects/miopen/src/kernels/gfx942130.HIP.fdb.txt.bz2
Binary file not shown.
Binary file modified projects/miopen/src/kernels/gfx942e4.HIP.fdb.txt.bz2
Binary file not shown.
Binary file modified projects/miopen/src/kernels/gfx950100.HIP.fdb.txt.bz2
Binary file not shown.
Original file line number Diff line number Diff line change
Expand Up @@ -1000,6 +1000,21 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::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().GetNumBytes() > max_int32 ||
problem.GetOut().GetNumBytes() > max_int32 ||
problem.GetWeights().GetNumBytes() > max_int32)
return false;
}
Comment thread
yassinsolim marked this conversation as resolved.
Resolved

if(!problem.IsFp32() && !problem.IsFp16() &&
!(problem.IsBfp16() &&
(device_name == "gfx90a" || device_name == "gfx942" || StartsWith(device_name, "gfx95"))))
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -923,6 +923,21 @@ 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().GetNumBytes() > max_int32 ||
problem.GetOut().GetNumBytes() > max_int32 ||
problem.GetWeights().GetNumBytes() > max_int32)
return false;
Comment thread
yassinsolim marked this conversation as resolved.
Resolved
}

if(!problem.IsFp32() && !problem.IsFp16() &&
!(problem.IsBfp16() &&
(device_name == "gfx90a" || device_name == "gfx942" || StartsWith(device_name, "gfx95"))))
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -892,6 +892,21 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::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().GetNumBytes() > max_int32 ||
problem.GetOut().GetNumBytes() > max_int32 ||
problem.GetWeights().GetNumBytes() > max_int32)
return false;
Comment thread
yassinsolim marked this conversation as resolved.
Resolved
}

if(!problem.IsFp32() && !problem.IsFp16() &&
!(problem.IsBfp16() &&
(device_name == "gfx90a" || device_name == "gfx942" || StartsWith(device_name, "gfx95"))))
Expand Down
144 changes: 144 additions & 0 deletions projects/miopen/test/gtest/conv_asm_gtc_nhwc_large_tensor.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,144 @@
/*******************************************************************************
*
* 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"

// The ASM-GTC NHWC solvers ConvAsmImplicitGemmGTCDynamic{Fwd,Bwd,Wrw}XdlopsNHWC index
// global tensor memory with 32-bit BYTE offsets and do not implement large-tensor
// support. Each solver's IsApplicable() therefore gates off any problem whose tensor
// size in bytes exceeds INT_MAX, so it cannot be selected for a shape it would silently
// compute incorrectly. Note the bound is on bytes, not elements: for fp16 the effective
// element ceiling is half of INT_MAX.
//
// The two cases below share geometry and differ only in batch size N, so the byte
// count crosses the INT_MAX boundary and isolates the gate. The gate checks the
// (direction-independent) in/out/weights descriptors; for this shape both the input and
// output tensors have C*H*W per sample = 1024*162*92 = 15,261,696 elements, and fp16
// doubles that to bytes:
// N=70 -> 2,136,637,440 bytes <= INT_MAX (2,147,483,647): int32-safe, applicable
// N=71 -> 2,167,160,832 bytes > INT_MAX : gated off, not applicable

namespace {

auto GetInRangeConvCase()
{
using miopen::unit_tests::ConvTestCase;
return ConvTestCase{{miopenHalf, miopenTensorNHWC, {70, 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, {71, 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;
}

// byte 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

// --- Forward -------------------------------------------------------------------------

using CPU_UnitTestConvSolverAsmGTCFwdNHWCLargeTensorDevApplicability_NONE =
CPU_UnitTestConvSolverDevApplicabilityFwd_NONE;

TEST_P(CPU_UnitTestConvSolverAsmGTCFwdNHWCLargeTensorDevApplicability_NONE,
ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC)
{
this->RunTest(miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC{});
};

INSTANTIATE_TEST_SUITE_P(SmokeInt32Safe,
CPU_UnitTestConvSolverAsmGTCFwdNHWCLargeTensorDevApplicability_NONE,
testing::Combine(testing::Values(GetInRangeParams()),
testing::Values(GetInRangeConvCase())));

INSTANTIATE_TEST_SUITE_P(SmokeOverInt32,
CPU_UnitTestConvSolverAsmGTCFwdNHWCLargeTensorDevApplicability_NONE,
testing::Combine(testing::Values(GetGatedParams()),
testing::Values(GetOverInt32ConvCase())));

// --- Backward-data -------------------------------------------------------------------

using CPU_UnitTestConvSolverAsmGTCBwdNHWCLargeTensorDevApplicability_NONE =
CPU_UnitTestConvSolverDevApplicabilityBwd_NONE;

TEST_P(CPU_UnitTestConvSolverAsmGTCBwdNHWCLargeTensorDevApplicability_NONE,
ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC)
{
this->RunTest(miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC{});
};

INSTANTIATE_TEST_SUITE_P(SmokeInt32Safe,
CPU_UnitTestConvSolverAsmGTCBwdNHWCLargeTensorDevApplicability_NONE,
testing::Combine(testing::Values(GetInRangeParams()),
testing::Values(GetInRangeConvCase())));

INSTANTIATE_TEST_SUITE_P(SmokeOverInt32,
CPU_UnitTestConvSolverAsmGTCBwdNHWCLargeTensorDevApplicability_NONE,
testing::Combine(testing::Values(GetGatedParams()),
testing::Values(GetOverInt32ConvCase())));

// --- Backward-weights ----------------------------------------------------------------

using CPU_UnitTestConvSolverAsmGTCWrwNHWCLargeTensorDevApplicability_NONE =
CPU_UnitTestConvSolverDevApplicabilityWrw_NONE;

TEST_P(CPU_UnitTestConvSolverAsmGTCWrwNHWCLargeTensorDevApplicability_NONE,
ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC)
{
this->RunTest(miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC{});
};

INSTANTIATE_TEST_SUITE_P(SmokeInt32Safe,
CPU_UnitTestConvSolverAsmGTCWrwNHWCLargeTensorDevApplicability_NONE,
testing::Combine(testing::Values(GetInRangeParams()),
testing::Values(GetInRangeConvCase())));

INSTANTIATE_TEST_SUITE_P(SmokeOverInt32,
CPU_UnitTestConvSolverAsmGTCWrwNHWCLargeTensorDevApplicability_NONE,
testing::Combine(testing::Values(GetGatedParams()),
testing::Values(GetOverInt32ConvCase())));
Loading