diff --git a/test/python/ep/mscclpp_ep_bench.cu b/test/python/ep/mscclpp_ep_bench.cu index 46147dad..997cd5d4 100644 --- a/test/python/ep/mscclpp_ep_bench.cu +++ b/test/python/ep/mscclpp_ep_bench.cu @@ -29,6 +29,7 @@ #include #include #include +#include #include #include "api.cuh" @@ -292,33 +293,42 @@ int main(int argc, char** argv) { CUDA_CHECK(cudaMalloc(&d_layout, (size_t)Elocal * W * sizeof(int64_t))); CUDA_CHECK(cudaMalloc(&d_count, (size_t)Elocal * sizeof(int))); - // Inputs (content is immaterial to timing). Route each rank independently to - // random distinct experts so top-k selections span destination ranks. + // Inputs. Token payloads are immaterial to timing, but the top-k routing is + // generated with the SAME scheme as NCCL-EP's ep_bench (generateRandomTopkIndicesLL): + // per-token abs(randn)+1 scores, take the top-k experts by score, then mask 10 + // random (token, slot) positions with -1 to simulate dropped tokens. CUDA_CHECK(cudaMemset(d_x, 0, (size_t)T * H * sizeof(Bf16))); std::vector h_topk((size_t)T * K); - std::vector h_weights((size_t)T * K); - std::mt19937 rng(args.seed + rank); - std::uniform_int_distribution expertDist(0, E - 1); - std::uniform_real_distribution weightDist(0.5f, 1.5f); - for (int t = 0; t < T; ++t) { - for (int j = 0; j < K; ++j) { - int expert; - bool duplicate; - do { - expert = expertDist(rng); - duplicate = false; - for (int previous = 0; previous < j; ++previous) { - duplicate |= h_topk[(size_t)t * K + previous] == expert; - } - } while (duplicate); - h_topk[(size_t)t * K + j] = expert; - h_weights[(size_t)t * K + j] = weightDist(rng); + std::vector h_weights((size_t)T * K, 1.0f); + { + std::mt19937 gen(args.seed + rank); + std::normal_distribution dist(0.0f, 1.0f); + std::vector> scoreIdx(E); + for (int t = 0; t < T; ++t) { + for (int e = 0; e < E; ++e) scoreIdx[e] = {std::abs(dist(gen)) + 1.0f, e}; + std::partial_sort(scoreIdx.begin(), scoreIdx.begin() + K, scoreIdx.end(), + [](const auto& a, const auto& b) { return a.first > b.first; }); + for (int j = 0; j < K; ++j) h_topk[(size_t)t * K + j] = scoreIdx[j].second; + } + // Randomly mask 10 positions with -1 (simulates dropped tokens); mirrors + // ep_bench. Guarded on T > 0 so the distribution bound is valid. + if (T > 0) { + std::uniform_int_distribution tokenDist(0, T - 1); + std::uniform_int_distribution topkDist(0, K - 1); + for (int i = 0; i < 10; ++i) { + int ti = tokenDist(gen), ki = topkDist(gen); + h_topk[(size_t)ti * K + ki] = -1; + } } } CUDA_CHECK(cudaMemcpy(d_topk, h_topk.data(), h_topk.size() * sizeof(int64_t), cudaMemcpyHostToDevice)); CUDA_CHECK(cudaMemcpy(d_weights, h_weights.data(), h_weights.size() * sizeof(float), cudaMemcpyHostToDevice)); - const long long num_valid_selections = (long long)T * K; + // Byte accounting counts only valid selections (topk >= 0), matching ep_bench's + // calculateLowLatencyBytes after the -1 masking above. + long long num_valid_selections = 0; + for (size_t i = 0; i < h_topk.size(); ++i) + if (h_topk[i] >= 0) ++num_valid_selections; const double dispatchBytesPerToken = fp8Dispatch ? H + (H / 128) * sizeof(float) : H * sizeof(Bf16); const double disp_bytes = (double)num_valid_selections * dispatchBytesPerToken; const double comb_bytes = (double)num_valid_selections * H * sizeof(Bf16);