Repository navigation
[CK] Disable test cases affected by compiler codegen bugs on gfx90a - #6343
Merged
Merged
Conversation
aosewski
approved these changes
Apr 10, 2026
aosewski
enabled auto-merge (squash)
April 10, 2026 13:51
The AMDGPU compiler generates incorrect code for fp16 d256 dropout kernels in batch mode on gfx90a. The compiler aliases Philox RNG VGPR parameters under high register pressure, causing corrupted dropout randval output (lanes 32-63) and NaN in output/LSE tensors. bf16 dropout is unaffected (different type conversion codegen path) and covers the same code paths, so disabling the fp16 variant does not reduce test coverage. The skip is version-gated and will be removed once the compiler fix is available.
poyenc
changed the title
3 tasks done
yraparti
reviewed
Apr 10, 2026
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Apr 13, 2026
[CK] Disable compilation of problematic bwd weight conv instances for gfx90a (#6343) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Motivation Due to compiler version update, there are test failures in the test suite `test_grouped_convnd_bwd_weight` when running on `gfx90a`. There are four failing tests for FP16/BF16 that arise from a single kernel instance. As the problem is in the current `develop` branch, the test failures are blocking any PR merges into `develop`. An example of a failed CI runs is here: [http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/develop/558/pipeline/](http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/develop/558/pipeline/). The underlying compiler problem is potentially the same as described in #6342 as tests are passing for clang compiler version 20.0 and failing for clang compiler version 22.0. ## Technical Details This PR disables the compilation of the problematic bwd weight conv instance for `gfx90a` by adding a new CMake flag `CK_USE_GFX90A` that allows us to detect when we are compiling for `gfx90a`. Using the new CMake flag, compilation of instance `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3<64, 128, 32, 32, Default, 8, 4, 1, 8, 8, 8, 8, 1, 1, 2>` is disabled for `gfx90a`. Co-authored-by: Ville Pietilä <>
brockhargreaves-amd
added a commit
that referenced
this pull request
Apr 14, 2026
…degen bugs on gfx90a" (#6400) Reverts #6343 This is causing failures in miopen, namely Dbsync gfx942 even though it shouldn't be affected so this needs to be investigated. Please add miopen as a label to the new PR for addressing the compiler codegen bug so that this can be addressed simultaneously.
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Apr 14, 2026
[MIOPEN] [CK] Revert "[CK] Disable test cases affected by compiler codegen bugs on gfx90a" (#6400) Reverts ROCm/rocm-libraries#6343 This is causing failures in miopen, namely Dbsync gfx942 even though it shouldn't be affected so this needs to be investigated. Please add miopen as a label to the new PR for addressing the compiler codegen bug so that this can be addressed simultaneously.
hyoon1
pushed a commit
to hyoon1/composable_kernel
that referenced
this pull request
Apr 17, 2026
…degen bugs on gfx90a" (#6400) Reverts ROCm/rocm-libraries#6343 This is causing failures in miopen, namely Dbsync gfx942 even though it shouldn't be affected so this needs to be investigated. Please add miopen as a label to the new PR for addressing the compiler codegen bug so that this can be addressed simultaneously.
aledudek
pushed a commit
that referenced
this pull request
May 20, 2026
… gfx90a (#6343) ## Motivation Due to compiler version update, there are test failures in the test suite `test_grouped_convnd_bwd_weight` when running on `gfx90a`. There are four failing tests for FP16/BF16 that arise from a single kernel instance. As the problem is in the current `develop` branch, the test failures are blocking any PR merges into `develop`. An example of a failed CI runs is here: [http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/develop/558/pipeline/](http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/develop/558/pipeline/). The underlying compiler problem is potentially the same as described in #6342 as tests are passing for clang compiler version 20.0 and failing for clang compiler version 22.0. ## Technical Details This PR disables the compilation of the problematic bwd weight conv instance for `gfx90a` by adding a new CMake flag `CK_USE_GFX90A` that allows us to detect when we are compiling for `gfx90a`. Using the new CMake flag, compilation of instance `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3<64, 128, 32, 32, Default, 8, 4, 1, 8, 8, 8, 8, 1, 1, 2>` is disabled for `gfx90a`. Co-authored-by: Ville Pietilä <>
aledudek
pushed a commit
that referenced
this pull request
May 20, 2026
…degen bugs on gfx90a" (#6400) Reverts #6343 This is causing failures in miopen, namely Dbsync gfx942 even though it shouldn't be affected so this needs to be investigated. Please add miopen as a label to the new PR for addressing the compiler codegen bug so that this can be addressed simultaneously.
shumway
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
May 27, 2026
[CK] Disable compilation of problematic bwd weight conv instances for gfx90a (#6343) ## Motivation Due to compiler version update, there are test failures in the test suite `test_grouped_convnd_bwd_weight` when running on `gfx90a`. There are four failing tests for FP16/BF16 that arise from a single kernel instance. As the problem is in the current `develop` branch, the test failures are blocking any PR merges into `develop`. An example of a failed CI runs is here: [http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/develop/558/pipeline/](http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/develop/558/pipeline/). The underlying compiler problem is potentially the same as described in #6342 as tests are passing for clang compiler version 20.0 and failing for clang compiler version 22.0. ## Technical Details This PR disables the compilation of the problematic bwd weight conv instance for `gfx90a` by adding a new CMake flag `CK_USE_GFX90A` that allows us to detect when we are compiling for `gfx90a`. Using the new CMake flag, compilation of instance `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3<64, 128, 32, 32, Default, 8, 4, 1, 8, 8, 8, 8, 1, 1, 2>` is disabled for `gfx90a`. Co-authored-by: Ville Pietilä <>
shumway
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
May 27, 2026
[MIOPEN] [CK] Revert "[CK] Disable test cases affected by compiler codegen bugs on gfx90a" (#6400) Reverts ROCm/rocm-libraries#6343 This is causing failures in miopen, namely Dbsync gfx942 even though it shouldn't be affected so this needs to be investigated. Please add miopen as a label to the new PR for addressing the compiler codegen bug so that this can be addressed simultaneously.
COrruDXC
pushed a commit
to COrruDXC/rocm-libraries
that referenced
this pull request
Jun 4, 2026
… gfx90a (ROCm#6343) ## Motivation Due to compiler version update, there are test failures in the test suite `test_grouped_convnd_bwd_weight` when running on `gfx90a`. There are four failing tests for FP16/BF16 that arise from a single kernel instance. As the problem is in the current `develop` branch, the test failures are blocking any PR merges into `develop`. An example of a failed CI runs is here: [http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/develop/558/pipeline/](http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/develop/558/pipeline/). The underlying compiler problem is potentially the same as described in ROCm#6342 as tests are passing for clang compiler version 20.0 and failing for clang compiler version 22.0. ## Technical Details This PR disables the compilation of the problematic bwd weight conv instance for `gfx90a` by adding a new CMake flag `CK_USE_GFX90A` that allows us to detect when we are compiling for `gfx90a`. Using the new CMake flag, compilation of instance `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3<64, 128, 32, 32, Default, 8, 4, 1, 8, 8, 8, 8, 1, 1, 2>` is disabled for `gfx90a`. Co-authored-by: Ville Pietilä <>
COrruDXC
pushed a commit
to COrruDXC/rocm-libraries
that referenced
this pull request
Jun 4, 2026
…degen bugs on gfx90a" (ROCm#6400) Reverts ROCm#6343 This is causing failures in miopen, namely Dbsync gfx942 even though it shouldn't be affected so this needs to be investigated. Please add miopen as a label to the new PR for addressing the compiler codegen bug so that this can be addressed simultaneously.
This was referenced Aug 15, 2026
Closed
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Aug 19, 2026
Adds an automatic per-PR `dbsync-rocjitsu` job to therock-ci.yml (matrixed over gfx942 + gfx950) that fetches this run's own TheRock source-CK build artifact and runs the MIOpen dbsync check (StaticFDBSync) under the rocjitsu KMD interposer on a GPU-less runner. dbsync is CPU-only (SKIP_KDB_TESTING=1, no kernel launch) except CK grouped-conv IsApplicable, which the arch-matched rocjitsu config services -- so this validates each arch's SystemDB against the PR's own CK without real hardware. This catches a PR's own CK change (the #6343 scenario), unlike the frozen-container CK path, and doesn't depend on scarce GPU test runners. Key wiring: - Gated on needs.setup.outputs.miopen_selected (computed in the setup job by parsing linux_projects' per-bundle comma-separated projects_to_test), so it runs only when TheRock's MIOpen CI would. - needs: setup only. The fetch step retry-polls (60s interval, 45 min budget) until this run's own MIOpen build has uploaded miopen_gtest, rather than depending on the whole therock-ci-linux matrix (which would also block on unrelated components' GPU test runs). The poll checks for miopen_gtest explicitly, not just the fetch script's exit code (which reports success as soon as any artifact -- e.g. another bundle's shared generic pieces -- lands). Also renames the standalone workflow_dispatch diagnostic to therock-dbsync-rocjitsu.yml (drop the "prototype" naming); it stays as a manual companion that takes an explicit artifact_run_id for pointing at arbitrary/pinned builds the automatic gate can't self-reference. ALMIOPEN-1779 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Aug 23, 2026
Adds an automatic per-PR `dbsync-rocjitsu` job to therock-ci.yml (matrixed over gfx942 + gfx950) that fetches this run's own TheRock source-CK build artifact and runs the MIOpen dbsync check (StaticFDBSync) under the rocjitsu KMD interposer on a GPU-less runner. dbsync is CPU-only (SKIP_KDB_TESTING=1, no kernel launch) except CK grouped-conv IsApplicable, which the arch-matched rocjitsu config services -- so this validates each arch's SystemDB against the PR's own CK without real hardware. This catches a PR's own CK change (the #6343 scenario), unlike the frozen-container CK path, and doesn't depend on scarce GPU test runners. Key wiring: - Gated on needs.setup.outputs.miopen_selected (computed in the setup job by parsing linux_projects' per-bundle comma-separated projects_to_test), so it runs only when TheRock's MIOpen CI would. - needs: setup only. The fetch step retry-polls (60s interval, 45 min budget) until this run's own MIOpen build has uploaded miopen_gtest, rather than depending on the whole therock-ci-linux matrix (which would also block on unrelated components' GPU test runs). The poll checks for miopen_gtest explicitly, not just the fetch script's exit code (which reports success as soon as any artifact -- e.g. another bundle's shared generic pieces -- lands). Also renames the standalone workflow_dispatch diagnostic to therock-dbsync-rocjitsu.yml (drop the "prototype" naming); it stays as a manual companion that takes an explicit artifact_run_id for pointing at arbitrary/pinned builds the automatic gate can't self-reference. ALMIOPEN-1779 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Aug 23, 2026
Adds an automatic per-PR `dbsync-rocjitsu` job to therock-ci.yml (matrixed over gfx942 + gfx950) that fetches this run's own TheRock source-CK build artifact and runs the MIOpen dbsync check (StaticFDBSync) under the rocjitsu KMD interposer on a GPU-less runner. dbsync is CPU-only (SKIP_KDB_TESTING=1, no kernel launch) except CK grouped-conv IsApplicable, which the arch-matched rocjitsu config services -- so this validates each arch's SystemDB against the PR's own CK without real hardware. This catches a PR's own CK change (the #6343 scenario), unlike the frozen-container CK path, and doesn't depend on scarce GPU test runners. Key wiring: - Gated on needs.setup.outputs.miopen_selected (computed in the setup job by parsing linux_projects' per-bundle comma-separated projects_to_test), so it runs only when TheRock's MIOpen CI would. - needs: setup only. The fetch step is build-aware: it polls therock_dbsync_build_leg_state.py for this run's MIOpen build-leg status, waiting while queued/running, failing fast if the build concluded without success, and erroring if it succeeded but the artifact never materializes -- rather than depending on the whole therock-ci-linux matrix (which would also block on unrelated components' GPU tests) or using a fixed wait budget. Readiness is keyed on the arch-specific CK plugin (libMIOpenCKGroupedConv_<arch>.so) being present in the fetched tree, not on miopen_gtest (which ships in a separate, arch-independent artifact and can appear before the arch-specific one). - The rocjitsu KMD config is CU-corrected to the arch's real CU count before each run (therock_dbsync_cu_correct_config.py rewrites the device-block CU fields; topology untouched, no emulation-cost change) -- rocjitsu's stock cdna configs declare a minimal (~4 CU) device, which makes CK's grouped-wrw occupancy-based auto-split-k miscompute and false-flag valid perf configs against real hardware CU counts. Also adds the standalone therock-dbsync-rocjitsu.yml: a manually-dispatchable companion that takes an explicit artifact_run_id for pointing at arbitrary/pinned builds the automatic gate can't self-reference (e.g. a ROCm/TheRock multi_arch_ci.yml dispatch, or bisecting an older run). ALMIOPEN-1779 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Aug 24, 2026
Adds an automatic per-PR `dbsync-rocjitsu` job to therock-ci.yml (matrixed over gfx942 + gfx950) that fetches this run's own TheRock source-CK build artifact and runs the MIOpen dbsync check (StaticFDBSync) under the rocjitsu KMD interposer on a GPU-less runner. dbsync is CPU-only (SKIP_KDB_TESTING=1, no kernel launch) except CK grouped-conv IsApplicable, which the arch-matched rocjitsu config services -- so this validates each arch's SystemDB against the PR's own CK without real hardware. This catches a PR's own CK change (the #6343 scenario), unlike the frozen-container CK path, and doesn't depend on scarce GPU test runners. Key wiring: - Gated on needs.setup.outputs.miopen_selected (computed in the setup job by parsing linux_projects' per-bundle comma-separated projects_to_test), so it runs only when TheRock's MIOpen CI would. - needs: setup only. The fetch step is build-aware: it polls therock_dbsync_build_leg_state.py for this run's MIOpen build-leg status, waiting while queued/running, failing fast if the build concluded without success, and erroring if it succeeded but the artifact never materializes -- rather than depending on the whole therock-ci-linux matrix (which would also block on unrelated components' GPU tests) or using a fixed wait budget. Readiness is keyed on the arch-specific CK plugin (libMIOpenCKGroupedConv_<arch>.so) being present in the fetched tree, not on miopen_gtest (which ships in a separate, arch-independent artifact and can appear before the arch-specific one). - The rocjitsu KMD config is CU-corrected to the arch's real CU count before each run (therock_dbsync_cu_correct_config.py rewrites the device-block CU fields; topology untouched, no emulation-cost change) -- rocjitsu's stock cdna configs declare a minimal (~4 CU) device, which makes CK's grouped-wrw occupancy-based auto-split-k miscompute and false-flag valid perf configs against real hardware CU counts. Also adds the standalone therock-dbsync-rocjitsu.yml: a manually-dispatchable companion that takes an explicit artifact_run_id for pointing at arbitrary/pinned builds the automatic gate can't self-reference (e.g. a ROCm/TheRock multi_arch_ci.yml dispatch, or bisecting an older run). ALMIOPEN-1779 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Aug 24, 2026
Adds an automatic per-PR `dbsync-rocjitsu` job to therock-ci.yml (matrixed over gfx942 + gfx950) that fetches this run's own TheRock source-CK build artifact and runs the MIOpen dbsync check (StaticFDBSync) under the rocjitsu KMD interposer on a GPU-less runner. dbsync is CPU-only (SKIP_KDB_TESTING=1, no kernel launch) except CK grouped-conv IsApplicable, which the arch-matched rocjitsu config services -- so this validates each arch's SystemDB against the PR's own CK without real hardware. This catches a PR's own CK change (the #6343 scenario), unlike the frozen-container CK path, and doesn't depend on scarce GPU test runners. Key wiring: - Gated on needs.setup.outputs.miopen_selected (computed in the setup job by parsing linux_projects' per-bundle comma-separated projects_to_test), so it runs only when TheRock's MIOpen CI would. - needs: setup only. The fetch step is build-aware: it polls therock_dbsync_build_leg_state.py for this run's MIOpen build-leg status, waiting while queued/running, failing fast if the build concluded without success, and erroring if it succeeded but the artifact never materializes -- rather than depending on the whole therock-ci-linux matrix (which would also block on unrelated components' GPU tests) or using a fixed wait budget. Readiness is keyed on the arch-specific CK plugin (libMIOpenCKGroupedConv_<arch>.so) being present in the fetched tree, not on miopen_gtest (which ships in a separate, arch-independent artifact and can appear before the arch-specific one). - The rocjitsu KMD config is CU-corrected to the arch's real CU count before each run (therock_dbsync_cu_correct_config.py rewrites the device-block CU fields; topology untouched, no emulation-cost change) -- rocjitsu's stock cdna configs declare a minimal (~4 CU) device, which makes CK's grouped-wrw occupancy-based auto-split-k miscompute and false-flag valid perf configs against real hardware CU counts. Also adds the standalone therock-dbsync-rocjitsu.yml: a manually-dispatchable companion that takes an explicit artifact_run_id for pointing at arbitrary/pinned builds the automatic gate can't self-reference (e.g. a ROCm/TheRock multi_arch_ci.yml dispatch, or bisecting an older run). ALMIOPEN-1779 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Due to compiler version update, there are test failures in the test suite
test_grouped_convnd_bwd_weightwhen running ongfx90a. There are four failing tests for FP16/BF16 that arise from a single kernel instance. As the problem is in the currentdevelopbranch, the test failures are blocking any PR merges intodevelop. An example of a failed CI runs is here: http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/develop/558/pipeline/. The underlying compiler problem is potentially the same as described in #6342 as tests are passing for clang compiler version 20.0 and failing for clang compiler version 22.0.Technical Details
This PR disables the compilation of the problematic bwd weight conv instance for
gfx90aby adding a new CMake flagCK_USE_GFX90Athat allows us to detect when we are compiling forgfx90a. Using the new CMake flag, compilation of instanceDeviceGroupedConvBwdWeight_Xdl_CShuffleV3<64, 128, 32, 32, Default, 8, 4, 1, 8, 8, 8, 8, 1, 1, 2>is disabled forgfx90a.Additionally, this PR skips fp16 FMHA forward dropout tests that use the d256 tile in batch mode. The AMDGPU compiler (clang 22) aliases Philox RNG VGPR parameters (
ph_seed,ph_head_offset) with other live data under high register pressure (383 VGPRs), causing corruptedbuffer_store_bytewrites for dropout randval on wave lanes 32-63 and NaN in output/LSE tensors. bf16 dropout is unaffected (different type conversion codegen path) and covers the same code paths.Changes:
include/ck_tile/core/config.hpp: Add version-gatedCK_TILE_WORKAROUND_ROCM_7_12_FP16_DROPOUT_MISCOMPILEmacrotest/ck_tile/fmha/test_fmha_fwd.cpp:GTEST_SKIPfor fp16 + d256 tile + batch mode inTEST_P(Dropout, ...)Verified: