zichguan-amd · Oct 16, 2026 · Oct 16, 2026 · Oct 19, 2026 · Oct 19, 2026 · Oct 19, 2026
@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -23,6 +23,11 @@
 cmake_minimum_required(VERSION 3.21 FATAL_ERROR)
 project(ComposableKernel_examples LANGUAGES CXX)
 
+# Set default build type to Release, as Composable Kernel's own build does
+if(NOT CMAKE_BUILD_TYPE)
+    set(CMAKE_BUILD_TYPE Release)
+endif()
+
 if(CMAKE_SYSTEM_NAME STREQUAL "Windows")
     message(STATUS "Composable Kernel examples are supported on Linux only.")
     return()
@@ -44,9 +49,6 @@ if(ROCM_EXAMPLES_GPU_LANGUAGE STREQUAL "CUDA")
     return()
 endif()
 
-file(RELATIVE_PATH folder_bin ${CMAKE_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR})
-set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin/${folder_bin})
-
 include("${CMAKE_CURRENT_LIST_DIR}/../../cmake/ROCmPath.cmake")
 
 find_package(composable_kernel)

@@ -32,3 +32,9 @@ independently.
 - `$ cd Libraries/ComposableKernel`
 - `$ cmake -S . -B build`
 - `$ cmake --build build`
+
+As in Composable Kernel's own build, some settings depend on the GPU targets an example is built for and apply to all
+of them: the examples use the OCP fp8 format when the targets include gfx950 or gfx12 (the FNUZ format otherwise),
+the kernel configurations for WMMA instructions when they include gfx11 or gfx12, and gfx950-specific features such
+as tf32 GEMMs and fp8 kernel configurations when they include gfx950. Build for gfx950 and gfx12 separately from
+other targets such as gfx942: in a build for both, the fp8, bf8 and tf32 results on the other targets are wrong.
@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -27,5 +27,4 @@ file(RELATIVE_PATH folder_bin ${CMAKE_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR})
 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin/${folder_bin})
 
 add_subdirectory(fmha)
-add_subdirectory(fmha_fwd_v3)
 add_subdirectory(topk_softmax)
@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -28,9 +28,6 @@ set(example_fwd_script test_ck_tile_fmha_fwd.sh)
 set(example_bwd ${example_name}_bwd)
 set(example_bwd_script test_ck_tile_fmha_bwd.sh)
 
-set(example_fwd_v3 ${example_name}_fwd_v3)
-set(example_fwd_v3_script test_ck_tile_fmha_fwd_v3.sh)
-
 cmake_minimum_required(VERSION 3.21 FATAL_ERROR)
 project(${example_name} LANGUAGES CXX)
 
@@ -105,7 +102,15 @@ add_dependencies(${example_fwd} fmha_fwd_gen_srcs)
 
 set_source_files_properties(fmha_fwd.cpp ${EXAMPLE_FWD_GEN_SRCS} PROPERTIES LANGUAGE ${ROCM_EXAMPLES_GPU_LANGUAGE})
 
