Skip to content

Commit a1e0e6f

Browse files
JonathanLichtnerAMDclaude
JonathanLichtnerAMD
and
claude
committed
[Conv] Add large 2D grouped conv coverage (ROCM-27526)
Adds NHWC 2D grouped convolution coverage for the grouped CK xdlops solvers across 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 3x3 with SAME padding, exercising the native NHWC path. - conv_api_solution_count_large2d.cpp: fast applicability/compile probe (fwd/bwd/wrw) for all shapes. - unit_conv_solver_ConvHipImplicitGemmGroupXdlops_Large2D.cpp: on-device execute+verify coverage; excluded from the default gtest filter as multi-GB (run explicitly). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
1 parent 512edb4 commit a1e0e6f

3 files changed

Lines changed: 336 additions & 2 deletions

File tree

‎projects/miopen/test/gtest/CMakeLists.txt‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -102,9 +102,10 @@ function(add_gtest TEST_NAME TEST_CPP)
102102
add_gtest_negative_filter("*GPU*")
103103
endif()
104104

105-
# Heavyweight (multi-GB) large-stride numerical tests (ROCM-23997); excluded from the
106-
# default filter so they don't run in standard/pre-checkin. Run explicitly when needed.
105+
# Heavyweight (multi-GB) numerical conv tests; excluded from the default
106+
# filter so they don't run in standard/pre-checkin. Run explicitly when needed.
107107
add_gtest_negative_filter("*UnitTestConvSolver*LargeStride*")
108+
add_gtest_negative_filter("*UnitTestConvSolver*Large2D*")
108109

109110
# NONE datatype is explicitly enabled for all the state, all the NONE tests take 4 seconds to complete.
110111
set(MIOPEN_GTEST_FILTER "${MIOPEN_GTEST_PREFIX}*${MIOPEN_GTEST_SUFFIX}*:${MIOPEN_GTEST_PREFIX}*NONE*${MIOPEN_GTEST_FILTER_NEGATIVE}")

‎projects/miopen/test/gtest/conv_api_solution_count_large2d.cpp‎

Lines changed: 120 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,120 @@
1+
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
2+
// SPDX-License-Identifier: MIT
3+
//
4+
// API-level applicability / compile probe for large 2D grouped convolutions in
5+
// NHWC layout. Shapes span three regimes: (a) sub-INT_MAX element counts,
6+
// (b) tensors above 2 GB whose element count still fits int32, and (c) tensors
7+
// whose element count exceeds INT_MAX (needs the CK large-tensor instances).
8+
// All geometries are 3x3 with SAME padding. Fast: CompileSolution only, no
9+
// allocation or execution.
10+
//
11+
// fwd/bwd/wrw all assert a compilable CK solution, so any shape/direction the
12+
// solver does not cover shows up as a FAILED case.
13+
14+
#include "conv_api_solution_count_large_stride_common.hpp"
15+
#include <vector>
16+
17+
namespace {
18+
19+
using miopen_test_large_stride::Descriptors;
20+
using miopen_test_large_stride::RunCompileBwdData;
21+
using miopen_test_large_stride::RunCompileFwd;
22+
using miopen_test_large_stride::RunCompileWrw;
23+
24+
// A 2D conv problem: input [N,C,H,W], weight [K,C,kH,kW], pad (pH,pW). Lengths
25+
// are in logical [N,C,H,W] order; all descriptors are created NHWC.
26+
struct Shape2D
27+
{
28+
const char* tag;
29+
int x[4];
30+
int w[4];
31+
int pad[2];
32+
};
33+
34+
std::vector<Shape2D> Shapes()
35+
{
36+
return {
37+
// ---- sub-INT_MAX elements, < 2 GB ----
38+
{"small_hc", {48, 2048, 64, 64}, {2048, 2048, 3, 3}, {1, 1}}, // 4.0e8 elem
39+
// ---- > 2 GB bytes, element count still fits int32 ----
40+
{"gt2gb", {96, 1024, 112, 112}, {1024, 1024, 3, 3}, {1, 1}}, // 1.23e9 elem
41+
{"gt2gb_kexp", {96, 2048, 64, 64}, {4096, 2048, 3, 3}, {1, 1}}, // 1.61e9 elem, K=2C
42+
// ---- element count > INT_MAX (needs CK large-tensor instances) ----
43+
{"over_intmax", {160, 1024, 140, 100}, {1024, 1024, 3, 3}, {1, 1}}, // 2.29e9 elem
44+
{"over_intmax_swap", {160, 1024, 100, 140}, {1024, 1024, 3, 3}, {1, 1}}, // 2.29e9, H/W swap
45+
{"over_intmax_spatial", {200, 512, 200, 200}, {512, 512, 3, 3}, {1, 1}}, // 4.10e9 elem
46+
{"over_intmax_kexp", {160, 1024, 120, 120}, {2048, 1024, 3, 3}, {1, 1}}, // 4.72e9, K=2C
47+
};
48+
}
49+
50+
::testing::AssertionResult SetupShape2D(const Shape2D& s, miopenDataType_t dtype, Descriptors& d)
51+
{
52+
constexpr int rank = 4;
53+
54+
if(miopenCreateWithStream(&d.handle, /*stream=*/nullptr) != miopenStatusSuccess)
55+
return ::testing::AssertionFailure() << "miopenCreateWithStream failed";
56+
57+
if(miopenCreateTensorDescriptor(&d.xDesc) != miopenStatusSuccess)
58+
return ::testing::AssertionFailure() << "create xDesc failed";
59+
if(miopenSetNdTensorDescriptorWithLayout(d.xDesc, dtype, miopenTensorNHWC, s.x, rank) !=
60+
miopenStatusSuccess)
61+
return ::testing::AssertionFailure() << "set xDesc failed";
62+
63+
if(miopenCreateTensorDescriptor(&d.wDesc) != miopenStatusSuccess)
64+
return ::testing::AssertionFailure() << "create wDesc failed";
65+
if(miopenSetNdTensorDescriptorWithLayout(d.wDesc, dtype, miopenTensorNHWC, s.w, rank) !=
66+
miopenStatusSuccess)
67+
return ::testing::AssertionFailure() << "set wDesc failed";
68+
69+
if(miopenCreateConvolutionDescriptor(&d.convDesc) != miopenStatusSuccess)
70+
return ::testing::AssertionFailure() << "create convDesc failed";
71+
{
72+
int pads[2] = {s.pad[0], s.pad[1]};
73+
int strides[2] = {1, 1};
74+
int dils[2] = {1, 1};
75+
if(miopenInitConvolutionNdDescriptor(
76+
d.convDesc, 2, pads, strides, dils, miopenConvolution) != miopenStatusSuccess)
77+
return ::testing::AssertionFailure() << "init convDesc failed";
78+
}
79+
80+
if(miopenCreateTensorDescriptor(&d.yDesc) != miopenStatusSuccess)
81+
return ::testing::AssertionFailure() << "create yDesc failed";
82+
{
83+
int yDim[rank] = {0};
84+
int yNbDims = 0;
85+
if(miopenGetConvolutionNdForwardOutputDim(d.convDesc, d.xDesc, d.wDesc, &yNbDims, yDim) !=
86+
miopenStatusSuccess)
87+
return ::testing::AssertionFailure() << "get yDim failed";
88+
if(yNbDims != rank)
89+
return ::testing::AssertionFailure() << "yNbDims != " << rank;
90+
if(miopenSetNdTensorDescriptorWithLayout(d.yDesc, dtype, miopenTensorNHWC, yDim, rank) !=
91+
miopenStatusSuccess)
92+
return ::testing::AssertionFailure() << "set yDesc failed";
93+
}
94+
return ::testing::AssertionSuccess();
95+
}
96+
97+
void RunFwd(const Shape2D& s, miopenDataType_t dtype)
98+
{
99+
RunCompileFwd(s, dtype, SetupShape2D, "ConvHipImplicitGemmGroupFwdXdlops");
100+
}
101+
void RunBwd(const Shape2D& s, miopenDataType_t dtype)
102+
{
103+
RunCompileBwdData(s, dtype, SetupShape2D, "ConvHipImplicitGemmGroupBwdXdlops");
104+
}
105+
void RunWrw(const Shape2D& s, miopenDataType_t dtype)
106+
{
107+
RunCompileWrw(s, dtype, SetupShape2D, "ConvHipImplicitGemmGroupWrwXdlops");
108+
}
109+
110+
class GPU_ConvApi_Large2D_FP16 : public ::testing::TestWithParam<Shape2D>
111+
{
112+
};
113+
114+
} // namespace
115+
116+
TEST_P(GPU_ConvApi_Large2D_FP16, FwdIncludesCK) { RunFwd(GetParam(), miopenHalf); }
117+
TEST_P(GPU_ConvApi_Large2D_FP16, BwdDataIncludesCK) { RunBwd(GetParam(), miopenHalf); }
118+
TEST_P(GPU_ConvApi_Large2D_FP16, WrwIncludesCK) { RunWrw(GetParam(), miopenHalf); }
119+
120+
INSTANTIATE_TEST_SUITE_P(Large2D, GPU_ConvApi_Large2D_FP16, ::testing::ValuesIn(Shapes()));

‎projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupXdlops_Large2D.cpp‎

Lines changed: 213 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,213 @@
1+
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
2+
// SPDX-License-Identifier: MIT
3+
//
4+
// Numerical (on-device execute + verify) coverage for large 2D grouped
5+
// convolutions in NHWC layout. Shapes span three regimes: (a) sub-INT_MAX
6+
// element counts, (b) tensors above 2 GB whose element count still fits int32,
7+
// and (c) tensors whose element count exceeds INT_MAX (needs the CK
8+
// large-tensor instances). All geometries are 3x3 with SAME padding.
9+
// Complements the API probe in conv_api_solution_count_large2d.cpp; catches any
10+
// silent int32 offset overflow on the regular (non-large-tensor) CK path.
11+
//
12+
// Excluded from the default gtest filter (test/gtest/CMakeLists.txt) as
13+
// multi-GB; run explicitly. A case skips when the solver is inapplicable or the
14+
// device lacks memory.
15+
16+
#include <algorithm>
17+
#include <cstddef>
18+
#include <vector>
19+
20+
#include "unit_conv_solver.hpp"
21+
#include "get_handle.hpp"
22+
23+
namespace {
24+
25+
using miopen::unit_tests::ConvolutionDescriptorParams;
26+
using miopen::unit_tests::ConvTestCase;
27+
using miopen::unit_tests::TensorDescriptorParams;
28+
29+
std::vector<ConvTestCase> GetTestCases(miopenDataType_t dt)
30+
{
31+
const auto L = miopenTensorNHWC;
32+
// Lengths are in logical [N, C, H, W] order; L selects NHWC memory layout.
33+
// Filters are [K, C, 3, 3], SAME padding. Element counts annotated per case.
34+
return {
35+
// ---- sub-INT_MAX elements, < 2 GB ----
36+
{TensorDescriptorParams{dt, L, {48, 2048, 64, 64}},
37+
TensorDescriptorParams{dt, L, {2048, 2048, 3, 3}},
38+
dt,
39+
ConvolutionDescriptorParams{{1, 1}, {1, 1}, {1, 1}, 1}}, // 4.0e8 elem
40+
// ---- > 2 GB bytes, element count still fits int32 ----
41+
{TensorDescriptorParams{dt, L, {96, 1024, 112, 112}},
42+
TensorDescriptorParams{dt, L, {1024, 1024, 3, 3}},
43+
dt,
44+
ConvolutionDescriptorParams{{1, 1}, {1, 1}, {1, 1}, 1}}, // 1.23e9 elem
45+
{TensorDescriptorParams{dt, L, {96, 2048, 64, 64}},
46+
TensorDescriptorParams{dt, L, {4096, 2048, 3, 3}},
47+
dt,
48+
ConvolutionDescriptorParams{{1, 1}, {1, 1}, {1, 1}, 1}}, // 1.61e9 elem, K=2C
49+
// ---- element count > INT_MAX (needs CK large-tensor instances) ----
50+
{TensorDescriptorParams{dt, L, {160, 1024, 140, 100}},
51+
TensorDescriptorParams{dt, L, {1024, 1024, 3, 3}},
52+
dt,
53+
ConvolutionDescriptorParams{{1, 1}, {1, 1}, {1, 1}, 1}}, // 2.29e9 elem
54+
{TensorDescriptorParams{dt, L, {160, 1024, 100, 140}},
55+
TensorDescriptorParams{dt, L, {1024, 1024, 3, 3}},
56+
dt,
57+
ConvolutionDescriptorParams{{1, 1}, {1, 1}, {1, 1}, 1}}, // 2.29e9 elem, H/W swap
58+
{TensorDescriptorParams{dt, L, {200, 512, 200, 200}},
59+
TensorDescriptorParams{dt, L, {512, 512, 3, 3}},
60+
dt,
61+
ConvolutionDescriptorParams{{1, 1}, {1, 1}, {1, 1}, 1}}, // 4.10e9 elem, large spatial
62+
{TensorDescriptorParams{dt, L, {160, 1024, 120, 120}},
63+
TensorDescriptorParams{dt, L, {2048, 1024, 3, 3}},
64+
dt,
65+
ConvolutionDescriptorParams{{1, 1}, {1, 1}, {1, 1}, 1}}, // 4.72e9 elem, K=2C
66+
};
67+
}
68+
69+
const miopen::unit_tests::UnitTestConvSolverParams& GetTestParams()
70+
{
71+
static const auto params = [] {
72+
#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
73+
Gpu supportedDevices = Gpu::gfx90A | Gpu::gfx94X | Gpu::gfx950;
74+
#else
75+
Gpu supportedDevices = Gpu::None;
76+
#endif
77+
auto p = miopen::unit_tests::UnitTestConvSolverParams(supportedDevices);
78+
p.Tunable(5);
79+
return p;
80+
}();
81+
return params;
82+
}
83+
84+
enum class Dir
85+
{
86+
Fwd,
87+
Bwd,
88+
Wrw
89+
};
90+
91+
struct RunGate
92+
{
93+
std::size_t required;
94+
std::size_t available;
95+
bool applicable;
96+
};
97+
98+
RunGate GateFor(const ConvTestCase& tc,
99+
const miopen::solver::conv::ConvSolverInterface& solver,
100+
Dir dir)
101+
{
102+
auto&& handle = get_handle();
103+
104+
const auto x_desc = tc.GetXTensorDescriptor();
105+
const auto w_desc = tc.GetWTensorDescriptor();
106+
const auto conv_desc = tc.GetConv();
107+
const auto y_desc = conv_desc.GetForwardOutputTensor(x_desc, w_desc, tc.GetYDataType());
108+
109+
const auto problem = [&] {
110+
switch(dir)
111+
{
112+
case Dir::Fwd:
113+
return miopen::conv::ProblemDescription(
114+
x_desc, w_desc, y_desc, conv_desc, miopen::conv::Direction::Forward);
115+
case Dir::Bwd:
116+
return miopen::conv::ProblemDescription(
117+
y_desc, w_desc, x_desc, conv_desc, miopen::conv::Direction::BackwardData);
118+
default:
119+
return miopen::conv::ProblemDescription(
120+
y_desc, w_desc, x_desc, conv_desc, miopen::conv::Direction::BackwardWeights);
121+
}
122+
}();
123+
124+
auto ctx = miopen::ExecutionContext{&handle};
125+
problem.SetupFloats(ctx);
126+
127+
if(!solver.IsApplicable(ctx, problem))
128+
return {0, 0, false};
129+
130+
const std::size_t ws_size =
131+
solver.MayNeedWorkspace() ? solver.GetWorkspaceSize(ctx, problem) : 0;
132+
const std::size_t x_bytes = x_desc.GetNumBytes();
133+
const std::size_t y_bytes = y_desc.GetNumBytes();
134+
const std::size_t w_bytes = w_desc.GetNumBytes();
135+
const std::size_t h_bytes = std::max(x_bytes, y_bytes);
136+
137+
const std::size_t raw_mem = ws_size + x_bytes + y_bytes + w_bytes + 4 * h_bytes;
138+
const std::size_t headroom = std::max<std::size_t>(1ULL << 30, raw_mem / 10);
139+
const std::size_t required_mem = raw_mem + headroom;
140+
const std::size_t device_mem = handle.GetGlobalMemorySize();
141+
142+
return {required_mem, device_mem, true};
143+
}
144+
145+
} // namespace
146+
147+
#define SKIP_IF_NOT_RUNNABLE(solver_expr, dir) \
148+
do \
149+
{ \
150+
miopen::unit_tests::UnitTestConvSolverParams _params; \
151+
miopenConvAlgorithm_t _algo; \
152+
ConvTestCase _tc; \
153+
std::tie(_params, _algo, _tc) = this->GetParam(); \
154+
const auto _g = GateFor(_tc, (solver_expr), (dir)); \
155+
if(!_g.applicable) \
156+
{ \
157+
GTEST_SKIP() << "solver not applicable to this shape"; \
158+
} \
159+
if(_g.available < _g.required) \
160+
{ \
161+
GTEST_SKIP() << "Insufficient device memory: need " << _g.required \
162+
<< " bytes, device has " << _g.available; \
163+
} \
164+
} while(0)
165+
166+
using GPU_UnitTestConvSolverImplicitGemmGroupFwdXdlops_Large2D_FP16 =
167+
GPU_UnitTestConvSolverFwd_FP16;
168+
using GPU_UnitTestConvSolverImplicitGemmGroupBwdXdlops_Large2D_FP16 =
169+
GPU_UnitTestConvSolverBwd_FP16;
170+
using GPU_UnitTestConvSolverImplicitGemmGroupWrwXdlops_Large2D_FP16 =
171+
GPU_UnitTestConvSolverWrw_FP16;
172+
173+
TEST_P(GPU_UnitTestConvSolverImplicitGemmGroupFwdXdlops_Large2D_FP16,
174+
ConvHipImplicitGemmGroupFwdXdlops)
175+
{
176+
const auto solver = miopen::solver::conv::ConvHipImplicitGemmGroupFwdXdlops{};
177+
SKIP_IF_NOT_RUNNABLE(solver, Dir::Fwd);
178+
this->RunTest(solver);
179+
};
180+
181+
TEST_P(GPU_UnitTestConvSolverImplicitGemmGroupBwdXdlops_Large2D_FP16,
182+
ConvHipImplicitGemmGroupBwdXdlops)
183+
{
184+
const auto solver = miopen::solver::conv::ConvHipImplicitGemmGroupBwdXdlops{};
185+
SKIP_IF_NOT_RUNNABLE(solver, Dir::Bwd);
186+
this->RunTest(solver);
187+
};
188+
189+
TEST_P(GPU_UnitTestConvSolverImplicitGemmGroupWrwXdlops_Large2D_FP16,
190+
ConvHipImplicitGemmGroupWrwXdlops)
191+
{
192+
const auto solver = miopen::solver::conv::ConvHipImplicitGemmGroupWrwXdlops{};
193+
SKIP_IF_NOT_RUNNABLE(solver, Dir::Wrw);
194+
this->RunTest(solver);
195+
};
196+
197+
INSTANTIATE_TEST_SUITE_P(Large2D,
198+
GPU_UnitTestConvSolverImplicitGemmGroupFwdXdlops_Large2D_FP16,
199+
testing::Combine(testing::Values(GetTestParams()),
200+
testing::Values(miopenConvolutionAlgoImplicitGEMM),
201+
testing::ValuesIn(GetTestCases(miopenHalf))));
202+
203+
INSTANTIATE_TEST_SUITE_P(Large2D,
204+
GPU_UnitTestConvSolverImplicitGemmGroupBwdXdlops_Large2D_FP16,
205+
testing::Combine(testing::Values(GetTestParams()),
206+
testing::Values(miopenConvolutionAlgoImplicitGEMM),
207+
testing::ValuesIn(GetTestCases(miopenHalf))));
208+
209+
INSTANTIATE_TEST_SUITE_P(Large2D,
210+
GPU_UnitTestConvSolverImplicitGemmGroupWrwXdlops_Large2D_FP16,
211+
testing::Combine(testing::Values(GetTestParams()),
212+
testing::Values(miopenConvolutionAlgoImplicitGEMM),
213+
testing::ValuesIn(GetTestCases(miopenHalf))));

0 commit comments

Comments
 (0)