ck: sync ck_tile examples with ROCm 10.1 by zichguan-amd · Pull Request #552 · ROCm/rocm-examples

zichguan-amd · 2026-10-16T22:55:31Z

Motivation

Align Libraries/ComposableKernel with the Composable Kernel shipped in ROCm 10.1. The stable CI image ghcr.io/rocm/rocm-examples-ubuntu-24.04-rocm:10.1 (#547, #548) ships CK 39d8d60365d2272a5366f76ccb3bfba24b219eaf, and the CK examples on amd-staging no longer compile against it: reduce, streamk_gemm, gemm_multi_abd and batched_contraction fail with errors such as unknown type name 'index_t' in ck_tile/ops/common/generic_2d_block_shape.hpp. CI only builds CK when Libraries/ComposableKernel/** changes (or on the nightly schedule), so this has not shown up on PRs.

Once the CK tests are registered with ctest, many of them fail on the GPUs they are built for (gfx942 and gfx950). Those failures were there before this PR, hidden because ctest reported every CK test as "Not Run". The later commits fix them, so every registered CK test now passes on gfx942 and gfx950.

Technical Details

ck: sync ck_tile examples with TheRock 10.1.0rc3 syncs the examples with ROCm/rocm-libraries@39d8d60, the rocm-libraries pin of TheRock 10.1.0rc3. This is the same CK commit as the ROCm 10.1 stable image.
- All 220 synced sources are upstream's files, formatted in rocm-examples style by this repo's pre-commit hooks (clang-format 18.1.4 with no .clang-format, i.e. LLVM style; whitespace; black). Upstream relies on its #include order (SortIncludes: false), so blank lines separate includes wherever LLVM sorting would reorder them.
- Synced files carry upstream's license header (Copyright (c) Advanced Micro Devices, Inc., or its affiliates. + SPDX-License-Identifier: MIT) instead of the full MIT text. The license is still MIT, as in LICENSE.md. rocm-examples-owned files (CMake, scripts, READMEs) keep their headers, with the copyright year updated to 2026 where they change.
- CMake and README changes that follow upstream:
  - attention/fmha_fwd_v3 is removed (not in upstream at this pin), along with its root README.md entry.
  - The large-tensor grouped convolution example and basic/reduce/reduce.hpp are removed.
  - block_scale_gemm builds gemm_quant.cpp and the new gemm_bquant_quantgrouped_* sources, with upstream's compile options.
  - grouped_gemm's quant example builds the new per-type quant_grouped_gemm_* sources.
  - fmha gains fmha_fwd_head_grouping.hpp and quant.hpp.
- batched_contraction and gemm_multi_abd are restricted to gfx908/gfx90a/gfx942/gfx950, and streamk_gemm to gfx90a/gfx942/gfx950, through cmake/FilterHIPArchitectures.cmake.
- basic/copy has no upstream counterpart at this pin and is kept as is.
ck: register the CK tests so they run under ctest fixes the CK test registration: on amd-staging, ctest reports every CK test as "Not Run".
- The tests used add_test(<name> COMMAND ...), so CTest looked for an executable named COMMAND. They now use add_test(NAME ... COMMAND ...), with the script and the executable passed as separate arguments.
- layernorm2d and add_rmsnorm2d_rdquant now point at test scripts that exist.
- The test scripts are now executable, use LF line endings and have a valid shebang.
- A failing case now fails its test.

The remaining commits fix the tests that then fail. Each commit message explains its change in detail.

CMake (cmake/FindCKTile.cmake, Libraries/ComposableKernel/CMakeLists.txt, attention/fmha/CMakeLists.txt):
- ck: use the OCP fp8 format in host code when building for gfx950 or gfx12 and ck: define Composable Kernel's architecture macros in FindCKTile: define CK_TILE_USE_OCP_FP8, CK_TILE_USE_WMMA, CK_GFX950_SUPPORT and CK_USE_GFX1250 from the HIP architectures, as Composable Kernel's own build does. Without CK_TILE_USE_OCP_FP8, the host references of the fp8/bf8 examples used the FNUZ format while gfx950 kernels use OCP, so their results did not match. As in CK's own build, the macros apply to all targets of a build, so the CK README now says to build gfx950 and gfx12 separately from other targets.
- ck: generate the FMHA backward kernels: fmha_bwd was built without any kernels ("no kernel found"). The backward generator takes one --filter pattern per kernel, separated by @.
- ck: build the CK examples in Release by default: Libraries/ComposableKernel configured on its own, as its README shows, had no build type.
- ck: set the CK output directory once: removes a duplicate CMAKE_RUNTIME_OUTPUT_DIRECTORY.
Test scripts and arguments:
- ck: fix the arguments of the CK normalization, quantization and FMHA tests: several tests passed options that the examples no longer accept at this CK version (-stride, -squant, -prec_i on the basic examples).
- ck: limit the FMHA tests to the kernels the examples build: the examples generate only fp16 group-mode kernels up to head dimension 64, but the scripts ran every precision, mode and head dimension.
- ck: test the GEMM examples with the sizes of Composable Kernel's GEMM tests: the scripts used sizes that the unpadded basic and weight-preshuffle kernels reject. They now use the sizes of CK's own test/ck_tile/gemm* tests, and weight preshuffle gets its own script.
- ck: run the LDS pipeline of the batched transpose test only on gfx950: the script ran -pipeline=1 on every GPU. That pipeline uses gfx950's LDS transpose loads and gives wrong results on gfx942. Composable Kernel's own test/ck_tile/batched_transpose test also skips it on other GPUs.
Synced example code. These two commits change upstream's C++, because the host check is wrong upstream too:
- ck: scale and check the fp8 outputs of the norm examples with the GPU's fp8 format: layernorm2d and rmsnorm2d scaled fp8 outputs to a hard-coded 240 (the FNUZ maximum) instead of numeric<YDataType>::max() (448 for OCP on gfx950/gfx12), and checked fp8 outputs for an exact match. The fp8 output check now allows one fp8 rounding step.
- ck: check the fused MoE example with gemm-style thresholds: the kernel rounds its intermediate result to bf16 and adds the topk expert results with bf16 atomics, while the reference stays in float, so about 20% of values missed the fixed 1e-2 tolerance. It now uses ck_tile::get_relative_threshold/get_absolute_threshold, as the GEMM examples do, with the topk sum treated like split-K.

Test Plan

CI: Build Libraries (build_libraries.yml, the pull_request run at 04342289). It runs in the ROCm 10.1 stable image with ROCM_EXAMPLES_ENABLE_CK=ON and GPU_TARGETS gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1200;gfx1201.
ctest of every registered CK test at eba28778, and of batched_transpose again at 04342289 on gfx942 and gfx950, in ghcr.io/rocm/rocm-examples-ubuntu-24.04-rocm:10.1, configured from the repository root like CI, for the local GPU only:
- gfx950 (MI350X): all 37 CK tests are registered, so every CK test runs on at least one GPU it supports.
- gfx942 (MI300A): 36 tests (fused_moe is gfx950-only).
- gfx1100: 9 tests (the others are filtered to gfx9 or gfx12).
pre-commit and the linting.yml checks.

Test Result

CI: Build Libraries run 40702009727 at 04342289 passed with ROCM_EXAMPLES_ENABLE_CK=ON. All 47 CK targets built for all nine GPU targets. Build ROCm Examples, Linting, pre-commit and the security scan also passed.
ctest at eba28778:

GPU Registered CK tests Passed

gfx950 (MI350X) 37 37

gfx942 (MI300A) 36 35 at eba28778, 36 at 04342289

gfx1100 9 9

At eba28778, batched_transpose failed on gfx942 because its script also ran the gfx950-only -pipeline=1. At 04342289 it passes on gfx942 (pipeline 0 only) and on gfx950 (pipelines 0 and 1).

gfx12 was not tested on hardware.
Checks: pre-commit and the linting checks pass.
On amd-staging, for comparison: the four examples listed under Motivation fail to compile in the 10.1 image, and ctest reports all CK tests as "Not Run".

Added/Updated documentation?

Yes
No, does not apply to this PR.

Included Visual Studio files?

Yes
No, does not apply to this PR.

Submission Checklist

New examples contain proper CMake and Make files.
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.
Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

🤖 Generated with Claude Code

Sync Libraries/ComposableKernel with ROCm/rocm-libraries@39d8d60 (projects/composablekernel/example/ck_tile), the rocm-libraries pin of the TheRock 10.1.0rc3 release candidate: https://rc.repo.amd.com/rocm/core/tarball/therock-dist-linux-gfx110X-all-tests-10.1.0rc3.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 - upstream example dir missing (kept): - basic/copy -> 39_copy - overlays (manual CMake/README adaptations): - applied 0001-ck-therock-10.1-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>

Copilot

🟡 Changes recommended

One or more issues must be addressed before approval.

4 open issues

What changed in this PR

Aligns CK Tile examples with ROCm 10.1 and repairs CTest integration.

Changes:

Updates upstream headers, API usage, code generation, and quantized GEMM instantiations.
Adjusts architecture restrictions and removes obsolete example references.
Corrects test registration and failure propagation; three basic-example registrations still use incompatible scripts.

File	Description
README.md	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/smoothquant.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/smoothquant.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/instances/smoothquant_instance_common.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/instances/smoothquant_fwd_api.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/instances/smoothquant_fp16_n768_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/instances/smoothquant_fp16_n64_n128_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/instances/smoothquant_fp16_n256_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/instances/smoothquant_bf16_n768_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/instances/smoothquant_bf16_n64_n128_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/instances/smoothquant_bf16_n512_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/instances/smoothquant_bf16_n256_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/quantization/smoothquant/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/pooling/pool3d/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/rmsnorm2d/rmsnorm2d_fwd.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/rmsnorm2d/rmsnorm2d_fwd.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/rmsnorm2d/generate.py	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/rmsnorm2d/example_rmsnorm2d_fwd.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/rmsnorm2d/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/layernorm2d/layernorm2d_fwd.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/layernorm2d/layernorm2d_fwd.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/layernorm2d/generate.py	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/layernorm2d/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_instance_common.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n8192_tp_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n8192_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n768_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n64_n128_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n512_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n4096_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n3072_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n256_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n2048_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_fp16_n1536_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n8192_tp_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n8192_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n768_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n64_n128_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n512_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n4096_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n3072_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n256_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n2048_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_bf16_n1536_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/instances/add_rmsnorm2d_rdquant_fwd_api.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/example_add_rmsnorm2d_rdquant_fwd.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/normalization/add_rmsnorm2d_rdquant/add_rmsnorm2d_rdquant_fwd.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_sorting/moe_sorting_api.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_sorting/moe_sorting_api.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_sorting/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/moe_smoothquant.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/moe_smoothquant.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_instance_common.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fwd_api.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n768_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n64_n128_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n512_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n4096_tp_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n4096_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n3072_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n256_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n2048_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n1536_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_fp16_n1024_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n768_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n64_n128_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n512_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n4096_tp_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n4096_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n3072_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n256_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n2048_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n1536_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/instances/moe_smoothquant_bf16_n1024_instance.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/moe_smoothquant/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/instances/fused_moesorting_api.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/instances/fused_moegemm_fp16_m32.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/instances/fused_moegemm_bf16_m32.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/instances/fused_moegemm_api.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/instances/fused_moegemm_api_traits.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/instances/fused_moegemm_api_internal.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/instances/fused_moe_api.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/fused_moesorting.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/fused_moegemm.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/fused_moe.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/moe/fused_moe/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/streamk_gemm/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/run_grouped_gemm_example.inc	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/README.md	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/quant_grouped_gemm_fp8_tensor.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/quant_grouped_gemm_fp8_rowcol.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/quant_grouped_gemm_fp8_bquant.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/quant_grouped_gemm_fp8_aquant.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/quant_grouped_gemm_bf8_tensor.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/quant_grouped_gemm_bf8_rowcol.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/quant_grouped_gemm_bf8_bquant.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/quant_grouped_gemm_bf8_aquant.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/grouped_gemm/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm/test_ck_tile_gemm_basic.sh	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm/run_gemm_example_common.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm/README.md	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm/gemm_weight_preshuffle.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm/gemm_splitk_two_stage.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm/gemm_basic.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm_multi_d/utils.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm_multi_d/run_gemm_multi_d_fp16_example.inc	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm_multi_d/README.md	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm_multi_d/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm_multi_abd/utils.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm_multi_abd/run_gemm_multi_abd_fp16_example.inc	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/gemm_multi_abd/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/flatmm/run_grouped_flatmm_example.inc	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/flatmm/run_flatmm_example.inc	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/flatmm/README.md	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/flatmm/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/block_scale_gemm/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/batched_gemm/run_batched_gemm_example.inc	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/batched_gemm/README.md	Updated as part of this pull request.
Libraries/ComposableKernel/gemm/batched_gemm/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/convolution/grouped_convolution/README.md	Updated as part of this pull request.
Libraries/ComposableKernel/convolution/grouped_convolution/grouped_convolution_forward.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/convolution/grouped_convolution/grouped_convolution_forward_bias_clamp.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/convolution/grouped_convolution/grouped_convolution_backward_weight.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/convolution/grouped_convolution/grouped_convolution_backward_weight_two_stage.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/convolution/grouped_convolution/grouped_convolution_backward_data.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/convolution/grouped_convolution/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/contraction/batched_contraction/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/basic/reduce/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/basic/permute/permute.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/permute/permute.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/permute/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/basic/permute/alternative_impl/matrix_core_swizzle.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/permute/alternative_impl/matrix_core_swizzle.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/permute/alternative_impl/matrix_core_swizzle_kernel.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/img2col/image_to_column.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/img2col/image_to_column.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/img2col/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/basic/elementwise/elementwise_example.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/elementwise/elementwise_example_transpose.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/elementwise/elementwise_example_add_4d.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/elementwise/elementwise_common.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/elementwise/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/basic/copy/test_tile_example.sh	Updated as part of this pull request.
Libraries/ComposableKernel/basic/copy/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/basic/batched_transpose/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/basic/batched_transpose/batched_transpose_example.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/basic/batched_transpose/batched_transpose_example.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/topk_softmax/topk_softmax.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/topk_softmax/topk_softmax_api.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/topk_softmax/topk_softmax_api.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/topk_softmax/CMakeLists.txt	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/utils.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/test_ck_tile_fmha_fwd.sh	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/rotary.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/README.md	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/quant.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/generate.py	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/codegen/utils.py	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/codegen/ops/fmha_fwd_appendkv.py	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/codegen/ops/__init__.py	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/codegen/cmake_config.py	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/codegen/arch.py	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/codegen/__init__.py	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha/bias.hpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha_fwd_v3/instances/fmha_fwd_v3_d128_fp16_nmask.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha_fwd_v3/instances/fmha_fwd_v3_d128_fp16_mask.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha_fwd_v3/instances/fmha_fwd_v3_d128_bf16_nmask.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha_fwd_v3/instances/fmha_fwd_v3_d128_bf16_mask.cpp	Updated as part of this pull request.
Libraries/ComposableKernel/attention/fmha_fwd_v3/.gitignore	Updated as part of this pull request.
Libraries/ComposableKernel/attention/CMakeLists.txt	Updated as part of this pull request.

_{Review effort: Balanced}

@@ -91,4 +91,4 @@ target_include_directories(${example_name} PRIVATE ${CMAKE_CURRENT_LIST_DIR})
 target_link_libraries(${example_name} PRIVATE composable_kernel::ck_tile hip::host)
 target_sources(${example_name} PRIVATE layernorm2d_fwd.cpp ${EXAMPLE_FWD_GEN_SRCS})

-add_test(${example_name} COMMAND "${CMAKE_CURRENT_LIST_DIR}/${example_script} $<TARGET_FILE:${example_name}>")
+add_test(NAME ${example_name} COMMAND "${CMAKE_CURRENT_LIST_DIR}/${example_script}" "$<TARGET_FILE:${example_name}>")


@@ -83,7 +83,7 @@ add_executable(${example_name}
    smoothquant.cpp
 )
 target_include_directories(${example_name} PRIVATE ${CMAKE_CURRENT_LIST_DIR})
-add_test(${example_name} COMMAND "${CMAKE_CURRENT_LIST_DIR}/${example_script}" "$<TARGET_FILE:${example_name}>")
+add_test(NAME ${example_name} COMMAND "${CMAKE_CURRENT_LIST_DIR}/${example_script}" "$<TARGET_FILE:${example_name}>")


…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>

…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?   - [ ] 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>

zichguan-amd and others added 2 commits October 16, 2026 17:27

zichguan-amd requested review from a team as code owners October 16, 2026 22:55

zichguan-amd requested review from j-stephan and a balanced review from Copilot October 19, 2026 13:19

Copilot started reviewing on behalf of zichguan-amd October 19, 2026 13:20

View session

Copilot AI reviewed Oct 19, 2026

View reviewed changes

zichguan-amd and others added 11 commits October 19, 2026 15:13

ck: set the CK output directory once

b8b2e10

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>

ck: build the CK examples in Release by default

4ab9844

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>

This was referenced Oct 30, 2026

ck: sync ck_tile examples with ROCm 10.2.0rc0 #565

Open

ci: add ROCm version and CK inputs to the TheRock workflow dispatch #567

Merged

Navigation Menu

Repository navigation

ck: sync ck_tile examples with ROCm 10.1 - #552

ck: sync ck_tile examples with ROCm 10.1#552
zichguan-amd wants to merge 13 commits into
amd-stagingfrom
zichguan/ck-sync-10.1-rc3

zichguan-amd commented Oct 16, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

GPU	Registered CK tests	Passed
gfx950 (MI350X)	37	37
gfx942 (MI300A)	36	35 at `eba28778`, 36 at `04342289`
gfx1100	9	9

Conversation

zichguan-amd commented Oct 16, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Added/Updated documentation?

Included Visual Studio files?

Submission Checklist

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

zichguan-amd commented Oct 16, 2026 •

edited

Loading