Skip to content

feat(miopen): Add gfx120x support to Winograd Rage - #6439

Merged
k-artem merged 22 commits into
developfrom
users/scerzh/winograd-rage-add-gfx12
Jul 17, 2026
Merged

feat(miopen): Add gfx120x support to Winograd Rage#6439
k-artem merged 22 commits into
developfrom
users/scerzh/winograd-rage-add-gfx12

Conversation

@scerzh

@scerzh scerzh commented Apr 14, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

JIRA ID : ALMIOPEN-1853

Motivation

Previous Rage update #2429 was reverted by #6196 due to memory faults found on several models. Support for gfx12 was lost in Rage solver leading to a performance drop on gfx12.

Technical Details

The reverted PR was split into dedicated steps. This step brings gfx12 support back.

Test Plan

  • madengine Tests
  • Unit Tests
  • Speedup Check for StableDiffusion
  • CI

Test Result

  • madengine Tests
Index Model Topology Launcher Deployment GPU Arch Performance Metric Status Duration Data Name Data Provider
1 pyt_detectron2_retinanet_singlenode 1N×8G docker local gfx942 2.03 samples_per_second ✅ Success 19.4m coco2017 minio
2 pyt_detectron2_retinanet_singlenode 1N×8G docker local gfx942 2.42 samples_per_second ✅ Success 16.9m coco2017 minio
  • unittest miopen_gtest --gtest_filter=*UnitTestConvSolverWinoRage2x3*
[==========] Running 17 tests from 4 test suites.
[----------] Global test environment set-up.
[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/0
warning: xnack 'Off' was requested for a processor that does not support it!
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/0 (2483 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/1
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/1 (1 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/2
warning: xnack 'Off' was requested for a processor that does not support it!
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/2 (37 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/3
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/3 (1 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/4
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/4 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/5
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/5 (0 ms)
[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16 (2526 ms total)

[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/0 (10 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/1
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/1 (1 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/2
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/2 (18 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/3
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/3 (2 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/4
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/4 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/5
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/5 (0 ms)
[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16 (34 ms total)

[----------] 4 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/0 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/1
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/1 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/2
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/2 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/3
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/3 (1 ms)
[----------] 4 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16 (2 ms total)

[----------] 1 test from Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE
[ RUN      ] Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE.ConvWinoRageRxSf2x3/0 (67 ms)
[----------] 1 test from Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE (67 ms total)

[----------] Global test environment tear-down
[==========] 17 tests from 4 test suites ran. (2632 ms total)
[  PASSED  ] 17 tests.
  • Speedup check for StableDiffusion
    The average gain on certain shapes is around 44%
N C H W K R S Time (ms) develop→PR GFLOPs develop→PR Speedup
2 1152 32 32 768 1 1 0.206 → 0.201 18K → 18K 1.02x
2 1152 32 32 768 3 3 0.317 → 0.205 103K → 159K 1.54x
2 1536 32 32 1536 3 3 1.221 → 0.485 71K → 179K 2.52x
2 1536 32 32 768 1 1 0.281 → 0.275 17K → 18K 1.02x
2 1536 32 32 768 3 3 0.405 → 0.264 107K → 165K 1.53x
2 2304 32 32 768 1 1 0.402 → 0.362 18K → 20K 1.11x
2 2304 32 32 768 3 3 0.592 → 0.396 110K → 165K 1.50x
2 384 32 32 768 1 1 0.080 → 0.077 15K → 16K 1.04x
2 384 32 32 768 3 3 0.125 → 0.089 87K → 123K 1.41x
2 768 32 32 768 3 3 0.219 → 0.142 99K → 153K 1.54x
2 768 64 64 768 3 3 0.777 → 0.497 112K → 175K 1.56x
Average 1.44x

Submission Checklist

@scerzh
scerzh force-pushed the users/scerzh/winograd-rage-add-gfx12 branch from a29097b to 49d14fb Compare April 14, 2026 19:59
@scerzh
scerzh force-pushed the users/scerzh/winograd-rage-add-gfx12 branch from 3ca253a to 154b76b Compare April 15, 2026 06:42
@k-artem

k-artem commented Apr 15, 2026

Copy link
Copy Markdown
Contributor

@scerzh shapes from this comment (see any of attached log files) can be used for testing.

@scerzh
scerzh marked this pull request as ready for review April 23, 2026 11:55
@scerzh
scerzh requested review from a team as code owners April 23, 2026 11:55
@scerzh scerzh changed the title [miopen][winograd][gfx12] Add gfx12 support to Winograd Rage [miopen][winograd][gfx12][ALMIOPEN-1853] Add gfx12 support to Winograd Rage Apr 23, 2026
k-artem
k-artem approved these changes Apr 23, 2026
View reviewed changes

@k-artem k-artem 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

SreecharanGundaboluAMD
SreecharanGundaboluAMD requested changes May 27, 2026
View reviewed changes

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

Please review all the changes where you are doing device checks for "gfx12".

Comment thread projects/miopen/src/solver/conv/conv_wino_rage_RxS.cpp Resolved
Comment thread projects/miopen/src/solver/conv/conv_wino_rage_RxS.cpp Resolved
Comment thread projects/miopen/src/solver/conv/conv_wino_rage_RxS.cpp Resolved
Comment thread projects/miopen/src/solver/conv/conv_wino_rage_RxS.cpp Resolved
Comment thread projects/miopen/src/solver/conv/conv_wino_rage_RxS.cpp Resolved
scerzh and others added 2 commits June 30, 2026 11:11
@scerzh
Merge remote-tracking branch 'origin/develop' into users/scerzh/winog…
897cd81 
…rad-rage-add-gfx12
@scerzh @claude
[miopen][winograd][gfx12] Re-register TransposedConvWinoRageRxS and a…
79a71ca 
…dd FP16 NHWC tests

Re-register the conv::TransposedConvWinoRageRxS<2, 3> solver (NHWC
transposing wrapper around ConvWinoRageRxS) that was disabled when the
Winograd Rage integration was reverted. The transpose wrapper inherits
the forward solver's applicability, which is FP16-only on this branch,
so it cannot reach the fp32 kernels implicated in the revert.

Restore the FP16 NHWC unit tests for the transposed solver, including the
g=1 and degenerate spatial-dim regression cases for the GetGroupConvLayout
NHWC WrW fix.

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

therock-pr-bot Bot commented Jul 1, 2026 •
edited
Loading

Copy link
Copy Markdown

❌ PR Check — Action Required

Check Status Details
🌿 Branch Name ✅ Pass —
📝 PR Title/Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ❌ Fail Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/miopen/src/env_debug.cpp, projects/miopen/src/solver.cpp, projects/miopen/src/solver/conv/conv_wino_rage_RxS.cpp, projects/miopen/test/gtest/unit_conv_solver_ConvWinoRageRxS.cpp; no test file found
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

⚠️ 1 policy check(s) failed. Please address the issues above before this PR can be Reviewed.

🚫 Please fix the failed policies

  • ❌ Unit Test

The Not ready to Review label was added to this PR. Once all policies pass, the label is removed automatically.

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

@therock-pr-bot

therock-pr-bot Bot commented Jul 1, 2026 •
edited
Loading

Copy link
Copy Markdown

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

  • ❌ Unit Test

The Not ready to Review label has been added to this PR.
Once all policies pass, the label will be removed automatically.

@scerzh scerzh changed the title [miopen][winograd][gfx12][ALMIOPEN-1853] Add gfx12 support to Winograd Rage feat(miopen) Add gfx120x support to Winograd Rage Jul 1, 2026
@scerzh scerzh changed the title feat(miopen) Add gfx120x support to Winograd Rage feat(miopen): Add gfx120x support to Winograd Rage Jul 1, 2026
@k-artem

k-artem commented Jul 7, 2026

Copy link
Copy Markdown
Contributor

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

  • ❌ Unit Test

The Not ready to Review label has been added to this PR. Once all policies pass, the label will be removed automatically.

false positive, tests are added.

@k-artem

k-artem commented Jul 7, 2026

Copy link
Copy Markdown
Contributor

@JonathanLichtnerAMD could you please review and help to merge?

@k-artem k-artem mentioned this pull request Jul 8, 2026
Open
@samathew121

samathew121 commented Jul 8, 2026

Copy link
Copy Markdown
Collaborator

Hi All, The therock-pr-bot is not blocking now you from merging this ,currently this is only showing the issues and raising the Not ready to review flag.We are collecting the requirements across and will be providing enhancements in the upcoming releases.

@k-artem
k-artem enabled auto-merge (squash) July 8, 2026 11:49
JonathanLichtnerAMD
JonathanLichtnerAMD reviewed Jul 8, 2026
View reviewed changes

@JonathanLichtnerAMD JonathanLichtnerAMD left a comment •
edited
Loading

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.

We have some MIOpenDriver out-of-bounds access checks that can be used:

  • env variable: MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=2
  • MIOpenDriver option: --gpubuffer_check=2

Both are needed, and both check different buffers. (You can also run with these set to 1, but we've never seen an issue before the start of the buffer--it has always been an illegal access after the buffer.) Note that ASAN would not help with the Winograd kernels, since these are assembly kernels.

Can we run a wide set of shapes using the above env variable and flag just to make sure we don't hit any new issues?

@JonathanLichtnerAMD

JonathanLichtnerAMD commented Jul 8, 2026

Copy link
Copy Markdown
Contributor

I asked claude to spit out some suggested shapes to run:

Suggested OOB test coverage for the gfx12 Winograd Rage kernels

Since these kernels are hand-written GCN assembly, ASan cannot instrument them and the solver hard-gates HSA_XNACK=0 (so xnack replay isn't available through the normal path). Instead we can use
MIOpen's page-guard buffer check, which is xnack-independent:

  • --gpubuffer_check=2 / MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=2 — Back: real buffer is right-aligned flush to the end of a 2 MB-rounded allocation, so any overrun hits an unmapped page →
    fault. This is the primary detector for the Revert "[miopen][winograd][gfx12] Integrate several Wino Rage kernels… #6196-class overrun.
  • --gpubuffer_check=1 / ...=1 — Front: buffer left-aligned to the page start, catches underflow (negative halo / input-transform under-indexing).

Both pass -V 1 as an orthogonal check (guard-band catches cross-page faults; verification catches silent in-pad corruption).

Unit tests

The gtest fixture doesn't disable the buffer check, so it can run directly under guard-banding — this gives OOB coverage on the unit suite without needing xnack:

Overrun (primary)
MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=2 HSA_XNACK=0
./bin/miopen_gtest --gtest_filter='WinoRage'

Underflow
MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=1 HSA_XNACK=0
./bin/miopen_gtest --gtest_filter='WinoRage'

(Note: WinoRage, not WinoRage2x3 — the latter misses the new TransposedWinoRageRxS suites.)

Driver shape sweep

Common env (forces Rage in isolation so the guard-band is attributable to these kernels):

export HSA_XNACK=0
export MIOPEN_DEBUG_CONV_WINOGRAD=0 # disable other winograd solvers...
export MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3=1 # ...then re-enable ONLY Rage
export MIOPEN_FIND_MODE=NORMAL

Applicability constraints these shapes satisfy: fp16, stride=1, dilation=1, packed/non-casted, gfx942 or gfx120x. Two kernel bodies must both be covered: R≤3 && S≤3 → v4_6_1_gfx12; R>3 || S>3 →
v4_9_1_gfx12.

Guard-band note: the check faults only when an access crosses the 2 MB page boundary, so shapes with larger per-tensor byte sizes leave a smaller rounding pad and catch tighter overruns. The
group-A production shapes are sized well for this (>4 MB tensors); tiny shapes have large pads and hide small overruns.

Run each line twice — once with --gpubuffer_check 2 (shown), then again with --gpubuffer_check 1:

--- A. Production StableDiffusion shapes (the #6196 OOB class — highest priority) ---
R=S=3 → v4_6_1_gfx12 body
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 2304 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 384 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2

--- B. v4_9_1_gfx12 body (R>3 or S>3) — NOT reachable by any group-A shape ---
MIOpenDriver convfp16 -n 2 -c 512 -H 32 -W 32 -k 512 -y 5 -x 5 -p 2 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 256 -H 64 -W 64 -k 256 -y 7 -x 7 -p 3 -q 3 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 128 -H 32 -W 32 -k 128 -y 1 -x 5 -p 0 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2 # asym R≠S

--- C. Direction coverage (bwd-data sets F_REVERSE_R/S; wrw is a distinct path) ---
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 2 -V 1 --gpubuffer_check 2 # bwd-data
MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 4 -V 1 --gpubuffer_check 2 # wrw

--- D. Grouped conv (F_GROUPED_CONVOLUTION; CgG / KgG near the 2^16 guard) ---
MIOpenDriver convfp16 -n 2 -c 1024 -H 32 -W 32 -k 1024 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -g 8 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 1 -c 4088 -H 8 -W 8 -k 4088 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -g 16 -F 1 -V 1 --gpubuffer_check 2

--- E. Boundary / limit edges (dimsFit16bit, paddedSizesFit16bits, out_w < 2^16-3) ---
MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2 # H near 2^16
MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2 # out_w near -3 margin

--- F. Small-channel path (C<=16 → V4_6 chosen for perf; the region gtest already covers) ---
MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2

Groups A and B are the must-haves: A reproduces the production shapes behind the original revert (which the gtest C/K≤16 cases can't), and B is the only way to reach the second new kernel body
(v4_9_1_gfx12). C/D/E exercise the distinct flag paths (F_REVERSE, F_GROUPED_CONVOLUTION) and the 2^16/2^31 element-limit edges that are the arithmetic origin of OOB in these kernels.

@JonathanLichtnerAMD

JonathanLichtnerAMD commented Jul 8, 2026

Copy link
Copy Markdown
Contributor

And just because it might be missed in the above comment, we can also run the unit tests with MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS to provide extra coverage.

@k-artem

k-artem commented Jul 10, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

on gfx1201:

  1. UT with MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=2 - passed
root@RSB-RORLAB-44:/remote_therock/dist/rocm# MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=2 HSA_XNACK=0 ./bin/miopen_gtest --gtest_filter=*UnitTestConvSolverWinoRage2x3*
PRNG seed: 12345678
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
Note: Google Test filter = *UnitTestConvSolverWinoRage2x3*
[==========] Running 17 tests from 4 test suites.
[----------] Global test environment set-up.
[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/0 (72 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/1
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/1 (1 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/2
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/2 (12 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/3
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/3 (1 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/4
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/4 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/5
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/5 (0 ms)
[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16 (89 ms total)

[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/0 (9 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/1
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/1 (1 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/2
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/2 (16 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/3
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/3 (2 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/4
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/4 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/5
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/5 (0 ms)
[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16 (30 ms total)

[----------] 4 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/0 (3 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/1
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/1 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/2
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/2 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/3
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/3 (0 ms)
[----------] 4 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16 (5 ms total)

[----------] 1 test from Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE
[ RUN      ] Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE.ConvWinoRageRxSf2x3/0 (42 ms)
[----------] 1 test from Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE (42 ms total)

[----------] Global test environment tear-down
[==========] 17 tests from 4 test suites ran. (167 ms total)
[  PASSED  ] 17 tests.
  1. UT with MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=1 - passed
root@RSB-RORLAB-44:/remote_therock/dist/rocm# MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=1 HSA_XNACK=0 ./bin/miopen_gtest --gtest_filter=*UnitTestConvSolverWinoRage2x3*
PRNG seed: 12345678
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
Note: Google Test filter = *UnitTestConvSolverWinoRage2x3*
[==========] Running 17 tests from 4 test suites.
[----------] Global test environment set-up.
[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/0 (71 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/1
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/1 (1 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/2
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/2 (13 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/3
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/3 (1 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/4
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/4 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/5
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16.ConvWinoRageRxSf2x3/5 (0 ms)
[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16 (88 ms total)

[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/0 (9 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/1
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/1 (1 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/2
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/2 (16 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/3
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/3 (2 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/4
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/4 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/5
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16.ConvWinoRageRxSf2x3/5 (0 ms)
[----------] 6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16 (31 ms total)

[----------] 4 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/0 (3 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/1
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/1 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/2
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/2 (0 ms)
[ RUN      ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/3
[       OK ] Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16.ConvWinoRageRxSf2x3/3 (0 ms)
[----------] 4 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16 (5 ms total)

[----------] 1 test from Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE
[ RUN      ] Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE.ConvWinoRageRxSf2x3/0
[       OK ] Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE.ConvWinoRageRxSf2x3/0 (42 ms)
[----------] 1 test from Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE (42 ms total)

[----------] Global test environment tear-down
[==========] 17 tests from 4 test suites ran. (167 ms total)
[  PASSED  ] 17 tests.
  1. A - passed
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 2304 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 384 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Forward Convolution Verifies OK on GPU reference (0.00023697 < 0.0082)
MIOpenDriver convfp16 -n 2 -c 2304 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Forward Convolution Verifies OK on GPU reference (0.000309294 < 0.0082)
MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Forward Convolution Verifies OK on GPU reference (0.000226055 < 0.0082)
MIOpenDriver convfp16 -n 2 -c 384 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Forward Convolution Verifies OK on GPU reference (0.000223925 < 0.0082)
  1. B - passed
 MIOpenDriver convfp16 -n 2 -c 512 -H 32 -W 32 -k 512 -y 5 -x 5 -p 2 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 256 -H 64 -W 64 -k 256 -y 7 -x 7 -p 3 -q 3 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 128 -H 32 -W 32 -k 128 -y 1 -x 5 -p 0 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 512 -H 32 -W 32 -k 512 -y 5 -x 5 -p 2 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Forward Convolution Verifies OK on GPU reference (0.000249951 < 0.0082)
MIOpenDriver convfp16 -n 2 -c 256 -H 64 -W 64 -k 256 -y 7 -x 7 -p 3 -q 3 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Forward Convolution Verifies OK on GPU reference (0.000252114 < 0.0082)
MIOpenDriver convfp16 -n 2 -c 128 -H 32 -W 32 -k 128 -y 1 -x 5 -p 0 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Forward Convolution Verifies OK on GPU reference (0.000221911 < 0.0082)
  1. C - passed
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 2 -V 1 --gpubuffer_check 2 && MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 4 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 2 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Backward Convolution Data Verifies OK on GPU reference (0.000235168 < 0.0082)
MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 4 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Backward Convolution Weights Verifies OK on GPU reference (0.00 < 0.02)
  1. D - passed
root@RSB-RORLAB-44:/remote_therock/dist/rocm/bin# MIOpenDriver convfp16 -n 2 -c 1024 -H 32 -W 32 -k 1024 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -g 8 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 1 -c 4088 -H 8 -W 8 -k 4088 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -g 16 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 1024 -H 32 -W 32 -k 1024 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -g 8 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
Forward Convolution Verifies OK on GPU reference (0.000263216 < 0.0082)
MIOpenDriver convfp16 -n 1 -c 4088 -H 8 -W 8 -k 4088 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -g 16 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
Error: RSB-RORLAB-44:/root/workspace/gfx1201/TheRock/rocm-libraries/projects/miopen/driver/conv_driver.hpp:1230: Invalid group number
  1. E - failed, but it only with --gpubuffer_check 2 and for MIOpenIm2d2Col kernel, so not related with this PR, passed with disabling GEMM MIOPEN_DEBUG_CONV_GEMM=0 and Wino kernels successfully executed.
root@RSB-RORLAB-44:/remote_therock/dist/rocm/bin# export MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=2
root@RSB-RORLAB-44:/remote_therock/dist/rocm/bin# MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
MIOpen(HIP): Command [LogCmdFindConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
Warning: Queue error - HSA_STATUS_ERROR_MEMORY_FAULT
VGPU=0x5fb586b0e3c0 SWq=0x7073ee252000, HWq=0x7073d4600000, id=1
        Dispatch Header =0xd02 (type=2, barrier=1, acquire=2, release=1), setup=3
        grid=[1048064, 1, 1], workgroup=[128, 1, 1]
        private_seg_size=0, group_seg_size=13312
        kernel_obj=0x706fa6d18380, kernarg_address=0x0x706fc7600100
        completion_signal=0x0
        rptr=3, wptr=5
:0:rocdevice.cpp            :3781: 1303743360194 us:  Memory Fault Error [host: RSB-RORLAB-44, GPU index: 0, faulting addr: 0x706fb5200000, kernel: Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT64x64x32_MI16x16x1_SN_LDSB0_AFC1_AG0_AGGSUA0_AGNTAB0_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR1_CLS0_CADS0_DTLA0_DTLB0_DTLM0_DTVA1_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_HPLR0_ISA1201_ICIW0_IU1_K1_LDSTI0_LBSPPA0_LBSPPB128_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA0_LPB16_LPMXSA0_LPMXSB0_LPM0_LRVW8_LWPMn1_MIAV1_MIWT2_2_MXLIBL_MXSFNS_MO40_MGRIPM1_NTn1_NTA0_NTB0_NTC0_NTD0_NTE0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS0_NLCA2_NLCB1_ONLL0_PAP0_PGL0_PGR2_PLR1_PKA0_SGROB0_SIA3_SS0_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKFDPO0_SKXCCM0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THMXSA0_THMXSB0_THM0_THWS0_TLDS1_TLDSMn1_ULSGRO0_USL1_USLMX0_UDFMAC0_UIOFGRO0_UPLRP0_USFGROn1_USI0_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS32_WG32_4_1_WGMXCC1]
Memory access fault by GPU node-1 (Agent handle: 0x5fb586a4ebd0) on address 0x706fb5200000. Reason: Page not present or supervisor privilege.
GPU coredump: Pipe handler '/usr/share/apport/apport' not found or not executable.
Falling back to file-based dump. Set HSA_COREDUMP_PATTERN to override.
Failed to create GPU coredump: Permission denied
GPU core dump failed
Aborted (core dumped)
root@RSB-RORLAB-44:/remote_therock/dist/rocm/bin# MIOPEN_DEBUG_CONV_GEMM=0 MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
MIOpen(HIP): Command [LogCmdFindConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
Forward Convolution Verifies OK on GPU reference (0.000194166 < 0.0082)
root@RSB-RORLAB-44:/remote_therock/dist/rocm/bin# MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
MIOpen(HIP): Command [LogCmdFindConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
Warning: Queue error - HSA_STATUS_ERROR_MEMORY_FAULT
VGPU=0x614c1abc93c0 SWq=0x7b44dd844000, HWq=0x7b44bde00000, id=1
        Dispatch Header =0xd02 (type=2, barrier=1, acquire=2, release=1), setup=3
        grid=[1048064, 1, 1], workgroup=[128, 1, 1]
        private_seg_size=0, group_seg_size=13312
        kernel_obj=0x7b4090b18380, kernarg_address=0x0x7b44bd813f00
        completion_signal=0x0
        rptr=639, wptr=641
:0:rocdevice.cpp            :3781: 1303776494358 us:  Memory Fault Error [host: RSB-RORLAB-44, GPU index: 0, faulting addr: 0x7b40a4c00000, kernel: Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT64x64x32_MI16x16x1_SN_LDSB0_AFC1_AG0_AGGSUA0_AGNTAB0_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR1_CLS0_CADS0_DTLA0_DTLB0_DTLM0_DTVA1_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_HPLR0_ISA1201_ICIW0_IU1_K1_LDSTI0_LBSPPA0_LBSPPB128_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA0_LPB16_LPMXSA0_LPMXSB0_LPM0_LRVW8_LWPMn1_MIAV1_MIWT2_2_MXLIBL_MXSFNS_MO40_MGRIPM1_NTn1_NTA0_NTB0_NTC0_NTD0_NTE0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS0_NLCA2_NLCB1_ONLL0_PAP0_PGL0_PGR2_PLR1_PKA0_SGROB0_SIA3_SS0_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKFDPO0_SKXCCM0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THMXSA0_THMXSB0_THM0_THWS0_TLDS1_TLDSMn1_ULSGRO0_USL1_USLMX0_UDFMAC0_UIOFGRO0_UPLRP0_USFGROn1_USI0_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS32_WG32_4_1_WGMXCC1]
Memory access fault by GPU node-1 (Agent handle: 0x614c1ab09bd0) on address 0x7b40a4c00000. Reason: Page not present or supervisor privilege.
GPU coredump: Pipe handler '/usr/share/apport/apport' not found or not executable.
Falling back to file-based dump. Set HSA_COREDUMP_PATTERN to override.
Failed to create GPU coredump: Permission denied
GPU core dump failed
Aborted (core dumped)
root@RSB-RORLAB-44:/remote_therock/dist/rocm/bin# MIOPEN_DEBUG_CONV_GEMM=0 MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
MIOpen(HIP): Command [LogCmdFindConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
Forward Convolution Verifies OK on GPU reference (0.000199525 < 0.0082)
  1. F - passed
root@RSB-RORLAB-44:/remote_therock/dist/rocm/bin# MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
MIOpen(HIP): Warning [OpenRuntimeLibraryForDevice] CK grouped conv library not found for device gfx1201: libMIOpenCKGroupedConv_gfx1201.so: cannot open shared object file: No such file or directory
PRNG seed: 12345678
MIOpen(HIP): Command [LogCmdFindConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
MIOpen(HIP): Command [LogCmdConvolution] ./bin/MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1
Forward Convolution Verifies OK on GPU reference (0.00 < 0.01)

CC @JonathanLichtnerAMD @scerzh

JonathanLichtnerAMD
JonathanLichtnerAMD approved these changes Jul 14, 2026
View reviewed changes

@JonathanLichtnerAMD JonathanLichtnerAMD 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. Thanks for doing the extra testing. I've created a jira issue for the unrelated failure

@k-artem
k-artem merged commit 29f6487 into develop Jul 17, 2026
70 of 73 checks passed
@k-artem
k-artem deleted the users/scerzh/winograd-rage-add-gfx12 branch July 17, 2026 14:16
@scerzh

scerzh commented Jul 17, 2026

Copy link
Copy Markdown
Contributor Author

on gfx942:

Same OOB / guard-band test plan as the gfx1201 run above, executed on gfx942 against this build's miopen_gtest and MIOpenDriver. Winograd Rage forced in isolation:

export HSA_XNACK=0
export MIOPEN_DEBUG_CONV_WINOGRAD=0                # disable other winograd solvers...
export MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3=1   # ...then re-enable ONLY Rage
export MIOPEN_FIND_MODE=NORMAL

Every driver line was run twice — once with --gpubuffer_check 2 (overrun / right-aligned guard page) and once with --gpubuffer_check 1 (underflow / left-aligned). No GPU memory faults were observed in any case.

  1. UT with MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=2 - passed
MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=2 HSA_XNACK=0 ./bin/miopen_gtest --gtest_filter='*WinoRage*'
PRNG seed: 12345678
[==========] Running 189 tests from 12 test suites.
...
[----------] 115 tests from Full/GPU_ConvBiasActivFind2Infer_FP16
[----------]   6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Fwd_FP16
[----------]   6 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Bwd_FP16
[----------]   4 tests from Smoke/GPU_UnitTestConvSolverWinoRage2x3Wrw_FP16
[----------]   2 tests from Smoke/GPU_UnitTestConvSolverTransposedWinoRageRxSFwd_FP16
[----------]   2 tests from Smoke/GPU_UnitTestConvSolverTransposedWinoRageRxSBwd_FP16
[----------]   3 tests from Smoke/GPU_UnitTestConvSolverTransposedWinoRageRxSWrw_FP16
[----------]   1 test  from Smoke/CPU_UnitTestConvSolverWinoRage2x3DevApplicabilityFwd_NONE
[----------]   1 test  from Smoke/CPU_UnitTestConvSolverTransposedWinoRageRxSDevApplicabilityFwd_NONE
...
[==========] 189 tests from 12 test suites ran. (39307 ms total)
[  PASSED  ] 122 tests.

(Using *WinoRage* rather than *WinoRage2x3* so the new TransposedWinoRageRxS Fwd/Bwd/Wrw suites are included. The remaining tests of the 189 are GTEST_SKIPped by arch/dtype gating — 0 failures.)

  1. UT with MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=1 - passed
MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS=1 HSA_XNACK=0 ./bin/miopen_gtest --gtest_filter='*WinoRage*'
[==========] Running 189 tests from 12 test suites.
...
[==========] 189 tests from 12 test suites ran. (28076 ms total)
[  PASSED  ] 122 tests.
  1. A. Production StableDiffusion shapes (R=S=3 → v4_6_1_gfx12 body) - passed
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
PRNG seed: 12345678
Forward Convolution Verifies OK on GPU reference (0.000424229 < 0.082)
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (0.000424106 < 0.082)
MIOpenDriver convfp16 -n 2 -c 2304 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (0.000503618 < 0.082)
MIOpenDriver convfp16 -n 2 -c 2304 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (0.000503345 < 0.082)
MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (1.08852e-05 < 0.082)
MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (1.08852e-05 < 0.082)
MIOpenDriver convfp16 -n 2 -c 384 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (0.000358348 < 0.082)
MIOpenDriver convfp16 -n 2 -c 384 -H 32 -W 32 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (0.000358539 < 0.082)
  1. B. v4_9_1_gfx12 body (R>3 or S>3) - passed
MIOpenDriver convfp16 -n 2 -c 512 -H 32 -W 32 -k 512 -y 5 -x 5 -p 2 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (0.00050697 < 0.082)
MIOpenDriver convfp16 -n 2 -c 512 -H 32 -W 32 -k 512 -y 5 -x 5 -p 2 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (0.00050726 < 0.082)
MIOpenDriver convfp16 -n 2 -c 256 -H 64 -W 64 -k 256 -y 7 -x 7 -p 3 -q 3 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (0.00042998 < 0.082)
MIOpenDriver convfp16 -n 2 -c 256 -H 64 -W 64 -k 256 -y 7 -x 7 -p 3 -q 3 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (0.000430241 < 0.082)
MIOpenDriver convfp16 -n 2 -c 128 -H 32 -W 32 -k 128 -y 1 -x 5 -p 0 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2   # asym R≠S
Forward Convolution Verifies OK on GPU reference (0.000327885 < 0.082)
MIOpenDriver convfp16 -n 2 -c 128 -H 32 -W 32 -k 128 -y 1 -x 5 -p 0 -q 2 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1   # asym R≠S
Forward Convolution Verifies OK on GPU reference (0.000327906 < 0.082)
  1. C. Direction coverage (bwd-data -F 2, wrw -F 4) - passed
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 2 -V 1 --gpubuffer_check 2
Backward Convolution Data Verifies OK on GPU reference (0.000365585 < 0.082)
MIOpenDriver convfp16 -n 2 -c 1536 -H 32 -W 32 -k 1536 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 2 -V 1 --gpubuffer_check 1
Backward Convolution Data Verifies OK on GPU reference (0.000365593 < 0.082)
MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 4 -V 1 --gpubuffer_check 2
Backward Convolution Weights Verifies OK on GPU reference (0.000201711 < 0.0164)
MIOpenDriver convfp16 -n 2 -c 768 -H 64 -W 64 -k 768 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 4 -V 1 --gpubuffer_check 1
Backward Convolution Weights Verifies OK on GPU reference (0.000201711 < 0.0164)
  1. D. Grouped conv - passed (g=8); the g=16 case is rejected by driver argument validation, not a kernel issue
MIOpenDriver convfp16 -n 2 -c 1024 -H 32 -W 32 -k 1024 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -g 8 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (7.13474e-06 < 0.082)
MIOpenDriver convfp16 -n 2 -c 1024 -H 32 -W 32 -k 1024 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -g 8 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (7.13474e-06 < 0.082)
MIOpenDriver convfp16 -n 1 -c 4088 -H 8 -W 8 -k 4088 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -g 16 -F 1 -V 1 --gpubuffer_check 2
Error: conv_driver.hpp:1230: Invalid group number

(c=4088 is not divisible by g=16, so the driver rejects the shape before any kernel runs — same behavior noted on gfx1201.)

  1. E. Boundary / limit edges (H/W near 2¹⁶) - passed
MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (3.31279e-06 < 0.082)
MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (3.31279e-06 < 0.082)
MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (0.000240668 < 0.082)
MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (0.000240668 < 0.082)

Note: unlike gfx1201, these pass on gfx942 with --gpubuffer_check 2 without needing MIOPEN_DEBUG_CONV_GEMM=0 — the im2col/GEMM memory fault seen on gfx1201 does not reproduce here. Re-running with MIOPEN_DEBUG_CONV_GEMM=0 (Rage in isolation) also passes:

MIOPEN_DEBUG_CONV_GEMM=0 MIOpenDriver convfp16 -n 1 -c 16 -H 65500 -W 8 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (3.31279e-06 < 0.082)
MIOPEN_DEBUG_CONV_GEMM=0 MIOpenDriver convfp16 -n 1 -c 16 -H 8 -W 65500 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (0.000240668 < 0.082)
  1. F. Small-channel path (C≤16 → v4_6 chosen for perf) - passed
MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 2
Forward Convolution Verifies OK on GPU reference (2.32193e-06 < 0.082)
MIOpenDriver convfp16 -n 2 -c 16 -H 32 -W 32 -k 16 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -F 1 -V 1 --gpubuffer_check 1
Forward Convolution Verifies OK on GPU reference (2.32193e-06 < 0.082)

Summary: all groups pass under both page-guard modes with no GPU memory faults; the unit suite (including the new transposed Rage kernels) is green (122 passed / 0 failed); correctness verified against the GPU reference for fwd / bwd-data / wrw. The #6196-class overrun does not reproduce on gfx942.

@JonathanLichtnerAMD JonathanLichtnerAMD mentioned this pull request Jul 31, 2026
Merged
JonathanLichtnerAMD added a commit that referenced this pull request Aug 13, 2026
@JonathanLichtnerAMD @claude
docs(miopen): add ROCm 10.0.0 changelog section (#10208)
baf3e19 
JIRA ID: ALMIOPEN-2364

## Motivation

ROCm 7.15 has been renamed to ROCm 10.0, and MIOpen 3.6.0 is the
corresponding
release. `projects/miopen/CHANGELOG.md` needs a new version section
documenting
what changed for downstream users, release-notes tooling, and packagers
between
MIOpen 3.5.2 (ROCm 7.14.0) and MIOpen 3.6.0 (ROCm 10.0.0), consistent
with the
changelog's existing per-release format.

## Technical Details

- Inserted a new `## MIOpen 3.6.0 for ROCm 10.0.0` section at the top of
  `projects/miopen/CHANGELOG.md`, directly above the existing
`## MIOpen 3.5.2 for ROCm 7.14.0` section, using the same `### Added` /
`### Changed` / `### Resolved Issues` / `### Removed` subheading
structure as
  every prior release section in the file.
- Entries were compiled by diffing `origin/release/therock-10.0` against
  `origin/release/therock-7.14` for `projects/miopen` and hand-selecting
changelog-worthy commits, excluding: test-only commits, CI-only commits,
build/version-bump commits, doc-only commits, reverted commit/revert
pairs, and
  commits already surfaced in the 3.5.2/7.14.0 section via cherry-pick.
- Resulting entries added:
- **Added:** gfx950 (MI350X/MI355X) 7x7 depthwise forward and
backward-data
convolution support, fp16/bf16 (#9168) — fixes a slow
fallback-to-naive-kernel
    regression in ConvNeXt-style depthwise convolutions.
  - **Changed:**
- Restored gfx12x support in the Winograd Rage solver (#6439),
recovering
      performance lost when earlier gfx12 support was reverted.
- Refreshed the gfx1100, gfx1102, and gfx1201 (Navi) SystemDBs (#9529)
with
      updated tuned find/perf-database entries.
- Refreshed the gfx950 SystemDB (#9111) with additional tuned entries.
  - **Resolved Issues:**
- Fixed an off-by-stride indexing bug in BatchNorm backward `CalcStats`
      mean/variance remainder loop causing a ~1% systematic bias in NCHW
      batchnorm backward results (#9031).
- Fixed a tensor-op integer overflow on gfx1250 for large allocations
that
      could cause memory access faults (#9049).
- Fixed naive convolution solvers silently failing to launch (and
leaving a
stale HIP error visible after Find returned success) when a global work
size reached 2^32 or more work-items; such launches are now rejected up
      front (#10188).
- **Removed:** the OpenCL (OCL) backend (#8433); MIOpen now supports the
HIP
    backend only.
- No code, build, or test files were touched — this is a pure
documentation
  addition; the underlying functional changes were already merged to
`release/therock-10.0` in their own PRs (referenced above) and are only
being
  documented here.

## Risk Assessment

Single file (`projects/miopen/CHANGELOG.md`), 17 lines added, 0 removed,
no code
or build-system changes.

**Risk Level: 🟢 Low**

### Impacted Components

- `projects/miopen/CHANGELOG.md` only — no source, kernel, solver,
build, or test
  files are touched.

### Potential Side Effects

- None functionally. The only observable effect is the new changelog
text itself;
  no build artifacts, APIs, or runtime behavior change.
- Cosmetic risk only: if any release-notes/packaging tooling parses this
file
programmatically, it should be checked for compatibility with the new
section,
  though the added section strictly follows the existing heading/bullet
  conventions used by all prior sections in the file.

### Mitigation Steps

- Followed the exact existing section format (`##` version header, `###`
category subheadings, `* ` bullet entries, `[Component]` prefix tags)
used by
the adjacent 3.5.2/7.14.0 and 3.5.1/7.13.0 sections, minimizing
formatting risk.
- Each entry was cross-referenced against a specific merged PR number
  (#9168, #6439, #9529, #9111, #9031, #9049, #10188, #8433) rather than
paraphrased from memory, so reviewers can verify each line against its
source
  commit.

## Test Plan

No functional tests apply — this is a documentation-only change to a
Markdown
changelog file. Verification consisted of manually reviewing the diff
against the
existing file's structure and confirming the new section's heading
hierarchy and
bullet formatting match every other version section in `CHANGELOG.md`.

## Test Result

- `git diff origin/develop...HEAD --stat` confirms the change is scoped
to exactly
  one file (`projects/miopen/CHANGELOG.md`, +17/-0).
- Visual/manual review confirms the new `## MIOpen 3.6.0 for ROCm
10.0.0` section
renders as valid Markdown consistent with surrounding sections. No build
or CI
  run is applicable to a docs-only change.

---------

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

Reviewers

@k-artem k-artem k-artem approved these changes
@iq136boy iq136boy iq136boy approved these changes
@AnaghaRaoAMD AnaghaRaoAMD AnaghaRaoAMD approved these changes
@JonathanLichtnerAMD JonathanLichtnerAMD JonathanLichtnerAMD approved these changes
@SreecharanGundaboluAMD SreecharanGundaboluAMD SreecharanGundaboluAMD approved these changes
@mattwill-amd mattwill-amd mattwill-amd approved these changes

Assignees

No one assigned

Labels

documentation Not ready to Review organization: ROCm project: miopen

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

8 participants

@scerzh @k-artem @samathew121 @JonathanLichtnerAMD @iq136boy @AnaghaRaoAMD @SreecharanGundaboluAMD @mattwill-amd