[amd-alavin]

JIRA ID : ALMIOPEN-2361

Takes the vendored src/hipconv snapshot from v0.3.1 to v0.3.2. Per-version detail is in the shipped src/hipconv/CHANGELOG.md.

Data types

ConvHipConv supports fp16, bf16 and TF32, and not fp32. This is unchanged from v0.2.7 and v0.3.1. hipconv emulates TF32 with three bf16 MFMAs on a (big, small) bf16 split of the stored fp32, so it needs no TF32 hardware, and its gfx950 and gfx1250 kernels both support it.

MIOpen decides when a problem is TF32: MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE set to miopenMathDefault or MIOPEN_TF32_OVERRIDE, on a device IsTF32Supported admits. That list is gfx942 and gfx95x today, so through MIOpen, TF32 reaches hipconv on gfx950 and not yet on gfx1250.

Additional TF32 support on gfx950 since v0.3.1:

• direct_l1 forward and backward-data, 2x2 to 5x5 filters (hipconv [rocrand] backport from release-staging-7.0 #226);
• depthwise_1d_toeplitz forward and backward-data (hipconv [rocprim] backport from release-staging-7.0 #223);
• direct_wgrad backward-weights at 2x2, 3x1, 3x3, 4x4 and 5x5 (hipconv [Discussion]: Cherry-pick from monorepo to original repositories #234).

Kernels

• gfx950 patch_embed: patch-embedding convolutions (stride equal to the filter, no padding), forward, backward-data and backward-weights, fp16/bf16.
• gfx1250 depthwise_wgrad_hankel: depthwise backward-weights, fp16/bf16, 3x3 to 11x11 odd square, stride 1 or 2.
• gfx1250 direct: fixes a crash when groups x batch tiles exceeded the 65535 grid z limit (ALMIOPEN-2624), and an LDS race between a TDM load and another wave's reads; the race fix lifts the odd-channel backward-data restriction v0.3.1 added. direct no longer claims depthwise layers (one channel per group).
• gfx950 depthwise_wgrad_hankel sums its partitions in a fold kernel instead of with atomics: geometric mean +88% over fifteen shapes (hipconv [hipCUB] add memory to benchmark_device_batch_memcpy (#219) #227).

Build

src/CMakeLists.txt renames HIPCONV_ENABLE_POINTWISE to HIPCONV_ENABLE_EXPLICIT_GEMM, the name hipconv v0.3.2 gives the option, which now enables explicit-GEMM patch embedding as well as pointwise. CMake ignores the old name without a warning. Both explicit-GEMM paths stay off, since the new option also defaults to OFF.

Downstream patches

On develop, projects/miopen/src/hipconv differed from hipconv v0.3.1 only by #12406, which patched src/arch/cdna4/direct/direct_wgrad/epilogue.h. hipconv v0.3.2 contains the same fix (hipconv e1777ac9), so this sync keeps the file byte-identical.

Validation

• hipconv's internal test layers (549 convolutions, each in all three directions) through MIOpen with ConvHipConv forced, on gfx950 (MI355X), with MIOpen built by clang 24 (ROCm 10.1) against this snapshot: every case ConvHipConv accepted passed, 1476 in fp16, 1476 in bf16 and 1283 in TF32. 18 of the cases it rejected (12 in TF32) are ones hipconv supports on its own. The cause, noted in [MIOpen] Sync hipconv v0.3.1 #12000, is that GetAdjustedConvolutionStride{H,W} hands hipconv a stride of 1 on an axis whose output extent is 1. They are 6 grouped stride-2 layers with a degenerate spatial extent, in forward and backward-data, and 2 patch-embedding layers with a 1x1 output, in all three directions.
• hipconv's internal test suite on MI455 A0 (gfx1250): 64/64, including about 14,000 cases that run every kernel config against the same test layers. On gfx950, 65/65. Both ran at hipconv f58e0771, from which v0.3.2 differs in formatting, its changelog, host warning flags and one constexpr expression.

clang-tidy over the consumer TU is clean with AMD LLVM 24.

ConvHipConv unit tests

unit_conv_solver_ConvHipConv.cpp gains NHWC tests that cover every hipconv kernel available
through ConvHipConv. The layers come from hipconv's internal tests.

List	gfx950	gfx1250
GetConvDenseTestCases	direct_l1, direct (Fwd, Bwd); direct_wgrad (Wrw); tf32 on direct_l1, direct_wgrad	direct (Fwd, Bwd)
GetConvDenseVerticalTestCases	direct_l1, direct_wgrad	—
GetConvPatchEmbedTestCases	patch_embed (Fwd, Bwd, Wrw)	—
GetConvOddChannelTestCases	tf32 only: direct_l1, direct_wgrad	direct (Fwd, Bwd)
GetConvPointwiseTestCases	—	direct 1×1 (Fwd, Bwd)
GetConvGroupedTestCases	grouped, grouped_wgrad at 4, 8, 16 and 32 channels per group; fp16, bf16, tf32	grouped_multi_g, grouped_multi_g_wgrad
GetConvGroupedStride2TestCases	grouped at 4, 16 and 32 channels per group (Fwd, Bwd)	grouped_multi_g (Fwd, Bwd)
GetConvGrouped8cStride2TestCases	grouped at 8 channels per group (Fwd)	grouped_multi_g (Fwd, Bwd)
GetConvDepthwiseTestCases	depthwise_1d_toeplitz (Fwd, Bwd; tf32 too), depthwise_wgrad_hankel (Wrw)	the same, fp16 and bf16

Each list is instantiated only on the architectures and directions whose kernels support all of its
layers, since the harness fails an inapplicable case. The new lists are NHWC only: the NCHW
staging does not depend on the kernel, and the existing smoke lists cover it.

Tunable(0). GenericSearch shuffles every valid perf config and benchmarks the first
Tunable(N), so Tunable(5) verifies the fastest of five random configs. On a layer several
kernels support, which kernel that is varies from run to run. The grouped stride-1 layers are such
layers (the direct kernels support them too), as are the tf32 depthwise layers. Their suites use
Tunable(0): the sample is empty, so the search runs only the default config, which is hipconv's
top-ranked, and the grouped or depthwise kernel is the one verified. Every other list keeps
Tunable(5), since one kernel supports each of its layers and the sampling still exercises its
non-default configs.

Testing. On gfx950 (MI355X): 399 tests, 375 pass and 24 skip (the gfx125X-only suites), the
same in each of five repeats. With MIOPEN_LOG_LEVEL=5, each new case's final GetSolution names
the kernel it was meant to probe. On MI455 A0 (gfx1250): 208 pass and 191 skip (the gfx950-only suites), each case on its intended kernel; CI skips the gfx1250 cases, and hipconv's own suite also passed on MI455 A0
(Validation above).

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[amd-alavin]

The ConvHipConv unit tests include several layers that hipconv does not yet support on gfx1250, so da19417e2bc restricts those tests to gfx950. This went undetected because MIOpen CI does not cover gfx1250 yet.

The same commit adjusts one test for an overflow. MIOpen sizes the test data's range by the forward convolution's reduction depth, but dgrad reduces over a different depth, which can overflow fp16. The odd-channel grouped case now uses fewer output channels, which keeps it on the same kernels.

With the commit, all ConvHipConv unit tests pass on MI455 (gfx1250) and gfx950.

[JonathanLichtnerAMD]

LGTM. Just a minor NIT on the miopen changelog: we don't usually mention the jira tickets in the changelog or the code, e.g., ALMIOPEN-2624.

[amd-alavin]

LGTM. Just a minor NIT on the miopen changelog: we don't usually mention the jira tickets in the changelog or the code, e.g., ALMIOPEN-2624.

Fixed, thanks.

[therock-pr-bot]

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

• python -m pip install pre-commit
• pre-commit install
• pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

[SreecharanGundaboluAMD]

LGTM