Skip to content

bench: rewrite latency perf test into a ping-pong (backend-generic support)#242

Draft
dssgabriel wants to merge 2 commits into
kokkos:developfrom
dssgabriel:bench/generic-osu-lat
Draft

bench: rewrite latency perf test into a ping-pong (backend-generic support)#242
dssgabriel wants to merge 2 commits into
kokkos:developfrom
dssgabriel:bench/generic-osu-lat

Conversation

@dssgabriel

@dssgabriel dssgabriel commented Jun 4, 2026

Copy link
Copy Markdown
Collaborator

Description

This PR is a first draft of a backend-generic ping-pong latency benchmark for Kokkos Comm, fixing #244.

It replaces bench.mpi.osu_latency with a new bench.core.latency perf test that compares the Kokkos Comm APIs against direct backend calls (both MPI and NCCL).

One important difference from the existing MPI-only implementation in perf_tests/mpi/test_osu_latency.cpp is that this version measures a true ping-pong: rank 0 sends to rank 1, then receives the response. The existing MPI benchmark measured only one-way latency.

Changes

  • Affected areas: perf tests, build
  • Breaking change(s)? no

List:

  • Add backend-generic latency benchmark in perf_tests/bench_latency.cpp.
  • Rewrite shared do_iteration helper in perf_tests/utils.hpp with high_resolution_clock instead of steady_clock, and remove MPI_Comm param.
  • Wire the new benchmark into perf_tests/CMakeLists.txt as bench.core.latency.
  • Update kc_add_perf_test so core perf tests can link the same Kokkos Comm/MPI/NCCL support used by core unit tests (ugly but avoids duplicating NCCL init logic across unit and perf tests).

Checklist

  • Tests are up-to-date
  • Documentation is up-to-date

@dssgabriel dssgabriel self-assigned this Jun 4, 2026
@dssgabriel dssgabriel added C-enhancement Category: an enhancement T-performance Topic: something related to performance A-perf-tests Area: KokkosComm performance tests labels Jun 4, 2026
@dssgabriel

Copy link
Copy Markdown
Collaborator Author

@nicoleavans This is the PR with the backend-agnostic OSU lat implementation. For some reason, I cannot add you as a reviewer 🤷‍♂️

And here are the results we've obtained on a GH200 + Bull eXascale Interconnect v2 (BXI) system:

image image

@dssgabriel
dssgabriel marked this pull request as draft June 4, 2026 17:43
@dssgabriel
dssgabriel force-pushed the bench/generic-osu-lat branch from bb2cb9d to c26ca59 Compare June 5, 2026 14:32
Signed-off-by: Gabriel Dos Santos <gabriel.dossantos@cea.fr>
@dssgabriel
dssgabriel force-pushed the bench/generic-osu-lat branch from c26ca59 to c81b0fd Compare June 5, 2026 15:50

@cedricchevalier19 cedricchevalier19 left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks ok, just a few comments.

Comment thread perf_tests/bench_osu_latency.cpp Outdated

using DES = Kokkos::DefaultExecutionSpace;
#ifdef KOKKOSCOMM_ENABLE_NCCL
// If NCCL is enabled, then `KOKKOS_ENABLE_CUDA` must be defined.

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You can add a check or change must be by is.

Comment on lines +254 to +261
BENCHMARK(bench_KC_MpiSpace)->UseManualTime()->Unit(benchmark::kMicrosecond)->RangeMultiplier(8)->Range(1, 1 << 28);
BENCHMARK(bench_KC_mpi)->UseManualTime()->Unit(benchmark::kMicrosecond)->RangeMultiplier(8)->Range(1, 1 << 28);
BENCHMARK(bench_MPI)->UseManualTime()->Unit(benchmark::kMicrosecond)->RangeMultiplier(8)->Range(1, 1 << 28);
BENCHMARK(bench_MPI_nb)->UseManualTime()->Unit(benchmark::kMicrosecond)->RangeMultiplier(8)->Range(1, 1 << 28);
#ifdef KOKKOSCOMM_ENABLE_NCCL
BENCHMARK(bench_KC_NcclSpace)->UseManualTime()->Unit(benchmark::kMicrosecond)->RangeMultiplier(8)->Range(1, 1 << 28);
BENCHMARK(bench_KC_nccl)->UseManualTime()->Unit(benchmark::kMicrosecond)->RangeMultiplier(8)->Range(1, 1 << 28);
BENCHMARK(bench_NCCL)->UseManualTime()->Unit(benchmark::kMicrosecond)->RangeMultiplier(8)->Range(1, 1 << 28);

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can we group a little bit the post treatment?
->Unit(benchmark::kMicrosecond)->RangeMultiplier(8)->Range(1, 1 << 28); is quite heavy to read.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree, will look into how we can achieve this.

Alternatively, we may consider porting all our test/benchmark infrastructure to a more modern framework, such as Catch2. We definitely need to rethink how we validate multi-rank tests anyway.

Comment thread perf_tests/CMakeLists.txt
Comment on lines +34 to +40
FetchContent_Declare(
fmt
GIT_REPOSITORY https://github.com/fmtlib/fmt.git
GIT_TAG
407c905e45ad75fc29bf0f9bb7c5c2fd3475976f #v12.1.0
)
FetchContent_MakeAvailable(fmt)

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am not a huge fan of using FetchContent. Can we test if it is available locally or can we avoid making it mandatory ?

Comment thread perf_tests/utils.hpp

