"...composable_kernel_rocm.git" did not exist on "2a0e5439e176fd5063c1c39fc1e14bd68e0f6796"
Commit 8f360206 authored by Po Yen Chen's avatar Po Yen Chen
Browse files

Only use vllm kcache layout in group mode

parent ead9c3cb
...@@ -770,9 +770,10 @@ bool run(const ck_tile::ArgParser& arg_parser) ...@@ -770,9 +770,10 @@ bool run(const ck_tile::ArgParser& arg_parser)
q_buf.ToDevice(q_host.data()); q_buf.ToDevice(q_host.data());
if(0 < page_block_size) if(mode == mode_enum::group && 0 < page_block_size)
{ {
if (!(i_perm && !is_v_rowmajor)) { if(!(i_perm && !is_v_rowmajor))
{
std::cerr << "make sure input layout is correct" << std::endl; std::cerr << "make sure input layout is correct" << std::endl;
return false; return false;
} }
......
...@@ -636,7 +636,7 @@ struct FmhaFwdSplitKVKernel ...@@ -636,7 +636,7 @@ struct FmhaFwdSplitKVKernel
const auto make_k_dram = [&](const KDataType* data, index_t height) { const auto make_k_dram = [&](const KDataType* data, index_t height) {
auto k_dram_naive = [&] { auto k_dram_naive = [&] {
if constexpr(kIsPagedKV) if constexpr(kIsGroupMode && kIsPagedKV)
{ {
constexpr index_t vector_size = 16 / sizeof(KDataType); constexpr index_t vector_size = 16 / sizeof(KDataType);
// (hdim_q/vector_size, page_block_size, vector_size) // (hdim_q/vector_size, page_block_size, vector_size)
......
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