Skip to content

fix(miopen): add naive reference kernel for 3D INT8 forward convolutions - #10051

Merged
amd-ahyttine merged 16 commits into
developfrom
users/bartgips/naive_conv_3d_int8
Aug 28, 2026
Merged

amd-ahyttine merged 16 commits into
developfrom
users/bartgips/naive_conv_3d_int8

Conversation

@amd-bartgips

@amd-bartgips amd-bartgips commented Jul 28, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID : ALMIOPEN-2275

Motivation

MIOpenDriver verification fails for 3D INT8 forward convolutions in both NCDHW and NDHWC
layouts (SILOTIGER-709 / ALMIOPEN-2275). The fast solver
(ConvHipImplicitGemm3DGroupFwdXdlops) runs successfully, but the naive reference
kernel — used for output verification — aborts with a HIP symbol error:

MIOpen Error: named symbol not found
Failed to get function: naive_conv_ab_nonpacked_fwd_ndhwc_int8_t_int32_t_int8_t_0
reference kernel fail to run ConvDirectNaiveConvFwd

The same failure occurs for the NCDHW variant
(naive_conv_ab_nonpacked_fwd_ncdhw_int8_t_int32_t_int8_t_0).

Technical Details

Root cause. The JIT-compiled reference kernel (src/kernels/gpu_reference_kernel/naive_conv.cpp)
instantiates typed kernel variants via the DEFINE_3D_NAIVE_CONV_KERNEL macro. The
int8_t-output variant (…, int8_t, int32_t, int8_t, 0) was present for 2D
(nchw/nhwc) but had never been added for 3D (ncdhw/ndhwc), even though the
int32_t- and float-output 3D variants already existed and the underlying template
logic for 3D INT8 is correct.

Fix. Two macro instantiation lines added; no logic change:

DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, int8_t, int32_t, int8_t, 0)
DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, int8_t, int32_t, int8_t, 0)

Regression test (test/gtest/test_naive_conv_3d_int8_fwd.cpp). The test runs
ConvDirectNaiveConvFwd (the naive solver itself) on two 3D INT8 cases — one NCDHW
(default layout) and one explicit NDHWC via TensorDescriptorParams{miopenTensorNDHWC}.
Both use type_y=miopenInt8 (the previously missing Tout=int8_t symbol path), and
UseCpuRef() so the CPU provides the reference values while the GPU naive kernel is the
solution-under-test. Channel count is kept small (C=16) to satisfy the test data
generator's num_add < int8_max=127 overflow guard for int8 accumulation.

Test Plan

Build with the existing projects/miopen/build config
(MIOPEN_BACKEND=HIP, GPU_TARGETS=gfx942;gfx950, MIOPEN_TEST_DISCRETE=ON):

# Re-run cmake to pick up new glob entry, then build
cmake -DMIOPEN_TEST_DISCRETE=ON .
ninja test_test_naive_conv_3d_int8_fwd \
      test_unit_conv_solver_ConvDirectNaiveConvFwd \
      test_unit_conv_solver_ConvDirectNaiveConvBwd \
      test_unit_conv_solver_ConvDirectNaiveConvWrw \
      test_unit_conv_solver_GemmFwd1x1_0_1_int8

# Run regression test (both NCDHW and NDHWC 3D INT8 int8_t-output cases)
./bin/test_test_naive_conv_3d_int8_fwd

# Run broader naive/int8 smoke suite for regression check
./bin/test_unit_conv_solver_ConvDirectNaiveConvFwd --gtest_filter="Smoke*"
./bin/test_unit_conv_solver_ConvDirectNaiveConvBwd --gtest_filter="Smoke*"
./bin/test_unit_conv_solver_ConvDirectNaiveConvWrw --gtest_filter="Smoke*"
./bin/test_unit_conv_solver_GemmFwd1x1_0_1_int8   --gtest_filter="Smoke*"

Test Result

All tests run on GPU MI300X (gfx942):

Test binary Filter Result
test_test_naive_conv_3d_int8_fwd all 2/2 PASSED (NCDHW 13.6s, NDHWC 13ms)
test_unit_conv_solver_ConvDirectNaiveConvFwd Smoke* 5/5 PASSED
test_unit_conv_solver_ConvDirectNaiveConvBwd Smoke* 4/4 PASSED
test_unit_conv_solver_ConvDirectNaiveConvWrw Smoke* 4/4 PASSED
test_unit_conv_solver_GemmFwd1x1_0_1_int8 Smoke* 2/2 PASSED

pre-commit (scoped to the two changed files, matching CI's --from-ref/--to-ref
behavior): all hooks pass (trailing-whitespace, end-of-file-fixer, clang-format v18.1.4,
gtest-format-miopen).

Risk Assessment

Risk level: Low

  • Change scope: two additive DEFINE_3D_NAIVE_CONV_KERNEL macro calls instantiating
    a template that is already used and proven for other 3D INT8 output variants. No
    changes to kernel logic, solvers, dispatch, APIs, or ABI.
  • Impact: only the 3D INT8 fwd naive reference/verification path, which was
    previously broken (crashes → works). All other paths are untouched.
  • Existing variants: the 2D nchw/nhwc int8_t-output variants and the 3D
    int32_t/float-output variants continue to pass their existing smoke tests.
  • Test coverage: new isolated regression test added; CI will catch any future
    regression on this symbol.

Checklist

  • I have added automated tests relevant to the introduced functionality
  • I have sufficient test coverage for the changes, and code coverage hasn't decreased as a result of my PR
  • I have ran the tests, and they are all passing locally
  • I have added relevant documentation for the changes
  • I have removed the stale documentation which is no longer relevant after this pull request
  • I have ran make format & make check_format to ensure the changes have been formatted

…d convolutions

Missing DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw/ndhwc, int8_t, int32_t, int8_t, 0)
macro calls caused "named symbol not found" HIP errors when MIOpenDriver tried to
verify 3D INT8 forward convolutions against the naive reference. The int8_t→int8_t
output variant existed for 2D (nchw/nhwc) but was never added for 3D layouts.

Fixes SILOTIGER-709 / ALMIOPEN-2275.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@therock-pr-bot

therock-pr-bot Bot commented Jul 28, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

  • python -m pip install pre-commit
  • pre-commit install
  • pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

… fwd regression

- Rename fixture alias GPU_UnitTestNaiveConv3dInt8Fwd -> GPU_UnitTestConvNaiveConv3dFwd_I8
  to satisfy the gtest-format-miopen pre-commit hook (suite name must end in a valid
  datatype token, e.g. _I8).
- Replace two identical NCDHW cases with one NCDHW + one explicit NDHWC case (via
  TensorDescriptorParams{miopenTensorNDHWC}) so both fixed kernel instantiations are
  actually exercised at runtime.
- Reduce tensor sizes (C=16, K=8) so num_add=16 < int8_max=127, satisfying the
  conv_tensor_gen.hpp overflow guard for int8_t accumulator test data generation.

All four tests pass on gfx942 (MI300X): 2/2 new 3D int8 cases + 4/4 smoke suites for
ConvDirectNaiveConvFwd/Bwd/Wrw and GemmFwd1x1_0_1_int8.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@therock-pr-bot

therock-pr-bot Bot commented Jul 28, 2026 •

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@amd-bartgips

amd-bartgips commented Jul 31, 2026 •

Copy link
Copy Markdown
Contributor Author

CI fails on the Dbsync test on GemmFwdRest solver checks
This is what Claude can tell me about this (i.e. it's unrelated to this PR, so it's something wrong in the develop branch)

Dbsync gfx942 failure is pre-existing gfx942 find-db drift, unrelated to this PR

The failing Dbsync gfx942 stage (StaticFDBSync, db_sync.cpp:816) is not caused by this change. All failures are GEMM forward solvers (GemmFwdRest, GemmFwd1x1_0_2) on FP32/BF16 2D NCHW keys, e.g. 768-2048-1024-3x3-256-2048-1024-1-1x1-1x1-1x1-0-NCHW-FP32-F.

Root cause: GemmFwdRest::IsApplicable returns false when the recorded workspace exceeds handle.GetMaxMemoryAllocSize() (src/solver/conv/gemm.cpp:903 — GetWorkspaceSize → 0). The failing key records a ~27 GB workspace. These entries were reintroduced into src/kernels/gfx942130.HIP.fdb.txt.bz2 by #8174 (gfx942 SystemDB refresh), which is already on develop — so this drift exists on develop independently of this PR. Dbsync gfx90a passes, confirming it is gfx942-specific.

