Skip to content

fix(miopen/hipconv) Fix build break with GCC 14+ system headers - #11989

Merged
BradPepersAMD merged 1 commit into
developfrom
users/benrichard-amd/miopen-hipconv-constexpr-steps
Sep 12, 2026
Merged

fix(miopen/hipconv) Fix build break with GCC 14+ system headers#11989
BradPepersAMD merged 1 commit into
developfrom
users/benrichard-amd/miopen-hipconv-constexpr-steps

Conversation

@benrichard-amd

@benrichard-amd benrichard-amd commented Sep 11, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Motivation

Addresses #11990

  • Add -fconstexpr-steps=2097152 to hipconv compile options to fix a build failure with GCC 14+ system headers.

Technical Details

The TransposeSwizzle class in lds_layout.h validates its LDS bank-conflict-free layout at compile time using static_assert. The rounds_are_rungs() check iterates through all swizzle combinations using std::array constexpr operations.

GCC 14 and later changed how std::array implements constexpr — the new implementation requires significantly more evaluation steps per operation. The same validation that completed within clang's default constexpr step limit (1048576) with GCC 13's libstdc++ now exceeds it with GCC 14+, producing:

lds_layout.h:287: error: static assertion expression is not an integral constant expression
note: use -fconstexpr-steps= to increase this limit (=0 will remove the limit)

This is not a code bug — the constexpr evaluation is finite and correct. It just needs a higher step budget with the newer std::array implementation.

Test Plan

  • MIOpen builds successfully on x86 with GCC 15 system headers (Ubuntu 26.04)
  • MIOpen builds successfully on ppc64le with GCC 14 system headers (AlmaLinux 10)
  • No impact on GCC 13 systems (flag raises the limit but doesn't change behavior when the default is sufficient)

Submission Checklist

@benrichard-amd
benrichard-amd requested a review from a team as a code owner September 11, 2026 17:11
@therock-pr-bot

therock-pr-bot Bot commented Sep 11, 2026 •
edited
Loading

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/miopen/src/hipconv/src/arch/cdna4/direct/direct_wgrad/lds_layout.h; no test file found
🔎 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

therock-pr-bot Bot commented Sep 11, 2026 •
edited
Loading

Copy link
Copy Markdown

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

@benrichard-amd benrichard-amd changed the title fix(miopen/hipconv) Fix build break with GCC 14+ fix(miopen/hipconv) Fix build break with GCC 14+ system headers Sep 11, 2026
@amd-alavin

amd-alavin commented Sep 11, 2026

Copy link
Copy Markdown
Contributor

@benrichard-amd I am curious if simply deleting static_assert(rounds_are_rungs(), ...) fixes the issue. This static_assert hurts compilation time also, so we should just get rid of it. Would be helpful if you can verify that fix on GCC 14+.

I can add unit tests that perform the equivalent checks.

@benrichard-amd

benrichard-amd commented Sep 12, 2026

Copy link
Copy Markdown
Contributor Author

@benrichard-amd I am curious if simply deleting static_assert(rounds_are_rungs(), ...) fixes the issue. This static_assert hurts compilation time also, so we should just get rid of it. Would be helpful if you can verify that fix on GCC 14+.

I can add unit tests that perform the equivalent checks.

Hi @amd-alavin your suggestion works. MIOpen builds successfully with just the static_assert removal and no -fconstexpr-steps flag.

@benrichard-amd
miopen/hipconv: remove rounds_are_rungs static_assert for GCC 14+ compat
e87cd14 
The rounds_are_rungs() constexpr validation exceeds clang's default
constexpr step limit when compiled with GCC 14+ libstdc++ headers, due
to changes in std::array's constexpr implementation. Remove the
static_assert; equivalent checks will be covered by unit tests.

Suggested-by: amd-alavin
@benrichard-amd
benrichard-amd force-pushed the users/benrichard-amd/miopen-hipconv-constexpr-steps branch from 78720a5 to e87cd14 Compare September 12, 2026 01:54
@amd-alavin

amd-alavin commented Sep 12, 2026

Copy link
Copy Markdown
Contributor

@benrichard-amd I am curious if simply deleting static_assert(rounds_are_rungs(), ...) fixes the issue. This static_assert hurts compilation time also, so we should just get rid of it. Would be helpful if you can verify that fix on GCC 14+.
I can add unit tests that perform the equivalent checks.

Hi @amd-alavin your suggestion works. MIOpen builds successfully with just the static_assert removal and no -fconstexpr-steps flag.

Good, that is my preferred fix. I will add it to PR #12000 for hipconv v0.3. Continue this PR too, because it might get merged first.

BradPepersAMD
BradPepersAMD approved these changes Sep 12, 2026
View reviewed changes

@BradPepersAMD BradPepersAMD left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Trivial fix at this point so approved.

@BradPepersAMD
BradPepersAMD merged commit 3792095 into develop Sep 12, 2026
139 checks passed
@BradPepersAMD
BradPepersAMD deleted the users/benrichard-amd/miopen-hipconv-constexpr-steps branch September 12, 2026 15:17
@assistant-librarian

assistant-librarian Bot commented Sep 14, 2026

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

@amd-alavin amd-alavin mentioned this pull request Sep 18, 2026
Merged
amd-alavin added a commit that referenced this pull request Sep 25, 2026
@amd-alavin Boris Prokopenko
@uphoffc @jfactory07 @petar-jovanovic @yingluAMD
[MIOpen] Sync hipconv v0.3.1 (#12000)
2ce4c67 
JIRA ID : ALMIOPEN-2361

Updates the vendored `src/hipconv` from v0.2.8 to v0.3.1 and the
ConvHipConv solver with it. Per-version detail is in the new
`src/hipconv/CHANGELOG.md`.

## Kernels

v0.3.0 adds `depthwise_wgrad_hankel` (gfx950 depthwise conv2d wgrad,
fp16/bf16, odd square 3x3 to 11x11, stride 1 or 2),
`depthwise_1d_toeplitz` (gfx1250 fprop/dgrad, fp16/bf16/tf32), tf32
across the gfx1250 kernels (not yet selectable: `IsTF32Supported`
admits only gfx942 and gfx95x), and conv3d via unfolding to conv2d. It
removes 1x1 from the gfx950 `direct` filter set.

v0.3.1 adds a gfx1250 dgrad dispatch guard, `-fno-offload-lto` on the
device targets, and the `static_assert` removal below.

## Solver

`IsSupportedProblem` admits 3D on the NHWC/NDHWC path and rejects
non-packed tensors for every layout. NCHW stays 2D-only because
`MakeTransposePlan` reads four extents positionally. The #11845 merge
removes the wrw staging helpers along with this PR's depth fix to them;
`GetWgradCastSize` counts depth through `GetElementSize()`.

## Downstream patches reconciled

- #12085 (`gfx1250-strict` in `src/arch/cdna5/CMakeLists.txt`) is now
  upstream in hipconv, with the `GPU_TARGETS` filter fix that keeps the
  `-strict` suffix. The file differs from #12085's only by the new
  `depthwise_1d_toeplitz` shard.
- #11989 (the `rounds_are_rungs()` `static_assert` removed for GCC 14+)
  matches upstream, which also removes the now-unused function.
- #12328 (`-fno-offload-lto`) matches v0.3.1.
- #12406 (direct_wgrad LDS ordering, hipconv `e1777ac9`, after v0.3.1)
  merges cleanly; `HIPCONV_VERSION` keeps its out-of-snapshot note
  until a sync at or after `e1777ac9`.

## Validation

gfx950, clang 24 (ROCm 10.1):

- ConvHipConv unit test: 144/144, including 60 conv3d cases and
  #12406's 12 dense wgrad cases.
- hipconv feature corpus through MIOpen, solver forced: 454 pass,
  0 fail, 9 no-match. Three are 1x1 (pointwise, off in MIOpen builds).
  Six are grouped stride-2 layers with an output extent of 1 on one
  axis, where `GetAdjustedConvolutionStride{H,W}` produces a mixed
  stride hipconv's config table does not cover. Pre-existing.
- hipconv standalone ctest: 54/54, including 11020 corpus cases.

gfx1250-strict: MIOpen configures with it on ROCm 10.1, and
`hipconv_arch_cdna5` compiles for both gfx1250 and gfx1250-strict.

Co-authored-by: Boris Prokopenko <boris.prokopenko@amd.com>
Co-authored-by: Carsten Uphoff <carsten.uphoff@amd.com>
Co-authored-by: Jin Zhou <jin.zhou@amd.com>
Co-authored-by: Petar Jovanovic <petar.jovanovic@amd.com>
Co-authored-by: Yingmao Lu <yingmao.lu@amd.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@BradPepersAMD BradPepersAMD BradPepersAMD approved these changes

Assignees

No one assigned

Labels

organization: ROCm project: miopen

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@benrichard-amd @amd-alavin @BradPepersAMD