Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
41 commits
Select commit Hold shift + click to select a range
ef93610
feat(ck_tile): add gfx1250 hardware named barrier producer/consumer r…
chris-tsiaousis-hpc Sep 14, 2026
21c6c5a
test(ck_tile): cover named barrier ring id allocation at compile time
chris-tsiaousis-hpc Sep 14, 2026
7aad6bb
test(ck_tile): add gfx1250-gated device test for the named barrier ring
chris-tsiaousis-hpc Sep 14, 2026
3ad9947
fix(ck_tile): wait for the TDM store in TdmEpilogue before returning
chris-tsiaousis-hpc Sep 24, 2026
12f6108
refactor(ck_tile): move the wavelet helpers into a shared header
chris-tsiaousis-hpc Sep 24, 2026
fad82d1
feat(ck_tile): add run() to drive a named barrier ring's whole K loop
chris-tsiaousis-hpc Sep 24, 2026
180d628
test(ck_tile): replay the named barrier ring protocol on the host
chris-tsiaousis-hpc Sep 24, 2026
3414816
test(ck_tile): cover named barrier ring run() on gfx1250
chris-tsiaousis-hpc Sep 24, 2026
037fd3c
feat(ck_tile): run wavelet and named barrier pipelines in UniversalGe…
chris-tsiaousis-hpc Sep 24, 2026
f58555a
feat(ck_tile): add the producer/consumer TDM GEMM pipeline
chris-tsiaousis-hpc Sep 24, 2026
5da5065
refactor(ck_tile): select TDM test behaviour by trait, not suite name
chris-tsiaousis-hpc Sep 24, 2026
a22c37d
test(ck_tile): cover the producer/consumer TDM GEMM pipeline
chris-tsiaousis-hpc Sep 24, 2026
18c4041
refactor(ck_tile): rename Lag to PublishLag and tidy the named barrie…
chris-tsiaousis-hpc Oct 7, 2026
40fcad7
test(ck_tile): tidy the named barrier device test
chris-tsiaousis-hpc Oct 7, 2026
8fa9d48
test(ck_tile): replay the GEMM test's named barrier ring on the host
chris-tsiaousis-hpc Oct 7, 2026
10a28aa
refactor(ck_tile): arm the named barriers once per launch, not once p…
chris-tsiaousis-hpc Oct 7, 2026
ff01e95
refactor(ck_tile): spell out the producer/consumer pipeline assumptions
chris-tsiaousis-hpc Oct 7, 2026
850c6a4
test(ck_tile): add the producer/consumer GEMM test to ck_tile_gemm_tests
chris-tsiaousis-hpc Oct 7, 2026
65ee20f
fix(ck_tile): restore the release/acquire fences in the named barriers
chris-tsiaousis-hpc Oct 9, 2026
496dc3d
fix(ck_tile): validate memory safety and split-K before the launch op…
chris-tsiaousis-hpc Oct 9, 2026
14dc151
fix(ck_tile): clear the whole split-K workspace between timed launches
chris-tsiaousis-hpc Oct 9, 2026
94a8a22
fix(ck_tile): accept gfx1250-strict in is_gfx125_supported
chris-tsiaousis-hpc Oct 9, 2026
4ee4ab6
refactor(ck_tile): generalise producer dispatch and assert PC pipelin…
chris-tsiaousis-hpc Oct 9, 2026
cebe5f0
test(ck_tile): add host-side argument validation and widen TDM coverage
chris-tsiaousis-hpc Oct 9, 2026
b7eedc3
fix(ck_tile): keep the producer/consumer gtest in the gfx1250 smoke lane
chris-tsiaousis-hpc Oct 9, 2026
2b23345
fix(ck_tile): split trait combos by name in the gemm benchmark parser
chris-tsiaousis-hpc Oct 9, 2026
444e1b1
feat(ck_tile): generate comp_tdm_pc instances in the gfx1250 tile engine
chris-tsiaousis-hpc Oct 9, 2026
30374fc
fix(ck_tile): Make the 4 slot PC pipeline the default
chris-tsiaousis-hpc Oct 19, 2026
7edfbc8
Merge branch 'develop' into users/chris-tsiaousis-hpc/ck/producer-con…
illsilin Oct 19, 2026
7e0ebb6
feat(ck-tile): comp_tdm_pc producer/consumer GEMM TE to dispatcher br…
ozturkosu Oct 19, 2026
0b57de0
Merge remote-tracking branch 'origin/develop' into work/pr13975-fix
ozturkosu Oct 20, 2026
90e54e1
Merge branch 'develop' into users/muozturk/ck/comp-tdm-pc-dispatcher-…
ozturkosu Oct 20, 2026
1c8168e
Merge branch 'develop' into users/muozturk/ck/comp-tdm-pc-dispatcher-…
ozturkosu Oct 20, 2026
4655280
fix(ck-tile): address Copilot review of the comp_tdm_pc bridge
ozturkosu Oct 21, 2026
7abbc28
Merge develop into comp_tdm_pc dispatcher bridge
ozturkosu Oct 22, 2026
e23a561
Merge branch 'develop' into users/muozturk/ck/comp-tdm-pc-dispatcher-…
ozturkosu Oct 22, 2026
97da23d
test(ck-tile): gate the comp_tdm_pc split-K argument test on gfx1250
ozturkosu Oct 22, 2026
75789c0
Merge remote-tracking branch 'origin/develop' into users/muozturk/ck/…
ozturkosu Oct 30, 2026
02a5244
style(ck-tile): format kernel_key.hpp Pipeline enum with clang-format-18
ozturkosu Oct 30, 2026
8e02a1a
Merge remote-tracking branch 'origin/develop' into users/muozturk/ck/…
ozturkosu Oct 30, 2026
90ce0fe
Merge branch 'develop' into users/muozturk/ck/comp-tdm-pc-dispatcher-…
ozturkosu Oct 30, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
42 changes: 30 additions & 12 deletions projects/composablekernel/dispatcher/codegen/arch_filter.py
Original file line number Diff line number Diff line change
Expand Up @@ -50,6 +50,8 @@
from codegen_common import (
CommonTypeMappings,
gfx1250_fp32_tile_reject_reason,
gfx1250_lds_operand_bytes,
gfx1250_lds_operand_transposed,
gfx1250_pipeline_reject_reason,
)

Expand Down Expand Up @@ -278,6 +280,8 @@ class OperatorType(Enum):
"preshuffle_tdm": 32768,
# Mandatory double buffering (num_lds_buffers = 2), so half the budget.
"comp_async": 32768,
# Ring of four LDS slots (producer/consumer TDM), so a quarter.
"comp_tdm_pc": 16384,
"wavelet": 65536,
"default": 65536,
}
Expand Down Expand Up @@ -314,6 +318,9 @@ def get_lds_limit(
("comp_tdm_v2", "tdm", "interwave"),
("comp_tdm_v2", "cshuffle", "interwave"),
("comp_tdm_v2", "default", "interwave"),
("comp_tdm_pc", "tdm", "interwave"),
("comp_tdm_pc", "cshuffle", "interwave"),
("comp_tdm_pc", "default", "interwave"),
}

DTYPE_COMBINATIONS = {
Expand All @@ -331,7 +338,7 @@ def get_lds_limit(

# Tensor Data Mover pipelines. They exist only on gfx1250; elsewhere the TDM
# instructions are no-ops and the kernel silently produces zeros.
TDM_PIPELINES = ("comp_tdm", "comp_tdm_v2")
TDM_PIPELINES = ("comp_tdm", "comp_tdm_v2", "comp_tdm_pc")
Comment thread
ozturkosu marked this conversation as resolved.
Resolved
TDM_ARCH = "gfx1250"


Expand Down Expand Up @@ -856,16 +863,18 @@ def _validate_trait_combo(self, config: KernelConfig, result: ValidationResult):
def _validate_gfx1250_pipeline(
self, config: KernelConfig, result: ValidationResult
):
"""Validate the gfx1250 pipelines (comp_async, comp_tdm, comp_tdm_v2)
"""Validate the gfx1250 pipelines (comp_async, comp_tdm*)
and the TDM epilogue.

The rules come from codegen_common.gfx1250_pipeline_reject_reason, the
single source of truth shared with unified_gemm_codegen and
python/gemm_utils: TDM pipelines are gfx1250-only (off gfx1250 the TDM
instructions compile to no-ops and the kernel silently writes zeros),
need the TDM epilogue (and vice versa), intrawave, unpadded tiles and,
for comp_tdm_v2, exactly four waves; only the plain and preshuffle GEMM
operators can use them. Non-MX comp_async GEMM on gfx1250 needs the cshuffle
for comp_tdm_v2, exactly four waves, for comp_tdm_pc, at most 256 fp32
accumulators per lane; only the plain and preshuffle GEMM operators can
use them (comp_tdm_pc is plain only). Non-MX comp_async GEMM on gfx1250
needs the cshuffle
epilogue, an rc A/B layout, pad_m=pad_n=pad_k=True and, for fp8/bf8,
warp_tile_k >= 128 and a 16x16 warp tile (32x32 does not compile).

Expand Down Expand Up @@ -902,6 +911,8 @@ def _validate_gfx1250_pipeline(
variant_supported=is_gemm,
variant_name=f"operator={config.operator.value}",
pads=None if None in pads else tuple(bool(p) for p in pads),
tile_mn=(config.tile_m, config.tile_n),
warp_k=config.warp_k,
)
if reason:
result.add_error(reason)
Expand All @@ -928,14 +939,21 @@ def _validate_lds_capacity(self, config: KernelConfig, result: ValidationResult)
and config.operator in (OperatorType.GEMM, OperatorType.GEMM_PRESHUFFLE)
)
if uses_gfx1250_base_lds:
# TDM and non-MX comp_async use the gfx1250 non-transposed base LDS
# descriptors: each group of rows spanning at least 256 bytes is
# followed by 16 padding bytes (same accounting as the Tile Engine
# validate_lds_capacity).
a_lds_layer = max(1, 256 // (config.tile_k * elem_size_a))
b_lds_layer = max(1, 256 // (config.tile_k * elem_size_b))
matrix_a_size += max(0, config.tile_m // a_lds_layer - 1) * 16
matrix_b_size += max(0, config.tile_n // b_lds_layer - 1) * 16
# TDM and non-MX comp_async use the gfx1250 base LDS descriptors,
# whose padding depends on whether the operand is transpose-loaded
# (same accounting as the Tile Engine validate_lds_capacity).
matrix_a_size = gfx1250_lds_operand_bytes(
config.tile_m,
config.tile_k,
elem_size_a,
gfx1250_lds_operand_transposed("a", config.datatype_a, config.layout),
)
matrix_b_size = gfx1250_lds_operand_bytes(
config.tile_n,
config.tile_k,
elem_size_b_staged,
gfx1250_lds_operand_transposed("b", config.datatype_b, config.layout),
)
total_lds = matrix_a_size + matrix_b_size

# The budget depends on the target, not just the pipeline: a tile that
Expand Down
16 changes: 15 additions & 1 deletion projects/composablekernel/dispatcher/codegen/arch_specs.json
Original file line number Diff line number Diff line change
Expand Up @@ -348,6 +348,17 @@
"pipeline_lds_budget_aliases."
]
},
"comp_tdm_pc": {
"basis": "fraction",
"value": 0.25,
"_reason": [
"Ring of NumSlots LDS buffers, no configuration choice:",
"GemmPipelineAgBgCrCompTDMProducerConsumer::GetSmemSize() returns",
"NumSlots * Policy::GetSmemSize() with the policy default NumSlots_ = 4.",
"Emitted to C++ as Pipeline::CompTDMPC. Matches the Tile Engine",
"PIPELINE_SMEM_SLOTS entry; raise both if NumSlots_ changes."
]
},
"wavelet": {
"basis": "fraction",
"value": 1.0,
Expand Down Expand Up @@ -410,7 +421,10 @@
["comp_tdm", "tdm", "interwave"],
["comp_tdm_v2", "cshuffle", "interwave"],
["comp_tdm_v2", "default", "interwave"],
["comp_tdm_v2", "tdm", "interwave"]
["comp_tdm_v2", "tdm", "interwave"],
["comp_tdm_pc", "cshuffle", "interwave"],
["comp_tdm_pc", "default", "interwave"],
["comp_tdm_pc", "tdm", "interwave"]
]
},

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
AUTO-GENERATED FILE - DO NOT EDIT DIRECTLY!

Generated from: arch_specs.json
Generated at: 2026-10-25T20:57:27.671022
Generated at: 2026-10-30T14:59:00.519447

To update this file:
1. Edit arch_specs.json
Expand Down Expand Up @@ -162,6 +162,7 @@
"preshufflev2": 32768,
"preshuffle_tdm": 32768,
"comp_async": 32768,
"comp_tdm_pc": 16384,
"wavelet": 65536,
"compv6": 32768,
"preshufflev1": 32768,
Expand All @@ -178,6 +179,7 @@
"preshufflev2": 32768,
"preshuffle_tdm": 32768,
"comp_async": 32768,
"comp_tdm_pc": 16384,
"wavelet": 65536,
"compv6": 32768,
"preshufflev1": 32768,
Expand All @@ -194,6 +196,7 @@
"preshufflev2": 32768,
"preshuffle_tdm": 32768,
"comp_async": 32768,
"comp_tdm_pc": 16384,
"wavelet": 65536,
"compv6": 32768,
"preshufflev1": 32768,
Expand All @@ -210,6 +213,7 @@
"preshufflev2": 81920,
"preshuffle_tdm": 81920,
"comp_async": 81920,
"comp_tdm_pc": 40960,
"wavelet": 163840,
"compv6": 81920,
"preshufflev1": 81920,
Expand All @@ -226,6 +230,7 @@
"preshufflev2": 32768,
"preshuffle_tdm": 32768,
"comp_async": 32768,
"comp_tdm_pc": 16384,
"wavelet": 65536,
"compv6": 32768,
"preshufflev1": 32768,
Expand All @@ -242,6 +247,7 @@
"preshufflev2": 32768,
"preshuffle_tdm": 32768,
"comp_async": 32768,
"comp_tdm_pc": 16384,
"wavelet": 65536,
"compv6": 32768,
"preshufflev1": 32768,
Expand All @@ -258,6 +264,7 @@
"preshufflev2": 32768,
"preshuffle_tdm": 32768,
"comp_async": 32768,
"comp_tdm_pc": 16384,
"wavelet": 65536,
"compv6": 32768,
"preshufflev1": 32768,
Expand All @@ -274,6 +281,7 @@
"preshufflev2": 163840,
"preshuffle_tdm": 163840,
"comp_async": 163840,
"comp_tdm_pc": 81920,
"wavelet": 327680,
"compv6": 163840,
"preshufflev1": 163840,
Expand Down Expand Up @@ -319,6 +327,9 @@
("comp_tdm_v2", "cshuffle", "interwave"),
("comp_tdm_v2", "default", "interwave"),
("comp_tdm_v2", "tdm", "interwave"),
("comp_tdm_pc", "cshuffle", "interwave"),
("comp_tdm_pc", "default", "interwave"),
("comp_tdm_pc", "tdm", "interwave"),
}

# Valid dtype combinations: (A_dtype, B_dtype) -> acc_dtype and notes
Expand Down
Loading
Loading