This PR only adds two 3D INT8 naive-kernel instantiations + a gtest; it cannot affect GEMM applicability. The fix belongs in a gfx942 DB re-clean (MIOPEN_DBSYNC_CLEAN=1) on develop; this branch will be rebased once that lands.

@amd-bartgips
amd-bartgips marked this pull request as ready for review July 31, 2026 08:35
@amd-bartgips
amd-bartgips requested a review from a team as a code owner July 31, 2026 08:35
amd-bartgips and others added 2 commits August 3, 2026 02:52
…ectNaiveConvFwd

Add GetConvTestCases3dInt8() with NCDHW and NDHWC 3D int8 cases (type_y=miopenInt8)
and a Smoke3dInt8 instantiation for GPU_UnitTestConvSolverDirectNaiveFwd_I8. Remove
the now-redundant standalone test_naive_conv_3d_int8_fwd.cpp.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@amd-ahyttine
amd-ahyttine merged commit 2029f4f into develop Aug 28, 2026
65 checks passed
@amd-ahyttine
amd-ahyttine deleted the users/bartgips/naive_conv_3d_int8 branch August 28, 2026 10:49
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

JonathanLichtnerAMD added a commit that referenced this pull request Sep 2, 2026
## Summary

Backfills the **MIOpen 3.6.1 / ROCm 10.1.0** changelog section with
externally-visible changes that landed since the `release/therock-10.0`
branch but weren't yet documented. All 43 MIOpen commits in that range
were reviewed; CI/build-infra, test-only, docs, version-bump, and
already-documented changes were excluded.

21 entries added:

### Added
* [RNN] Fused single-layer LSTM forward-inference kernel, gated behind
`MIOPEN_DEBUG_RNN_FUSED_INFERENCE` (#8915)
* [Conv] Naive reference kernel enabling 3D INT8 forward convolution
(#10051)
* [Conv] `ConvHipConv` solver backed by vendored hipconv kernels (#9746)
* [Conv] float8 vector type support + deduplicated math kernels (#10130)
* [Conv] Validated gfx1151 and gfx1200 (Navi) SystemDBs (#10050)
* [Conv] `MIOPEN_DEBUG_DISABLE_SYSTEM_DB` /
`MIOPEN_DEBUG_DISABLE_USER_DB` env vars (#10728)
* [Conv] CK depthwise convolution on RDNA wave32 GPUs (#9373)

### Changed
* [Conv] Winograd Rage RxS f2x3 on gfx950 + gfx942 v4_6_1/v4_9_1 kernels
(#9014)
* [Conv] Updated 3D AI heuristics for gfx942 & gfx950 (#8954)

### Optimized
* [Conv] Single GEMM for point (1x1) output convolution, all directions,
2D/3D (#10912)
* [Conv] Naive backward-data kernel work group widened to 1024 (#10684)

### Resolved Issues
* [Conv] Batch-split miscalculation on large WrW tensors — ROCM-26011
(#11188)
* [Conv] Conv failure on XNACK-unsupported gfx9 APUs (gfx90c) (#11211)
* [BatchNorm] Device heap-overflow in fwd spatial `FinalMeanVariance` —
ALMIOPEN-2424 (#11093)
* [Conv] Deterministic `SetNextValue` polarity for grouped Bwd/Wrw CK
solvers (#10240)
* [LRN][LayerNorm] Grid alignment + OOB threads on gfx1250 (#10251)
* [Conv] Stale gfx950 grouped-wrw SystemDB entries (#11130)
* [Conv][Windows] System-DB path resolution from the MIOpen module
(#10838)
* [Conv] System-DB fallback so MI308X (gfx942) resolves a usable DB
(#8802)
* JSON perf logs (`MIOPEN_PERFORMANCE_LOGS`) dropping the last solver
during Find (#11166)
* Build failure against GCC 15 libstdc++ under C++20 (#9642)

## Notes

Deliberately excluded: internal/default-off changes (#10219, #10088,
#10649, #10231), a near-duplicate of an existing 10.0.0 entry (#10188),
and CI/build-infra/test-only/docs commits.

Docs-only change; no code impact.

🤖 Generated with [Claude Code](https://claude.com/claude-code)

---------

Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
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