diff --git a/example/ck_tile/01_unified_attention/example_unified_attention.cpp b/example/ck_tile/01_unified_attention/example_unified_attention.cpp index 401dd90496..d4313d79d5 100644 --- a/example/ck_tile/01_unified_attention/example_unified_attention.cpp +++ b/example/ck_tile/01_unified_attention/example_unified_attention.cpp @@ -35,7 +35,10 @@ auto parse_cmd_args(int argc, char* argv[]) -> std::pair std::pair std::pair std::pair& query_lens_input, + const std::vector& kv_lens_input, + bool varlen) -> std::pair, std::vector> +{ + // If both query_lens and kv_lens are provided, return them directly + if(!query_lens_input.empty() && !kv_lens_input.empty()) + { + return std::make_pair(query_lens_input, kv_lens_input); + } + + std::vector query_lens; + std::vector kv_lens; + + if(!varlen) + { + // Fixed length mode: fill with max seqlen + query_lens.assign(batch, max_seqlen_q); + kv_lens.assign(batch, max_seqlen_kv); + } + else + { + // Variable length mode: generate random lengths up to max + std::random_device rd; + std::mt19937 gen(rd()); + std::uniform_int_distribution q_dist(1, max_seqlen_q); + std::uniform_int_distribution kv_dist(1, max_seqlen_kv); + + query_lens.resize(batch); + kv_lens.resize(batch); + + for(ck_tile::index_t i = 0; i < batch; ++i) + { + query_lens[i] = q_dist(gen); + kv_lens[i] = kv_dist(gen); + } + } + + return std::make_pair(query_lens, kv_lens); +} + struct Problem { explicit Problem(const ck_tile::ArgParser& args) @@ -82,10 +129,30 @@ struct Problem // TODO: support other GQA/MQA cases than just 4x nhead_q = nhead_kv * num_queries_per_kv; + ck_tile::index_t max_seqlen_q = args.get_int("s"); + ck_tile::index_t max_seqlen_kv = args.get_int("s_k"); + + if (max_seqlen_kv == -1) { + max_seqlen_kv = max_seqlen_q; + } + hdim = args.get_int("d"); query_lens = args.get_int_vec("query_lens"); kv_lens = args.get_int_vec("kv_lens"); assert(query_lens.size() == kv_lens.size() && "query_lens and kv_lens must have the same length b"); + batch = args.get_int("b"); + + bool varlen = args.get_bool("varlen"); + auto [query_lens_, kv_lens_] = seqlen_preprocess( + batch, + max_seqlen_q, + max_seqlen_kv, + query_lens, + kv_lens, + varlen); + + query_lens = query_lens_; + kv_lens = kv_lens_; batch = query_lens.size(); // Calculate scale_s @@ -436,7 +503,18 @@ bool run_impl(const Problem& problem, const RunConfig& run_config) std::cout << "[" << problem.data_type << "|"; std::cout << "] b:" << problem.batch << ", h:" << problem.nhead_q << "/" << problem.nhead_kv - << ", d:" << problem.hdim << ", mask:" << "causal mask" << std::fixed << ", " + << ", d:" << problem.hdim << ", scale_s:" << problem.scale_s + << ", query_lens:["; + for (size_t i = 0; i < problem.query_lens.size(); ++i) { + std::cout << problem.query_lens[i]; + if (i < problem.query_lens.size() - 1) std::cout << ","; + } + std::cout << "], kv_lens:["; + for (size_t i = 0; i < problem.kv_lens.size(); ++i) { + std::cout << problem.kv_lens[i]; + if (i < problem.kv_lens.size() - 1) std::cout << ","; + } + std::cout << "], mask:" << "causal mask" << std::fixed << ", " << std::setprecision(8) << time << " ms, " << std::setprecision(2) << tflops << " TFlops, " << std::setprecision(2) << (static_cast(mem) / 1e12 / (time / 1e3)) << " TB/s" << std::endl;