double max_elapsed_second;
double elapsed_seconds = elapsed.count();
MPI_Allreduce(&elapsed_seconds, &max_elapsed_second, 1, MPI_DOUBLE, MPI_MAX, MPI_COMM_WORLD);

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ideally can we keep all ? Or can we at least compute the average too ?

@dssgabriel dssgabriel Jun 10, 2026

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Not sure if we can keep all (I don't know if it's possible to attach an array of results to the custom counters), but we can certainly compute the average (as well as min, median, stddev, etc.).

Note that the average isn't necessarily meaningful here, since it is the average latency of all ranks. I had tried to report all these metrics at some point, but the results varied so much across ranks that I didn't see the point in keeping them, so I reverted everything back to how it is here.

Comment thread perf_tests/utils.hpp
@dssgabriel

Copy link
Copy Markdown
Collaborator Author

Thanks for the review @cedricchevalier19.

I think this still needs quite a bit of work before we can merge it into Kokkos Comm, and I am not sure the "OSU latency" naming fits.
Specifically, OSU latency measures ping-pong, so all calls should have blocking semantics: rank 1 needs to actually receive rank 0's message before responding. The time for the round trip divided by two is the latency. That is not what the benchmarks in this PR implement, nor what the existing KokkosComm::mpi-based benchmarks do.
We should have a proper latency benchmark implementation and also consider adding a suite of bandwidth benchmarks (for P2P, AG, AR and A2A).

@dssgabriel
dssgabriel force-pushed the bench/generic-osu-lat branch 3 times, most recently from 4b8bd79 to 175faf4 Compare June 10, 2026 22:17
@dssgabriel dssgabriel changed the title wip(bench): backend-generic OSU latency benchmark wip(bench): backend-generic ping-pong latency benchmark Jun 10, 2026
Signed-off-by: Gabriel Dos Santos <gabriel.dossantos@cea.fr>
@dssgabriel
dssgabriel force-pushed the bench/generic-osu-lat branch from 175faf4 to 9ed7be1 Compare June 10, 2026 23:21
@dssgabriel

dssgabriel commented Jun 11, 2026

Copy link
Copy Markdown
Collaborator Author

Commands to run config/build/run this benchmark:

# Load dependencies via `module` or `spack`
spack load cmake cuda mpi nccl

# Fetch + configure + build + install Kokkos (if not available on your machine, else load via module/spack directly)
export KOKKOS_VERSION=5.1.1
export KOKKOS_DOWNLOAD_URL=https://github.com/kokkos/kokkos/releases/download/${KOKKOS_VERSION}
mkdir -p _deps && cd _deps && curl -sLO ${KOKKOS_DOWNLOAD_URL}/kokkos-${KOKKOS_VERSION}.tar.gz && cd ..
tar xzf _deps/kokkos-${KOKKOS_VERSION}.tar.gz -C _deps/
cmake -S _deps/kokkos-${KOKKOS_VERSION} -B _deps/kokkos-${KOKKOS_VERSION}-build -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_FLAGS="-mcpu=native" -DKokkos_ENABLE_CUDA=ON -DKokkos_ARCH_<SYSTEM_ARCH>=ON
cmake --build _deps/kokkos-${KOKKOS_VERSION}-build -j
cmake --install _deps/kokkos-${KOKKOS_VERSION}-build --prefix _deps/kokkos-${KOKKOS_VERSION}-install

# Configure Kokkos Comm with just-installed Kokkos
cmake -B build -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_FLAGS="-mcpu=native" -DCMAKE_CXX_COMPILER="$(pwd)/_deps/kokkos-${KOKKOS_VERSION}/bin/nvcc_wrapper" -DKokkos_ROOT="$(pwd)/_deps/kokkos-${KOKKOS_VERSION}-install" -DKokkosComm_ENABLE_NCCL=ON -DKokkosComm_ENABLE_PERFTESTS=ON

# Build
cmake --build build --target bench.core.latency -j

# Run the benchmark (via ctest)
ctest --test-dir build -R bench.core.latency
# (or manually)
mpirun -n 2 build/perf_tests/bench.core.latency
# (or via a job scheduler for internode/intranode; here assuming GH200 core counts to see the GPUs correctly)
srun -p <PARTITION> -N2 -n2 -c72 build/perf_tests/bench.core.latency
srun -p <PARTITION> -N1 -n2 -c144 build/perf_tests/bench.core.latency

Comment on lines +32 to +41
auto s_req = KC::send(comm, sv, 1);
auto r_req = KC::recv(comm, rv, 1);
s_req.wait();
r_req.wait();
} else {
auto r_req = KC::recv(comm, rv, 0);
// Actually wait for the ping reception before sending the pong
r_req.wait();
auto s_req = KC::send(comm, sv, 0);
s_req.wait();

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Perhaps we should add another comm to measure a ping from 1 :

sequenceDiagram
    0->>1: send1
    activate 0
    1->>0: send2
    deactivate 0
    activate 1
     0->>1: send 3
     deactivate 1
Loading

We time the grey boxes.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Good idea. How do we report that, though? Currently, only rank 0 writes the results (hence we reported "iteration" time as the max on all ranks).

@dssgabriel dssgabriel changed the title wip(bench): backend-generic ping-pong latency benchmark bench: rewrite latency perf test into a ping-pong (backend-generic support) Jun 11, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

A-perf-tests Area: KokkosComm performance tests C-enhancement Category: an enhancement T-performance Topic: something related to performance

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants