Skip to content

GitLab

Commit ec1b65e7 authored by Aleksander Dudek's avatar Aleksander Dudek
Browse files

[CK_TILE] Update the file structure

parent aa30ef56
Hide whitespace changes
Inline Side-by-side
Showing with 184 additions and 29 deletions
example/ck_tile/13_moe_sorting/moe_sorting_api.hpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include <string>
#include "ck_tile/core.hpp"
#include "ck_tile/host.hpp"
#include "ck_tile/ops/fused_moe.hpp"
#include "ck_tile/device/ops/fused_moe.hpp"
struct moe_sorting_trait
{
......
example/ck_tile/14_moe_smoothquant/moe_smoothquant.hpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "ck_tile/core.hpp"
#include "ck_tile/host/kernel_launch.hpp"
#include "ck_tile/ops/smoothquant.hpp"
#include "ck_tile/host/util/kernel_launch.hpp"
#include "ck_tile/device/ops/smoothquant.hpp"
#include <string>
template <typename DataType>
......
example/ck_tile/15_fused_moe/fused_moegemm.hpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "ck_tile/core.hpp"
#include "ck_tile/host/kernel_launch.hpp"
#include "ck_tile/ops/fused_moe.hpp"
#include "ck_tile/host/util/kernel_launch.hpp"
#include "ck_tile/device/ops/fused_moe.hpp"
#include <string>
// this is only a convenient structure for creating an example
......
example/ck_tile/15_fused_moe/fused_moesorting.hpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include <string>
#include "ck_tile/core.hpp"
#include "ck_tile/host.hpp"
#include "ck_tile/ops/fused_moe.hpp"
#include "ck_tile/device/ops/fused_moe.hpp"
struct fused_moesorting_trait
{
......
example/ck_tile/15_fused_moe/instances/fused_moegemm_api_internal.hpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "fused_moegemm_api_traits.hpp"
#include "ck_tile/ops/fused_moe.hpp"
#include "ck_tile/device/ops/fused_moe.hpp"
#include <iostream>
template <ck_tile::index_t... Is>
......
example/ck_tile/16_batched_gemm/batched_gemm.cpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <hip/hip_runtime.h>
......@@ -10,8 +10,8 @@
#include <tuple>
#include "ck_tile/core.hpp"
#include "ck_tile/ops/epilogue.hpp"
#include "ck_tile/ops/gemm.hpp"
#include "ck_tile/device/ops/epilogue.hpp"
#include "ck_tile/device/ops/gemm.hpp"
#include "ck_tile/host.hpp"
#include "batched_gemm.hpp"
......
example/ck_tile/16_batched_gemm/batched_gemm.hpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include <string>
#include "ck_tile/core.hpp"
#include "ck_tile/host/kernel_launch.hpp"
#include "ck_tile/ops/gemm/kernel/batched_gemm_kernel.hpp"
#include "ck_tile/host/util/kernel_launch.hpp"
#include "ck_tile/device/ops/gemm/kernel/batched_gemm_kernel.hpp"
template <typename DataType>
struct BatchedGemmTypeConfig;
......
example/ck_tile/17_grouped_gemm/grouped_gemm.cpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <hip/hip_runtime.h>
......@@ -11,8 +11,8 @@
#include <memory>
#include "ck_tile/core.hpp"
#include "ck_tile/ops/epilogue.hpp"
#include "ck_tile/ops/gemm.hpp"
#include "ck_tile/device/ops/epilogue.hpp"
#include "ck_tile/device/ops/gemm.hpp"
#include "ck_tile/host.hpp"
#include "grouped_gemm.hpp"
#include "utils.hpp"
......
example/ck_tile/17_grouped_gemm/grouped_gemm.hpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include <string>
#include "ck_tile/core.hpp"
#include "ck_tile/host/kernel_launch.hpp"
#include "ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp"
#include "ck_tile/host/util/kernel_launch.hpp"
#include "ck_tile/device/ops/gemm/kernel/grouped_gemm_kernel.hpp"
template <typename DataType>
struct GemmBasicTypeConfig;
......
include/ck_tile/README.md
View file @ ec1b65e7
......@@ -8,7 +8,7 @@
`ck_tile` is independently from the old ck, located under [/include/ck_tile](/include/ck_tile). You don't need to include anything from old CK, `ck_tile` has similiar (indeed almost the same) implementations for users to build operators. We will have a transition period to pull everything from old ck into `ck_tile`, stay tuned.
## component
`ck_tile` is splitted into several componenets including `core`, `host`, `ops/gemm`, `ops/fmha`... each component you only need to include a single header (e.g `#include "ck_tile/core.hpp"`, `#include "ck_tile/ops/fmha.hpp"`) then you are able to use the function/structure inside (different from old `ck`)
`ck_tile` is splitted into several componenets including `core`, `host`, `ops/gemm`, `ops/fmha`... each component you only need to include a single header (e.g `#include "ck_tile/core.hpp"`, `#include "ck_tile/device/ops/fmha.hpp"`) then you are able to use the function/structure inside (different from old `ck`)
**[core]**
`ck_tile/core` contains all the basic data structure and function to build the kernel, you can only include this header and build your own operators that utilizing all the basic building blocks introduced in ck.
......
include/ck_tile/algorithm.hpp 0 → 100644
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "ck_tile/core/algorithm/cluster_descriptor.hpp"
#include "ck_tile/core/algorithm/coordinate_transform.hpp"
#include "ck_tile/core/algorithm/indexing_adaptor.hpp"
#include "ck_tile/core/algorithm/space_filling_curve.hpp"
include/ck_tile/arch.hpp 0 → 100644
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "ck_tile/core/arch/amd_buffer_addressing.hpp"
#include "ck_tile/core/arch/arch.hpp"
#include "ck_tile/core/arch/generic_memory_space_atomic.hpp"
#include "ck_tile/core/arch/utility.hpp"
include/ck_tile/ops/common.hpp → include/ck_tile/config.hpp
View file @ ec1b65e7
......@@ -3,5 +3,4 @@
#pragma once
#include "ck_tile/ops/common/generic_2d_block_shape.hpp"
#include "ck_tile/ops/common/tensor_layout.hpp"
#include "ck_tile/core/config.hpp"
include/ck_tile/container.hpp 0 → 100644
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "ck_tile/core/container/array.hpp"
#include "ck_tile/core/container/container_helper.hpp"
#include "ck_tile/core/container/map.hpp"
#include "ck_tile/core/container/meta_data_buffer.hpp"
#include "ck_tile/core/container/multi_index.hpp"
#include "ck_tile/core/container/sequence.hpp"
#include "ck_tile/core/container/span.hpp"
#include "ck_tile/core/container/statically_indexed_array.hpp"
#include "ck_tile/core/container/thread_buffer.hpp"
#include "ck_tile/core/container/tuple.hpp"
include/ck_tile/device.hpp 0 → 100644
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "ck_tile/device/ops/common/generic_2d_block_shape.hpp"
#include "ck_tile/device/ops/common/tensor_layout.hpp"
#include "ck_tile/device/ops/common.hpp"
#include "ck_tile/device/ops/epilogue/cshuffle_epilogue.hpp"
#include "ck_tile/device/ops/epilogue/default_2d_epilogue.hpp"
#include "ck_tile/device/ops/epilogue/dynamic_quant_epilogue.hpp"
#include "ck_tile/device/ops/epilogue.hpp"
#include "ck_tile/device/ops/fmha/block/block_attention_bias_enum.hpp"
#include "ck_tile/device/ops/fmha/block/block_dropout.hpp"
#include "ck_tile/device/ops/fmha/block/block_masking.hpp"
#include "ck_tile/device/ops/fmha/block/block_position_encoding.hpp"
#include "ck_tile/device/ops/fmha/block/block_rotary_embedding.hpp"
#include "ck_tile/device/ops/fmha/block/page_block_navigator.hpp"
#include "ck_tile/device/ops/fmha/kernel/fmha_bwd_kernel.hpp"
#include "ck_tile/device/ops/fmha/kernel/fmha_fwd_appendkv_kernel.hpp"
#include "ck_tile/device/ops/fmha/kernel/fmha_fwd_appendkv_tile_partitioner.hpp"
#include "ck_tile/device/ops/fmha/kernel/fmha_fwd_kernel.hpp"
#include "ck_tile/device/ops/fmha/kernel/fmha_fwd_splitkv_combine_kernel.hpp"
#include "ck_tile/device/ops/fmha/kernel/fmha_fwd_splitkv_kernel.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_bwd_convert_dq.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_bwd_dot_do_o.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr_iglp.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_bwd_pipeline_default_policy.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_bwd_pipeline_enum.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_bwd_pipeline_problem.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_fwd_appendkv_pipeline.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_fwd_appendkv_pipeline_default_policy.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_fwd_splitkv_combine_pipeline.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_fwd_splitkv_combine_pipeline_default_policy.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs_default_policy.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_qr_ks_vs.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_qr_ks_vs_default_policy.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_enum.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_problem.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_async.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_async_default_policy.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_default_policy.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_fp8.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_qs_ks_vs.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_qs_ks_vs_default_policy.hpp"
#include "ck_tile/device/ops/fmha/pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy.hpp"
#include "ck_tile/device/ops/fmha/pipeline/tile_fmha_shape.hpp"
#include "ck_tile/device/ops/fmha/pipeline/tile_fmha_traits.hpp"
#include "ck_tile/device/ops/fmha.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_bgmem_creg_v1.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_bgmem_creg_v1_default_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_breg_creg_v1.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_breg_creg_v1_custom_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_breg_creg_v1_default_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_bsmem_creg_one_warp_v1.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_bsmem_creg_v1.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_bsmem_creg_v1_custom_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_bsmem_creg_v1_default_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_bsmem_creg_v2.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_bsmem_creg_v2_custom_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_areg_bsmem_creg_v2_default_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_asmem_breg_creg_v1.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_asmem_breg_creg_v1_custom_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_asmem_breg_creg_v1_default_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_asmem_bsmem_creg_v1.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_asmem_bsmem_creg_v1_custom_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_asmem_bsmem_creg_v1_default_policy.hpp"
#include "ck_tile/device/ops/gemm/block/block_gemm_problem.hpp"
#include "ck_tile/device/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp"
#include "ck_tile/device/ops/gemm/kernel/batched_gemm_kernel.hpp"
#include "ck_tile/device/ops/gemm/kernel/gemm_kernel.hpp"
#include "ck_tile/device/ops/gemm/kernel/gemm_tile_partitioner.hpp"
#include "ck_tile/device/ops/gemm/kernel/grouped_gemm_kernel.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_base.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_mem.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_pipeline_agmem_bgmem_creg_v1.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_pipeline_agmem_bgmem_creg_v1_default_policy.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_pipeline_agmem_bgmem_creg_v2.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_pipeline_agmem_bgmem_creg_v2_default_policy.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_pipeline_problem.hpp"
#include "ck_tile/device/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp"
#include "ck_tile/device/ops/gemm/pipeline/tile_gemm_shape.hpp"
#include "ck_tile/device/ops/gemm/pipeline/tile_gemm_traits.hpp"
#include "ck_tile/device/ops/gemm/warp/warp_gemm.hpp"
#include "ck_tile/device/ops/gemm/warp/warp_gemm_attribute_mfma.hpp"
#include "ck_tile/device/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp"
#include "ck_tile/device/ops/gemm/warp/warp_gemm_dispatcher.hpp"
#include "ck_tile/device/ops/gemm/warp/warp_gemm_impl.hpp"
#include "ck_tile/device/ops/gemm.hpp"
#include "ck_tile/device/ops/layernorm2d/kernel/layernorm2d_fwd_kernel.hpp"
#include "ck_tile/device/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_default_policy.hpp"
#include "ck_tile/device/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_one_pass.hpp"
#include "ck_tile/device/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_problem.hpp"
#include "ck_tile/device/ops/layernorm2d/pipeline/layernorm2d_fwd_pipeline_two_pass.hpp"
#include "ck_tile/device/ops/layernorm2d/pipeline/layernorm2d_fwd_traits.hpp"
#include "ck_tile/device/ops/layernorm2d.hpp"
#include "ck_tile/device/ops/reduce/block/block_reduce.hpp"
#include "ck_tile/device/ops/reduce/block/block_reduce2d.hpp"
#include "ck_tile/device/ops/reduce/block/block_reduce2d_default_policy.hpp"
#include "ck_tile/device/ops/reduce/block/block_reduce2d_problem.hpp"
#include "ck_tile/device/ops/reduce.hpp"
include/ck_tile/device/ops/common.hpp 0 → 100644
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "ck_tile/device/ops/common/generic_2d_block_shape.hpp"
#include "ck_tile/device/ops/common/tensor_layout.hpp"
include/ck_tile/ops/common/generic_2d_block_shape.hpp → include/ck_tile/device/ops/common/generic_2d_block_shape.hpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
......
include/ck_tile/ops/common/tensor_layout.hpp → include/ck_tile/device/ops/common/tensor_layout.hpp
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
......
include/ck_tile/device/ops/epilogue.hpp 0 → 100644
View file @ ec1b65e7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "ck_tile/device/ops/epilogue/cshuffle_epilogue.hpp"
#include "ck_tile/device/ops/epilogue/default_2d_epilogue.hpp"
#include "ck_tile/device/ops/epilogue/dynamic_quant_epilogue.hpp"
#include "ck_tile/device/ops/common/generic_2d_block_shape.hpp"
#include "ck_tile/device/ops/common/tensor_layout.hpp"
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment