Skip to content

[MIOpen] Remove OCL solver ConvOclDirectFwd1x1 - #6491

Merged
iq136boy merged 19 commits into
developfrom
dfeng/remove_ConvOclDirectFwd1x1
May 7, 2026
Merged

iq136boy merged 19 commits into
developfrom
dfeng/remove_ConvOclDirectFwd1x1

Conversation

@iq136boy

@iq136boy iq136boy commented Apr 16, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

ConvOclDirectFwd1x1 is a legacy OpenCL-based solver for 1x1 forward convolution. As MIOpen continues to migrate away from OpenCL toward HIP, this solver has become a maintenance burden: it relies on three OpenCL kernel files (~1200 lines of OpenCL C), has entries in the find-db for several GPU architectures, and carries the overhead of the OpenCL compilation path. Removing it reduces technical debt and aligns with the broader effort to standardize on HIP-based kernels.

1×1 convolution is mathematically equivalent to a matrix multiply (GEMM), and MIOpen's GEMM-based and CK-based solvers are specifically optimized for exactly this. In practice, those solvers win the find-db tuning for virtually all 1×1 problem shapes. ConvOclDirectFwd1x1 was already rarely — if ever — selected as the best solver, which is precisely why its find-db entries are being cleaned up in this PR.

No HIP replacement solver is introduced. ConvDirectNaiveConvFwd already covers the same problem space (and more — NHWC layout, groups, padding, dilation, INT8/FP8) by dispatching to the same conv_internal::GetConv2DFWDSolution() function that any replacement would use. Since there is no new HIP kernel and no performance difference, a separate solver wrapper for 1×1 problems adds complexity with no benefit.

Additionally, the old OCL solver carried accumulated tech debt that makes it not worth porting:

  • It was restricted to a subset of GPU architectures (gfx8, gfx90a, gfx103), requiring explicit device guards
  • It was disabled on gfx10/gfx11 via WORKAROUND_SWDEV_271887 due to precision issues
  • Padding support was disabled as a workaround for correctness bugs
  • It had a history of FP16 precision fixes

Technical Details

  • Removed ConvOclDirectFwd1x1 solver (src/solver/conv/conv_ocl_dir2Dfwd1x1.cpp, ~433 lines) and its three OpenCL kernel files (MIOpenConv1x1J1.cl, MIOpenConv1x1J1_stride.cl, MIOpenConv1x1S.cl, ~1200 lines total).
  • Skipped the solver ID slot in solver.cpp with ++id; // removed ConvOclDirectFwd1x1 to preserve backward compatibility with existing find-db entries.
  • ConvDirectNaiveConvFwd acts as the fallback for any 1×1 problem shapes not claimed by higher-performance solvers.
  • Updated find-db binary files for gfx803, gfx900, gfx906, gfx908, and gfx90a to remove entries referencing the old solver.

Test Plan

  • Existing unit tests for ConvDirectNaiveConvFwd cover the problem space previously handled by ConvOclDirectFwd1x1.
  • CI test

Test Result

CI test pending

Submission Checklist

  • Removed legacy OpenCL solver and all associated kernel files.
  • Updated find-db, CMake, fin interface, and documentation.
  • No new solver introduced; ConvDirectNaiveConvFwd serves as the fallback.

Comment thread projects/miopen/src/mlo_dir_conv.cpp Resolved

@AnaghaRaoAMD AnaghaRaoAMD 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.

LGTM, minor suggestions posted in comments, needs rebase and ensure there is no accuracy/performance regressions on any gfx platform.

@AnaghaRaoAMD

Copy link
Copy Markdown
Contributor

Rethinking on this some more, we are not introducing any new hip kernel here, there's no real reason for ConvHipDirectFwd1x1 to exist as a separate solver.

  1. Same kernel — Both ConvHipDirectFwd1x1 and ConvDirectNaiveConvFwd call the exact same function: conv_internal::GetConv2DFWDSolution()
  2. ConvDirectNaiveConvFwd is a superset — It accepts all layouts (NCHW + NHWC), all filter sizes, all strides, all padding, all dilation, groups, more data types (FP8, Int8 in addition to FP16/BFP16/FP32). Every problem that ConvHipDirectFwd1x1 would accept, ConvDirectNaiveConvFwd already accepts.
  3. No performance difference — Since they dispatch to the same kernel, there's zero performance advantage to having the 1x1-specific wrapper.

If there's no specific reason to keep a separate solver ID for 1x1 cases, the simpler approach is to just delete ConvOclDirectFwd1x1, skip the ID slot with ++id; in solver.cpp, and let ConvDirectNaiveConvFwd handle the problem space. That would eliminate the new solver file, new test file, and the added complexity entirely.

@iq136boy

Copy link
Copy Markdown
Contributor Author

Rethinking on this some more, we are not introducing any new hip kernel here, there's no real reason for ConvHipDirectFwd1x1 to exist as a separate solver.

  1. Same kernel — Both ConvHipDirectFwd1x1 and ConvDirectNaiveConvFwd call the exact same function: conv_internal::GetConv2DFWDSolution()
  2. ConvDirectNaiveConvFwd is a superset — It accepts all layouts (NCHW + NHWC), all filter sizes, all strides, all padding, all dilation, groups, more data types (FP8, Int8 in addition to FP16/BFP16/FP32). Every problem that ConvHipDirectFwd1x1 would accept, ConvDirectNaiveConvFwd already accepts.
  3. No performance difference — Since they dispatch to the same kernel, there's zero performance advantage to having the 1x1-specific wrapper.

If there's no specific reason to keep a separate solver ID for 1x1 cases, the simpler approach is to just delete ConvOclDirectFwd1x1, skip the ID slot with ++id; in solver.cpp, and let ConvDirectNaiveConvFwd handle the problem space. That would eliminate the new solver file, new test file, and the added complexity entirely.

Yes. You are right! The new hip solver is not needed. I will address it. Thanks.

iq136boy and others added 4 commits April 30, 2026 13:54
The previous merge incorrectly took develop's version of the FindDB
binary files. This PR modified these files, so our version should be
kept.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
The previous merge resolution was incorrect. Develop's FindDB changes
(PR #6164) removed ConvOclDirectFwd11x11 entries, while our branch
removed ConvOclDirectFwd1x1 entries. These are independent changes
that both need to be present. Rebuild the FindDB files by starting
from develop's version and stripping ConvOclDirectFwd1x1 entries.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@iq136boy iq136boy changed the title [MIOpen] Remove OCL solver ConvOclDirectFwd1x1 from MIOpen and add the replacement hip solver [MIOpen] Remove OCL solver ConvOclDirectFwd1x1 Apr 30, 2026
@iq136boy
iq136boy merged commit 126c720 into develop May 7, 2026
36 checks passed
@iq136boy
iq136boy deleted the dfeng/remove_ConvOclDirectFwd1x1 branch May 7, 2026 13:36
aledudek pushed a commit that referenced this pull request May 20, 2026
## Motivation
ConvOclDirectFwd1x1 is a legacy OpenCL-based solver for 1x1 forward
convolution. As MIOpen continues to migrate away from OpenCL toward HIP,
this solver has become a maintenance burden: it relies on three OpenCL
kernel files (~1200 lines of OpenCL C), has entries in the find-db for
several GPU architectures, and carries the overhead of the OpenCL
compilation path. Removing it reduces technical debt and aligns with the
broader effort to standardize on HIP-based kernels.

1×1 convolution is mathematically equivalent to a matrix multiply
(GEMM), and MIOpen's GEMM-based and CK-based solvers are specifically
optimized for exactly this. In practice, those solvers win the find-db
tuning for virtually all 1×1 problem shapes. ConvOclDirectFwd1x1 was
already rarely — if ever — selected as the best solver, which is
precisely why its find-db entries are being cleaned up in this PR.

No HIP replacement solver is introduced. ConvDirectNaiveConvFwd already
covers the same problem space (and more — NHWC layout, groups, padding,
dilation, INT8/FP8) by dispatching to the same
`conv_internal::GetConv2DFWDSolution()` function that any replacement
would use. Since there is no new HIP kernel and no performance
difference, a separate solver wrapper for 1×1 problems adds complexity
with no benefit.

Additionally, the old OCL solver carried accumulated tech debt that
makes it not worth porting:
- It was restricted to a subset of GPU architectures (gfx8, gfx90a,
gfx103), requiring explicit device guards
- It was disabled on gfx10/gfx11 via WORKAROUND_SWDEV_271887 due to
precision issues
  - Padding support was disabled as a workaround for correctness bugs
  - It had a history of FP16 precision fixes

## Technical Details
- Removed ConvOclDirectFwd1x1 solver
(src/solver/conv/conv_ocl_dir2Dfwd1x1.cpp, ~433 lines) and its three
OpenCL kernel files (MIOpenConv1x1J1.cl, MIOpenConv1x1J1_stride.cl,
MIOpenConv1x1S.cl, ~1200 lines total).
- Skipped the solver ID slot in solver.cpp with `++id; // removed
ConvOclDirectFwd1x1` to preserve backward compatibility with existing
find-db entries.
- ConvDirectNaiveConvFwd acts as the fallback for any 1×1 problem shapes
not claimed by higher-performance solvers.
- Updated find-db binary files for gfx803, gfx900, gfx906, gfx908, and
gfx90a to remove entries referencing the old solver.

## Test Plan
- Existing unit tests for ConvDirectNaiveConvFwd cover the problem space
previously handled by ConvOclDirectFwd1x1.
- CI test

## Test Result
CI test passed

## Submission Checklist
- Removed legacy OpenCL solver and all associated kernel files.
- Updated find-db, CMake, fin interface, and documentation.
- No new solver introduced; ConvDirectNaiveConvFwd serves as the
fallback.

---------

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.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.
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.

2 participants