Skip to content

[MIOpen] Convert the OpenCL conv solver conv_ocl_dir2D11x11 to hip solver - #6164

Merged
iq136boy merged 25 commits into
developfrom
dfeng/remove_conv2dfwdCL
Apr 28, 2026
Merged

iq136boy merged 25 commits into
developfrom
dfeng/remove_conv2dfwdCL

Conversation

@iq136boy

@iq136boy iq136boy commented Apr 3, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Convert the OpenCL kernel and solver for the 11×11 direct forward convolution (ConvOclDirectFwd11x11) to hip solver as part of the ongoing effort to replace OpenCL solvers/kernels with hip solvers/kernels in MIOpen.

Technical Details

  • Removed the OpenCL solver conv_ocl_dir2D11x11.cpp and its kernel MIOpenConvFwd_LxL_11.cl (1466 lines).
  • Added a replacement HIP solver conv_hip_dir2D11x11.cpp and HIP kernel MIOpenConvFwd_LxL_11.cpp, which is a direct port of the OpenCL kernel supporting the same 11×11, stride-4 tiled forward convolution
    algorithm.
  • Replaced ConvOclDirectFwd11x11 with ConvHipDirectFwd11x11 in the solver registry (solver.cpp) and solver list (mlo_dir_conv.cpp).
  • Updated all find-db files (HIP and OpenCL) that referenced the old solver.
  • Extended GPU support to gfx90x, gfx94x, gfx95x, gfx103x, gfx110x, gfx115x, and gfx120x.
  • Updated docs to reflect the renamed environment variable (MIOPEN_DEBUG_CONV_DIRECT_HIP_FWD11X11).

Test Plan

  • Run GPU_UnitTestConvSolverHipDirectFwd11x11 for FP16, BFP16, and FP32 on supported hardware.
  • Run CPU_UnitTestConvSolverHipDirectFwd11x11DevApplicability_NONE for device applicability checks.

Test Result

  • All unit tests pass on gfx94x (FP32 tolerance set to 3.0× ULP to account for accumulation differences across architectures).
  • CI test pending

Submission Checklist

  • Removed OpenCL kernel and solver
  • Added replacement HIP kernel and solver
  • Updated solver registry with ID slot preserved
  • Updated find-db files
  • Added unit tests for the new HIP solver
  • Updated documentation

@iq136boy iq136boy changed the title Deprecate the opencl conv solver conv_ocl_dir2D11x11 [MIOpen] Convert the OpenCL conv solver conv_ocl_dir2D11x11 to hip solver Apr 6, 2026

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

I would like a bit more rigor given how frequently broken kernels have gotten through. Could you please do some elementwise comparisons against the original code and share the results?

@iq136boy

iq136boy commented Apr 9, 2026 •

Copy link
Copy Markdown
Contributor Author

I would like a bit more rigor given how frequently broken kernels have gotten through. Could you please do some elementwise comparisons against the original code and share the results?

I understand the concern. While there is no direct element-wise comparison between the OCL and HIP outputs, there is an indirect one — both solvers are independently verified against the same CPU reference in the unit test suite, which is mathematically equivalent. Here is why the unit tests are sufficient:
Both solvers are tested against the same CPU reference (computed in FP64 accumulation), using the same fixed PRNG seed and identical input tensors. By the triangle inequality, for any output element at index [i]:
|OCL_output[i] - HIP_output[i]|≤ |OCL_output[i] - CPU_ref[i]| + |HIP_output[i] - CPU_ref[i]| ≤ 2 × threshold
where threshold = epsilon(T) × tolerance_multiplier:

  • FP32: 3 × 1.19e-7 ≈ 3.58e-7 (tolerance multiplier = 3.0, set explicitly in both tests)
  • FP16 / BF16: 1 × 9.77e-4 (default tolerance multiplier = 1.0)

Since both unit tests pass on all three datatypes with the same test cases, this provides a mathematically rigorous bound on the deviation between the two implementations — and also illustrates exactly why having these unit tests matters.

Comment thread projects/miopen/src/solver.cpp Resolved
Comment thread projects/miopen/src/CMakeLists.txt 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, some CI failures need attention

@iq136boy

iq136boy commented Apr 13, 2026 •

Copy link
Copy Markdown
Contributor Author

LGTM, some CI failures need attention

Thanks so much for the comments. I have checked about the CI failure. It seems that the failure is unrelated to this PR. It refers to ConvHipImplicitGemmGroupWrwXdlops on gfx942, which is neither touched by this PR nor present in any of the modified fdb files. The invalid perf config in gfx942130.db.txt should be a pre-existing issue in the performance database.

@randyspauldingamd

Copy link
Copy Markdown
Contributor

I would like a bit more rigor given how frequently broken kernels have gotten through. Could you please do some elementwise comparisons against the original code and share the results?

I understand the concern. While there is no direct element-wise comparison between the OCL and HIP outputs, there is an indirect one — both solvers are independently verified against the same CPU reference in the unit test suite, which is mathematically equivalent. Here is why the unit tests are sufficient: Both solvers are tested against the same CPU reference (computed in FP64 accumulation), using the same fixed PRNG seed and identical input tensors. By the triangle inequality, for any output element at index [i]: |OCL_output[i] - HIP_output[i]|≤ |OCL_output[i] - CPU_ref[i]| + |HIP_output[i] - CPU_ref[i]| ≤ 2 × threshold where threshold = epsilon(T) × tolerance_multiplier:

  • FP32: 3 × 1.19e-7 ≈ 3.58e-7 (tolerance multiplier = 3.0, set explicitly in both tests)
  • FP16 / BF16: 1 × 9.77e-4 (default tolerance multiplier = 1.0)

Since both unit tests pass on all three datatypes with the same test cases, this provides a mathematically rigorous bound on the deviation between the two implementations — and also illustrates exactly why having these unit tests matters.

While you're correct in regard to elementwise comparisons, I remind everyone that these tests do not do that. They use RRMSE, which can hide many significant elementwise inaccuracies. In fact, this is the purpose of using it. It is dangerous to use as a standalone metric though, and what I'm asking for is a sanity check to gauge how bad this is in practice.

You may have overgeneralized in your example. I worked through what I think you had in mind. First, go back to the base definition of triangle inequality by defining:

x[i] = OCL_output[i] - CPU_ref[i]
y[i] = HIP_output[i] - CPU_ref[i]
=> OCL_output[i] - HIP_output[i] = x[i] - y[i]

You're trying to relate RRMSE(|x - y|) to RRMSE(|x|) and RRMSE(|y|) since we have information about the latter two. Lets try it! Note, M[v] refers to the 'max' term. For readability I'll move the M's to the right and just look at RMSE, and also skip writing unneeded absolute value symbols:

RMSE(x) = sqrt(1/N * sum(x[i]^2) <= M[x]*T
RMSE(y) = sqrt(1/N * sum(y[i]^2) <= M[y]*T

Now apply RMSE to the triangle inequality:

LHS: RMSE(|x - y|)   = sqrt(1/N*sum((x[i] - y[i])^2))     = sqrt(1/N*(sum(x[i]^2) + sum(y[i]^2) - 2 * sum(x[i] * y[i])))     <=
RHS: RMSE(|x| + |y|) = sqrt(1/N*sum((|x[i]| + |y[i]|)^2)) = sqrt(1/N*(sum(x[i]^2) + sum(y[i]^2) + 2 * sum(|x[i] * y[i]|)))

After squaring both sides and multiplying by N, all the sum-of-squares terms cancel and we are left with:

-sum(x[i] * y[i]) <= sum(|x[i] * y[i]|)

Which is true by inspection and, assuming my method and algebra are correct, shows that the triangle inequality does not help when comparing RRMSE.

We should really be looking at the histograms for all three sets, OCL vs HIP, OCL vs CPU, HIP vs CPU.

@iq136boy

iq136boy commented Apr 14, 2026 •

Copy link
Copy Markdown
Contributor Author

Hi @randyspauldingamd, you're right about the proof. The triangle inequality does not transfer to RRMSE the way I described — thank you for working through the math carefully.
I'd like to reframe the core question though. This PR is a part of the effort of retiring OCL solvers from MIOpen. The relevant question is not "are OCL and HIP element-wise equivalent?" but rather "after removing OCL, is the system still correct?" That question is answered by the HIP solver's own independent validation against the FP64 CPU reference, which is what the unit tests verify.
A direct OCL vs HIP element-wise comparison would tell us how similar the two implementations are to each other, but not which one is more correct — and since we're retiring OCL, that comparison has limited value going forward. The HIP solver's correctness stands on its own test coverage.
If the broader concern is that RRMSE can mask large element-wise errors in the HIP solver itself, I agree that is a legitimate concern — but it is a property of the existing test infrastructure, not something introduced by this PR. That would be worth addressing separately.

@iq136boy

iq136boy commented Apr 21, 2026 •

Copy link
Copy Markdown
Contributor Author

Hi @randyspauldingamd , please check the attacked results of element comparisons.

Comparison Analysis: OCL vs HIP Solver vs CPU Reference

Total output elements: 81 (I dumped the result to a text file and each line contains one element, datatype fp32)
--- OCL solver vs CPU reference ---
Differing values : 12 / 81
Max abs diff : 1.000000e+14
RMS diff : 2.757795e+13
Line 6: -5.4061e+17 vs -5.40608e+17
Line 10: -4.97138e+19 vs -4.97139e+19
Line 13: 5.68387e+18 vs 5.68386e+18
Line 19: -6.08214e+18 vs -6.08213e+18
Line 35: -8.9369e+19 vs -8.93691e+19
Line 39: 3.95324e+19 vs 3.95323e+19
Line 46: 9.94519e+19 vs 9.94518e+19
Line 60: 2.48852e+19 vs 2.48853e+19
Line 66: 6.01315e+19 vs 6.01314e+19
Line 70: 4.54517e+18 vs 4.54518e+18
Line 73: -6.29509e+18 vs -6.29506e+18
Line 78: 3.7039e+18 vs 3.70388e+18

--- HIP solver vs CPU reference ---
Differing values : 11 / 81
Max abs diff : 1.000000e+14
RMS diff : 2.524840e+13
Line 6: -5.4061e+17 vs -5.40604e+17
Line 10: -4.97138e+19 vs -4.97139e+19
Line 13: 5.68387e+18 vs 5.68386e+18
Line 19: -6.08214e+18 vs -6.08213e+18
Line 35: -8.9369e+19 vs -8.93691e+19
Line 46: 9.94519e+19 vs 9.94518e+19
Line 60: 2.48852e+19 vs 2.48853e+19
Line 66: 6.01315e+19 vs 6.01314e+19
Line 70: 4.54517e+18 vs 4.54518e+18
Line 73: -6.29509e+18 vs -6.29506e+18
Line 78: 3.7039e+18 vs 3.70388e+18

The differences are all at the last displayed significant digit (e.g. 5.68386e+18 vs 5.68387e+18). The one extra diff for OCL is at line 39 (3.95323e+19 vs 3.95324e+19) — the HIP solver happens to match the CPU reference exactly on that value while OCL doesn't. Overall both solvers agree with CPU to within floating-point rounding tolerance, and all tests pass. The small divergences are normal for GPU convolutions due to non-associative FP accumulation order.

--- OCL solver vs HIP solver ---
Differing values : 2 / 81
Max abs diff : 1.000000e+14
RMS diff : 1.112000e+13
Line 6: -5.40608e+17 vs -5.40604e+17
Line 39: 3.95323e+19 vs 3.95324e+19
For the comparison of OCL solver vs Hip solver, the differences are tiny (last significant digit), consistent with floating-point non-associativity — the OCL and HIP kernels likely accumulate in slightly different order or use different rounding, resulting in ULP-level divergence. Both are within tolerance of the CPU reference (which is why both tests pass).

Summary

OCL vs CPU RMS: 2.757795e+13
HIP vs CPU RMS: 2.524840e+13
OCL vs HIP RMS: 1.112000e+13

@iq136boy
iq136boy dismissed randyspauldingamd’s stale review April 28, 2026 14:44

The accuracy and performance comparison tests of the ocl and hip solver show that they are equivalent to each other. About the concern of the current test infrastructure, that worth to address separately.

@iq136boy
iq136boy merged commit 7dfb005 into develop Apr 28, 2026
36 checks passed
@iq136boy
iq136boy deleted the dfeng/remove_conv2dfwdCL branch April 28, 2026 14:45
iq136boy added a commit that referenced this pull request Apr 30, 2026
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>
aledudek pushed a commit that referenced this pull request May 20, 2026
…lver (#6164)

## Motivation
Convert the OpenCL kernel and solver for the 11×11 direct forward
convolution (ConvOclDirectFwd11x11) to hip solver as part of the ongoing
effort to replace OpenCL solvers/kernels with hip solvers/kernels in
MIOpen.

## Technical Details
- Removed the OpenCL solver conv_ocl_dir2D11x11.cpp and its kernel
MIOpenConvFwd_LxL_11.cl (1466 lines).
- Added a replacement HIP solver conv_hip_dir2D11x11.cpp and HIP kernel
MIOpenConvFwd_LxL_11.cpp, which is a direct port of the OpenCL kernel
supporting the same 11×11, stride-4 tiled forward convolution
  algorithm.
- Replaced ConvOclDirectFwd11x11 with ConvHipDirectFwd11x11 in the
solver registry (solver.cpp) and solver list (mlo_dir_conv.cpp).
- Updated all find-db files (HIP and OpenCL) that referenced the old
solver.
- Extended GPU support to gfx90x, gfx94x, gfx95x, gfx103x, gfx110x,
gfx115x, and gfx120x.
- Updated docs to reflect the renamed environment variable
(MIOPEN_DEBUG_CONV_DIRECT_HIP_FWD11X11).

## Test Plan
- Run GPU_UnitTestConvSolverHipDirectFwd11x11 for FP16, BFP16, and FP32
on supported hardware.
- Run CPU_UnitTestConvSolverHipDirectFwd11x11DevApplicability_NONE for
device applicability checks.

## Test Result
- All unit tests pass on gfx94x (FP32 tolerance set to 3.0× ULP to
account for accumulation differences across architectures).
 - CI test pending

## Submission Checklist
  - Removed OpenCL kernel and solver
  - Added replacement HIP kernel and solver
  - Updated solver registry with ID slot preserved
  - Updated find-db files
  - Added unit tests for the new HIP solver
  - Updated documentation

---------

Co-authored-by: BradPepersAMD <Brad.Pepers@amd.com>
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
JH-Leon-KIM-AMD added a commit that referenced this pull request May 20, 2026
…hip rename

PR #6164 on develop renamed solver/conv/conv_ocl_dir2D11x11.cpp to
conv_hip_dir2D11x11.cpp (OpenCL solver converted to HIP). The uplift
branch picked up the renamed source file but not the corresponding
CMakeLists.txt update, causing:

  CMake Error at src/CMakeLists.txt:870 (add_library):
    Cannot find source file:
      solver/conv/conv_ocl_dir2D11x11.cpp

This patch updates CMakeLists.txt to reference the new filename,
matching the state on develop.
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.

5 participants