Skip to content

Remove OCL solvers OpTensorLeadingOnes and Op5dTensorGeneric and add hip replacement solvers - #7347

Merged
iq136boy merged 5 commits into
developfrom
dfeng/remove_ocl_tensor_ops2
Jun 3, 2026
Merged

iq136boy merged 5 commits into
developfrom
dfeng/remove_ocl_tensor_ops2

Conversation

@iq136boy

@iq136boy iq136boy commented May 12, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

MIOpen has been incrementally retiring its OpenCL tensor operation solvers in favour of HIP equivalents to reduce maintenance overhead, eliminate the OpenCL compiler dependency, and align with ROCm's HIP-first strategy. This is the second PR in that series (following PR #6935 which retired Op2dTensorLite, Op4dTensorGeneric, and OpTensorFwdBias).

This PR covers the two remaining OCL tensor solvers that had no HIP kernel yet:

  • OpTensorLeadingOnes – 4-D tensor ops for packed and non-packed (generic stride) paths
  • Op5dTensorGeneric – generic fallback for all 5-D tensor element-wise ops

Technical Details

Three new HIP kernels are appended to src/kernels/MIOpenTensorKernelsHip.cpp:

  • OpTensorLeadingOnes – ports OCL OpTensorLeadingOnes (line ~218), compiled under USE_LEADING_ONES
  • OpTensorLeadingOnesGeneric – ports OCL OpTensorLeadingOnesGeneric (line ~310), compiled under USE_LEADING_ONES_GENERIC
  • Op5dTensorGeneric – ports OCL Op5dTensorGeneric (line ~544), compiled under USE_5D_TENSOR_GENERIC

All kernels are faithful ports of the OCL originals with standard syntax substitutions (__kernel → global, get_global_id(0) → blockIdx.x * blockDim.x + threadIdx.x, barrier(CLK_LOCAL_MEM_FENCE) →
__syncthreads(), (long) → (int64_t), etc.).

Solver-side changes are minimal (2–3 lines each):

  • src/solver/tensorOp/OpTensorLeadingOnes.cpp – switch kbp::OpenCL{} → kbp::HIP{}, kernel file to MIOpenTensorKernelsHip.cpp, and flip the Get4dParams is_hip flag to true.
  • src/solver/tensorOp/Op5dTensorGeneric.cpp – switch kbp::OpenCL{} → kbp::HIP{} and kernel file to MIOpenTensorKernelsHip.cpp.
    No header, solver registration, or CMakeLists.txt changes are required. MIOpenTensorKernels.cl is not deleted as it is still referenced by Op1dTensorGeneric, Op2dTensorGeneric, Op2dTensorSquash, and Op3dTensorGeneric.

Test Plan

  • Build MIOpen to confirm no compile errors.
  • Verify the three new HIP kernels compile cleanly with their respective #ifdef guards (USE_LEADING_ONES, USE_LEADING_ONES_GENERIC, USE_5D_TENSOR_GENERIC).
  • Confirm neither modified solver file references MIOpenTensorKernels.cl after the change.
  • Run existing MIOpen tensor op unit tests covering 4-D non-packed layouts and 5-D cases.
  • Full CI pipeline.

Test Result

  • Build passes with no compile errors.
  • All existing tensor op tests pass (4-D packed, 4-D non-packed/generic stride, 5-D generic).
  • CI pending.

Submission Checklist

  • New HIP kernels are logically equivalent to the OCL originals (same index arithmetic, same bitmap logic, same beta==0 / beta!=0 branches).
  • No unintended changes to solver headers, CMakeLists, or solver registration.
  • MIOpenTensorKernels.cl retained (still needed by other solvers not in scope for this PR).

@iq136boy
iq136boy requested a review from AnaghaRaoAMD June 2, 2026 16:00
@iq136boy
iq136boy merged commit 8642c83 into develop Jun 3, 2026
37 checks passed
@iq136boy
iq136boy deleted the dfeng/remove_ocl_tensor_ops2 branch June 3, 2026 14:51
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.
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