Repository navigation
[MIOpen] Remove ocl solver ConvOclBwdWrW1x1 and ConvOclBwdWrW2 - #6859
Merged
Merged
Conversation
…ocl_dir2D_bwdWrW_1x1
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
iq136boy
changed the title
iq136boy
changed the title
DarylHawkinsAMD
approved these changes
May 6, 2026
Resolve finddb conflicts by taking develop's latest fdb files and re-applying ConvOclBwdWrW1x1 and ConvOclBwdWrW2 removal on top. Doc conflicts in debug-log.rst and env_variables.rst resolved by keeping our removal of WRW2 entries (replaced by FWD1X1 entries). Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
aledudek
pushed a commit
that referenced
this pull request
May 20, 2026
## Motivation
This PR removes two legacy OpenCL-based backward-weight convolution
solvers as part of an ongoing OCL→HIP migration effort:
- **ConvOclBwdWrW1x1**: ~548 lines of OCL solver code + 26 KB kernel
file. 1×1 convolutions are mathematically equivalent to matrix
multiplication, and existing HIP-based solvers already provide superior
coverage.
- **ConvOclBwdWrW2** (variants N=1,2,4,8,16 and NonTunable): ~784 lines
of OCL solver code + two kernel files (~26 KB each). Find-db data
confirms these solvers were never selected as optimal on modern hardware
(MI100, MI200, RDNA2), only appearing on legacy gfx8/gfx9 hardware.
- No HIP replacement solvers are added — existing HIP-based solvers
already provide equivalent and superior coverage for both problem types.
## Technical Details
**Removed:**
- OCL solver implementations, associated kernel files, unit tests,
solver
registrations, fin interface entries, and related find-db entries for
both
`ConvOclBwdWrW1x1` and the `ConvOclBwdWrW2` family.
**Updated:**
- Build configuration, solver list, and debug environment variable
documentation.
## Test Plan
- CI pipeline
## Test Result
- CI pending
## Submission Checklist
- Removed OpenCL kernel and solver
- Updated solver registry with ID slot preserved
- Updated find-db files
- Updated documentation
---------
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
iq136boy
added a commit
that referenced
this pull request
Jul 8, 2026
## Motivation MIOpen is retiring its OpenCL backend. All OCL solvers have already been converted to HIP in prior PRs (#6859, #7281, #7637, #7047, #7347, #7575). This PR comprehensively removes the OpenCL backend — conditional compilation blocks, OCL-only source files, OpenCL compiler infrastructure, and all remaining OCL residue — keeping only the HIP code paths. This is dead-code removal with no behavior changes for HIP builds. ## Summary of Changes 118 files changed, 214 insertions, 16,129 deletions ### Phase 1: Remove `#if MIOPEN_BACKEND_OPENCL` conditionals - **Core headers (4):** Remove OCL includes (`CL/cl.h`), OCL typedefs (`cl_command_queue`, `cl_mem`), and OCL kernel aliases from `miopen.h`, `common.hpp`, `kernel.hpp`, `mlo_internal.hpp` - **Core source (5):** Remove OCL-specific logic (`IsAmdRocmOpencl`, `CalculateIsAmdRocmOpencl`, `IsTokenWithin`), OCL backend string in logger, OCL warning prefix, and OCL subbuffer alignment query from `execution_context.cpp/.hpp`, `logger.cpp`, `kernel_warnings.cpp`, `subbuffers.cpp` - **MLIR invoker (1):** Remove `EXPAND_MLIR_CONV_ARGS` macros and 4 OCL kernel invocation blocks from `mlir_impl_gemm.cpp` - **Handle header (1):** Remove `#if MIOPEN_BACKEND_HIP` guard from `CreateSubBuffer` in `handle.hpp` - **Env var cleanup (3):** Remove `MIOPEN_DEBUG_OPENCL_CONVOLUTIONS` env var declaration, registry entry, and documentation - **Driver files (16):** Remove OCL typedefs, `GPUMem`/`miopenCreate`/`GetStream` OCL paths, and `clGetCommandQueueInfo` blocks - **Test files (9):** Remove OCL skip lists, OCL `Device`/`DevMem` implementations, OCL kernel type dispatch, OCL error test, OCL backend defines ### Phase 2: Delete OCL-only source files and build infrastructure - **Delete 10 OCL source files:** `src/ocl/activ_ocl.cpp`, `batchnormocl.cpp`, `convolutionocl.cpp`, `dropoutocl.cpp`, `fusionopbiasbnactivocl.cpp`, `fusionopconvocl.cpp`, `lrn_ocl.cpp`, `mloNorm.cpp`, `rnnocl.cpp`, `utilocl.cpp` - **Delete OpenCL find database files:** `src/kernels/find_db/gfx*.fdb.txt` OCL entries - **Move `gcn_asm_utils.cpp`** from `src/ocl/` to `src/hip/` (used by 27 HIP solver files, not OCL-specific) - **Rename `ocl_wrw_rdc`** invoker to `wrw_rdc` (used by HIP solver `ConvHipBwdWrW53`, not OCL-specific) - **Remove OpenCL compiler infrastructure:** `BuildOcl()` from `comgr.cpp`/`comgr.hpp`, `getOpenCLCompilerPath()` from `hip_build_utils.cpp`, `HIP_OC_COMPILER` from `config.h.in` and `CMakeLists.txt` - **Remove `kbp::OpenCL`** kernel build parameter struct (identical output to `kbp::HIP`) - **Remove OCL debug env vars:** `MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1`, `MIOPEN_DEBUG_OPENCL_*` env vars from `hipoc_program.cpp` - **Clean up OCL code paths** in `hipoc_program.cpp`: remove `DetectCodeObjectOptionSyntax()`, `DetectCodeObjectVersion()`, offline OCL compiler invocation ### Phase 3: Rename OCL-named symbols and clean up residual references - **Rename batchnorm functions:** `IsOCLInferTypeValid` → `IsInferTypeValid`, `IsOCLFwdTrainTypeValid` → `IsFwdTrainTypeValid`, `IsOCLBwdTypeValid` → `IsBwdTypeValid` - **Rename variables:** `bot_ocl_buf` → `bot_buf`, `top_ocl_buf` → `top_buf`, `wei_ocl_buf` → `wei_buf`, etc. across solver files - **Rename enum:** `OpenclText` → `KernelText` in fusion.hpp - **Rename function:** `IsAmdRocmOpencl` → `DetectAmdRocmMetadata` in execution_context.cpp - **Remove stale OpenCL comments** across ~20 files - **Remove OpenCL search paths** from CMakeLists.txt ### Phase 4: Remove `#pragma OPENCL EXTENSION` and dead non-HIP code paths - **Remove `#pragma OPENCL EXTENSION`** from `batchnorm_functions.h` and `float_types.h` - **Remove dead `#else` (non-HIP) branches** from `float_types.h` — eliminates `_FLOAT`/`_FLOAT_ACCUM` OpenCL naming - **Remove `-Wno-pedantic-core-features`** warning suppression (no longer needed) - **Remove commented-out `opencl_unroll_hint`** attributes from `MIOpenBatchNormActivFwdTrainSpatial.cpp` ## Test Plan - HIP backend build compiles cleanly (`libMIOpen.so` verified) - `grep -rn "MIOPEN_BACKEND_OPENCL" src/ include/ driver/ test/` returns zero matches (only `config.h.in` remains for build-system use) - `grep -rn "#pragma OPENCL\|BuildOcl\|HIP_OC_COMPILER\|kbp::OpenCL\|MIOPEN_DEBUG_OPENCL\|MIOPEN_DEBUG_CONV_DIRECT_OCL" src/ include/` returns zero matches - Existing gtests pass without regressions - CI pending --------- Co-authored-by: Claude Opus 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.
3 tasks done
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Nov 4, 2026
…OCM-23567) (#14905) JIRA ID: ROCM-23567 ## Summary `Smoke/GPU_UnitTestConvSolverHipImplicitGemmV4R1WrW_FP32.ConvHipImplicitGemmV4R1WrW/0` has been failing in TheRock's nightly gfx103X run since 2026-04 (ROCM-23567). Root cause is a test calibration gap, not a product defect: - `GetTestParams()` declares `Gpu::gfx103X` in `supported_gpus` but only calls `SetTolerance` for `gfx908`/`gfx90A` (at `40.0f`). - gfx1030 therefore falls back to the default tolerance of `1.0`, making the threshold exactly `FLT_EPSILON` = `1.1920928955078125e-07` — matching the reported threshold bit-for-bit. - The verification metric is `rms_range` (RMS relative error), so the reported `2.2468e-07` is **1.88 ULP** of accumulation-order divergence on a backward-weights reduction — not a correctness problem. This adds `p.SetTolerance(Gpu::gfx103X, miopenFloat, 40.0f);`, matching the existing `gfx908`/`gfx90A` value. An initial attempt used `4.0f` (~2.1x headroom over the ticket's reported 1.88x), but real-hardware testing (see below) showed ~30 ULP error on this device — well over `4.0f` — so the value was raised to match gfx9 rather than pick another unverified intermediate number. This solver's `static_ck` numerics already carry that same documented looseness on gfx9 (see the existing `/// \todo 40.0f is too much` comment above the gfx9 lines). Of the four tests originally reported in ROCM-23567, this is the only one still live — the two `ConvOclBwdWrW2` tests were removed with their solvers (#6859), and `MIOpenDriverRegressionBigTensorTest_FP32` is now excluded from every `test_categories.yaml` category (#6198). ## Test plan - [x] `cmake --build projects/miopen/build --target miopen_gtest` — builds clean on gfx1151. - [x] `./miopen_gtest.exe --gtest_filter='*UnitTestConvSolverHipImplicitGemmV4R1WrW*'` on local gfx1151 — GPU FP32/FP16/BFP16 variants report `SKIPPED` (gfx1151 isn't in `supported_gpus`), CPU applicability tests still `OK`. Confirms the `SetTolerance` call is inert off gfx103X. - [x] **Validated on real gfx1030 hardware** (AMD Radeon Pro V620, the same SKU named in the original defect report) via an AMD-internal Slurm cluster, since rocm-libraries pre-merge CI does not cover gfx103X (gfx942-only) and TheRock's gfx1030 CI runners are currently disabled cluster-wide (`therock-ci-config#57`, known-bad-runner exclusion, unrelated to this change). Built MIOpen from this branch inside `rocm/miopen-dev:latest` and ran `test_unit_conv_solver_ConvHipImplicitGemmV4R1WrW --gtest_filter='*ConvHipImplicitGemmV4R1WrW*'`: - At the default tolerance (no override): FP32 failed, `3.59e-06` actual vs `1.19e-07` threshold — reproduces the ticket. - At `4.0f`: FP32 still failed, `3.59e-06` actual vs `4.77e-07` threshold (~30 ULP, not the ~1.88 ULP the ticket reported) — confirms `4.0f` was insufficient and motivated the revision to `40.0f`. - At `40.0f` (final): **all 5 tests pass** — `[ PASSED ] 5 tests.` 🤖 Generated with [Claude Code](https://claude.com/claude-code)
cycheng
pushed a commit
that referenced
this pull request
Nov 6, 2026
…OCM-23567) (#14905) JIRA ID: ROCM-23567 ## Summary `Smoke/GPU_UnitTestConvSolverHipImplicitGemmV4R1WrW_FP32.ConvHipImplicitGemmV4R1WrW/0` has been failing in TheRock's nightly gfx103X run since 2026-04 (ROCM-23567). Root cause is a test calibration gap, not a product defect: - `GetTestParams()` declares `Gpu::gfx103X` in `supported_gpus` but only calls `SetTolerance` for `gfx908`/`gfx90A` (at `40.0f`). - gfx1030 therefore falls back to the default tolerance of `1.0`, making the threshold exactly `FLT_EPSILON` = `1.1920928955078125e-07` — matching the reported threshold bit-for-bit. - The verification metric is `rms_range` (RMS relative error), so the reported `2.2468e-07` is **1.88 ULP** of accumulation-order divergence on a backward-weights reduction — not a correctness problem. This adds `p.SetTolerance(Gpu::gfx103X, miopenFloat, 40.0f);`, matching the existing `gfx908`/`gfx90A` value. An initial attempt used `4.0f` (~2.1x headroom over the ticket's reported 1.88x), but real-hardware testing (see below) showed ~30 ULP error on this device — well over `4.0f` — so the value was raised to match gfx9 rather than pick another unverified intermediate number. This solver's `static_ck` numerics already carry that same documented looseness on gfx9 (see the existing `/// \todo 40.0f is too much` comment above the gfx9 lines). Of the four tests originally reported in ROCM-23567, this is the only one still live — the two `ConvOclBwdWrW2` tests were removed with their solvers (#6859), and `MIOpenDriverRegressionBigTensorTest_FP32` is now excluded from every `test_categories.yaml` category (#6198). ## Test plan - [x] `cmake --build projects/miopen/build --target miopen_gtest` — builds clean on gfx1151. - [x] `./miopen_gtest.exe --gtest_filter='*UnitTestConvSolverHipImplicitGemmV4R1WrW*'` on local gfx1151 — GPU FP32/FP16/BFP16 variants report `SKIPPED` (gfx1151 isn't in `supported_gpus`), CPU applicability tests still `OK`. Confirms the `SetTolerance` call is inert off gfx103X. - [x] **Validated on real gfx1030 hardware** (AMD Radeon Pro V620, the same SKU named in the original defect report) via an AMD-internal Slurm cluster, since rocm-libraries pre-merge CI does not cover gfx103X (gfx942-only) and TheRock's gfx1030 CI runners are currently disabled cluster-wide (`therock-ci-config#57`, known-bad-runner exclusion, unrelated to this change). Built MIOpen from this branch inside `rocm/miopen-dev:latest` and ran `test_unit_conv_solver_ConvHipImplicitGemmV4R1WrW --gtest_filter='*ConvHipImplicitGemmV4R1WrW*'`: - At the default tolerance (no override): FP32 failed, `3.59e-06` actual vs `1.19e-07` threshold — reproduces the ticket. - At `4.0f`: FP32 still failed, `3.59e-06` actual vs `4.77e-07` threshold (~30 ULP, not the ~1.88 ULP the ticket reported) — confirms `4.0f` was insufficient and motivated the revision to `40.0f`. - At `40.0f` (final): **all 5 tests pass** — `[ PASSED ] 5 tests.` 🤖 Generated with [Claude Code](https://claude.com/claude-code)
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
This PR removes two legacy OpenCL-based backward-weight convolution solvers as part of an ongoing OCL→HIP migration effort:
ConvOclBwdWrW1x1: ~548 lines of OCL solver code + 26 KB kernel file. 1×1 convolutions are mathematically equivalent to matrix multiplication, and existing HIP-based solvers already provide superior coverage.
ConvOclBwdWrW2 (variants N=1,2,4,8,16 and NonTunable): ~784 lines of OCL solver code + two kernel files (~26 KB each). Find-db data confirms these solvers were never selected as optimal on modern hardware (MI100, MI200, RDNA2), only appearing on legacy gfx8/gfx9 hardware.
No HIP replacement solvers are added — existing HIP-based solvers already provide equivalent and superior coverage for both problem types.
Technical Details
Removed:
registrations, fin interface entries, and related find-db entries for both
ConvOclBwdWrW1x1and theConvOclBwdWrW2family.Updated:
documentation.
Test Plan
Test Result
Submission Checklist