From 823ceed7fb12c4852f406a8f59be3f9d2af3f8b3 Mon Sep 17 00:00:00 2001 From: huizzhan Date: Wed, 6 Aug 2025 06:47:17 +0000 Subject: [PATCH] refine code --- .../CMakeLists.txt | 4 ++-- .../README.md | 0 .../block_gemm_asmem_bsmem_creg.hpp | 0 ...ock_gemm_asmem_bsmem_creg_default_policy.hpp | 0 ...pipeline_agmem_bgmem_creg_default_policy.hpp | 0 ..._grouped_topk_pipeline_agmem_bgmem_creg.hpp} | 2 +- .../config.h | 0 .../gemm_softmax_grouped_topk.cpp} | 17 ++++------------- .../gemm_softmax_grouped_topk.hpp} | 6 +++--- .../grid_gemm_softmax_grouped_topk.hpp} | 0 .../reference_gemm_softmax_grouped_topk.hpp} | 2 +- .../stream_config.hpp | 0 12 files changed, 11 insertions(+), 20 deletions(-) rename example/ck_tile/{39_gemm_softmax => 39_gemm_softmax_grouped_topk}/CMakeLists.txt (88%) rename example/ck_tile/{39_gemm_softmax => 39_gemm_softmax_grouped_topk}/README.md (100%) rename example/ck_tile/{39_gemm_softmax => 39_gemm_softmax_grouped_topk}/block_gemm_asmem_bsmem_creg.hpp (100%) rename example/ck_tile/{39_gemm_softmax => 39_gemm_softmax_grouped_topk}/block_gemm_asmem_bsmem_creg_default_policy.hpp (100%) rename example/ck_tile/{39_gemm_softmax => 39_gemm_softmax_grouped_topk}/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp (100%) rename example/ck_tile/{39_gemm_softmax/block_gemm_pipeline_agmem_bgmem_creg.hpp => 39_gemm_softmax_grouped_topk/block_gemm_softmax_grouped_topk_pipeline_agmem_bgmem_creg.hpp} (99%) rename example/ck_tile/{39_gemm_softmax => 39_gemm_softmax_grouped_topk}/config.h (100%) rename example/ck_tile/{39_gemm_softmax/gemm_softmax.cpp => 39_gemm_softmax_grouped_topk/gemm_softmax_grouped_topk.cpp} (94%) rename example/ck_tile/{39_gemm_softmax/gemm.hpp => 39_gemm_softmax_grouped_topk/gemm_softmax_grouped_topk.hpp} (97%) rename example/ck_tile/{39_gemm_softmax/grid_gemm.hpp => 39_gemm_softmax_grouped_topk/grid_gemm_softmax_grouped_topk.hpp} (100%) rename example/ck_tile/{39_gemm_softmax/reference_gemm.hpp => 39_gemm_softmax_grouped_topk/reference_gemm_softmax_grouped_topk.hpp} (95%) rename example/ck_tile/{39_gemm_softmax => 39_gemm_softmax_grouped_topk}/stream_config.hpp (100%) diff --git a/example/ck_tile/39_gemm_softmax/CMakeLists.txt b/example/ck_tile/39_gemm_softmax_grouped_topk/CMakeLists.txt similarity index 88% rename from example/ck_tile/39_gemm_softmax/CMakeLists.txt rename to example/ck_tile/39_gemm_softmax_grouped_topk/CMakeLists.txt index b250157523..e6b94ebd26 100755 --- a/example/ck_tile/39_gemm_softmax/CMakeLists.txt +++ b/example/ck_tile/39_gemm_softmax_grouped_topk/CMakeLists.txt @@ -1,9 +1,9 @@ -set(EXAMPLE_REDUCE "tile_example_basic_gemm_softmax") +set(EXAMPLE_REDUCE "tile_example_basic_gemm_softmax_grouped_topk") # not using add_example_executable() to add this target, since we don't want this to have # to be included in "make all/install/check" message("adding example ${EXAMPLE_REDUCE}") -add_executable(${EXAMPLE_REDUCE} EXCLUDE_FROM_ALL gemm_softmax.cpp) +add_executable(${EXAMPLE_REDUCE} EXCLUDE_FROM_ALL gemm_softmax_grouped_topk.cpp) target_include_directories(${EXAMPLE_REDUCE} PRIVATE ${CMAKE_CURRENT_LIST_DIR}) set(EXAMPLE_REDUCE_COMPILE_OPTIONS) diff --git a/example/ck_tile/39_gemm_softmax/README.md b/example/ck_tile/39_gemm_softmax_grouped_topk/README.md similarity index 100% rename from example/ck_tile/39_gemm_softmax/README.md rename to example/ck_tile/39_gemm_softmax_grouped_topk/README.md diff --git a/example/ck_tile/39_gemm_softmax/block_gemm_asmem_bsmem_creg.hpp b/example/ck_tile/39_gemm_softmax_grouped_topk/block_gemm_asmem_bsmem_creg.hpp similarity index 100% rename from example/ck_tile/39_gemm_softmax/block_gemm_asmem_bsmem_creg.hpp rename to example/ck_tile/39_gemm_softmax_grouped_topk/block_gemm_asmem_bsmem_creg.hpp diff --git a/example/ck_tile/39_gemm_softmax/block_gemm_asmem_bsmem_creg_default_policy.hpp b/example/ck_tile/39_gemm_softmax_grouped_topk/block_gemm_asmem_bsmem_creg_default_policy.hpp similarity index 100% rename from example/ck_tile/39_gemm_softmax/block_gemm_asmem_bsmem_creg_default_policy.hpp rename to example/ck_tile/39_gemm_softmax_grouped_topk/block_gemm_asmem_bsmem_creg_default_policy.hpp diff --git a/example/ck_tile/39_gemm_softmax/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp b/example/ck_tile/39_gemm_softmax_grouped_topk/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp similarity index 100% rename from example/ck_tile/39_gemm_softmax/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp rename to example/ck_tile/39_gemm_softmax_grouped_topk/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp diff --git a/example/ck_tile/39_gemm_softmax/block_gemm_pipeline_agmem_bgmem_creg.hpp b/example/ck_tile/39_gemm_softmax_grouped_topk/block_gemm_softmax_grouped_topk_pipeline_agmem_bgmem_creg.hpp similarity index 99% rename from example/ck_tile/39_gemm_softmax/block_gemm_pipeline_agmem_bgmem_creg.hpp rename to example/ck_tile/39_gemm_softmax_grouped_topk/block_gemm_softmax_grouped_topk_pipeline_agmem_bgmem_creg.hpp index b56169373e..6f219def57 100755 --- a/example/ck_tile/39_gemm_softmax/block_gemm_pipeline_agmem_bgmem_creg.hpp +++ b/example/ck_tile/39_gemm_softmax_grouped_topk/block_gemm_softmax_grouped_topk_pipeline_agmem_bgmem_creg.hpp @@ -18,7 +18,7 @@ namespace ck_tile { // B Tile Window: global memory // C Distributed tensor: register template -struct BlockGemmSoftmaxPipelineAGmemBGmemCReg +struct BlockGemmSoftmaxGroupedTopkPipelineAGmemBGmemCReg { using ADataType = remove_cvref_t; using BDataType = remove_cvref_t; diff --git a/example/ck_tile/39_gemm_softmax/config.h b/example/ck_tile/39_gemm_softmax_grouped_topk/config.h similarity index 100% rename from example/ck_tile/39_gemm_softmax/config.h rename to example/ck_tile/39_gemm_softmax_grouped_topk/config.h diff --git a/example/ck_tile/39_gemm_softmax/gemm_softmax.cpp b/example/ck_tile/39_gemm_softmax_grouped_topk/gemm_softmax_grouped_topk.cpp similarity index 94% rename from example/ck_tile/39_gemm_softmax/gemm_softmax.cpp rename to example/ck_tile/39_gemm_softmax_grouped_topk/gemm_softmax_grouped_topk.cpp index 2adbb907f8..babd502f3b 100755 --- a/example/ck_tile/39_gemm_softmax/gemm_softmax.cpp +++ b/example/ck_tile/39_gemm_softmax_grouped_topk/gemm_softmax_grouped_topk.cpp @@ -5,8 +5,8 @@ #include "config.h" #include "ck_tile/host.hpp" -#include "gemm.hpp" -#include "reference_gemm.hpp" +#include "gemm_softmax_grouped_topk.hpp" +#include "reference_gemm_softmax_grouped_topk.hpp" /* * Toy code of GEMM @@ -227,21 +227,13 @@ int main(int argc, char* argv[]) Ldb, Ldout, CElementFunction{})); - // auto pass = true; + bool rtn = true; if(verification) { - // reference gemm - // ck_tile::HostTensor c_host_ref(c_lengths, c_strides); - // reference_basic_gemm_softmax( - // a_host, b_host, c_host_ref); - // c_buf.FromDevice(c_host_dev.mData.data()); - // pass &= ck_tile::check_err(c_host_dev, c_host_ref); - // std::cout << "valid:" << (pass ? "y" : "n") << std::endl; - ck_tile::HostTensor value_ref(out_lengths, out_strides); ck_tile::HostTensor index_ref(out_lengths, out_strides); - reference_basic_gemm_softmax_topk( + reference_basic_gemm_softmax_grouped_topk( a_host, b_host, value_ref, index_ref, topk); value_buf.FromDevice(value_host_dev.mData.data()); index_buf.FromDevice(index_host_dev.mData.data()); @@ -285,6 +277,5 @@ int main(int argc, char* argv[]) std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s" << std::endl; - // return !pass; return rtn; } diff --git a/example/ck_tile/39_gemm_softmax/gemm.hpp b/example/ck_tile/39_gemm_softmax_grouped_topk/gemm_softmax_grouped_topk.hpp similarity index 97% rename from example/ck_tile/39_gemm_softmax/gemm.hpp rename to example/ck_tile/39_gemm_softmax_grouped_topk/gemm_softmax_grouped_topk.hpp index a58533eefb..9f5d5bf436 100755 --- a/example/ck_tile/39_gemm_softmax/gemm.hpp +++ b/example/ck_tile/39_gemm_softmax_grouped_topk/gemm_softmax_grouped_topk.hpp @@ -8,9 +8,9 @@ #include "ck_tile/ops/common.hpp" #include "ck_tile/ops/gemm/warp/warp_gemm.hpp" -#include "block_gemm_pipeline_agmem_bgmem_creg.hpp" +#include "block_gemm_softmax_grouped_topk_pipeline_agmem_bgmem_creg.hpp" #include "config.h" -#include "grid_gemm.hpp" +#include "grid_gemm_softmax_grouped_topk.hpp" namespace ck_tile { @@ -171,7 +171,7 @@ struct Gemm IndexType, kBlockSize, TileGemmShape>; - return BlockGemmSoftmaxPipelineAGmemBGmemCReg{}; + return BlockGemmSoftmaxGroupedTopkPipelineAGmemBGmemCReg{}; } }; diff --git a/example/ck_tile/39_gemm_softmax/grid_gemm.hpp b/example/ck_tile/39_gemm_softmax_grouped_topk/grid_gemm_softmax_grouped_topk.hpp similarity index 100% rename from example/ck_tile/39_gemm_softmax/grid_gemm.hpp rename to example/ck_tile/39_gemm_softmax_grouped_topk/grid_gemm_softmax_grouped_topk.hpp diff --git a/example/ck_tile/39_gemm_softmax/reference_gemm.hpp b/example/ck_tile/39_gemm_softmax_grouped_topk/reference_gemm_softmax_grouped_topk.hpp similarity index 95% rename from example/ck_tile/39_gemm_softmax/reference_gemm.hpp rename to example/ck_tile/39_gemm_softmax_grouped_topk/reference_gemm_softmax_grouped_topk.hpp index fd01780a53..a4544cdb03 100755 --- a/example/ck_tile/39_gemm_softmax/reference_gemm.hpp +++ b/example/ck_tile/39_gemm_softmax_grouped_topk/reference_gemm_softmax_grouped_topk.hpp @@ -7,7 +7,7 @@ #include "ck_tile/host/host_tensor.hpp" template -void reference_basic_gemm_softmax_topk(const ck_tile::HostTensor& a_m_k, +void reference_basic_gemm_softmax_grouped_topk(const ck_tile::HostTensor& a_m_k, const ck_tile::HostTensor& b_n_k, ck_tile::HostTensor& y_values, ck_tile::HostTensor& y_indices, diff --git a/example/ck_tile/39_gemm_softmax/stream_config.hpp b/example/ck_tile/39_gemm_softmax_grouped_topk/stream_config.hpp similarity index 100% rename from example/ck_tile/39_gemm_softmax/stream_config.hpp rename to example/ck_tile/39_gemm_softmax_grouped_topk/stream_config.hpp