-target_compile_definitions(${example_fwd} PRIVATE CK_TILE_FMHA_FWD_FAST_EXP2=1)
+target_compile_definitions(
+    ${example_fwd}
+    PRIVATE
+        CK_TILE_FMHA_FWD_FAST_EXP2=1
+        CK_TILE_FMHA_FWD_SPLITKV_API=0
+        CK_TILE_FMHA_FWD_APPENDKV_API=0
+        CK_TILE_FMHA_FWD_PAGEDKV_API=0
+        CK_TILE_FMHA_FWD_BATCH_PREFILL_API=0
+)
 target_compile_features(${example_fwd} PRIVATE cxx_std_17 hip_std_17)
 target_compile_options(${example_fwd} PRIVATE -Wno-float-equal
                                               -Wno-undefined-func-template
@@ -114,19 +119,21 @@ target_include_directories(${example_fwd} PRIVATE ${CMAKE_CURRENT_LIST_DIR})
 target_link_libraries(${example_fwd} PRIVATE composable_kernel::ck_tile hip::host)
 target_sources(${example_fwd} PRIVATE fmha_fwd.cpp ${EXAMPLE_FWD_GEN_SRCS})
 
-add_test(${example_fwd} COMMAND "${CMAKE_CURRENT_LIST_DIR}/${example_fwd_script}" "$<TARGET_FILE:${example_fwd}>")
+add_test(NAME ${example_fwd} COMMAND "${CMAKE_CURRENT_LIST_DIR}/${example_fwd_script}" "$<TARGET_FILE:${example_fwd}>")
 
 ########################################################################################################################
 # BWD
 ########################################################################################################################
 add_executable(${example_bwd})
 
+# The backward pass needs three kernels per configuration, and --filter takes one pattern for each, separated by
+# "@": dot_do_o@convert_dq@dq_dk_dv. A configuration is generated only if all three kernel names match.
 execute_process(
     COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/generate.py
         --api bwd
         --targets ${GENERATOR_TARGETS}
         --optdim 64
-        --filter *fp16_group*
+        --filter *fp16*group*@*fp16*group*@*fp16*group*
         --receipt 3
         --list_blobs ${CMAKE_CURRENT_BINARY_DIR}/bwd_src_list.txt
     RESULT_VARIABLE ret
@@ -139,7 +146,7 @@ add_custom_target(fmha_bwd_gen_srcs
         --api bwd
         --targets ${GENERATOR_TARGETS}
         --optdim 64
-        --filter *fp16_group*
+        --filter *fp16*group*@*fp16*group*@*fp16*group*
         --receipt 3
         --output_dir ${CMAKE_CURRENT_BINARY_DIR}
     DEPENDS "${CMAKE_CURRENT_BINARY_DIR}/bwd_src_list.txt" ${EXAMPLE_CODE_GEN_SCRIPTS}
@@ -151,7 +158,7 @@ add_custom_target(fmha_bwd_gen_srcs
 add_dependencies(${example_bwd} fmha_bwd_gen_srcs)
 target_compile_definitions(${example_bwd} PRIVATE CK_TILE_FLOAT_TO_BFLOAT16_DEFAULT=3)
 target_compile_features(${example_bwd} PRIVATE cxx_std_17 hip_std_17)
-target_compile_options(${example_fwd} PRIVATE -Wno-float-equal
+target_compile_options(${example_bwd} PRIVATE -Wno-float-equal
                                               -Wno-undefined-func-template
                                               -fgpu-flush-denormals-to-zero
 )
@@ -162,4 +169,4 @@ target_sources(${example_bwd} PRIVATE fmha_bwd.cpp ${EXAMPLE_BWD_GEN_SRCS})
 set_source_files_properties(fmha_bwd.cpp ${EXAMPLE_BWD_GEN_SRCS}
                             PROPERTIES LANGUAGE ${ROCM_EXAMPLES_GPU_LANGUAGE})
 
-add_test(${example_bwd} COMMAND "${CMAKE_CURRENT_LIST_DIR}/${example_bwd_script}" "$<TARGET_FILE:${example_bwd}>")
+add_test(NAME ${example_bwd} COMMAND "${CMAKE_CURRENT_LIST_DIR}/${example_bwd_script}" "$<TARGET_FILE:${example_bwd}>")
@@ -69,7 +69,6 @@ For more information on CK Tile terminology, refer to the
 * `ck_tile::bf8_t`
 * `ck_tile::BlockAttentionBiasEnum`
 * `ck_tile::BlockFmhaPipelineEnum`
-* `ck_tile::composes`
 * `ck_tile::DeviceMem`
 * `ck_tile::FillNormalDistribution`
 * `ck_tile::FillNormalDistributionIntegerValue`
@@ -98,6 +97,7 @@ For more information on CK Tile terminology, refer to the
 * `ck_tile::get_alibi_slopes`
 * `ck_tile::integer_divide_ceil`
 * `ck_tile::make_alibi_from_lr_mask`
+* `ck_tile::make_composes`
 * `ck_tile::make_generic_attention_mask_coordinates_from_lr_window`
 * `ck_tile::make_tuple`
 * `ck_tile::numeric<T>::max`

@@ -1,32 +1,15 @@
-// MIT License
-//
-// Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
-//
-// Permission is hereby granted, free of charge, to any person obtaining a copy
-// of this software and associated documentation files (the "Software"), to deal
-// in the Software without restriction, including without limitation the rights
-// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-// copies of the Software, and to permit persons to whom the Software is
-// furnished to do so, subject to the following conditions:
-//
-// The above copyright notice and this permission notice shall be included in
-// all copies or substantial portions of the Software.
-//
-// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
-// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
-// SOFTWARE.
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
 
 #pragma once
 
-#include "ck_tile/core.hpp"
-#include "ck_tile/ops/fmha.hpp"
 #include <ostream>
+
 #include <string>
 
+#include "ck_tile/core.hpp"
+#include "ck_tile/ops/fmha.hpp"
+
 // keep sync with BlockAttentionBiasEnum
 enum class bias_enum {
   no_bias = 0,
@@ -102,7 +85,8 @@ struct bias_info {
     return info;
   }
 
-  friend std::ostream &operator<<(std::ostream &os, const bias_info &bi) {
+  friend std::ostream &operator<<([[clang::lifetimebound]] std::ostream &os,
+                                  const bias_info &bi) {
     bi.serialize(os);
     return os;
   }

@@ -0,0 +1,2 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
@@ -1,5 +1,5 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 # SPDX-License-Identifier: MIT
-# Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
 
 from dataclasses import dataclass, field
 from typing import Any, List, Callable

@@ -1,5 +1,4 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 # SPDX-License-Identifier: MIT
-# Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
 # generate kernel instances to speed up compilation
-
 GEN_DIR = ""  # in Cmake, have to generate files in same folder
@@ -1,7 +1,6 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 # SPDX-License-Identifier: MIT
-# Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 # generate kernel instances to speed up compilation
-
 FWD_DTYPE_MAP = {
     "fp32": "FmhaFwdFp32",
     "fp16": "FmhaFwdFp16",
@@ -10,6 +9,8 @@
     "fp8fp16": "FmhaFwdFp8Fp16",
     "fp8bf16": "FmhaFwdFp8Bf16",
     "fp8fp32": "FmhaFwdFp8Fp32",
+    "mxfp8": "FmhaFwdMxFp8",
+    "mxfp4": "FmhaFwdMxFp4",
 }
 
 BWD_DTYPE_MAP = {"fp32": "FmhaBwdFp32", "fp16": "FmhaBwdFp16", "bf16": "FmhaBwdBf16"}
@@ -31,16 +32,24 @@
 }
 
 
-def get_mask_map(mask: str):
-    if mask == "generic":
+def get_mask_map(mask_impl: str):
+    if mask_impl == "generic":
         return _MASK_MAP
-    elif mask == "simplified":
+    elif mask_impl == "simplified":
         return _MASK_SIMPLIFIED_MAP
     else:
         assert False
         return None
 
 
+def get_mask_impl(mask: str) -> str:
+    return "simplified" if mask.startswith("s_") else "generic"
+
+
+def get_mask_cpp_type(mask: str) -> str:
+    return get_mask_map(get_mask_impl(mask))[mask]
+
+
 _MASK_CHECK_MAP = {
     "no": "t.mask_type == mask_enum::no_mask",
     "causal": "t.mask_type == mask_enum::mask_top_left || t.mask_type == mask_enum::mask_bottom_right",
@@ -63,6 +72,26 @@ def get_mask_check_map(mask: str):
         return None
 
 
+def get_mask_cpp_check_expr(mask: str) -> str:
+    return get_mask_check_map(get_mask_impl(mask))[mask]
+
+
+QSCALE_MAP = {
+    "no": "ck_tile::BlockAttentionQuantScaleEnum::NO_SCALE",
+    "pertensor": "ck_tile::BlockAttentionQuantScaleEnum::PERTENSOR",
+    "blockscale": "ck_tile::BlockAttentionQuantScaleEnum::BLOCKSCALE",
+    "kv_blockscale": "ck_tile::BlockAttentionQuantScaleEnum::KV_BLOCKSCALE",
+    "mx": "ck_tile::BlockAttentionQuantScaleEnum::MX",
+}
+
+QSCALE_CHECK_MAP = {
+    "no": "quant_scale_enum::no_scale",
+    "pertensor": "quant_scale_enum::pertensor",
+    "blockscale": "quant_scale_enum::blockscale",
+    "kv_blockscale": "quant_scale_enum::kv_blockscale",
+    "mx": "quant_scale_enum::mx",
+}
+
 BIAS_MAP = {
     "no": "ck_tile::BlockAttentionBiasEnum::NO_BIAS",
     "bias": "ck_tile::BlockAttentionBiasEnum::ELEMENTWISE_BIAS",
@@ -110,18 +139,24 @@ def get_mask_check_map(mask: str):
 
 PIPELINE_MAP = {
     "qr": "ck_tile::BlockFmhaPipelineQRKSVS",
+    "qr_hpad": "ck_tile::BlockFmhaPipelineQRKSVSHpad",
     "qr_async": "ck_tile::BlockFmhaPipelineQRKSVSAsync",
     "qs": "ck_tile::BlockFmhaPipelineQSKSVS",
     "qr_async_trload": "ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload",
+    "qr_async_trload_v3": "ck_tile::BlockFmhaFwdV3Pipeline",
+    "qr_tdm": "ck_tile::BlockFmhaPipelineQRKSVSTdm",
 }
 
 PIPELINE_ENUM_MAP = {
     "qr": "ck_tile::BlockFmhaPipelineEnum::QRKSVS",
+    "qr_hpad": "ck_tile::BlockFmhaPipelineEnum::QRKSVS_HPAD",
     "qr_async": "ck_tile::BlockFmhaPipelineEnum::QRKSVS_ASYNC",
     "qr_nwarp_sshuffle": "ck_tile::BlockFmhaPipelineEnum::QRKSVS",
     "qs": "ck_tile::BlockFmhaPipelineEnum::QSKSVS",
     "qr_pagedkv": "ck_tile::BlockFmhaPipelineEnum::QRKSVS",
     "qr_async_trload": "ck_tile::BlockFmhaPipelineEnum::QRKSVS_ASYNC_TRLOAD",
+    "qr_async_trload_v3": "ck_tile::BlockFmhaPipelineEnum::QRKSVS_ASYNC_TRLOAD_V3",
+    "qr_tdm": "ck_tile::BlockFmhaPipelineEnum::QRKSVS_TDM",
 }
 
 BOOL_MAP = {

@@ -0,0 +1,2 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
		# SPDX-License-Identifier: MIT