Skip to content

[MIOpen] Remove opencl convolution solver conv_ocl_dir2Dfwdgen - #6633

Merged
iq136boy merged 17 commits into
developfrom
dfeng/remove_conv_ocl_dir2Dfwdgen
May 15, 2026
Merged

iq136boy merged 17 commits into
developfrom
dfeng/remove_conv_ocl_dir2Dfwdgen

Conversation

@iq136boy

@iq136boy iq136boy commented Apr 21, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Remove the OpenCL solver ConvOclDirectFwdGen and its kernel MIOpenConvDirGenFwd.cl as part of the ongoing effort to eliminate OpenCL solvers/kernels from MIOpen. No HIP replacement solver is required here. The problem space covered by ConvOclDirectFwdGen (non-1x1 kernels with stride > 1, and kernels larger than 11x11) is already fully
covered by existing HIP-based solvers:

  • ConvHipImplicitGemmForwardV4R4Xdlops and ConvHipImplicitGemmForwardV4R5Xdlops
    handle the common strided convolution cases (e.g., 3x3/stride=2) with better
    performance via XDLOPs acceleration.
  • ConvDirectNaiveConvFwd provides a HIP fallback for any remaining edge cases
    (e.g., small or oddly-shaped channel counts) with no shape restrictions.

Technical Details

  • Removed the OpenCL solver conv_ocl_dir2Dfwdgen.cpp and its kernel
    MIOpenConvDirGenFwd.cl (588 lines).
  • Removed ConvOclDirectFwdGen struct declaration from solvers.hpp.
  • Replaced the RegisterWithSolver call for ConvOclDirectFwdGen in solver.cpp
    with a ++id tombstone comment to preserve solver ID numbering for database
    compatibility.
  • Removed ConvOclDirectFwdGen from GetDirectSolvers() in mlo_dir_conv.cpp.
  • Updated all find-db files (HIP and OpenCL) to remove entries referencing ConvOclDirectFwdGen.
    Lines where ConvOclDirectFwdGen was the sole solver for a problem config
    were dropped entirely; lines with multiple solvers had only the
    ConvOclDirectFwdGen segment removed.
  • Removed the unit test file unit_conv_solver_ConvOclDirectFwdGen.cpp.
  • Removed ConvOclDirectFwdGen entry from unit_FinInterface.cpp.

Test Plan

  • Verify the build compiles without errors after the removal.
  • Run existing HIP solver unit tests to confirm coverage of the problem space previously handled by ConvOclDirectFwdGen:
    • GPU_UnitTestConvSolverHipImplicitGemmV4R1Fwd (FP16, BFP16, FP32)
    • GPU_UnitTestConvSolverHipImplicitGemmForwardV4R4Xdlops (FP16, BFP16, FP32)
    • GPU_UnitTestConvSolverHipImplicitGemmForwardV4R5Xdlops (FP16, BFP16, FP32)
  • Confirm no remaining references to ConvOclDirectFwdGen or MIOpenConvDirGenFwd.cl in src/ and test/.

Test Result

  • Local tests passed
  • pending CI

Submission Checklist

  • Removed OpenCL solver conv_ocl_dir2Dfwdgen.cpp
  • Removed OpenCL kernel MIOpenConvDirGenFwd.cl
  • Removed solver declaration from solvers.hpp
  • Preserved solver ID slot in solver.cpp with tombstone comment
  • Removed solver from GetDirectSolvers() in mlo_dir_conv.cpp
  • Updated CMakeLists.txt to remove source and kernel file entries
  • Cleaned up all find-db files (.fdb.txt.bz2) across all GPU architectures
  • Removed unit test file unit_conv_solver_ConvOclDirectFwdGen.cpp
  • No HIP replacement solver added (existing HIP solvers provide full coverage)

iq136boy and others added 9 commits April 22, 2026 02:53
Resolved conflicts in find-db .bz2 files by taking develop's latest
versions and re-applying ConvOclDirectFwdGen removal on top.

Resolved text conflicts in mlo_dir_conv.cpp, solver.cpp, and
unit_FinInterface.cpp by taking develop's ConvHipDirectFwd11x11
(OCL->HIP replacement for ConvOclDirectFwd11x11) while keeping
this branch's removal of ConvOclDirectFwdGen.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
iq136boy and others added 5 commits May 7, 2026 18:53
Resolve finddb conflicts by taking develop's latest fdb files and
re-applying ConvOclDirectFwdGen removal on top. Text conflicts in
CMakeLists.txt and mlo_dir_conv.cpp resolved by keeping our removal
of ConvOclDirectFwdGen (replaced by ConvOclDirectFwd1x1).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Resolve fdb binary conflicts: take develop's version and remove
ConvOclDirectFwdGen entries. Where ConvOclDirectFwdGen was the winning
(first) solver, drop the entire fdb line to avoid promoting an arbitrary
next solver that may lack a KDB entry.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@iq136boy
iq136boy force-pushed the dfeng/remove_conv_ocl_dir2Dfwdgen branch from d1d4d25 to 5d9b525 Compare May 14, 2026 02:12
@iq136boy
iq136boy merged commit 2d499dd into develop May 15, 2026
37 checks passed
@iq136boy
iq136boy deleted the dfeng/remove_conv_ocl_dir2Dfwdgen branch May 15, 2026 01:26
aledudek pushed a commit that referenced this pull request May 20, 2026
## Motivation
Remove the OpenCL solver ConvOclDirectFwdGen and its kernel
MIOpenConvDirGenFwd.cl as part of the ongoing effort to eliminate OpenCL
solvers/kernels from MIOpen. No HIP replacement solver is required here.
The problem space covered by ConvOclDirectFwdGen (non-1x1 kernels with
stride > 1, and kernels larger than 11x11) is already fully
covered by existing HIP-based solvers:
- ConvHipImplicitGemmForwardV4R4Xdlops and
ConvHipImplicitGemmForwardV4R5Xdlops
handle the common strided convolution cases (e.g., 3x3/stride=2) with
better
  performance via XDLOPs acceleration.
- ConvDirectNaiveConvFwd provides a HIP fallback for any remaining edge
cases
(e.g., small or oddly-shaped channel counts) with no shape restrictions.

## Technical Details
- Removed the OpenCL solver conv_ocl_dir2Dfwdgen.cpp and its kernel
  MIOpenConvDirGenFwd.cl (588 lines).
- Removed ConvOclDirectFwdGen struct declaration from solvers.hpp.
- Replaced the RegisterWithSolver call for ConvOclDirectFwdGen in
solver.cpp
with a ++id tombstone comment to preserve solver ID numbering for
database
  compatibility.
- Removed ConvOclDirectFwdGen from GetDirectSolvers() in
mlo_dir_conv.cpp.
- Updated all find-db files (HIP and OpenCL) to remove entries
referencing ConvOclDirectFwdGen.
Lines where ConvOclDirectFwdGen was the sole solver for a problem config
  were dropped entirely; lines with multiple solvers had only the
  ConvOclDirectFwdGen segment removed.
- Removed the unit test file unit_conv_solver_ConvOclDirectFwdGen.cpp.
- Removed ConvOclDirectFwdGen entry from unit_FinInterface.cpp.

## Test Plan
- Verify the build compiles without errors after the removal.
- Run existing HIP solver unit tests to confirm coverage of the problem
space previously handled by ConvOclDirectFwdGen:
	- GPU_UnitTestConvSolverHipImplicitGemmV4R1Fwd (FP16, BFP16, FP32)
- GPU_UnitTestConvSolverHipImplicitGemmForwardV4R4Xdlops (FP16, BFP16,
FP32)
- GPU_UnitTestConvSolverHipImplicitGemmForwardV4R5Xdlops (FP16, BFP16,
FP32)
- Confirm no remaining references to ConvOclDirectFwdGen or
MIOpenConvDirGenFwd.cl in src/ and test/.

## Test Result
- Local tests passed
- pending CI

