Skip to content
Merged
Show file tree
Hide file tree
Changes from all 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
2 changes: 2 additions & 0 deletions projects/miopen/CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
* [Conv] Enabled Composable Kernel (CK) depthwise convolution on RDNA wave32 GPUs.

### Changed
* [Conv] Pruned ASM-GTC NHWC entries from the gfx908, gfx90a, gfx942, and gfx950 system find-databases for shapes the new large-tensor guard rejects, so immediate-mode lookups no longer resolve to a solver that is gated off at runtime. This removes 12,907 entries across the six databases, 4,668 of them rank-1, and empties 81 keys.
* [Conv] Enabled the Winograd Rage RxS f2x3 solver on gfx950 and updated the gfx942 kernels to v4_6_1/v4_9_1.
* [Conv] Updated the 3D AI heuristics (solver selection) for gfx942 and gfx950.

Expand Down Expand Up @@ -78,6 +79,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 size in bytes exceeds the int32 addressable 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
19 changes: 19 additions & 0 deletions projects/miopen/src/include/miopen/conv/asm_implicit_gemm.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,7 @@
#define CK_ASM_IMPLICITGEMM_HPP_

#include <miopen/config.h>
#include <miopen/conv/problem_description.hpp>

#include <string>
#include <cmath>
Expand Down Expand Up @@ -234,6 +235,24 @@ HeuristicInitMacroTileNoPadGemmK(size_t gemm_m,
return std::make_tuple(0, 0, 0);
}

// The ASM-GTC NHWC solvers originate in the (now unmaintained) MISA project and address
// global tensor memory with 32-bit BYTE offsets. AllTensorsDimsFitIntoInt() only validates
// that each individual length/stride fits in int32; it does NOT bound the flattened tensor
// size. A tensor whose total size in bytes exceeds the addressable range therefore overflows
// the kernel's 32-bit addressing and silently returns wrong results, so gate the solver off
// such shapes and let a large-tensor-capable solver be selected instead.
//
// The bound is expressed in BYTES, not elements: these kernels index memory by byte offset,
// so an element-count bound would admit shapes up to sizeof(T) times the addressable size.
static inline bool
igemm_tensor_size_exceeds_asm_gtc_addressing(const miopen::conv::ProblemDescription& problem)
{
constexpr std::size_t max_int32 = static_cast<std::size_t>(std::numeric_limits<int>::max());
return problem.GetIn().GetNumBytes() > max_int32 ||
problem.GetOut().GetNumBytes() > max_int32 ||
problem.GetWeights().GetNumBytes() > max_int32;
}

// This is to support big tensor > 4G. Need to decide how many splits needed.
// Return the number of splits.
static inline int igemm_split_batch_size(const int hi,
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,9 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsApplicable(
if(!problem.AllTensorsDimsFitIntoInt())
return false;

if(igemm_tensor_size_exceeds_asm_gtc_addressing(problem))
return false;

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,9 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsApplicable(
if(!problem.AllTensorsDimsFitIntoInt())
return false;

if(igemm_tensor_size_exceeds_asm_gtc_addressing(problem))
return false;

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,9 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsApplicable(
if(!problem.AllTensorsDimsFitIntoInt())
return false;

if(igemm_tensor_size_exceeds_asm_gtc_addressing(problem))
return false;

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