From 612bedd3ab652e2e77c4ece6ad25bec305740bee Mon Sep 17 00:00:00 2001 From: Gino Lu Date: Fri, 21 Nov 2025 02:48:13 -0600 Subject: [PATCH] test if CI failed on hdim=160 --- example/ck_tile/01_fmha/CMakeLists.txt | 4 ++-- example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py | 4 ++-- .../ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py | 4 ++-- example/ck_tile/01_fmha/script/smoke_test_fwd.sh | 10 +++++----- 4 files changed, 11 insertions(+), 11 deletions(-) diff --git a/example/ck_tile/01_fmha/CMakeLists.txt b/example/ck_tile/01_fmha/CMakeLists.txt index ce914b92af..f11ed9b496 100644 --- a/example/ck_tile/01_fmha/CMakeLists.txt +++ b/example/ck_tile/01_fmha/CMakeLists.txt @@ -44,7 +44,7 @@ set(FMHA_FWD_CODE_GEN_COMMON_ARGS ${CMAKE_CURRENT_LIST_DIR}/generate.py --targets ${FMHA_TARGETS_ARG} --api ${FMHA_FWD_APIS} - --optdim 32,64,128,256 + --optdim 32,64,128,160,256 # --filter fmha_fwd... ) set(FMHA_BWD_CODE_GEN_COMMON_ARGS @@ -52,7 +52,7 @@ set(FMHA_BWD_CODE_GEN_COMMON_ARGS --targets ${FMHA_TARGETS_ARG} --api bwd --receipt 3 - --optdim 32,64,96,128,256 + --optdim 32,64,96,128,160,256 # --filter fmha_bwd_dot...@fmha_bwd_convert...@fmha_bwd... ) diff --git a/example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py b/example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py index 2acc467410..f9a06f0c08 100644 --- a/example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py +++ b/example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py @@ -30,7 +30,7 @@ from codegen.utils import check_duplicates_and_paddings, if_, indent, update_fil DTYPE_BITS = {"fp32": 32, "fp16": 16, "bf16": 16, "fp8": 8, "bf8": 8} -K0_MAX_SUBMAX_MAP = {32: 32, 48: 48, 64: 64, 96: 128, 128: 128, 192: 192, 256: 256} +K0_MAX_SUBMAX_MAP = {32: 32, 48: 48, 64: 64, 96: 128, 128: 128, 160: 160, 192: 192, 256: 256} FMHA_FWD_KERNEL_HEADER = """// SPDX-License-Identifier: MIT // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.\n @@ -667,7 +667,7 @@ class KernelComponentFactoryGfx9: FmhaFwdTileSize( 32, 32, 128, 128, 32, 128, 1, 1, 1, 1, 1, 1, 32, 32, 16, 32, 32, 16, -1), FmhaFwdTileSize(128, 64, 32, 128, 16, 128, 4, 1, 1, 4, 1, 1, 32, 32, 16, 32, 32, 16, -1), FmhaFwdTileSize(128, 128, 32, 128, 32, 128, 4, 1, 1, 4, 1, 1, 32, 32, 16, 32, 32, 16, -1)], - # (160, 160) : [FmhaFwdTileSize(128, 128 , 32, 160, 32, 160, 4, 1, 1, 4, 1, 1, 32, 32, 16, 32, 32, 16, 1)], + (160, 160) : [FmhaFwdTileSize(128, 128 , 32, 160, 32, 160, 4, 1, 1, 4, 1, 1, 32, 32, 16, 32, 32, 16, 1)], (192, 128) : [FmhaFwdTileSize(128, 128, 32, 128, 32, 192, 4, 1, 1, 4, 1, 1, 32, 32, 16, 32, 32, 16, -1)], (192, 192) : [FmhaFwdTileSize(128, 128, 32, 192, 32, 192, 4, 1, 1, 4, 1, 1, 32, 32, 16, 32, 32, 16, 1)], (256, 256) : [FmhaFwdTileSize(128, 128, 32, 256, 32, 256, 4, 1, 1, 4, 1, 1, 32, 32, 16, 32, 32, 16, -1)], diff --git a/example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py b/example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py index 85c25561ea..7a1bf42c35 100644 --- a/example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py +++ b/example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py @@ -788,7 +788,7 @@ class KernelComponentFactoryBase: "64": FmhaFwdSplitKVCombineTileSize(32, -1), "96": FmhaFwdSplitKVCombineTileSize(32, -1), "128": FmhaFwdSplitKVCombineTileSize(32, -1), - # "160" : FmhaFwdSplitKVCombineTileSize(32, -1), + "160" : FmhaFwdSplitKVCombineTileSize(32, -1), "256": FmhaFwdSplitKVCombineTileSize(32, -1), } elif dtype in ["fp8", "bf8"]: @@ -812,7 +812,7 @@ class KernelComponentFactoryGfx9(KernelComponentFactoryBase): "64" : FmhaFwdTileSize( 64, 64, 32, 64, 32, 64, 4, 1, 1, 4, 1, 1, 16, 16, 16, 16, 16, 16, -1), "96" : FmhaFwdTileSize( 64, 128, 32, 128, 32, 96, 4, 1, 1, 4, 1, 1, 16, 16, 16, 16, 16, 16, -1), "128": FmhaFwdTileSize( 64, 128, 32, 128, 32, 128, 4, 1, 1, 4, 1, 1, 16, 16, 16, 16, 16, 16, -1), - # "160" : FmhaFwdTileSize(64, 128, 32, 160, 32, 160, 4, 1, 1, 4, 1, 1, 16, 16, 16, 16, 16, 16, -1), + "160" : FmhaFwdTileSize(64, 128, 32, 160, 32, 160, 4, 1, 1, 4, 1, 1, 16, 16, 16, 16, 16, 16, -1), "256": FmhaFwdTileSize( 64, 128, 32, 256, 32, 256, 4, 1, 1, 4, 1, 1, 16, 16, 16, 16, 16, 16, -1), } # fmt: skip elif dtype in ["fp8", "bf8"]: diff --git a/example/ck_tile/01_fmha/script/smoke_test_fwd.sh b/example/ck_tile/01_fmha/script/smoke_test_fwd.sh index 02bc5476fa..0d7de64b3a 100755 --- a/example/ck_tile/01_fmha/script/smoke_test_fwd.sh +++ b/example/ck_tile/01_fmha/script/smoke_test_fwd.sh @@ -65,7 +65,7 @@ run_fp16_bf16_tests() { for prec in "fp16" "bf16" ; do for mode in 1 0 ; do for perm in 0 1 ; do - for hdim in 32 64 128 256 ; do + for hdim in 32 64 128 160 256 ; do for lse in 0 1 ; do for bias in "n" "e" "a" ; do for p_drop in 0.0 0.2 ; do @@ -92,7 +92,7 @@ run_fp8_tests() { for perm in 0 1 ; do for bias in "n" "e" "a" ; do for b in 1 2 ; do - for hdim in 64 128 256 ; do + for hdim in 64 128 160 256 ; do $EXE -prec=fp8 -init=0 -b=$b -h=1 -d=$hdim -s=128 -bias=$bias -iperm=$perm -operm=$perm -vlayout=r -squant=1 -kname=$KNAME $COMMON_ARGS @@ -103,7 +103,7 @@ run_fp8bf16_tests() { for perm in 0 1 ; do for bias in "n" "e" "a" ; do for b in 1 2 ; do - for hdim in 64 128 256 ; do + for hdim in 64 128 160 256 ; do $EXE -prec=fp8bf16 -init=0 -b=$b -h=1 -d=$hdim -s=128 -bias=$bias -iperm=$perm -operm=$perm -vlayout=r -squant=1 -kname=$KNAME $COMMON_ARGS @@ -114,7 +114,7 @@ run_fp8fp32_tests() { for perm in 0 1 ; do for bias in "n" "e" "a" ; do for b in 1 2 ; do - for hdim in 128 ; do + for hdim in 128 160; do $EXE -prec=fp8fp32 -init=0 -b=$b -h=1 -d=$hdim -s=128 -bias=$bias -iperm=$perm -operm=$perm -vlayout=r -squant=1 -kname=$KNAME $COMMON_ARGS @@ -125,7 +125,7 @@ run_fp16_appendkv_tests() { for s in $(seq 63 1 65) ; do for s_k in 65 129 ; do for s_knew in 0 64 $s_k ; do - for hdim in 32 64 128 256 ; do + for hdim in 32 64 128 160 256 ; do for ri in 0 1 ; do for rdim in 0 16 32 $hdim ; do for page_block_size in 0 128 ; do