From 90b950b492f00990b175c99a453543830c557aa2 Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Mon, 12 Jan 2026 09:42:03 +0000 Subject: [PATCH 01/15] Rename the parameters of load_interleaved_pk_type and load_and_convert_tile --- .../ck_tile/ops/common/load_and_convert_tile.hpp | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/projects/composablekernel/include/ck_tile/ops/common/load_and_convert_tile.hpp b/projects/composablekernel/include/ck_tile/ops/common/load_and_convert_tile.hpp index 0748c5fb49e8..9fb26e7638f8 100644 --- a/projects/composablekernel/include/ck_tile/ops/common/load_and_convert_tile.hpp +++ b/projects/composablekernel/include/ck_tile/ops/common/load_and_convert_tile.hpp @@ -14,11 +14,11 @@ template struct ConverterLoader { template - CK_TILE_DEVICE static void load_interleaved_pk_type(WarpTile& dst, const WarpWindow& src) + CK_TILE_DEVICE static void load_interleaved_pk_type(WarpTile& dst, const WarpWindow& src_window) { static_assert(WarpTile::get_thread_buffer_size() % UnaryOpSize == 0); constexpr index_t thread_buffer_size = WarpTile::get_thread_buffer_size() / UnaryOpSize; - const auto tmp = load_tile(src); + const auto src = load_tile(src_window); // NOTE: we rely on types packing neatly here using RawSrcType = typename SrcDataType::type; @@ -30,13 +30,13 @@ struct ConverterLoader const element_wise::PassThroughPack8 elementwise_op{}; elementwise_op(dst.get_thread_buffer().template get_as()(i), - tmp.get_thread_buffer().template get_as()[i]); + src.get_thread_buffer().template get_as()[i]); }); } }; template -CK_TILE_DEVICE void load_and_convert_tile(WarpTile& dst, const WarpWindow& src) +CK_TILE_DEVICE void load_and_convert_tile(WarpTile& dst, const WarpWindow& src_window) { using SrcDataType = typename WarpWindow::Base::DataType; using DstDataType = typename WarpTile::DataType; @@ -44,15 +44,16 @@ CK_TILE_DEVICE void load_and_convert_tile(WarpTile& dst, const WarpWindow& src) if constexpr(is_packed_type_v && !is_packed_type_v) { static_assert(!LoadTranspose, "LoadTranspose not supported with pk_int4_t or pk_fp4_t"); - ConverterLoader::load_interleaved_pk_type(dst, src); + ConverterLoader::load_interleaved_pk_type( + dst, src_window); } else if constexpr(LoadTranspose) { - load_tile_transpose(dst, src); + load_tile_transpose(dst, src_window); } else { - load_tile(dst, src); + load_tile(dst, src_window); } } From 999b4936dcd9becc2831f37b8337e62f72eff4eb Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Mon, 26 Jan 2026 09:26:59 +0000 Subject: [PATCH 02/15] Add load_tile_transpose_convert for mixed precision transpose loading --- .../core/tensor/load_tile_transpose.hpp | 116 ++++++++++++++++++ .../unary_element_wise_operation.hpp | 38 ++++++ 2 files changed, 154 insertions(+) diff --git a/projects/composablekernel/include/ck_tile/core/tensor/load_tile_transpose.hpp b/projects/composablekernel/include/ck_tile/core/tensor/load_tile_transpose.hpp index 5f73d4934a1f..82c49261b2aa 100644 --- a/projects/composablekernel/include/ck_tile/core/tensor/load_tile_transpose.hpp +++ b/projects/composablekernel/include/ck_tile/core/tensor/load_tile_transpose.hpp @@ -14,6 +14,7 @@ #include "ck_tile/core/container/statically_indexed_array.hpp" #include "ck_tile/core/numeric/math.hpp" #include "ck_tile/core/utility/type_traits.hpp" +#include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp" namespace ck_tile { @@ -529,4 +530,119 @@ load_tile_transpose(const tile_window_with_static_distribution, + typename = std::enable_if_t::distr_encoding_valid, + Policy>> +CK_TILE_DEVICE void load_tile_transpose_convert_with_offset( + DistributedTensor_& out_tensor, + const tile_window_with_static_distribution& __restrict__ tile_window, + const index_t offset, + number = {}) +{ + using SrcDataType = typename BottomTensorView_::DataType; + using DstDataType = typename DistributedTensor_::DataType; + + auto trans_tensor = tile_window.template load_transpose_with_offset(offset); + constexpr auto input_distr = TileDistribution_{}; + constexpr auto output_distr = typename DistributedTensor_::StaticTileDistribution{}; + + constexpr auto y_in_desc = input_distr.get_ys_to_d_descriptor(); + constexpr auto y_out_desc = output_distr.get_ys_to_d_descriptor(); + + constexpr auto y_in_lengths = to_sequence(y_in_desc.get_lengths()); + constexpr auto y_out_lengths = to_sequence(y_out_desc.get_lengths()); + + constexpr auto y_in_element_space_size = y_in_desc.get_element_space_size(); + constexpr auto y_out_element_space_size = y_out_desc.get_element_space_size(); + + // For mixed precision: element space size must be the same (total bytes match) + static_assert(y_in_element_space_size == y_out_element_space_size, + "For mixed precision transpose, input and output element space size must match!"); + + // Ensure total element counts are consistent and divisible by the input vector length. + constexpr index_t total_elems_in = reduce_on_sequence(y_in_lengths, multiplies<>{}, number<1>{}); + constexpr index_t total_elems_out = reduce_on_sequence(y_out_lengths, multiplies<>{}, number<1>{}); + static_assert(total_elems_in == total_elems_out, + "For mixed precision transpose, input/output element counts must match!"); + static_assert(total_elems_in % number{} == 0, + "Input vector length must evenly divide total elements."); + + constexpr index_t num_of_access = total_elems_in / number{}; + + // Read as input type, convert to output type + using SrcDataVec = ext_vector_t{}>; + using DstDataVec = ext_vector_t{}>; + static_for<0, num_of_access, 1>{}([&](auto i) { + static_assert(number{} == 8, "Only PassThroughPack8 is supported for now."); + const element_wise::PassThroughPack8 elementwise_op{}; + + elementwise_op( + out_tensor.get_thread_buffer().template get_as()(i), + trans_tensor.get_thread_buffer().template get_as()[i] + ); + }); +} + +/** + * @brief Mixed-precision transpose load with zero offset. + * + * Convenience wrapper for load_tile_transpose_convert_with_offset with offset=0. + */ +template < + typename DistributedTensor_, + typename BottomTensorView_, + typename WindowLengths_, + typename TileDistribution_, + index_t NumCoord_, + index_t UnaryOpSize_, + typename Policy = DefaultTranspose, + typename = std::enable_if_t::distr_encoding_valid, + Policy>> +CK_TILE_DEVICE void +load_tile_transpose_convert( + DistributedTensor_& out_tensor, + const tile_window_with_static_distribution& __restrict__ tile_window, + number = {}) +{ + load_tile_transpose_convert_with_offset(out_tensor, tile_window, 0, number{}); +} + } // namespace ck_tile diff --git a/projects/composablekernel/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp b/projects/composablekernel/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp index 4ad699629c02..eea1eb6acd62 100644 --- a/projects/composablekernel/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp +++ b/projects/composablekernel/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp @@ -447,6 +447,15 @@ CK_TILE_HOST_DEVICE bf16x8_t fp8x8_to_bf16x8_scale(const fp8x8_t& src, const flo return y; } +CK_TILE_HOST_DEVICE fp8x8_t bf16x8_to_fp8x8_scale(const bf16x8_t& src, const float& scale) +{ + fp8x8_t y; + static_for<0, 8, 1>{}([&](auto i) { + y[i.value] = type_convert(type_convert(src[i.value]) * scale); + }); + return y; +} + CK_TILE_HOST_DEVICE fp16x8_t fp8x8_to_fp16x8_scale(const fp8x8_t& src, const float& scale) { fp16x8_t y; @@ -491,6 +500,15 @@ CK_TILE_HOST_DEVICE fp16x8_t fp8x8_to_fp16x8_scale(const fp8x8_t& src, const flo return y; } +CK_TILE_HOST_DEVICE fp8x8_t fp16x8_to_fp8x8_scale(const fp16x8_t& src, const float& scale) +{ + fp8x8_t y; + static_for<0, 8, 1>{}([&](auto i) { + y[i.value] = type_convert(type_convert(src[i.value]) * scale); + }); + return y; +} + CK_TILE_HOST_DEVICE fp16x8_t bf8x8_to_fp16x8_scale(const bf8x8_t& src, const float& scale) { fp16x8_t y; @@ -620,12 +638,32 @@ struct PassThroughPack8 template CK_TILE_HOST_DEVICE void operator()(Y& y, const X& x) const; + CK_TILE_HOST_DEVICE constexpr void operator()(fp16x8_t& y, const fp8x8_t& x) const + { + y = fp8x8_to_fp16x8_scale(x, 1.0f); + } + + CK_TILE_HOST_DEVICE constexpr void operator()(fp8x8_t& y, const fp16x8_t& x) const + { + y = fp16x8_to_fp8x8_scale(x, 1.0f); + } + CK_TILE_HOST_DEVICE constexpr void operator()(fp16x8_t& y, const pk_int4x4_t& x) const { y.lo = i4_to_half4(bit_cast(x)); y.hi = i4_to_half4(bit_cast(x) >> 8); } + CK_TILE_HOST_DEVICE constexpr void operator()(bf16x8_t& y, const fp8x8_t& x) const + { + y = fp8x8_to_bf16x8_scale(x, 1.0f); + } + + CK_TILE_HOST_DEVICE constexpr void operator()(fp8x8_t& y, const bf16x8_t& x) const + { + y = bf16x8_to_fp8x8_scale(x, 1.0f); + } + CK_TILE_HOST_DEVICE constexpr void operator()(bf16x8_t& y, const pk_int4x4_t& x) const { y.lo = i4_to_bhalf4(bit_cast(x)); From afba075a2ad03894551fbb8aa1eb1bc6a3d15ec5 Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Wed, 12 Nov 2025 09:04:15 +0000 Subject: [PATCH 03/15] Add and use load_with_type_convert --- .../ops/common/load_and_convert_tile.hpp | 49 +++++++++++++++---- 1 file changed, 39 insertions(+), 10 deletions(-) diff --git a/projects/composablekernel/include/ck_tile/ops/common/load_and_convert_tile.hpp b/projects/composablekernel/include/ck_tile/ops/common/load_and_convert_tile.hpp index 9fb26e7638f8..6dc2ff35c0a1 100644 --- a/projects/composablekernel/include/ck_tile/ops/common/load_and_convert_tile.hpp +++ b/projects/composablekernel/include/ck_tile/ops/common/load_and_convert_tile.hpp @@ -10,12 +10,16 @@ namespace ck_tile { -template +template struct ConverterLoader { template CK_TILE_DEVICE static void load_interleaved_pk_type(WarpTile& dst, const WarpWindow& src_window) { + static_assert(!LoadTranspose, "LoadTranspose not supported with pk_int4_t or pk_fp4_t"); static_assert(WarpTile::get_thread_buffer_size() % UnaryOpSize == 0); constexpr index_t thread_buffer_size = WarpTile::get_thread_buffer_size() / UnaryOpSize; const auto src = load_tile(src_window); @@ -33,6 +37,36 @@ struct ConverterLoader src.get_thread_buffer().template get_as()[i]); }); } + + template + CK_TILE_DEVICE static void load_with_type_convert(WarpTile& dst, const WarpWindow& src_window) + { + if constexpr(LoadTranspose) + { + if constexpr(std::is_same_v) + { + load_tile_transpose(dst, src_window); + } + else + { + load_tile_transpose_convert(dst, src_window, number{}); + } + } + else + { + if constexpr(std::is_same_v) + { + load_tile(dst, src_window); + } + else + { + auto tmp = load_tile(src_window); + sweep_tile([&](auto i) { + dst(i) = type_convert(type_convert(tmp(i))); + }); + } + } + } }; template @@ -43,18 +77,13 @@ CK_TILE_DEVICE void load_and_convert_tile(WarpTile& dst, const WarpWindow& src_w if constexpr(is_packed_type_v && !is_packed_type_v) { - static_assert(!LoadTranspose, "LoadTranspose not supported with pk_int4_t or pk_fp4_t"); - ConverterLoader::load_interleaved_pk_type( - dst, src_window); - } - else if constexpr(LoadTranspose) - { - load_tile_transpose(dst, src_window); + ConverterLoader:: + load_interleaved_pk_type(dst, src_window); } else { - load_tile(dst, src_window); + ConverterLoader:: + load_with_type_convert(dst, src_window); } } - } // namespace ck_tile From a9c52489facd48949de0ab40a06fbea28fb66283 Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Tue, 13 Jan 2026 09:04:21 +0000 Subject: [PATCH 04/15] Add test_load_tile_transpose --- .../test/ck_tile/CMakeLists.txt | 1 + .../load_and_convert_tile/CMakeLists.txt | 9 + .../ck_tile/load_and_convert_tile/kernel.hpp | 230 ++++++++++++++++++ .../test_load_and_convert_tile.cpp | 133 ++++++++++ 4 files changed, 373 insertions(+) create mode 100644 projects/composablekernel/test/ck_tile/load_and_convert_tile/CMakeLists.txt create mode 100644 projects/composablekernel/test/ck_tile/load_and_convert_tile/kernel.hpp create mode 100644 projects/composablekernel/test/ck_tile/load_and_convert_tile/test_load_and_convert_tile.cpp diff --git a/projects/composablekernel/test/ck_tile/CMakeLists.txt b/projects/composablekernel/test/ck_tile/CMakeLists.txt index 320e5b1e91c9..a0e35261fdef 100644 --- a/projects/composablekernel/test/ck_tile/CMakeLists.txt +++ b/projects/composablekernel/test/ck_tile/CMakeLists.txt @@ -70,3 +70,4 @@ add_subdirectory(gemm_tile_engine) add_subdirectory(pooling) add_subdirectory(grouped_conv) add_subdirectory(gemm_streamk_tile_engine) +add_subdirectory(load_and_convert_tile) diff --git a/projects/composablekernel/test/ck_tile/load_and_convert_tile/CMakeLists.txt b/projects/composablekernel/test/ck_tile/load_and_convert_tile/CMakeLists.txt new file mode 100644 index 000000000000..e07fc05fc814 --- /dev/null +++ b/projects/composablekernel/test/ck_tile/load_and_convert_tile/CMakeLists.txt @@ -0,0 +1,9 @@ +# Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +set(LOAD_TILE_TRANSPOSE_COMPILE_OPTIONS) +if(GPU_TARGETS MATCHES "gfx9") + add_gtest_executable(test_load_and_convert_tile test_load_and_convert_tile.cpp) + list(APPEND LOAD_TILE_TRANSPOSE_COMPILE_OPTIONS -DCK_TILE_USE_OCP_FP8 -fverbose-asm --save-temps -Wno-gnu-line-marker) + target_compile_options(test_load_and_convert_tile PRIVATE ${LOAD_TILE_TRANSPOSE_COMPILE_OPTIONS}) +endif() diff --git a/projects/composablekernel/test/ck_tile/load_and_convert_tile/kernel.hpp b/projects/composablekernel/test/ck_tile/load_and_convert_tile/kernel.hpp new file mode 100644 index 000000000000..14d1722e96ef --- /dev/null +++ b/projects/composablekernel/test/ck_tile/load_and_convert_tile/kernel.hpp @@ -0,0 +1,230 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include "ck_tile/core.hpp" +#include "ck_tile/core/algorithm/coordinate_transform.hpp" +#include "ck_tile/ops/common.hpp" +#include "ck_tile/ops/gemm/warp/warp_gemm.hpp" + +namespace ck_tile { + +template +struct LoadAndConvertShape +{ + static constexpr index_t Block_M = BlockTile::at(number<0>{}); + static constexpr index_t Block_N = BlockTile::at(number<1>{}); + static constexpr index_t Block_K = BlockTile::at(number<2>{}); + + static constexpr index_t Warp_M = WarpTile::at(number<0>{}); + static constexpr index_t Warp_N = WarpTile::at(number<1>{}); + static constexpr index_t Warp_K = WarpTile::at(number<2>{}); + + static constexpr index_t Vector_N = Vector::at(number<1>{}); + + static constexpr index_t WarpPerBlock_M = BlockWarps::at(number<0>{}); + static constexpr index_t WarpPerBlock_N = BlockWarps::at(number<1>{}); + static constexpr index_t WarpPerBlock_K = BlockWarps::at(number<2>{}); + + static constexpr index_t Repeat_M = Block_M / (WarpPerBlock_M * Warp_M); + static constexpr index_t Repeat_N = Block_N / (WarpPerBlock_N * Warp_N); + static constexpr index_t Repeat_K = Block_K / (WarpPerBlock_K * Warp_K); + + static constexpr index_t BlockSize = + ck_tile::get_warp_size() * reduce_on_sequence(BlockWarps{}, multiplies<>{}, number<1>{}); +}; + +template +struct LoadAndConvertProblem +{ + using XDataType = remove_cvref_t; + using YDataType = remove_cvref_t; + using BlockShape = remove_cvref_t; + using LoadTranspose = remove_cvref_t; +}; + +template +struct LoadAndConvertKernel +{ + using Problem = ck_tile::remove_cvref_t; + using XDataType = ck_tile::remove_cvref_t; + using YDataType = ck_tile::remove_cvref_t; + using LoadTranspose = ck_tile::remove_cvref_t; + + static constexpr index_t kBlockSize = Problem::BlockShape::BlockSize; + + template + static constexpr auto get_warp_dstr_encoding() + { + using S = typename Problem::BlockShape; + + if constexpr(NumAccess == 1) + return tile_distribution_encoding, + tuple, sequence<2, S::Vector_N>>, + tuple>, + tuple>, + sequence<2>, + sequence<1>>{}; + else + return tile_distribution_encoding< + sequence<>, + tuple, sequence>, + tuple>, + tuple>, + sequence<2, 2>, + sequence<0, 2>>{}; + } + + template + CK_TILE_DEVICE static constexpr auto GetVectorSize() + { + return DS_READ_TR_SIZE() / sizeof(DataType); + } + + template + CK_TILE_DEVICE static constexpr auto MakeDRAMDistribution() + { + using S = typename Problem::BlockShape; + constexpr index_t thread_elements = S::Warp_N * S::Warp_K / get_warp_size(); + constexpr index_t NumAccess = LoadTranspose::value ? thread_elements / GetVectorSize() : 1; + + constexpr auto a_block_outer_dstr_encode = tile_distribution_encoding< + sequence, + tuple, sequence>, + tuple>, + tuple>, + sequence<1, 2>, + sequence<0, 0>>{}; + + constexpr auto a_block_dstr_encode = detail::make_embed_tile_distribution_encoding( + a_block_outer_dstr_encode, get_warp_dstr_encoding()); + + return make_static_tile_distribution(a_block_dstr_encode); + } + + template + CK_TILE_DEVICE static constexpr auto MakeDRAMTransposedDistribution() + { + return make_static_tile_distribution( + typename InputTileDistributionTraits< + typename decltype(MakeDRAMDistribution())::DstrEncode, + DataType>::TransposedDstrEncode{}); + } + + CK_TILE_DEVICE void + operator()(const XDataType* a, YDataType* c, index_t M, index_t N, index_t K) const + { + using S = typename Problem::BlockShape; + + const index_t kMPerBlock = S::WarpPerBlock_M * S::Repeat_M * S::Block_M; + const index_t kNPerBlock = S::WarpPerBlock_N * S::Repeat_N * S::Block_N; + + constexpr auto block_dims = make_tuple(number{}, number{}); + constexpr auto block_strides = make_tuple(number<1>{}, number{}); + const index_t num_blocks_n = N / kNPerBlock; + const index_t block_m = get_block_id() / num_blocks_n; + + const index_t m_block_base = block_m * kMPerBlock; + + // LDS buffer + __shared__ XDataType a_lds[kMPerBlock * S::Block_K]; + + auto a_lds_write_view = make_naive_tensor_view( + a_lds, block_dims, block_strides, number<1>{}, number<1>{}); + + auto a_block_lds_write_window = make_tile_window(a_lds_write_view, block_dims, {0, 0}); + + auto a_block_lds_read_window = [&] { + if constexpr(LoadTranspose::value) + { + constexpr auto block_dims_t = + make_tuple(number{}, number{}); + constexpr auto block_strides_t = make_tuple(number{}, number<1>{}); + + auto view = make_naive_tensor_view( + a_lds, + block_dims_t, + block_strides_t, + number()>{}, + number<1>{}); + + return make_tile_window( + view, block_dims_t, {0, 0}, MakeDRAMTransposedDistribution()); + } + else + { + auto view = make_naive_tensor_view( + a_lds, block_dims, block_strides, number<1>{}, number<1>{}); + + return make_tile_window( + view, block_dims, {0, 0}, MakeDRAMDistribution()); + } + }(); + + // Input tensor + const auto a_tensor = make_naive_tensor_view( + a, make_tuple(M, K), make_tuple(1, M), number<1>{}, number<1>{}); + + auto a_block_window = make_tile_window( + a_tensor, block_dims, {m_block_base, 0}, MakeDRAMDistribution()); + + // Output tensor + auto c_tensor = [&]() { + if constexpr(LoadTranspose::value && !std::is_same_v) + { + // Similar to QuantGemmKernel::PermuteB: reinterpret the output logical layout via + // descriptor transform so YDataType distribution can be used without row-group + // permutation artifacts in mixed-precision transpose loads. + const auto c_m_n_desc = make_naive_tensor_descriptor( + make_tuple(M, N), make_tuple(1, M), number<1>{}, number<1>{}); + + const auto c_m_n0_b0_b1_n4_desc = transform_tensor_descriptor( + c_m_n_desc, + make_tuple( + make_pass_through_transform(M), + make_unmerge_transform(make_tuple(N / 16, number<2>{}, number<2>{}, number<4>{}))), + make_tuple(sequence<0>{}, sequence<1>{}), + make_tuple(sequence<0>{}, sequence<1, 2, 3, 4>{})); + + const auto c_perm_m_n_desc = transform_tensor_descriptor( + c_m_n0_b0_b1_n4_desc, + make_tuple(make_pass_through_transform(M), + make_merge_transform( + make_tuple(N / 16, number<2>{}, number<2>{}, number<4>{}))), + make_tuple(sequence<0>{}, sequence<1, 3, 2, 4>{}), + make_tuple(sequence<0>{}, sequence<1>{})); + + return make_tensor_view(c, c_perm_m_n_desc); + } + else + { + return make_naive_tensor_view( + c, make_tuple(M, N), make_tuple(1, M), number<1>{}, number<1>{}); + } + }(); + + auto c_block_window = make_tile_window( + c_tensor, block_dims, {m_block_base, 0}, MakeDRAMDistribution()); + + const index_t num_k_loops = K / S::Block_K; + for(index_t k_iter = 0; k_iter < num_k_loops; ++k_iter) + { + auto dram_tile = load_tile(a_block_window); + store_tile(a_block_lds_write_window, dram_tile); + block_sync_lds(); + + decltype(load_tile(c_block_window)) c_tile; + load_and_convert_tile<8, LoadTranspose::value>(c_tile, a_block_lds_read_window); + store_tile(c_block_window, c_tile); + + if(k_iter < num_k_loops - 1) + { + move_tile_window(a_block_window, {0, S::Block_K}); + move_tile_window(c_block_window, {0, S::Block_K}); + } + } + } +}; + +} // namespace ck_tile diff --git a/projects/composablekernel/test/ck_tile/load_and_convert_tile/test_load_and_convert_tile.cpp b/projects/composablekernel/test/ck_tile/load_and_convert_tile/test_load_and_convert_tile.cpp new file mode 100644 index 000000000000..24b61b8f4bf0 --- /dev/null +++ b/projects/composablekernel/test/ck_tile/load_and_convert_tile/test_load_and_convert_tile.cpp @@ -0,0 +1,133 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include "ck_tile/host.hpp" +#include "ck_tile/ops/common.hpp" +#include "kernel.hpp" + +#define USE_MONOTONIC_SEQUENCE 0 +#define PRINT_MATRICES 0 + +// Helper to print matrix (for debugging) +template +void print_matrix(const ck_tile::HostTensor& mat, const std::string& name = "Matrix") +{ + const auto lens = mat.get_lengths(); + assert(len(lens) == 2); + const ck_tile::index_t rows = lens[0]; + const ck_tile::index_t cols = lens[1]; + const ck_tile::index_t limit = 32; + + std::cout << name << " (" << rows << "×" << cols << "):\n"; + for(ck_tile::index_t i = 0; i < std::min(rows, ck_tile::index_t(limit)); ++i) + { + for(ck_tile::index_t j = 0; j < std::min(cols, ck_tile::index_t(limit)); ++j) + { +#if !USE_MONOTONIC_SEQUENCE + std::cout << std::setw(3) << std::setprecision(3) +#else + std::cout << std::fixed << std::setprecision(2) << std::setw(6) +#endif + << ck_tile::type_convert(mat(i, j)) << " "; + } + if(cols > limit) + std::cout << "..."; + std::cout << "\n"; + } + if(rows > limit) + std::cout << "...\n"; + std::cout << "\n"; +} + +template +class TestLoadAndConvert : public ::testing::Test +{ + public: + using XDataType = std::tuple_element_t<0, Tuple>; + using YDataType = std::tuple_element_t<1, Tuple>; + using LoadTranspose = std::tuple_element_t<2, Tuple>; + + protected: + void RunTest() + { + constexpr ck_tile::index_t M = 32; + constexpr ck_tile::index_t N = 32; + constexpr ck_tile::index_t K = 32; + + ck_tile::HostTensor h_a({M, K}); + ck_tile::HostTensor h_c({M, K}); +#if !USE_MONOTONIC_SEQUENCE + ck_tile::FillUniformDistributionIntegerValue{-5.0, 5.0, 11939}(h_a); +#else + ck_tile::HostTensor h_a_tmp({M, K}); + ck_tile::FillMonotonicSeq{0.0, 0.1}(h_a_tmp); + ck_tile::reference_unary_elementwise( + h_a_tmp, h_a, [](const auto& x) { return x; }); +#endif + + ck_tile::DeviceMem d_a(h_a.get_element_space_size_in_bytes()); + ck_tile::DeviceMem d_c(h_c.get_element_space_size_in_bytes()); + + d_a.ToDevice(h_a.data()); + d_c.ToDevice(h_c.data()); + + using BlockWarps = ck_tile::sequence<1, 1, 1>; + using BlockTile = ck_tile::sequence<32, 32, 16>; + using WarpTile = ck_tile::sequence<32, 32, 16>; + using Vector = ck_tile::sequence<1, 8>; + + using Shape = ck_tile::LoadAndConvertShape; + using Problem = ck_tile::LoadAndConvertProblem; + using Kernel = ck_tile::LoadAndConvertKernel; + + constexpr ck_tile::index_t block_size = Kernel::kBlockSize; + const ck_tile::index_t grid_size = ck_tile::integer_divide_ceil(M, Shape::Block_M) * + ck_tile::integer_divide_ceil(N, Shape::Block_N); + + launch_kernel(ck_tile::stream_config{nullptr, true}, + make_kernel(Kernel{}, + dim3(grid_size), + dim3(block_size), + 0, + static_cast(d_a.GetDeviceBuffer()), + static_cast(d_c.GetDeviceBuffer()), + M, + N, + K)); + + ck_tile::hip_check_error(hipDeviceSynchronize()); + d_c.FromDevice(h_c.data()); + ck_tile::HostTensor h_a_ref({M, K}); + ck_tile::reference_unary_elementwise( + h_a, h_a_ref, [](const auto& x) { return x; }); + bool pass = ck_tile::check_err(h_c, h_a_ref); + +#if PRINT_MATRICES + print_matrix(h_a, "Matrix A"); + print_matrix(h_c, "Matrix C"); +#endif + + EXPECT_TRUE(pass); + } +}; + +using TestTypes = ::testing::Types, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple, + std::tuple>; + +TYPED_TEST_SUITE(TestLoadAndConvert, TestTypes); + +TYPED_TEST(TestLoadAndConvert, Test) { this->RunTest(); } From b381791150536005230ac09637475e65bdea76e7 Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Mon, 16 Mar 2026 12:38:20 +0000 Subject: [PATCH 05/15] When defining c_tensor, replace hard coded values with values derived from kernel parameters --- .../ck_tile/load_and_convert_tile/kernel.hpp | 22 +++++++++++++++++-- 1 file changed, 20 insertions(+), 2 deletions(-) diff --git a/projects/composablekernel/test/ck_tile/load_and_convert_tile/kernel.hpp b/projects/composablekernel/test/ck_tile/load_and_convert_tile/kernel.hpp index 14d1722e96ef..8549e7b64ac9 100644 --- a/projects/composablekernel/test/ck_tile/load_and_convert_tile/kernel.hpp +++ b/projects/composablekernel/test/ck_tile/load_and_convert_tile/kernel.hpp @@ -176,6 +176,20 @@ struct LoadAndConvertKernel // Similar to QuantGemmKernel::PermuteB: reinterpret the output logical layout via // descriptor transform so YDataType distribution can be used without row-group // permutation artifacts in mixed-precision transpose loads. + using TransposeGroupType = std::conditional_t<(sizeof(XDataType) >= sizeof(YDataType)), + XDataType, + YDataType>; + constexpr index_t thread_elements = S::Warp_N * S::Warp_K / get_warp_size(); + constexpr index_t n_group_0 = + thread_elements / GetVectorSize(); + constexpr index_t n_group_1 = 2; + constexpr index_t n_group_2 = S::Vector_N / n_group_0; + constexpr index_t n_perm_group = n_group_0 * n_group_1 * n_group_2; + + static_assert(n_group_0 > 0, "Invalid derived transpose grouping factor"); + static_assert(S::Vector_N % n_group_0 == 0, + "Vector_N must be divisible by derived grouping factor"); + const auto c_m_n_desc = make_naive_tensor_descriptor( make_tuple(M, N), make_tuple(1, M), number<1>{}, number<1>{}); @@ -183,7 +197,8 @@ struct LoadAndConvertKernel c_m_n_desc, make_tuple( make_pass_through_transform(M), - make_unmerge_transform(make_tuple(N / 16, number<2>{}, number<2>{}, number<4>{}))), + make_unmerge_transform( + make_tuple(N / n_perm_group, number{}, number{}, number{}))), make_tuple(sequence<0>{}, sequence<1>{}), make_tuple(sequence<0>{}, sequence<1, 2, 3, 4>{})); @@ -191,7 +206,10 @@ struct LoadAndConvertKernel c_m_n0_b0_b1_n4_desc, make_tuple(make_pass_through_transform(M), make_merge_transform( - make_tuple(N / 16, number<2>{}, number<2>{}, number<4>{}))), + make_tuple(N / n_perm_group, + number{}, + number{}, + number{}))), make_tuple(sequence<0>{}, sequence<1, 3, 2, 4>{}), make_tuple(sequence<0>{}, sequence<1>{})); From 7a8c4489390b05bcdb32ee3f768c509a5ef5cc93 Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Mon, 2 Mar 2026 15:14:12 +0000 Subject: [PATCH 06/15] Add AttrNumAccessV_ as a template parameter to WarpGemmAttributeMfma::get_warp_dstr_encoding --- .../block/block_universal_gemm_as_bs_cr.hpp | 16 +++- .../gemm/warp/warp_gemm_attribute_mfma.hpp | 89 ++++++++++++++----- .../gemm/warp/warp_gemm_attribute_wmma.hpp | 22 ++++- .../ck_tile/ops/gemm/warp/warp_gemm_impl.hpp | 4 +- 4 files changed, 105 insertions(+), 26 deletions(-) diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp index f7f5cd33dbb9..9751f97c3608 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp @@ -140,6 +140,7 @@ struct BlockUniversalGemmAsBsCr using I0 = number<0>; using I1 = number<1>; + template CK_TILE_DEVICE static constexpr auto MakeABlockDistributionEncode() { constexpr index_t KPerThread = Traits::KPerThread; @@ -159,12 +160,18 @@ struct BlockUniversalGemmAsBsCr tuple>, sequence<1, 2>, sequence<0, 0>>{}; + using Attr = typename WarpGemm::WarpGemmAttribute; + constexpr auto NumAccessA = + convert ? Attr::AttrNumAccessAV * sizeof(ADataType) / sizeof(ComputeDataType) + : Attr::AttrNumAccessAV; constexpr auto a_block_dstr_encode = detail::make_embed_tile_distribution_encoding( - a_block_outer_dstr_encoding, typename WarpGemm::AWarpDstrEncoding{}); + a_block_outer_dstr_encoding, + WarpGemm::WarpGemmAttribute::template get_awarp_dstr_encoding()); return a_block_dstr_encode; } + template CK_TILE_DEVICE static constexpr auto MakeBBlockDistributionEncode() { constexpr index_t KPerThread = Traits::KPerThread; @@ -184,8 +191,13 @@ struct BlockUniversalGemmAsBsCr tuple>, sequence<1, 2>, sequence<0, 0>>{}; + using Attr = typename WarpGemm::WarpGemmAttribute; + constexpr auto NumAccessB = + convert ? Attr::AttrNumAccessBV * sizeof(BDataType) / sizeof(ComputeDataType) + : Attr::AttrNumAccessBV; constexpr auto b_block_dstr_encode = detail::make_embed_tile_distribution_encoding( - b_block_outer_dstr_encoding, typename WarpGemm::BWarpDstrEncoding{}); + b_block_outer_dstr_encoding, + WarpGemm::WarpGemmAttribute::template get_bwarp_dstr_encoding()); return b_block_dstr_encode; } diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp index f79741ea9606..8ab685d06c03 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp @@ -132,8 +132,23 @@ struct WarpGemmAttributeMfma } } } - using AWarpDstrEncoding = decltype(get_warp_dstr_encoding()); - using BWarpDstrEncoding = decltype(get_warp_dstr_encoding()); + template + static constexpr auto get_awarp_dstr_encoding() + { + return get_warp_dstr_encoding(); + } + + template + static constexpr auto get_bwarp_dstr_encoding() + { + return get_warp_dstr_encoding(); + } + + template + using AWarpDstrEncoding = decltype(get_awarp_dstr_encoding()); + + template + using BWarpDstrEncoding = decltype(get_bwarp_dstr_encoding()); using CWarpDstrEncoding = tile_distribution_encoding< sequence<>, @@ -297,6 +312,18 @@ struct WarpGemmAttributeMfmaIterateK } } + template + CK_TILE_DEVICE static constexpr auto get_awarp_dstr_encoding() + { + return get_warp_dstr_encoding(); + } + + template + CK_TILE_DEVICE static constexpr auto get_bwarp_dstr_encoding() + { + return get_warp_dstr_encoding(); + } + CK_TILE_DEVICE static constexpr auto get_cwarp_dstr_encoding() { if constexpr(Impl::kAMBlock == 1 && Impl::kBNBlock == 1) @@ -335,14 +362,12 @@ struct WarpGemmAttributeMfmaIterateK } } - using AWarpDstrEncoding = decltype(get_warp_dstr_encoding()); - using BWarpDstrEncoding = decltype(get_warp_dstr_encoding()); + template + using AWarpDstrEncoding = decltype(get_awarp_dstr_encoding()); + + template + using BWarpDstrEncoding = decltype(get_bwarp_dstr_encoding()); + using CWarpDstrEncoding = decltype(get_cwarp_dstr_encoding()); // c_vec += a_vec * b_vec @@ -421,10 +446,23 @@ struct WarpGemmAttributeMfmaTransposedCDistribution static_assert(Impl::kAMBlock == 1 && Impl::kBNBlock == 1, "Multi-block WarpGemmAttributeMfmaImpl is not supported"); - using AWarpDstrEncoding = - typename WarpGemmAttributeMfma::BWarpDstrEncoding; - using BWarpDstrEncoding = - typename WarpGemmAttributeMfma::AWarpDstrEncoding; + template + CK_TILE_DEVICE static constexpr auto get_awarp_dstr_encoding() + { + return WarpGemmAttributeMfma::template get_bwarp_dstr_encoding(); + } + + template + CK_TILE_DEVICE static constexpr auto get_bwarp_dstr_encoding() + { + return WarpGemmAttributeMfma::template get_awarp_dstr_encoding(); + } + + template + using AWarpDstrEncoding = decltype(get_awarp_dstr_encoding()); + + template + using BWarpDstrEncoding = decltype(get_bwarp_dstr_encoding()); using CWarpDstrEncoding = tile_distribution_encoding< sequence<>, @@ -491,6 +529,7 @@ struct WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB static_assert(Impl::kAMBlock == 1 && Impl::kBNBlock == 1, "Multi-block WarpGemmAttributeMfmaImpl is not supported"); + template using AWarpDstrEncoding = tile_distribution_encoding< sequence<>, tuple, sequence>, @@ -499,6 +538,7 @@ struct WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB sequence<2>, sequence<1>>; #if 0 + template using BWarpDstrEncoding = tile_distribution_encoding< sequence<>, tuple>; #else // TODO: more test not only 32x32 + template using BWarpDstrEncoding = tile_distribution_encoding< sequence<>, tuple struct WarpGemmAttributeMfmaIterateKAndTransposedCDistribution { - using Impl = remove_cvref_t; - static constexpr auto AttrNumAccess = AttrNumAccess_; + using Impl = remove_cvref_t; + static constexpr auto AttrNumAccess = AttrNumAccess_; + static constexpr auto AttrNumAccessV = static_cast(AttrNumAccess); // swap A and B using ADataType = typename Impl::BDataType; @@ -641,10 +683,12 @@ struct WarpGemmAttributeMfmaIterateKAndTransposedCDistribution } } - using AWarpDstrEncoding = - typename WarpGemmAttributeMfmaIterateK::BWarpDstrEncoding; - using BWarpDstrEncoding = - typename WarpGemmAttributeMfmaIterateK::AWarpDstrEncoding; + template + using AWarpDstrEncoding = typename WarpGemmAttributeMfmaIterateK:: + template BWarpDstrEncoding; + template + using BWarpDstrEncoding = typename WarpGemmAttributeMfmaIterateK:: + template AWarpDstrEncoding; using CWarpDstrEncoding = decltype(get_cwarp_dstr_encoding()); // c_vec += a_vec * b_vec @@ -723,6 +767,7 @@ struct WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB static_assert(Impl::kAMBlock == 1 && Impl::kBNBlock == 1, "Multi-block WarpGemmAttributeMfmaImpl is not supported"); + template using AWarpDstrEncoding = tile_distribution_encoding< sequence<>, tuple, sequence>, @@ -731,6 +776,7 @@ struct WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB sequence<2>, sequence<1>>; #if 0 + template using BWarpDstrEncoding = tile_distribution_encoding< sequence<>, tuple>; #else // TODO: more test not only 32x32 + template using BWarpDstrEncoding = tile_distribution_encoding< sequence<>, tuple using AWarpDstrEncoding = tile_distribution_encoding< sequence<>, tuple, sequence<1>>; + template using BWarpDstrEncoding = tile_distribution_encoding< sequence<>, tuple, sequence>, diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma.hpp index ef31d06c9c27..4a275848b2b6 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma.hpp @@ -67,6 +67,10 @@ template struct WarpGemmAttributeWmma { using Impl = remove_cvref_t; + // AttrNumAccessV is required for compatibility with the block GEMM, and is currently ignored + // within WarpGemmAttributeWmma + static constexpr auto AttrNumAccess = WGAttrNumAccessEnum::Single; + static constexpr auto AttrNumAccessV = static_cast(AttrNumAccess); // When kTransC is true and A/B types differ, we need an impl with swapped types using TransposedImpl = @@ -99,8 +103,22 @@ struct WarpGemmAttributeWmma // 16 bit input, kAMLane = 16, kABK0PerLane = 4, kABKLane = 2, kABK1PerLane = 2 // 8 bit input, kAMLane = 16, kABK0PerLane = 2, kABKLane = 2, kABK1PerLane = 4 - using AWarpDstrEncoding = typename AWarpDstrEncodingTrait::type; - using BWarpDstrEncoding = typename BWarpDstrEncodingTrait::type; + template + static constexpr auto get_awarp_dstr_encoding() + { + return typename AWarpDstrEncodingTrait::type{}; + } + + template + static constexpr auto get_bwarp_dstr_encoding() + { + return typename BWarpDstrEncodingTrait::type{}; + } + + template + using AWarpDstrEncoding = decltype(get_awarp_dstr_encoding()); + template + using BWarpDstrEncoding = decltype(get_bwarp_dstr_encoding()); // kCM0PerLane = 1, kCMLane = 2, kCM1PerLane = 2, kCNLane = 16 using CWarpDstrEncoding = diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_impl.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_impl.hpp index ca7c32b6af55..5ff0660f49c4 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_impl.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_impl.hpp @@ -25,8 +25,8 @@ struct WarpGemmImpl using BDataType = typename WarpGemmAttribute::BDataType; using CDataType = typename WarpGemmAttribute::CDataType; - using AWarpDstrEncoding = typename WarpGemmAttribute::AWarpDstrEncoding; - using BWarpDstrEncoding = typename WarpGemmAttribute::BWarpDstrEncoding; + using AWarpDstrEncoding = typename WarpGemmAttribute::template AWarpDstrEncoding<>; + using BWarpDstrEncoding = typename WarpGemmAttribute::template BWarpDstrEncoding<>; using CWarpDstrEncoding = typename WarpGemmAttribute::CWarpDstrEncoding; using AWarpDstr = remove_cvref_t; From 6ae8db7f778263583155ffb5ee9bd1807b63db75 Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Thu, 9 Oct 2025 08:07:04 +0000 Subject: [PATCH 07/15] Introduce DetermineWarpPrecType for determining warp GEMM precision types --- .../include/ck_tile/ops/common.hpp | 1 + .../ops/common/determine_warp_prec_type.hpp | 41 +++++++++++++++++++ .../ops/epilogue/cshuffle_epilogue.hpp | 14 ++----- .../block/block_universal_gemm_as_bs_cr.hpp | 9 +--- .../gemm_pipeline_ag_bg_cr_comp_v3.hpp | 9 ++-- ...emm_universal_pipeline_ag_bg_cr_policy.hpp | 13 ++---- 6 files changed, 56 insertions(+), 31 deletions(-) create mode 100644 projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp diff --git a/projects/composablekernel/include/ck_tile/ops/common.hpp b/projects/composablekernel/include/ck_tile/ops/common.hpp index ad7da5c18339..0113d8c9a280 100644 --- a/projects/composablekernel/include/ck_tile/ops/common.hpp +++ b/projects/composablekernel/include/ck_tile/ops/common.hpp @@ -2,6 +2,7 @@ // SPDX-License-Identifier: MIT #pragma once +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp b/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp new file mode 100644 index 000000000000..2e28073c9c8d --- /dev/null +++ b/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp @@ -0,0 +1,41 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include "ck_tile/core.hpp" + +// DetermineWarpPrecType is a set of rules to determine the right precision type to use +// for the warp GEMM, given the other precision type. This gives rise to a type conversion: +// type conversions are sometimes needed to obtain a pair of types that are compatible with +// the hardware matrix operations available. A typical use case is mixed precision GEMMs. + +namespace ck_tile { +// For the most general case, we default to no conversion. +template +struct DetermineWarpPrecType +{ + using prec_type = PrecType; +}; + +// For pk_int4_t, we convert to the other precision type. +template +struct DetermineWarpPrecType +{ + using prec_type = OtherPrecType; +}; + +// For pk_fp4_t, we convert to the other precision type. +template +struct DetermineWarpPrecType +{ + using prec_type = OtherPrecType; +}; + +// For pk_fp4_raw_t, we convert to the other precision type. +template +struct DetermineWarpPrecType +{ + using prec_type = OtherPrecType; +}; +}; // namespace ck_tile diff --git a/projects/composablekernel/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp b/projects/composablekernel/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp index 8dad775ecfb1..6dcf66528795 100644 --- a/projects/composablekernel/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp +++ b/projects/composablekernel/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp @@ -7,7 +7,7 @@ #include "ck_tile/core.hpp" #include "ck_tile/ops/common/utils.hpp" #include "ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp" -#include "ck_tile/ops/common/tensor_layout.hpp" +#include "ck_tile/ops/common.hpp" #include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp" #include @@ -92,16 +92,8 @@ struct CShuffleEpilogue using ADataType = remove_cvref_t{}, AsDataTypeTuple>>; using BDataType = remove_cvref_t{}, BsDataTypeTuple>>; - using ATypeToUse = std::conditional_t || - std::is_same_v, - BDataType, - ADataType>; - // Used for weight-only quantization kernel, B would be dequantized to the same data type as A - using BTypeToUse = std::conditional_t || - std::is_same_v || - sizeof(BDataType) < sizeof(ADataType), - ADataType, - BDataType>; + using ATypeToUse = typename DetermineWarpPrecType::prec_type; + using BTypeToUse = typename DetermineWarpPrecType::prec_type; using ELayout = remove_cvref_t; using CDElementwise = remove_cvref_t; diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp index 9751f97c3608..9151607aa8e3 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp @@ -94,13 +94,8 @@ struct BlockUniversalGemmAsBsCr using ComputeDataType = remove_cvref_t; using CDataType = remove_cvref_t; - using ATypeToUse = - std::conditional_t, BDataType, ADataType>; - using BTypeToUse = std::conditional_t || - std::is_same_v || - sizeof(BDataType) < sizeof(ADataType), - ADataType, - BDataType>; + using ATypeToUse = typename DetermineWarpPrecType::prec_type; + using BTypeToUse = typename DetermineWarpPrecType::prec_type; using WarpGemm = remove_cvref_t; diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp index 83f7f8082410..e03e069aac56 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp @@ -441,10 +441,11 @@ struct GemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 auto&& [a_lds_block, b_lds_block] = Base::GetABLdsTensorViews(p_smem); // Tile distribution for load from lds - constexpr auto a_lds_load_tile_distr = - make_static_tile_distribution(BlockGemm::MakeABlockDistributionEncode()); - constexpr auto b_lds_load_tile_distr = - make_static_tile_distribution(BlockGemm::MakeBBlockDistributionEncode()); + constexpr bool is_load_tr = is_a_load_tr_v || is_b_load_tr_v; + constexpr auto a_lds_load_tile_distr = make_static_tile_distribution( + BlockGemm::template MakeABlockDistributionEncode()); + constexpr auto b_lds_load_tile_distr = make_static_tile_distribution( + BlockGemm::template MakeBBlockDistributionEncode()); // A DRAM tile window for load // A LDS tile window for store diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp index 4e06ac8f7be9..1365136be263 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp @@ -909,15 +909,10 @@ struct UniversalGemmPipelineAgBgCrPolicy : vector_size * 4 == thread_elements ? WGAttrNumAccessEnum::Quad : WGAttrNumAccessEnum::Invalid; - using ADataType = remove_cvref_t; - using BDataType = remove_cvref_t; - using ATypeToUse = - std::conditional_t, BDataType, ADataType>; - using BTypeToUse = std::conditional_t || - std::is_same_v || - sizeof(BDataType) < sizeof(ADataType), - ADataType, - BDataType>; + using ADataType = remove_cvref_t; + using BDataType = remove_cvref_t; + using ATypeToUse = typename DetermineWarpPrecType::prec_type; + using BTypeToUse = typename DetermineWarpPrecType::prec_type; using WarpGemm = WarpGemmDispatcher Date: Wed, 28 Jan 2026 10:05:18 +0000 Subject: [PATCH 08/15] Switch to an implementation of DetermineWarpPrecType that explicitly defines the A and B types - This is for improved clarity and finer control of the datatypes to use --- .../ops/common/determine_warp_prec_type.hpp | 66 ++++++++++++++----- .../ops/epilogue/cshuffle_epilogue.hpp | 4 +- .../block/block_universal_gemm_as_bs_cr.hpp | 4 +- ...emm_universal_pipeline_ag_bg_cr_policy.hpp | 4 +- 4 files changed, 57 insertions(+), 21 deletions(-) diff --git a/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp b/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp index 2e28073c9c8d..5387b4e1c393 100644 --- a/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp +++ b/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp @@ -11,31 +11,67 @@ // the hardware matrix operations available. A typical use case is mixed precision GEMMs. namespace ck_tile { -// For the most general case, we default to no conversion. -template +// For the most general case, default to no conversion. +template struct DetermineWarpPrecType { - using prec_type = PrecType; + using a_prec_type = APrecType; + using b_prec_type = BPrecType; }; -// For pk_int4_t, we convert to the other precision type. -template -struct DetermineWarpPrecType +// For pk_fp4_t x pk_fp4_t, keep pk_fp4_t +template <> +struct DetermineWarpPrecType { - using prec_type = OtherPrecType; + using a_prec_type = ck_tile::pk_fp4_t; + using b_prec_type = ck_tile::pk_fp4_t; }; -// For pk_fp4_t, we convert to the other precision type. -template -struct DetermineWarpPrecType +// For pk_int4_t x B, use the B type. +template +struct DetermineWarpPrecType { - using prec_type = OtherPrecType; + using a_prec_type = BPrecType; + using b_prec_type = BPrecType; }; -// For pk_fp4_raw_t, we convert to the other precision type. -template -struct DetermineWarpPrecType +// For A x pk_int4_t, use the A type. +template +struct DetermineWarpPrecType { - using prec_type = OtherPrecType; + using a_prec_type = APrecType; + using b_prec_type = APrecType; +}; + +// For pk_fp4_t x B, use the B type. +template +struct DetermineWarpPrecType +{ + using a_prec_type = BPrecType; + using b_prec_type = BPrecType; +}; + +// For A x pk_fp4_t, use the A type. +template +struct DetermineWarpPrecType +{ + using a_prec_type = APrecType; + using b_prec_type = APrecType; +}; + +// For B x pk_fp4_raw_t, use the B type. +template +struct DetermineWarpPrecType +{ + using a_prec_type = BPrecType; + using b_prec_type = BPrecType; +}; + +// For A x pk_fp4_raw_t, use the A type. +template +struct DetermineWarpPrecType +{ + using a_prec_type = APrecType; + using b_prec_type = APrecType; }; }; // namespace ck_tile diff --git a/projects/composablekernel/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp b/projects/composablekernel/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp index 6dcf66528795..162bd7cc2576 100644 --- a/projects/composablekernel/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp +++ b/projects/composablekernel/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp @@ -92,8 +92,8 @@ struct CShuffleEpilogue using ADataType = remove_cvref_t{}, AsDataTypeTuple>>; using BDataType = remove_cvref_t{}, BsDataTypeTuple>>; - using ATypeToUse = typename DetermineWarpPrecType::prec_type; - using BTypeToUse = typename DetermineWarpPrecType::prec_type; + using ATypeToUse = typename DetermineWarpPrecType::a_prec_type; + using BTypeToUse = typename DetermineWarpPrecType::b_prec_type; using ELayout = remove_cvref_t; using CDElementwise = remove_cvref_t; diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp index 9151607aa8e3..1c5d5ede3aff 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp @@ -94,8 +94,8 @@ struct BlockUniversalGemmAsBsCr using ComputeDataType = remove_cvref_t; using CDataType = remove_cvref_t; - using ATypeToUse = typename DetermineWarpPrecType::prec_type; - using BTypeToUse = typename DetermineWarpPrecType::prec_type; + using ATypeToUse = typename DetermineWarpPrecType::a_prec_type; + using BTypeToUse = typename DetermineWarpPrecType::b_prec_type; using WarpGemm = remove_cvref_t; diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp index 1365136be263..1ee334404bba 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp @@ -911,8 +911,8 @@ struct UniversalGemmPipelineAgBgCrPolicy using ADataType = remove_cvref_t; using BDataType = remove_cvref_t; - using ATypeToUse = typename DetermineWarpPrecType::prec_type; - using BTypeToUse = typename DetermineWarpPrecType::prec_type; + using ATypeToUse = typename DetermineWarpPrecType::a_prec_type; + using BTypeToUse = typename DetermineWarpPrecType::b_prec_type; using WarpGemm = WarpGemmDispatcher Date: Wed, 12 Nov 2025 12:38:04 +0000 Subject: [PATCH 09/15] Add MFMA warp gemm for float, float, float, 32, 32, 16 --- .../include/ck_tile/ops/gemm/warp/warp_gemm.hpp | 6 ++++++ .../include/ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp | 2 ++ 2 files changed, 8 insertions(+) diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm.hpp index f393526de10c..8d64420e6e41 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm.hpp @@ -17,6 +17,12 @@ namespace ck_tile { using WarpGemmMfmaF32F32F32M16N16K4 = WarpGemmImpl< WarpGemmAttributeMfma>>; +template +using WarpGemmMfmaF32F32F32M32N32K16 = WarpGemmImpl, + 8, + AttrNumAccess>>; + template using WarpGemmMfmaF32F32F32M16N16K16 = WarpGemmImpl, diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp index 081ff5150d53..0f4f2c61c461 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp @@ -40,6 +40,8 @@ template<> struct Dispatcher { using Typ template<> struct Dispatcher { using Type = WarpGemmMfmaF32F32F32M32N32K8<>; }; template<> struct Dispatcher { using Type = WarpGemmMfmaF32F32F32M32N32K8; }; template<> struct Dispatcher { using Type = WarpGemmMfmaF32F32F32M16N16K16TransposedCDistribution<>; }; +template<> struct Dispatcher { using Type = WarpGemmMfmaF32F32F32M32N32K16<>; }; +template<> struct Dispatcher { using Type = WarpGemmMfmaF32F32F32M32N32K16; }; // fp16 // ADataType, BDataType, AccDataType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity template<> struct Dispatcher { using Type = WarpGemmMfmaF16F16F32M32N32K8; }; From 9475ac1fb1006d56052751ed4e2cf67b682d699a Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Thu, 9 Oct 2025 09:04:13 +0000 Subject: [PATCH 10/15] Add functionality and tests for bf16 x fp8 and fp8 x bf16 --- .../ops/common/determine_warp_prec_type.hpp | 16 ++++++++++++++++ .../gemm/test_gemm_pipeline_kernel_types.hpp | 8 ++++++++ 2 files changed, 24 insertions(+) diff --git a/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp b/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp index 5387b4e1c393..b8b5606360fe 100644 --- a/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp +++ b/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp @@ -74,4 +74,20 @@ struct DetermineWarpPrecType using a_prec_type = APrecType; using b_prec_type = APrecType; }; + +// For fp8 x bf16, use fp8 +template <> +struct DetermineWarpPrecType +{ + using a_prec_type = ck_tile::fp8_t; + using b_prec_type = ck_tile::fp8_t; +}; + +// For bf16 x fp8, use bf16 +template <> +struct DetermineWarpPrecType +{ + using a_prec_type = ck_tile::bf16_t; + using b_prec_type = ck_tile::bf16_t; +}; }; // namespace ck_tile diff --git a/projects/composablekernel/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp b/projects/composablekernel/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp index 7874c6f58886..55dcf7a16e50 100644 --- a/projects/composablekernel/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp +++ b/projects/composablekernel/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp @@ -82,9 +82,11 @@ using KernelTypesCompV3 = ::testing::Types< std::tuple< Row, Row, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, F16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, BF16, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Row, Row, Row, BF16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, BF16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Row, Row, Row, F8, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, F8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, F8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, BF8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, @@ -92,9 +94,11 @@ using KernelTypesCompV3 = ::testing::Types< std::tuple< Row, Col, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, BF16, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Row, Col, Row, BF16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, BF16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Row, Col, Row, F8, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, BF8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, @@ -102,9 +106,11 @@ using KernelTypesCompV3 = ::testing::Types< std::tuple< Col, Row, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, BF16, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Col, Row, Row, BF16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, BF16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Col, Row, Row, F8, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, BF8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, @@ -112,9 +118,11 @@ using KernelTypesCompV3 = ::testing::Types< std::tuple< Col, Col, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, BF16, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Col, Col, Row, BF16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, BF16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Col, Col, Row, F8, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, BF8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, From abaa06bb559f5e1fb309efea8ce011bcd8756df5 Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Wed, 12 Nov 2025 15:09:01 +0000 Subject: [PATCH 11/15] Add functionality and tests for fp16 x fp8 and fp8 x fp16 --- .../ops/common/determine_warp_prec_type.hpp | 16 ++++++++++++++++ .../gemm/test_gemm_pipeline_kernel_types.hpp | 8 ++++++++ 2 files changed, 24 insertions(+) diff --git a/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp b/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp index b8b5606360fe..ae11ff131461 100644 --- a/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp +++ b/projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp @@ -90,4 +90,20 @@ struct DetermineWarpPrecType using a_prec_type = ck_tile::bf16_t; using b_prec_type = ck_tile::bf16_t; }; + +// For fp8 x fp16, use fp8 +template <> +struct DetermineWarpPrecType +{ + using a_prec_type = ck_tile::fp8_t; + using b_prec_type = ck_tile::fp8_t; +}; + +// For fp16 x fp8, use fp16 +template <> +struct DetermineWarpPrecType +{ + using a_prec_type = ck_tile::half_t; + using b_prec_type = ck_tile::half_t; +}; }; // namespace ck_tile diff --git a/projects/composablekernel/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp b/projects/composablekernel/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp index 55dcf7a16e50..48206f06e192 100644 --- a/projects/composablekernel/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp +++ b/projects/composablekernel/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp @@ -80,48 +80,56 @@ using KernelTypesMemWmma = ::testing::Types< using KernelTypesCompV3 = ::testing::Types< std::tuple< Row, Row, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Row, Row, Row, F16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, F16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, BF16, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, BF16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, BF16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Row, Row, Row, F8, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, F8, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, F8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, F8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, BF8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Row, Row, BF8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Row, Col, Row, F16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, BF16, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, BF16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, BF16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Row, Col, Row, F8, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F8, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, F8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, BF8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Row, Col, Row, BF8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Col, Row, Row, F16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, BF16, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, BF16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, BF16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Col, Row, Row, F8, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F8, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, F8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, BF8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Row, Row, BF8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F16, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Col, Col, Row, F16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, BF16, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, BF16, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, BF16, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, INT8, INT8, INT32, INT32, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F8, F8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, + std::tuple< Col, Col, Row, F8, F16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F8, BF16, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F8, BF8, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, std::tuple< Col, Col, Row, F8, I4, F32, F16, I256, I256, I64, I32, I32, I16, Intrawave, CompV3>, From 3c4553a31e79878cc371855e867ea1e01c45871c Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Fri, 10 Oct 2025 08:39:04 +0000 Subject: [PATCH 12/15] Add type conversions to V4 pipeline, WIP! --- ...emm_pipeline_ag_bg_cr_comp_v4_default_policy.hpp | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v4_default_policy.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v4_default_policy.hpp index 777537a83a0a..583d94a9f122 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v4_default_policy.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v4_default_policy.hpp @@ -5,6 +5,7 @@ #include "ck_tile/core.hpp" #include "ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/tensor_layout.hpp" #include "ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp" @@ -22,6 +23,10 @@ struct GemmPipelineAgBgCrCompV4DefaultPolicy { using BlockWarps = typename Problem::BlockGemmShape::BlockWarps; using WarpTile = typename Problem::BlockGemmShape::WarpTile; + using ATypeToUse = typename DetermineWarpPrecType::a_prec_type; + using BTypeToUse = typename DetermineWarpPrecType::b_prec_type; constexpr index_t vector_size = DS_READ_TR_SIZE() / sizeof(typename Problem::ComputeDataType); @@ -33,8 +38,8 @@ struct GemmPipelineAgBgCrCompV4DefaultPolicy : vector_size * 4 == thread_elements ? WGAttrNumAccessEnum::Quad : WGAttrNumAccessEnum::Invalid; - using WarpGemm = WarpGemmDispatcher; - using BlockGemmPolicy = BlockGemmARegBRegCRegV1CustomPolicy; From d168d2f93fffe06eb0f1f3409bdd10a4f7f22db5 Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Thu, 18 Dec 2025 09:14:11 +0000 Subject: [PATCH 13/15] Refactor type conversions out of MakeBLdsBlockDescriptor, WIP! --- .../gemm_universal_pipeline_ag_bg_cr_policy.hpp | 11 ++++------- 1 file changed, 4 insertions(+), 7 deletions(-) diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp index 1ee334404bba..04f75fe97adf 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp @@ -321,15 +321,12 @@ struct UniversalGemmBasePolicy * @tparam Problem Gemm pipeline problem. * @return B tensor LDS block descriptor. */ - template + template > CK_TILE_DEVICE static constexpr auto MakeBLdsBlockDescriptor() { - using BLayout = remove_cvref_t; - constexpr bool IsBCastPolicyBeforeLDSWrite = IsBCastPolicyBeforeLDSWrite_v; - using BDataType = std::conditional_t; - + using BLayout = remove_cvref_t; + using BDataType = OverrideBDataType; constexpr index_t NPerBlock = Problem::BlockGemmShape::kN; constexpr index_t KPerBlock = Problem::BlockGemmShape::kK; From 1abad5a93bc9e0ce273724a669ac323a22ebdd39 Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Wed, 28 Jan 2026 14:24:47 +0000 Subject: [PATCH 14/15] Add a changelog entry --- projects/composablekernel/CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/composablekernel/CHANGELOG.md b/projects/composablekernel/CHANGELOG.md index 370e9e4243ec..b26f91a20f29 100644 --- a/projects/composablekernel/CHANGELOG.md +++ b/projects/composablekernel/CHANGELOG.md @@ -30,6 +30,7 @@ Documentation for Composable Kernel available at [https://rocm.docs.amd.com/proj ## Composable Kernel 1.2.0 for ROCm 7.2.0 ### Added +* Added support for fp16 x fp8, bf16 x fp8, fp8 x fp16, and fp8 x bf16 for the V3 pipeline * Added tests for f8 x bf8 on CompV3, and f8 x bf8 with K_BlockSize 32 on CompV4 * Added CK-Tile dispatcher - a unified kernel dispatch, code generation and architecture-based kernel filtering system with with C++ and Python frontends starting with GEMM support. * Added support for bf16 data type to grouped_gemm and grouped_gemm_preshuffle. From 5a304ff36658045755fb9b81bb0f2725cb676e2f Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Thu, 29 Jan 2026 08:03:57 +0000 Subject: [PATCH 15/15] Add include statements added by remod.py --- .../include/ck_tile/ops/add_rmsnorm2d_rdquant.hpp | 1 + .../composablekernel/include/ck_tile/ops/batched_contraction.hpp | 1 + .../composablekernel/include/ck_tile/ops/batched_transpose.hpp | 1 + projects/composablekernel/include/ck_tile/ops/elementwise.hpp | 1 + projects/composablekernel/include/ck_tile/ops/epilogue.hpp | 1 + projects/composablekernel/include/ck_tile/ops/flatmm.hpp | 1 + projects/composablekernel/include/ck_tile/ops/fmha.hpp | 1 + projects/composablekernel/include/ck_tile/ops/fused_moe.hpp | 1 + projects/composablekernel/include/ck_tile/ops/gemm.hpp | 1 + projects/composablekernel/include/ck_tile/ops/gemm_quant.hpp | 1 + .../composablekernel/include/ck_tile/ops/grouped_convolution.hpp | 1 + .../composablekernel/include/ck_tile/ops/image_to_column.hpp | 1 + projects/composablekernel/include/ck_tile/ops/layernorm2d.hpp | 1 + projects/composablekernel/include/ck_tile/ops/norm_reduce.hpp | 1 + projects/composablekernel/include/ck_tile/ops/permute.hpp | 1 + projects/composablekernel/include/ck_tile/ops/pooling.hpp | 1 + projects/composablekernel/include/ck_tile/ops/reduce.hpp | 1 + projects/composablekernel/include/ck_tile/ops/rmsnorm2d.hpp | 1 + projects/composablekernel/include/ck_tile/ops/smoothquant.hpp | 1 + projects/composablekernel/include/ck_tile/ops/softmax.hpp | 1 + projects/composablekernel/include/ck_tile/ops/sparse_attn.hpp | 1 + projects/composablekernel/include/ck_tile/ops/topk.hpp | 1 + projects/composablekernel/include/ck_tile/ops/topk_softmax.hpp | 1 + 23 files changed, 23 insertions(+) diff --git a/projects/composablekernel/include/ck_tile/ops/add_rmsnorm2d_rdquant.hpp b/projects/composablekernel/include/ck_tile/ops/add_rmsnorm2d_rdquant.hpp index aa0f632c2169..a62bbe981cca 100644 --- a/projects/composablekernel/include/ck_tile/ops/add_rmsnorm2d_rdquant.hpp +++ b/projects/composablekernel/include/ck_tile/ops/add_rmsnorm2d_rdquant.hpp @@ -7,6 +7,7 @@ #include "ck_tile/ops/add_rmsnorm2d_rdquant/pipeline/add_rmsnorm2d_rdquant_fwd_pipeline_one_pass.hpp" #include "ck_tile/ops/add_rmsnorm2d_rdquant/pipeline/add_rmsnorm2d_rdquant_fwd_pipeline_problem.hpp" #include "ck_tile/ops/add_rmsnorm2d_rdquant/pipeline/add_rmsnorm2d_rdquant_fwd_pipeline_three_pass.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/batched_contraction.hpp b/projects/composablekernel/include/ck_tile/ops/batched_contraction.hpp index 9c90db67eddc..71919b61873d 100644 --- a/projects/composablekernel/include/ck_tile/ops/batched_contraction.hpp +++ b/projects/composablekernel/include/ck_tile/ops/batched_contraction.hpp @@ -5,6 +5,7 @@ #include "ck_tile/ops/batched_contraction/kernel/batched_contraction_kernel.hpp" #include "ck_tile/ops/batched_contraction/pipeline/batched_contraction_problem.hpp" #include "ck_tile/ops/batched_contraction/utils/tensor_descriptor_utils.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/batched_transpose.hpp b/projects/composablekernel/include/ck_tile/ops/batched_transpose.hpp index 9cac035c4457..924db5fb60e4 100644 --- a/projects/composablekernel/include/ck_tile/ops/batched_transpose.hpp +++ b/projects/composablekernel/include/ck_tile/ops/batched_transpose.hpp @@ -10,6 +10,7 @@ #include "ck_tile/ops/batched_transpose/pipeline/batched_transpose_pipeline.hpp" #include "ck_tile/ops/batched_transpose/pipeline/batched_transpose_policy.hpp" #include "ck_tile/ops/batched_transpose/pipeline/batched_transpose_problem.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/elementwise.hpp b/projects/composablekernel/include/ck_tile/ops/elementwise.hpp index bc72f3b0ba1f..2c0ae4ad093f 100644 --- a/projects/composablekernel/include/ck_tile/ops/elementwise.hpp +++ b/projects/composablekernel/include/ck_tile/ops/elementwise.hpp @@ -8,6 +8,7 @@ #include "ck_tile/ops/elementwise/pipeline/elementwise_pipeline_problem.hpp" #include "ck_tile/ops/elementwise/pipeline/elementwise_shape.hpp" #include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/epilogue.hpp b/projects/composablekernel/include/ck_tile/ops/epilogue.hpp index d1b38a8bca6f..0eb9e59e723c 100644 --- a/projects/composablekernel/include/ck_tile/ops/epilogue.hpp +++ b/projects/composablekernel/include/ck_tile/ops/epilogue.hpp @@ -10,6 +10,7 @@ #include "ck_tile/ops/epilogue/default_2d_and_dynamic_quant_epilogue.hpp" #include "ck_tile/ops/epilogue/default_2d_epilogue.hpp" #include "ck_tile/ops/epilogue/dynamic_quant_epilogue.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/flatmm.hpp b/projects/composablekernel/include/ck_tile/ops/flatmm.hpp index e08fac48c7e9..2e71957ac774 100644 --- a/projects/composablekernel/include/ck_tile/ops/flatmm.hpp +++ b/projects/composablekernel/include/ck_tile/ops/flatmm.hpp @@ -21,6 +21,7 @@ #include "ck_tile/ops/flatmm/pipeline/mx_flatmm_pipeline_agmem_bgmem_creg_v1.hpp" #include "ck_tile/ops/flatmm/pipeline/mx_flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp" #include "ck_tile/ops/flatmm/pipeline/tile_flatmm_shape.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/fmha.hpp b/projects/composablekernel/include/ck_tile/ops/fmha.hpp index 8a5d77bf462e..633a36b7ff68 100644 --- a/projects/composablekernel/include/ck_tile/ops/fmha.hpp +++ b/projects/composablekernel/include/ck_tile/ops/fmha.hpp @@ -61,6 +61,7 @@ #include "ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy.hpp" #include "ck_tile/ops/fmha/pipeline/tile_fmha_shape.hpp" #include "ck_tile/ops/fmha/pipeline/tile_fmha_traits.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/fused_moe.hpp b/projects/composablekernel/include/ck_tile/ops/fused_moe.hpp index 60f5bd1c4e35..2eb4abd64117 100644 --- a/projects/composablekernel/include/ck_tile/ops/fused_moe.hpp +++ b/projects/composablekernel/include/ck_tile/ops/fused_moe.hpp @@ -14,6 +14,7 @@ #include "ck_tile/ops/fused_moe/pipeline/fused_moegemm_traits.hpp" #include "ck_tile/ops/fused_moe/pipeline/moe_sorting_pipeline.hpp" #include "ck_tile/ops/fused_moe/pipeline/moe_sorting_policy.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/gemm.hpp b/projects/composablekernel/include/ck_tile/ops/gemm.hpp index a0ed2fe9dd60..8a81cb7f5244 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm.hpp @@ -80,6 +80,7 @@ #include "ck_tile/ops/gemm/warp/warp_gemm_smfmac_impl.hpp" #include "ck_tile/ops/gemm/warp/warp_wmma_gemm.hpp" #include "ck_tile/ops/gemm/warp/warp_wmma_gemm_gfx11_utils.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/gemm_quant.hpp b/projects/composablekernel/include/ck_tile/ops/gemm_quant.hpp index 891ab2e6fbd7..5ab7911941d6 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm_quant.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm_quant.hpp @@ -33,6 +33,7 @@ #include "ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_base_policy.hpp" #include "ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp" #include "ck_tile/ops/gemm_quant/pipeline/tile_gemm_quant_traits.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/grouped_convolution.hpp b/projects/composablekernel/include/ck_tile/ops/grouped_convolution.hpp index 5bc4f0c6a042..1a115f9489e6 100644 --- a/projects/composablekernel/include/ck_tile/ops/grouped_convolution.hpp +++ b/projects/composablekernel/include/ck_tile/ops/grouped_convolution.hpp @@ -12,6 +12,7 @@ #include "ck_tile/ops/grouped_convolution/utils/transform_conv_bwd_data_to_gemm.hpp" #include "ck_tile/ops/grouped_convolution/utils/transform_conv_bwd_weight_to_gemm.hpp" #include "ck_tile/ops/grouped_convolution/utils/transform_conv_fwd_to_gemm.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/image_to_column.hpp b/projects/composablekernel/include/ck_tile/ops/image_to_column.hpp index 07d99890869e..faa165a8b001 100644 --- a/projects/composablekernel/include/ck_tile/ops/image_to_column.hpp +++ b/projects/composablekernel/include/ck_tile/ops/image_to_column.hpp @@ -5,6 +5,7 @@ #include "ck_tile/ops/image_to_column/kernel/image_to_column_kernel.hpp" #include "ck_tile/ops/image_to_column/pipeline/block_image_to_column_problem.hpp" #include "ck_tile/ops/image_to_column/pipeline/tile_image_to_column_shape.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/layernorm2d.hpp b/projects/composablekernel/include/ck_tile/ops/layernorm2d.hpp index 8f9ab205ac45..2266c138729a 100644 --- a/projects/composablekernel/include/ck_tile/ops/layernorm2d.hpp +++ b/projects/composablekernel/include/ck_tile/ops/layernorm2d.hpp @@ -8,6 +8,7 @@ #include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_problem.hpp" #include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp" #include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_traits.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/norm_reduce.hpp b/projects/composablekernel/include/ck_tile/ops/norm_reduce.hpp index eae0ea14a337..9f572ff5cbb2 100644 --- a/projects/composablekernel/include/ck_tile/ops/norm_reduce.hpp +++ b/projects/composablekernel/include/ck_tile/ops/norm_reduce.hpp @@ -5,6 +5,7 @@ #include "ck_tile/ops/norm_reduce/block/block_norm_reduce.hpp" #include "ck_tile/ops/norm_reduce/block/block_norm_reduce_problem.hpp" #include "ck_tile/ops/norm_reduce/thread/thread_welford.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/permute.hpp b/projects/composablekernel/include/ck_tile/ops/permute.hpp index 4d37f4fbc12a..c7747a67e70c 100644 --- a/projects/composablekernel/include/ck_tile/ops/permute.hpp +++ b/projects/composablekernel/include/ck_tile/ops/permute.hpp @@ -4,6 +4,7 @@ #include "ck_tile/ops/permute/kernel/generic_permute_kernel.hpp" #include "ck_tile/ops/permute/pipeline/generic_petmute_problem.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/pooling.hpp b/projects/composablekernel/include/ck_tile/ops/pooling.hpp index faa77d53273e..43b24c7f8caf 100644 --- a/projects/composablekernel/include/ck_tile/ops/pooling.hpp +++ b/projects/composablekernel/include/ck_tile/ops/pooling.hpp @@ -6,6 +6,7 @@ #include "ck_tile/ops/pooling/pipeline/pool_default_policy.hpp" #include "ck_tile/ops/pooling/pipeline/pool_problem.hpp" #include "ck_tile/ops/pooling/pipeline/pool_shape.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/reduce.hpp b/projects/composablekernel/include/ck_tile/ops/reduce.hpp index b5e53283e485..e680d2574538 100644 --- a/projects/composablekernel/include/ck_tile/ops/reduce.hpp +++ b/projects/composablekernel/include/ck_tile/ops/reduce.hpp @@ -13,6 +13,7 @@ #include "ck_tile/ops/reduce/pipeline/reduce2d_default_policy.hpp" #include "ck_tile/ops/reduce/pipeline/reduce2d_problem.hpp" #include "ck_tile/ops/reduce/pipeline/reduce2d_shape.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/rmsnorm2d.hpp b/projects/composablekernel/include/ck_tile/ops/rmsnorm2d.hpp index f271be50068c..7ee67334d6b9 100644 --- a/projects/composablekernel/include/ck_tile/ops/rmsnorm2d.hpp +++ b/projects/composablekernel/include/ck_tile/ops/rmsnorm2d.hpp @@ -9,6 +9,7 @@ #include "ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_problem.hpp" #include "ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_pipeline_two_pass.hpp" #include "ck_tile/ops/rmsnorm2d/pipeline/rmsnorm2d_fwd_traits.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/smoothquant.hpp b/projects/composablekernel/include/ck_tile/ops/smoothquant.hpp index 4c2fe9bee434..ad984c033f02 100644 --- a/projects/composablekernel/include/ck_tile/ops/smoothquant.hpp +++ b/projects/composablekernel/include/ck_tile/ops/smoothquant.hpp @@ -8,6 +8,7 @@ #include "ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_one_pass.hpp" #include "ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_problem.hpp" #include "ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_two_pass.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/softmax.hpp b/projects/composablekernel/include/ck_tile/ops/softmax.hpp index c79ba06abfea..b810a57dda0d 100644 --- a/projects/composablekernel/include/ck_tile/ops/softmax.hpp +++ b/projects/composablekernel/include/ck_tile/ops/softmax.hpp @@ -4,6 +4,7 @@ #include "ck_tile/ops/softmax/block/block_softmax_2d.hpp" #include "ck_tile/ops/softmax/block/block_softmax_2d_problem.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/sparse_attn.hpp b/projects/composablekernel/include/ck_tile/ops/sparse_attn.hpp index c7c4171874aa..56e074794ddf 100644 --- a/projects/composablekernel/include/ck_tile/ops/sparse_attn.hpp +++ b/projects/composablekernel/include/ck_tile/ops/sparse_attn.hpp @@ -6,6 +6,7 @@ #include "ck_tile/ops/sparse_attn/kernel/fmha_fwd_vsa_kernel.hpp" #include "ck_tile/ops/sparse_attn/pipeline/block_fmha_pipeline_qr_ks_vs_async_jenga.hpp" #include "ck_tile/ops/sparse_attn/pipeline/block_fmha_pipeline_qr_ks_vs_async_vsa.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/topk.hpp b/projects/composablekernel/include/ck_tile/ops/topk.hpp index 474ba932270c..13d818174e62 100644 --- a/projects/composablekernel/include/ck_tile/ops/topk.hpp +++ b/projects/composablekernel/include/ck_tile/ops/topk.hpp @@ -4,6 +4,7 @@ #include "ck_tile/ops/topk/block/block_topk_stream_2d.hpp" #include "ck_tile/ops/topk/block/block_topk_stream_2d_problem.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp" diff --git a/projects/composablekernel/include/ck_tile/ops/topk_softmax.hpp b/projects/composablekernel/include/ck_tile/ops/topk_softmax.hpp index 066fbf5feea2..b7219511faa3 100644 --- a/projects/composablekernel/include/ck_tile/ops/topk_softmax.hpp +++ b/projects/composablekernel/include/ck_tile/ops/topk_softmax.hpp @@ -6,6 +6,7 @@ #include "ck_tile/ops/topk_softmax/pipeline/topk_softmax_warp_per_row_pipeline.hpp" #include "ck_tile/ops/topk_softmax/pipeline/topk_softmax_warp_per_row_policy.hpp" #include "ck_tile/ops/topk_softmax/pipeline/topk_softmax_warp_per_row_problem.hpp" +#include "ck_tile/ops/common/determine_warp_prec_type.hpp" #include "ck_tile/ops/common/generic_2d_block_shape.hpp" #include "ck_tile/ops/common/load_and_convert_tile.hpp" #include "ck_tile/ops/common/streamk_common.hpp"