Repository navigation
ck: sync ck_tile examples with ROCm 10.2.0rc0 - #565
Open
zichguan-amd wants to merge 16 commits into
Open
zichguan-amd wants to merge 16 commits into
zichguan-amd wants to merge 16 commits into
Conversation
Sync Libraries/ComposableKernel with ROCm/rocm-libraries@8a29d2e (projects/composablekernel/example/ck_tile), the rocm-libraries pin of the TheRock 10.2.0rc0 release candidate: https://rc.repo.amd.com/rocm/core/tarball/therock-dist-linux-gfx110X-all-tests-10.2.0rc0.tar.gz Synced files are upstream's, license headers included, formatted in rocm-examples style by this repo's pre-commit hooks (clang-format 18.1.4 without a .clang-format, i.e. LLVM style; whitespace; black). Changed rocm-examples-owned files keep their headers, with the copyright year brought up to date. - updated: 220 files, unchanged: 0 - added: - gemm/gemm_common/vector_size_fallback_dispatch.hpp (included by gemm/batched_gemm/run_batched_gemm_example.inc) - upstream example dir missing (kept): - basic/copy -> 39_copy - overlays (manual CMake/README adaptations): - applied 0001-ck-therock-10.1-adaptations.patch - applied 0002-ck-therock-10.2-adaptations.patch Validated locally at the branch tip (with any follow-up fixup commits): pre-commit, linting.yml checks, CK build + ctest on gfx1100 and CK build for build_libraries.yml GPU_TARGETS in ghcr.io/rocm/rocm-examples-ubuntu-24.04-multiarch:latest with the RC tarball. Co-Authored-By: Claude <noreply@anthropic.com>
The CK add_test() calls used the legacy `add_test(<name> COMMAND ...)` signature. That signature takes the word after the name as the command, so CTest looked for an executable literally named "COMMAND" and reported every CK test as "Not Run". It also does not expand generator expressions, and the script tests passed "<script> $<TARGET_FILE:...>" as one argument. CI builds the CK examples but never runs their tests, so this went unnoticed. - Register every CK test with add_test(NAME ... COMMAND ...), passing the script (by absolute path) and the executable as separate arguments. - Point layernorm2d and add_rmsnorm2d_rdquant at scripts that exist: layernorm2d/test_ck_tile_add_rmsnorm2d_rdquant.sh (layernorm cases) becomes test_ck_tile_layernorm2d.sh, and add_rmsnorm2d_rdquant/smoke_test.sh becomes test_ck_tile_add_rmsnorm2d_rdquant.sh. - Make the scripts executable by CTest: +x, LF line endings, and a valid shebang (batched_transpose, moe_sorting; bash for rmsnorm2d, which uses arrays). - Make a failing case fail the test: `set -e` in the command-list scripts, `exit 1` on the first failure in the gemm/flatmm scripts, and a nonzero exit when any rmsnorm2d case is not valid or any copy run fails. - copy and fmha_fwd take the executable from $1, as fmha_bwd already does, instead of a hard-coded path / searching the working directory; fmha_fwd no longer aborts under `set -u` when GPU_arch is unset. Co-Authored-By: Claude <noreply@anthropic.com>
…tests Now that the CK tests run, several of them fail on arguments the examples do not accept at the synced CK version: - rmsnorm2d_basic, smoothquant_basic and add_rmsnorm2d_rdquant_tile ran the script of the generated-kernel example, whose options and precisions the basic examples do not have (-prec_i, -fadd, -stride, bf16). They now run the example with its defaults. - layernorm2d and smoothquant passed -stride, which the examples replaced with per-tensor strides. They now pass -x_stride/-xr_stride/-y_stride/ -yr_stride (layernorm2d) and -x_stride/-y_stride (smoothquant). - layernorm2d no longer runs the n > 8192 cases with fused quant: generate.py builds the two-pass kernels without it. - fmha_fwd passed -squant, which the example replaced with -qscale. Its fp8 cases now follow upstream's smoke_test_fwd.sh: no plain fp8 cases, and the fp8bf16/fp8fp32 cases use -init=3 and -qscale without -bias. Co-Authored-By: Claude <noreply@anthropic.com>
…fx12 ck_tile picks the fp8 format of device code per target: OCP on gfx950 and gfx12, FNUZ elsewhere. Host code gets FNUZ unless CK_TILE_USE_OCP_FP8 is defined, so on gfx950 the host reference results of the fp8/bf8 examples did not match the device results (e.g. gemm_universal, block_scale_gemm and moe_smoothquant failed their tests). Define CK_TILE_USE_OCP_FP8 when the HIP architectures include gfx950 or gfx12, as Composable Kernel's own CMake does. FindCKTile.cmake is included by every CK example after its architecture filtering, so the check uses the architectures the example is actually built for, also when an example is built on its own. The README notes that a build for both gfx950/gfx12 and FNUZ targets such as gfx942 gives wrong fp8 results on the latter. Co-Authored-By: Claude <noreply@anthropic.com>
Libraries/ComposableKernel/CMakeLists.txt set folder_bin and CMAKE_RUNTIME_OUTPUT_DIRECTORY twice with the same values. Keep the first. Co-Authored-By: Claude <noreply@anthropic.com>
The fmha_bwd example was built without any kernels, so every run reported "no kernel found". The backward code generator needs three kernels per configuration (dot_do_o, convert_dq, dq_dk_dv) and takes one --filter pattern for each, separated by "@". It generates a configuration only if all three names match, and the dot_do_o and convert_dq names do not match *fp16_group* (e.g. fmha_bwd_dot_do_o_d64_fp16_b64_group_...). Use *fp16*group* for each of the three kernels. This generates the fp16 group-mode kernels with head dimension 64, as the forward pass does (76 kernels and the API for gfx950). Co-Authored-By: Claude <noreply@anthropic.com>
When Libraries/ComposableKernel is configured on its own, as its README shows, CMAKE_BUILD_TYPE was empty. Default it to Release, as the repository root and Composable Kernel's own build do. Co-Authored-By: Claude <noreply@anthropic.com>
The FMHA examples generate only the fp16 group mode kernels for head dimensions up to 64 (forward: with LSE, without elementwise bias, splitkv and appendkv; backward: without dbias and deterministic mode), but the test scripts ran all precisions, modes and head dimensions. The unsupported cases fail with "not supported yet", and the forward script stopped at the first batch mode padding test, so its padding, boundary and fp8 tests never ran. The scripts now run only the supported cases: - fmha_fwd: fp16, group mode, -d/-d_v up to 64, -lse=1, -bias=n/a. The group mode padding and boundary tests use these settings too and record failures like the other tests. The batch mode, fp8bf16/fp8fp32 and appendkv tests and the -s/-a options are removed. - fmha_bwd: fp16, group mode, -d up to 64. Co-Authored-By: Claude <noreply@anthropic.com>
Composable Kernel's own build defines CK_TILE_USE_WMMA, CK_TILE_USE_OCP_FP8, CK_GFX950_SUPPORT and CK_USE_GFX1250 depending on its GPU targets. ck_tile and the examples use them in host code too, so define them in the same way. This enables e.g. the tf32 GEMMs and the gfx950 fp8 kernel configurations on gfx950. Co-Authored-By: Claude <noreply@anthropic.com>
… tests The basic GEMM example does not pad its 256x256x64 tiles, so its test failed for the sizes that are not multiples of them, e.g. m=128. The weight preshuffle example used the same script, whose k=64 is less than its 128-wide fp8 tiles. Use the sizes of Composable Kernel's own tests instead: the SingleTile, Regular and LargeMatrix cases of test/ck_tile/gemm for the basic and universal examples, and the cases of test/ck_tile/gemm_weight_preshuffle in a script of its own for the weight preshuffle example. Co-Authored-By: Claude <noreply@anthropic.com>
…'s fp8 format The host references of the layernorm2d and rmsnorm2d examples scale fused-quant outputs to a hard-coded 240.0, the maximum of the FNUZ fp8 format. gfx950 and gfx12 use the OCP fp8 format, whose maximum is 448, as does the kernel, so all fp8 cases failed there. Use ck_tile::numeric<YDataType>::max() like the add_rmsnorm2d_rdquant example does. The output check also used rtol = atol = 1e-2 for fp8, which check_err takes as an exact match. One fp8 rounding step is about 12% of the value, so fp8 cases failed at random on gfx942 too. Allow one rounding step and atol = 1, like int8, for the output check only. Co-Authored-By: Claude <noreply@anthropic.com>
The example compared its bf16 output with a float host reference within rtol = atol = 1e-2, and about 20% of the values failed on every run, on gfx950 and gfx942. The kernel rounds the intermediate result to bf16 for the second gemm's MFMA and adds the topk expert results into the output with bf16 atomics, while the reference keeps everything in float. Even a reference that models these roundings cannot match the kernel within 1e-2, because outputs cancel near zero. Use ck_tile::get_relative_threshold and get_absolute_threshold like the gemm examples do, treating the topk accumulation like split-k with kbatch = topk. Co-Authored-By: Claude <noreply@anthropic.com>
The test script ran every case with -pipeline=0 and -pipeline=1 on every GPU. The LDS pipeline (-pipeline=1) uses the LDS transpose loads of gfx950, and on gfx942 most of its output values were wrong, so the test failed there. Composable Kernel's own test/ck_tile/batched_transpose test skips this pipeline on GPUs other than gfx950 and gfx125x, and the example's README describes it as a gfx950 mode. Run -pipeline=1 only on gfx950, the one of these GPUs the example is built for. Detect the GPU like the FMHA test scripts do, with GPU_arch to override it. Co-Authored-By: Claude <noreply@anthropic.com>
Composable Kernel builds example/ck_tile/03_gemm when GPU_TARGETS matches "gfx94|gfx95|gfx90a|gfx125|gfx11|gfx12", but the example here was built only for the gfx9 GPUs. CK_TILE_USE_WMMA selects the kernel configurations for all GPU targets of the example, and the MFMA configurations do not compile for gfx11 and gfx12, nor the WMMA ones for gfx9. Build the example for the gfx9 GPUs if the targets include any of them, and for the gfx11 and gfx12 GPUs otherwise. Define CK_USE_WMMA_FP8, which enables the fp8 and bf8 kernel configurations for the WMMA instructions, like Composable Kernel does for gfx12, but only if the targets do not also include gfx11: gfx11 has no fp8 WMMA instructions, and with the macro defined the examples do not compile for it. The test scripts skip these precisions on gfx11, where the examples report them as unsupported. On gfx11, the weight preshuffle example's results are wrong: its kernel configuration permutes N, and the host-side ck_tile::shuffle_b_permuteN preshuffles B for gfx12 or the MFMA GPUs only. Composable Kernel's own test of the pipeline does not permute N for WMMA. Build the example only for the GPUs other than gfx11. With ROCm 10.2.0rc0, compiling the weight preshuffle example for gfx12 with CK_USE_WMMA_FP8 defined can crash the device-code link step: the LLVM Attributor pass recurses deeply enough to overflow the default 8 MiB thread stack, and clang-linker-wrapper's call to ld.lld segfaults. Confirmed this is a toolchain bug and not specific to this repo's build or to containers, by reproducing the identical crash bare metal on gfx1201 hardware with the default stack and building clean with an unlimited one; the fix is applied only to this target via HIP_COMPILER_LAUNCHER, which does not require any environment or operator-level configuration since the hard stack limit is unlimited by default. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
FindCKTile.cmake defines CK_GFX950_SUPPORT for all GPU targets of an example when they include gfx950, and the basic GEMM example then also builds its tf32 GEMMs for the other targets. With ROCm 10.2.0rc0, ck_tile does not compile them for gfx908 and gfx90a (no viable overloaded '=' in block_universal_gemm_as_bs_cr.hpp: the warp tile of the fallback is float), so a build for the GPU targets of the Libraries CI, which include gfx908, gfx90a and gfx950, fails. Undefine the macro in the device code for the targets other than gfx950 with -Xarch_<target> -UCK_GFX950_SUPPORT. The host code and the gfx950 device code keep it, so the example still offers tf32 and runs it on gfx950, and the code objects of the other targets have no tf32 GEMMs. A build for gfx950 alone does not change. Co-Authored-By: Claude <noreply@anthropic.com>
zichguan-amd
marked this pull request as ready for review
October 30, 2026 18:15
4 of 8 tasks
zichguan-amd
added a commit
that referenced
this pull request
Oct 30, 2026
…567) ## Motivation Build ROCm Examples (`ci_therock.yml`) and Nightly CI (`ci_nightly.yml`) can only be dispatched against the latest TheRock nightly, and they never build the Composable Kernel examples: the tarball and wheel setups leave `ENABLE_CK` unset and the preinstalled one sets it to `OFF`. The only workflow that builds CK, Build Libraries, is pinned to the ROCm 10.1 image. So a CK change cannot be checked in CI against a specific ROCm, e.g. a release candidate. #565 (CK sync with 10.2.0rc0) is an example: its Build Libraries run compiles the 10.2 examples against 10.1's CK and fails for that reason alone. ## Technical Details - **`build-rocm-examples-reusable.yml`** - New `enable_ck` input (boolean, default `false`). A new "Set CK build flag" step writes `ENABLE_CK=ON|OFF` for every install method, replacing the hard-coded `ENABLE_CK=OFF` of the preinstalled setup. - `rocm_version` already pins a nightly (Quartz uses it). An rc version such as `10.2.0rc0`, without an explicit `wheels_url`/`tarballs_url`, now installs from the release candidate index (`https://rc.repo.amd.com/rocm/core/tarball/`, `https://rc.repo.amd.com/rocm/core/whl-next/`) instead of the nightly one. - **`ci_therock.yml`, `ci_nightly.yml`**: new `rocm_version` (string, empty = latest nightly) and `enable_ck` (boolean) dispatch inputs, passed to the reusable workflow. With a `rocm_version`, `configure_ci.py` gets `--exclude-preinstalled`, as in `quartz_test.yml`, because the `stable_release` image cannot install another ROCm. Push, pull_request, schedule and Quartz runs pass neither input and are unchanged. **Note:** TheRock's Python wheels currently do not include Composable Kernel. The 10.2.0rc0 `rocm-sdk-devel` wheel has no `ck`/`ck_tile` headers and no `composable_kernel` CMake package, and the `rocm` meta-package has no other component that provides it. With `install_method=whl-multi-arch`, `enable_ck=true` turns the flag on, but `Libraries/ComposableKernel` skips itself ("Composable Kernel could not be found"). To build CK, use `install_method=tarball-multi-arch`; the tarballs ship CK. Example, for a CK branch that contains this change: ``` gh workflow run ci_therock.yml --ref <branch> \ -f rocm_version=10.2.0rc0 -f enable_ck=true \ -f install_method=tarball-multi-arch -f distro=ubuntu-24.04 -f gpu_config=gfx1100 ``` ## Test Plan - Dispatch of Build ROCm Examples on this branch with `rocm_version=10.2.0rc0`, `enable_ck=true`, gfx1100, Ubuntu 24.04, both install methods: [run 41123640840](https://github.com/ROCm/rocm-examples/actions/runs/41123640840). - pre-commit (including the GitHub Actions workflow lint) and zizmor on the three workflows. ## Test Result - **whl-multi-arch:** passed. It installed `rocm==10.2.0rc0` from `https://rc.repo.amd.com/rocm/core/whl-next/`, set `ROCM_EXAMPLES_ENABLE_CK=ON`, and ctest passed 371/371. As expected without CK in the wheels (see the note above), `Libraries/ComposableKernel` reported "Composable Kernel could not be found" and built nothing. - **tarball-multi-arch:** queued at the time of writing (waiting for a gfx1100 runner). On this branch the CK examples are still `amd-staging`'s, which do not compile against 10.2.0rc0's CK (#552, #565), so this job is expected to fail in the CK build; what it shows is that the rc tarball installs and the CK examples are configured. - **Checks:** pre-commit passes. zizmor reports the same findings before and after the change (36, none new). ## Added/Updated documentation? <!-- Make sure each new example has a README and the root-level README contains all new examples. --> <!-- New Libraries examples should have a library-level README explaining the dependencies, build process, and supported platforms. --> - [ ] Yes - [x] No, does not apply to this PR. ## Included Visual Studio files? - [ ] Yes - [x] No, does not apply to this PR. ## Submission Checklist - [ ] New examples contain proper CMake and Make files. - [x] I have updated relevant CI workflows and the new examples are being built and tested in CI. - [ ] Check and guard against unsupported ASICs if relevant dependent libraries have limited support. - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude <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
Align
Libraries/ComposableKernelwith the Composable Kernel of the TheRock 10.2.0rc0 release candidate, and carry over the fixes of #552 (the ROCm 10.1 sync). This PR also builds the GEMM examples for gfx11 and gfx12, as Composable Kernel itself does, and fixes a build failure that is new with 10.2.0rc0.This PR is based on
amd-staging, not on #552. It contains #552's fix commits, re-applied on top of the 10.2.0rc0 sync, so it replaces #552's sync commit.Technical Details
ck: sync ck_tile examples with TheRock 10.2.0rc0syncs the examples with ROCm/rocm-libraries@8a29d2e, the rocm-libraries pin of TheRock 10.2.0rc0. It is generated the same way as in ck: sync ck_tile examples with ROCm 10.1 #552 (upstream's files and license headers, formatted by this repo's pre-commit hooks).gemm/gemm_common/vector_size_fallback_dispatch.hppis added, because thegemmandbatched_gemmexamples now include it, and those targets getgemm/on their include path.basic/copystill has no upstream counterpart and is kept as is.ck: register the CK tests so they run under ctestck: fix the arguments of the CK normalization, quantization and FMHA testsck: use the OCP fp8 format in host code when building for gfx950 or gfx12ck: set the CK output directory onceck: generate the FMHA backward kernelsck: build the CK examples in Release by defaultck: limit the FMHA tests to the kernels the examples buildck: define Composable Kernel's architecture macros in FindCKTileck: test the GEMM examples with the sizes of Composable Kernel's GEMM testsck: scale and check the fp8 outputs of the norm examples with the GPU's fp8 formatck: check the fused MoE example with gemm-style thresholdsck: run the LDS pipeline of the batched transpose test only on gfx950ck: build the GEMM examples for gfx11 and gfx12. Composable Kernel buildsexample/ck_tile/03_gemmfor gfx11 and gfx12 too, but here it was built only for gfx9.CK_TILE_USE_WMMAapplies to all targets of a build, so the example is built for the gfx9 targets if the targets include any, and for the gfx11/gfx12 targets otherwise.CK_USE_WMMA_FP8(fp8/bf8 WMMA configurations) is defined for gfx12 builds that do not also include gfx11, which has no fp8 WMMA instructions. The test scripts skip fp8/bf8 on gfx11.shuffle_b_permuteNpreshuffles B for gfx12 and MFMA only, so its gfx11 results are wrong.ld.lldsegfaults. This happens on bare metal too, not only in containers. The target'sHIP_COMPILER_LAUNCHERnow runs the compiler withulimit -s unlimited, which only raises the soft limit, so no environment or operator setup is needed.ck: build the tf32 GEMMs of the GEMM example only for gfx950.CK_GFX950_SUPPORTapplies to all targets of a build that includes gfx950, sogemm_basicalso builds its tf32 GEMMs for the other targets. With 10.2.0rc0's ck_tile, these no longer compile for gfx908 and gfx90a (no viable overloaded '='inblock_universal_gemm_as_bs_cr.hpp). 10.1.0rc3 built the same targets without errors. As a result, a build forbuild_libraries.yml'sGPU_TARGETSfails.-Xarch_<target> -UCK_GFX950_SUPPORT. Only the host code and gfx950's device code keep it.FindCKTile.cmakeno longer say tf32 results are wrong on the other targets of a mixed build.Test Plan
ck-synccontainer test (ghcr.io/rocm/rocm-examples-ubuntu-24.04-multiarch:latestwith the TheRock 10.2.0rc0 gfx110X tarball). Like CI's tarball setup, it does two things:build_libraries.yml'sGPU_TARGETS(gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1200;gfx1201).ghcr.io/rocm/rocm-examples-ubuntu-24.04-rocm:10.1with the TheRock 10.2.0rc0 tarball of each GPU family, configured from the repository root like CI, for the local GPU only.gemm_basicbuilt for gfx908;gfx90a;gfx942;gfx950, and the tf32 kernels in each target's code object.ci_therock.yml, inputs from ci: add ROCm version and CK inputs to the TheRock workflow dispatch #567) atac8c207bwithrocm_version=10.2.0rc0,enable_ck=true,install_method=tarball-multi-arch, Ubuntu 24.04, gfx1100 and gfx1151: run 41126656536. It installs the 10.2.0rc0 tarball fromrc.repo.amd.comand builds and tests all examples, including Composable Kernel, for each GPU.linting.ymlchecks.Test Result
ctest (10.2.0rc0):
gfx1100 and gfx1201 now also run
gemm_basicandgemm_universal, and gfx1201 runsgemm_weight_preshuffle. These ran at the previous tip of this branch, which had the same commits except the tf32 commit and was based ona817f653. The tf32 commit does not change builds for a single target.Weight preshuffle on gfx1201: the full GEMM example set builds using only the target's launcher (no
ulimitin the job). In the test scripts, weight preshuffle passed 36 cases, basic 12 and universal 21, with no errors.Build for CI's
GPU_TARGETS: at the previous tip,gemm_basicwas the only target that failed, with the tf32 errors. With the tf32 commit,gemm_basicbuilds for gfx908;gfx90a;gfx942;gfx950, and only gfx950's code object has tf32 kernels. The full container test atd9916574(the last commit before mergingamd-staging, which only brought in workflow changes) passed: Composable Kernel built for all nine targets (47 executables, no errors), and ctest on gfx1100 passed 11/11.Checks: pre-commit and the linting checks pass.
Not tested: tf32 from a multi-target binary on gfx950 hardware; gfx11 GPUs other than gfx1100; gfx12 GPUs other than gfx1201.
CI dispatch against 10.2.0rc0 (run 41126656536) passed on both GPUs, with the 10.2.0rc0 tarballs from
rc.repo.amd.comandROCM_EXAMPLES_ENABLE_CK=ON. 13 CK targets built on each:gfx110X-allgfx1151This is the first run of these examples on gfx1151, including
gemm_basicandgemm_universal.PR CI's Build Libraries uses the ROCm 10.1 stable image, so it compiles these examples against 10.1's Composable Kernel and fails (e.g.
gemm_universalandbatched_gemmin 10.1'swarp_gemm_dispatcher.hpp). This is expected for examples synced to 10.2.0rc0, until Build Libraries moves to a 10.2 image.Added/Updated documentation?
Included Visual Studio files?
Submission Checklist
🤖 Generated with Claude Code