## Submission Checklist
- [x] Removed OpenCL solver conv_ocl_dir2Dfwdgen.cpp
- [x] Removed OpenCL kernel MIOpenConvDirGenFwd.cl
- [x] Removed solver declaration from solvers.hpp
- [x] Preserved solver ID slot in solver.cpp with tombstone comment
- [x] Removed solver from GetDirectSolvers() in mlo_dir_conv.cpp
- [x] Updated CMakeLists.txt to remove source and kernel file entries
- [x] Cleaned up all find-db files (.fdb.txt.bz2) across all GPU
architectures
- [x] Removed unit test file unit_conv_solver_ConvOclDirectFwdGen.cpp
- [ ] No HIP replacement solver added (existing HIP solvers provide full
coverage)

---------

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
Co-authored-by: BradPepersAMD <Brad.Pepers@amd.com>
JonathanLichtnerAMD added a commit that referenced this pull request Jul 15, 2026
JIRA ID: ALMIOPEN-2312

## Motivation

`projects/miopen/CHANGELOG.md`'s "MIOpen 3.5.2 for ROCm 7.14.0" section
only had 4 entries, but 68 commits touched `projects/miopen` between the
`therock-7.13` tag and `release/therock-7.14`. Auditing that commit
range (via `git log --oneline
therock-7.13..remotes/origin/release/therock-7.14 -- projects/miopen`)
found 18 additional user-facing changes — new gfx950/gfx1250
architecture support, SystemDB refreshes, AI-heuristic improvements,
several correctness/crash bug fixes, and OpenCL solver removals — that
were never documented. This PR backfills those entries so the 7.14.0
changelog accurately reflects what shipped.

## Technical Details

Added 24 lines to the existing "MIOpen 3.5.2 for ROCm 7.14.0" section of
`projects/miopen/CHANGELOG.md`, merging new bullets into the existing
`Added`/`Changed`/`Resolved Issues` sub-sections (no new entries were
duplicated or reworded) and adding a new `Removed` sub-section for
solver removals. No existing content was altered.

Each new bullet was traced to a specific commit + PR:

| Entry | Commit | PR |
|---|---|---|
| Added gfx950 (MI350X/MI355X) SystemDB | `60d8e1e5175` | #8081 |
| Added initial gfx1250 architecture support | `d3f057b2f22` | #7587 |
| Added JSON-formatted performance logs (`MIOPEN_PERFORMANCE_LOGS`) |
`ffbde871260` | #5736 |
| Refreshed gfx942 SystemDB | `faf754a9ef6` | #8174 |
| Refined 2D AI heuristics (Tunanet/KTN) for gfx942 & gfx950 |
`c24ad415884` | #7109 |
| Fixed thread-safety issue in AI-heuristic model caches | `cd957402309`
| #9149 (cherry-pick of #9131) |
| Fixed CK grouped-conv solvers not always registered in host |
`c971d54d7ae` | #9105 |
| Fixed legacy CK buffer resource descriptor (gfx1250) | `1dee011b8b0` |
#8758 |
| Removed stale gfx942 SystemDB naive-solver entries | `f42668f6316` |
#7028 |
| Fixed softmax NaN when `beta=0` | `fb8157c8512` | #8426 |
| Fixed BN backward spatial variant 1 vectorized memory access bug |
`1351a7e8b10` | #6765 |
| Fixed use-after-free of `kernel_id` in `InitInvokerFactoryNCHW` |
`394ebffb50d` | #7922 |
| Fixed SoftMarginLoss forward reduce-kernel kernarg corruption |
`1e65a31570f` | #8065, #8135 |
| Fixed build with `MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF` | `e6288c3ca61`
| #6800 |
| Fixed install RPATH for MIOpenDriver / CK backend libs | `ef8a4cf2f93`
| #7485 |
| Fixed dangling lambda causing log-level 6 segfault | `190f9791737` |
#6953 |
| Enabled Windows long-path support for `addkernels` | `1f9d69724cf` |
#7387 |
| Removed OCL solver `ConvOclDirectFwd1x1` | `126c720cadd` | #6491 |
| Removed OCL solver `conv_ocl_dir2Dfwdgen` | `2d499ddf709` | #6633 |
| Removed OCL solvers `ConvOclBwdWrW1x1`, `ConvOclBwdWrW2` |
`b5d43a33dc6` | #6859 |
| Removed OCL solvers `OpTensorLeadingOnes`, `Op5dTensorGeneric` |
`8642c83b7de` | #7347 |

The remaining ~40 commits in the 68-commit range were deliberately
excluded (tests/CI/build/docs/internal refactors/1:1 OCL→HIP swaps with
no behavior change, per each commit's own stated scope) and are not
represented here.

## Risk Assessment

**Risk Level: 🟢 Low**

This is a documentation-only change to a single Markdown file. No code,
build configuration, tests, or public API surface is touched.

### Impacted Components

- `projects/miopen/CHANGELOG.md` only.

### Potential Side Effects

None expected. This does not change any runtime behavior; it only
documents changes that already shipped on `release/therock-7.14`.

### Mitigation Steps

None required beyond normal review — a reviewer familiar with the 7.14
release can spot-check a few entries against the linked commit hashes/PR
numbers for accuracy.

## Test Plan

Not applicable (no code changes). Verification was done by
cross-referencing every added bullet against a specific commit hash and
PR number pulled from `git log`, and by diffing the edited section
against `git show
remotes/origin/release/therock-7.14:projects/miopen/CHANGELOG.md` to
confirm no existing entry was duplicated or altered.

## Test Result

All 21 new bullets map 1:1 to a real commit/PR in the
`therock-7.13..release/therock-7.14` range for `projects/miopen`. No
pre-existing changelog text was modified — `git diff` shows a pure
addition (24 insertions, 0 deletions). Two borderline commits
(`a67834413e8` "Standalone build", `7c7cd5774a9` "enable warnings about
deprecated API/features") were deliberately left out pending author
judgment on whether they warrant an entry.
pschang-phy pushed a commit that referenced this pull request Jul 17, 2026
JIRA ID: ALMIOPEN-2312

## Motivation

`projects/miopen/CHANGELOG.md`'s "MIOpen 3.5.2 for ROCm 7.14.0" section
only had 4 entries, but 68 commits touched `projects/miopen` between the
`therock-7.13` tag and `release/therock-7.14`. Auditing that commit
range (via `git log --oneline
therock-7.13..remotes/origin/release/therock-7.14 -- projects/miopen`)
found 18 additional user-facing changes — new gfx950/gfx1250
architecture support, SystemDB refreshes, AI-heuristic improvements,
several correctness/crash bug fixes, and OpenCL solver removals — that
were never documented. This PR backfills those entries so the 7.14.0
changelog accurately reflects what shipped.

## Technical Details

Added 24 lines to the existing "MIOpen 3.5.2 for ROCm 7.14.0" section of
`projects/miopen/CHANGELOG.md`, merging new bullets into the existing
`Added`/`Changed`/`Resolved Issues` sub-sections (no new entries were
duplicated or reworded) and adding a new `Removed` sub-section for
solver removals. No existing content was altered.

Each new bullet was traced to a specific commit + PR:

| Entry | Commit | PR |
|---|---|---|
| Added gfx950 (MI350X/MI355X) SystemDB | `60d8e1e5175` | #8081 |
| Added initial gfx1250 architecture support | `d3f057b2f22` | #7587 |
| Added JSON-formatted performance logs (`MIOPEN_PERFORMANCE_LOGS`) |
`ffbde871260` | #5736 |
| Refreshed gfx942 SystemDB | `faf754a9ef6` | #8174 |
| Refined 2D AI heuristics (Tunanet/KTN) for gfx942 & gfx950 |
`c24ad415884` | #7109 |
| Fixed thread-safety issue in AI-heuristic model caches | `cd957402309`
| #9149 (cherry-pick of #9131) |
| Fixed CK grouped-conv solvers not always registered in host |
`c971d54d7ae` | #9105 |
| Fixed legacy CK buffer resource descriptor (gfx1250) | `1dee011b8b0` |
#8758 |
| Removed stale gfx942 SystemDB naive-solver entries | `f42668f6316` |
#7028 |
| Fixed softmax NaN when `beta=0` | `fb8157c8512` | #8426 |
| Fixed BN backward spatial variant 1 vectorized memory access bug |
`1351a7e8b10` | #6765 |
| Fixed use-after-free of `kernel_id` in `InitInvokerFactoryNCHW` |
`394ebffb50d` | #7922 |
| Fixed SoftMarginLoss forward reduce-kernel kernarg corruption |
`1e65a31570f` | #8065, #8135 |
| Fixed build with `MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF` | `e6288c3ca61`
| #6800 |
| Fixed install RPATH for MIOpenDriver / CK backend libs | `ef8a4cf2f93`
| #7485 |
| Fixed dangling lambda causing log-level 6 segfault | `190f9791737` |
#6953 |
| Enabled Windows long-path support for `addkernels` | `1f9d69724cf` |
#7387 |
| Removed OCL solver `ConvOclDirectFwd1x1` | `126c720cadd` | #6491 |
| Removed OCL solver `conv_ocl_dir2Dfwdgen` | `2d499ddf709` | #6633 |
| Removed OCL solvers `ConvOclBwdWrW1x1`, `ConvOclBwdWrW2` |
`b5d43a33dc6` | #6859 |
| Removed OCL solvers `OpTensorLeadingOnes`, `Op5dTensorGeneric` |
`8642c83b7de` | #7347 |

The remaining ~40 commits in the 68-commit range were deliberately
excluded (tests/CI/build/docs/internal refactors/1:1 OCL→HIP swaps with
no behavior change, per each commit's own stated scope) and are not
represented here.

## Risk Assessment

**Risk Level: 🟢 Low**

This is a documentation-only change to a single Markdown file. No code,
build configuration, tests, or public API surface is touched.

### Impacted Components

- `projects/miopen/CHANGELOG.md` only.

### Potential Side Effects

None expected. This does not change any runtime behavior; it only
documents changes that already shipped on `release/therock-7.14`.

### Mitigation Steps

None required beyond normal review — a reviewer familiar with the 7.14
release can spot-check a few entries against the linked commit hashes/PR
numbers for accuracy.

## Test Plan

Not applicable (no code changes). Verification was done by
cross-referencing every added bullet against a specific commit hash and
PR number pulled from `git log`, and by diffing the edited section
against `git show
remotes/origin/release/therock-7.14:projects/miopen/CHANGELOG.md` to
confirm no existing entry was duplicated or altered.

## Test Result

All 21 new bullets map 1:1 to a real commit/PR in the
`therock-7.13..release/therock-7.14` range for `projects/miopen`. No
pre-existing changelog text was modified — `git diff` shows a pure
addition (24 insertions, 0 deletions). Two borderline commits
(`a67834413e8` "Standalone build", `7c7cd5774a9` "enable warnings about
deprecated API/features") were deliberately left out pending author
judgment on whether they warrant an entry.
amd-alavin added a commit that referenced this pull request Sep 3, 2026
The three SmokeNhwc bf16 instantiations excluded gfx90a through a copy of
GetTestParamsNoGfx90A(), carried from #6633 with the comment "rocBLAS does not
support BF16->BF16 GEMM on gfx90a".

That does not hold. rocblas_gemm_ex_kernels.cpp dispatches bf16/bf16 with f32
compute under no architecture predicate, aldebaran's *_BBS_BH Tensile logic
covers all four transpose variants plus three _GB files, one more than gfx942,
and GemmFwdRest's own bf16 Smoke suite runs on gfx90a unmasked today.

Nothing in IsApplicable gated bf16 on gfx90a either, so these solvers were
selectable there while the tests that would have caught a failure were the only
thing skipped.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants