Repository navigation
Use Kokkos instead of OpenMP - #11440
ArturBieniek4 wants to merge 22 commits into
Conversation
Build Kokkos 4.7.02 from source with a local Bazel integration and link it into GPL by default. Use OpenMP for host execution by default while allowing Threads and Serial backends. Keep CUDA builds connected to the external GPU-enabled Kokkos installation. Signed-off-by: Artur Bieniek <abieniek@antmicro.com>
Add Kokkos 4.7.02 to the dependency installer and export its installation prefix for CMake package discovery. Support OpenMP, Threads, and Serial through -kokkos-backend, preserving an existing Threads backend unless explicitly overridden. Reuse matching host installations and configure libomp discovery on macOS. Signed-off-by: Artur Bieniek <abieniek@antmicro.com>
Move Kokkos initialization and finalization into a common GPL runtime available to both host and GPU backends. Preserve externally managed Kokkos initialization and finalize only when GPL initialized it. Require Kokkos in CMake and link the runtime in both build systems. Signed-off-by: Artur Bieniek <abieniek@antmicro.com>
Convert native GPL parallel loops to Kokkos::parallel_for using the default host execution space and a shared HostRange policy alias. Remove direct OpenMP API usage from the placement code. Centralize execution-space creation in the Kokkos runtime. Honor requested thread counts with the OpenMP backend while retaining initialization-time pool sizing for Threads. Compile Kokkos consumers in CUDA mode when using a CUDA-enabled Kokkos installation. Signed-off-by: Artur Bieniek <abieniek@antmicro.com>
Compile MBFF for the host only under Clang CUDA to avoid unsupported device-side Boost.Math operations. Enable Boost empty-base optimization consistently across C++ and Clang CUDA translation units to prevent container layout mismatches. Use a supported string allocation label in the GPU wirelength test and clarify comments about Kokkos host execution. Signed-off-by: Artur Bieniek <abieniek@antmicro.com>
Pin the documented KokkosFFT checkout to v2.0.0. Use static-library CMake compiler probes with the hermetic toolchain and correct the benchmark option to KokkosFFT_ENABLE_BENCHMARK. Signed-off-by: Artur Bieniek <abieniek@antmicro.com>
Update the Bazel dependency, CMake requirement, dependency installer, and documented GPU build to Kokkos 5.2.2. Match the generated Bazel host configuration to the new release while retaining OpenMP, Threads, and Serial backend selection. Replace View::HostMirror with host_mirror_type for Kokkos 5 compatibility and remove the obsolete deprecated-code option from the build instructions. Signed-off-by: Artur Bieniek <abieniek@antmicro.com>
Remove wrapper scopes introduced during the Kokkos loop conversion to reduce indentation and simplify the code. Reuse execution-space variables for loops in the same scope with matching thread counts. Signed-off-by: Artur Bieniek <abieniek@antmicro.com>
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
There was a problem hiding this comment.
Welcome to OpenROAD! Thanks for opening your first PR.
Before we review:
- Contribution Guide: https://openroad.readthedocs.io/en/latest/contrib/GettingInvolved.html
- Build Instructions: https://openroad.readthedocs.io/en/latest/contrib/BuildWithCMake.html
Please ensure:
- CI passes
- Code is properly formatted
- Tests are included where applicable
A maintainer will review shortly!
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Code Review
This pull request replaces OpenMP parallel loops and reductions with Kokkos-based implementations across the gpl tool, integrating Kokkos 5.2.2 into both the CMake and Bazel build configurations. It also updates the dependency installer script to automate Kokkos installation. The review feedback identifies several critical issues: a compilation error in kokkosRuntime.cpp because the OpenMP execution space constructor does not accept a thread count argument, a Bazel analysis error in bazel/kokkos/BUILD due to an undeclared @kokkos repository reference, a type mismatch compilation error in nesterovBase.cpp where a float reduction is bound to a double variable, and a potential failure in DependencyInstaller.sh if the kokkos directory already exists.
| Kokkos::DefaultHostExecutionSpace hostExecutionSpace( | ||
| [[maybe_unused]] int num_threads) | ||
| { | ||
| ensureKokkosInitialized(); | ||
| #ifdef KOKKOS_ENABLE_OPENMP | ||
| return Kokkos::DefaultHostExecutionSpace(num_threads); | ||
| #else | ||
| return Kokkos::DefaultHostExecutionSpace(); | ||
| #endif | ||
| } |
There was a problem hiding this comment.
Kokkos::OpenMP (and thus Kokkos::DefaultHostExecutionSpace when OpenMP is enabled) does not have a constructor that accepts an int for the number of threads. This will cause a compilation error when KOKKOS_ENABLE_OPENMP is defined.
In Kokkos, the number of threads for the OpenMP backend is configured globally at initialization time (via Kokkos::InitializationSettings::set_num_threads()) or via the OMP_NUM_THREADS environment variable. It cannot be set per execution space instance.
To fix this, the function should return Kokkos::DefaultHostExecutionSpace() without any arguments.
Kokkos::DefaultHostExecutionSpace hostExecutionSpace(
[[maybe_unused]] int num_threads)
{
ensureKokkosInitialized();
return Kokkos::DefaultHostExecutionSpace();
}| alias( | ||
| name = "kokkos", | ||
| actual = select({ | ||
| "//:gpu_cuda": "@kokkos", | ||
| "//conditions:default": "@kokkos_host//:kokkos", | ||
| }), | ||
| ) |
There was a problem hiding this comment.
In MODULE.bazel, the kokkos dependency is declared with repo_name = "kokkos_host". This renames the repository to kokkos_host within the module's scope, meaning the name kokkos is no longer bound to any repository.
However, in bazel/kokkos/BUILD, the alias target references @kokkos for the GPU build:
alias(
name = "kokkos",
actual = select({
"//:gpu_cuda": "@kokkos",
"//conditions:default": "@kokkos_host//:kokkos",
}),
)This will result in a Bazel analysis error: Repository '@kokkos' is not visible from this module because @kokkos is not declared in MODULE.bazel.
If the GPU build of Kokkos is intended to be provided by a local/system installation or a separate module extension, it must be explicitly declared and imported in MODULE.bazel (e.g., via use_repo). Otherwise, this reference will break the build when //:gpu_cuda is enabled.
| sumPhi_ = 0; | ||
| #pragma omp parallel for num_threads(nbc_->getNumThreads()) \ | ||
| reduction(+ : sumPhi_) | ||
| for (auto it = getBins().begin(); it < getBins().end(); ++it) { | ||
| auto& bin = *it; // old-style loop for old OpenMP | ||
| auto eFieldPair = fft_->getElectroField(bin.x(), bin.y()); | ||
| bin.setElectroField(eFieldPair.first, eFieldPair.second); | ||
|
|
||
| float electroPhi = fft_->getElectroPhi(bin.x(), bin.y()); | ||
| bin.setElectroPhi(electroPhi); | ||
|
|
||
| sumPhi_ += electroPhi | ||
| Kokkos::parallel_reduce( | ||
| "gpl::updateElectrostaticField", | ||
| HostRange(space, 0, getBins().size()), | ||
| [&](std::size_t index, float& phi_sum) { | ||
| auto it = getBins().begin() + index; | ||
| auto& bin = *it; // Host iteration over the placement data | ||
| auto eFieldPair = fft_->getElectroField(bin.x(), bin.y()); | ||
| bin.setElectroField(eFieldPair.first, eFieldPair.second); | ||
|
|
||
| float electroPhi = fft_->getElectroPhi(bin.x(), bin.y()); | ||
| bin.setElectroPhi(electroPhi); | ||
|
|
||
| phi_sum | ||
| += electroPhi | ||
| * static_cast<float>(bin.getNonPlaceArea() + bin.instPlacedArea() | ||
| + bin.getFillerArea()); | ||
| } | ||
| }, | ||
| Kokkos::Sum<float>(sumPhi_)); |
There was a problem hiding this comment.
In NesterovBase, sumPhi_ is typically declared as a double to maintain precision for the electrostatic potential sum.
Using float& phi_sum and Kokkos::Sum<float>(sumPhi_) will cause a compilation error due to type mismatch (binding a double variable to a float& reference). Even if it compiles, it introduces unnecessary precision loss by performing the reduction in single-precision float.
We should use double for the reduction to match the type of sumPhi_.
sumPhi_ = 0;
Kokkos::parallel_reduce(
"gpl::updateElectrostaticField",
HostRange(space, 0, getBins().size()),
[&](std::size_t index, double& phi_sum) {
auto it = getBins().begin() + index;
auto& bin = *it; // Host iteration over the placement data
auto eFieldPair = fft_->getElectroField(bin.x(), bin.y());
bin.setElectroField(eFieldPair.first, eFieldPair.second);
float electroPhi = fft_->getElectroPhi(bin.x(), bin.y());
bin.setElectroPhi(electroPhi);
phi_sum
+= electroPhi
* static_cast<double>(bin.getNonPlaceArea() + bin.instPlacedArea()
+ bin.getFillerArea());
},
Kokkos::Sum<double>(sumPhi_));| ( | ||
| cd "${BASE_DIR}" | ||
| _execute "Cloning Kokkos..." git clone --depth=1 -b "${KOKKOS_VERSION}" https://github.com/kokkos/kokkos.git |
There was a problem hiding this comment.
If the installer script is run multiple times, or if a previous run failed, the kokkos directory might already exist in ${BASE_DIR}. In this case, git clone will fail with fatal: destination path 'kokkos' already exists and is not an empty directory., causing the script to abort.
To make the installation idempotent and robust, we should remove any existing kokkos directory before cloning.
| ( | |
| cd "${BASE_DIR}" | |
| _execute "Cloning Kokkos..." git clone --depth=1 -b "${KOKKOS_VERSION}" https://github.com/kokkos/kokkos.git | |
| ( | |
| cd "${BASE_DIR}" | |
| rm -rf kokkos | |
| _execute "Cloning Kokkos..." git clone --depth=1 -b "${KOKKOS_VERSION}" https://github.com/kokkos/kokkos.git |
Default CC and CXX to the already installed GCC Toolset 13 when building common dependencies on RHEL 8 and compatible distributions. Preserve explicit compiler selections and cover both -all and -common invocations. This avoids selecting the system GCC 8.5, which Kokkos 5.2.2 does not support. Verified an OpenMP/Serial Kokkos build on Rocky Linux 8 with GCC 13.3.1. Signed-off-by: Artur Bieniek <abieniek@antmicro.com>
|
This looks good to me. I think the config select for different backends makes sense to me. Is this how the cmake version works? |
Yes, it's the same idea - in CMake the backend is chosen when Kokkos itself is built and Kokkos generates its config headers from that. Bazel builds Kokkos from source with no CMake step, so it writes the same headers based on the flag. For GPU, both builds use a prebuilt Kokkos, so the backend is defined by how that prebuilt Kokkos was configured. |
calewis
left a comment
calewis
left a comment
There was a problem hiding this comment.
I mainly wonder if we can take this opportunity to set threads once and stop caching and passing the thread count around in so many locations. I'm not just being pedantic and trying to tack on work to this PR, the strategy used in kokkosRuntime, doesn't work for the Thread backend and is inefficient for the OpenMP one. Also as proof that the current strategy is fragile, BinGrid doesn't seem to even cache the number of threads properly (an existing issue), unless I missed where it is being set.
| } | ||
| Kokkos::InitializationSettings settings; | ||
| settings.set_disable_warnings(true); | ||
| Kokkos::initialize(settings); |
There was a problem hiding this comment.
I think we should be setting the number of threads here and not in hostExecutionSpace. I realize that turns off dynamic thread sizing from tcl, but doing it this way without hwloc makes it so Kokkos::Threads backend will always silently be 1 thread.
I'll talk to @maliberty about whether we can lock in the number of threads before the first call to Kokkos::initialize or not.
There was a problem hiding this comment.
@maliberty and I discussed that the number of threads should be dynamic so setting it once here won't work.
There was a problem hiding this comment.
I'm afraid that dynamic changing of thread count in case of Threads backed won't work. Looks like the only way to set it is during Kokkos initialization via Kokkos::InitializationSettings and Kokkos cannot be reinitialized (both according to the docs: https://kokkos.org/kokkos-core-wiki/API/core/initialize_finalize/initialize.html Kokkos can be initialized at most once; subsequent calls are erroneous. and by inspecting the sources - it just aborts upon second call. The Threads backed implementation also doesn't expose any way to change thread count after initialization)
There was a problem hiding this comment.
I think that is fine, I think we expect 99% of users to be on OpenMP that is the best supported CPU backend for Kokkos anyways and it use to be the fastest CPU one (might still be). @maliberty okay with users of a not default backend losing the ability to dynamically change the number of threads (can still dynamically set it once, just can't change it once set)?
| { | ||
| ensureKokkosInitialized(); | ||
| #ifdef KOKKOS_ENABLE_OPENMP | ||
| return Kokkos::DefaultHostExecutionSpace(num_threads); |
There was a problem hiding this comment.
This is going to allocate new objects every time and be more expensive than strictly needed for small loops. I'm of the opinion we should be using return Kokkos::DefaultHostExecutionSpace(); every time, but we'll need to discuss with other maintainers first.
| } | ||
| } | ||
| } | ||
| const auto space = hostExecutionSpace(num_threads_); |
There was a problem hiding this comment.
BinGrid appears to never set this value (defaulting to 1), making all the code that uses it in BinGrid actually serial, this is an example of where I think each class tracking threads is leading to difficulty and a single set one global solves some problems. @maliberty WDYT?
There was a problem hiding this comment.
Sorry I forgot to leave this comment yesterday. I talked to @maliberty and @QuantamHD last night and we were in agreement that a better solution than reconstructing a host execution space via hostExecutionSpace(num_threads_); every single time you need a space would be to have GPL/nesterov construct a space once with the correct number of threads (ideally by reading the value that was most recently set from TCL) once and then reuse that space throughout.
AFAICT once global_placement has launched the user can no longer change the number of threads until after it returns so we shouldn't need to be reconstructing theses spaces or passing around num_threads_ as member variables at all. What do you think?
There was a problem hiding this comment.
I discussed this with @ArturBieniek4 and there wasn't any particular reason to create new execution space each time other than to keep it on par with OpenMP code this PR replaces that was setting number of threads for every loop
Will modify it to create execution space once and reuse it
There was a problem hiding this comment.
Let me know when that goes in please.
|
Please resolve the merge conflicts |
vvbandeira
left a comment
vvbandeira
left a comment
There was a problem hiding this comment.
Two build issues, both inline:
MODULE.bazel: thekokkosarchive_overrideonly applies in the root module, so Bazel builds that depend onopenroad(for example bazel-orfs) can no longer resolve it.etc/DependencyInstaller.sh: the Kokkos step runscmakefromPATHinstead of the installed${PREFIX:-/usr/local}/bin/cmake.
| bazel_dep(name = "kokkos", version = "5.2.2", repo_name = "kokkos_host") | ||
| archive_override( | ||
| module_name = "kokkos", | ||
| patch_strip = 1, | ||
| patches = ["//bazel/kokkos:0001-add-bazel-build.patch"], | ||
| sha256 = "491fe02cb127559dd3c996f6693686f1eab00a53fd5c184e18e663896884175e", | ||
| strip_prefix = "kokkos-5.2.2", | ||
| urls = ["https://github.com/kokkos/kokkos/archive/refs/tags/5.2.2.tar.gz"], | ||
| ) |
There was a problem hiding this comment.
This breaks Bazel builds that use openroad as a dependency instead of as the root module.
archive_override applies only in the root module (see the note near the freetype dep in this file). BCR has no kokkos module. So a downstream root such as bazel-orfs, which pulls openroad in through archive_override, now fails module resolution with kokkos@5.2.2 not found. To fix it there, every downstream root must repeat this override and copy the patch into its own repo, because patches must be labels in the root module. test/downstream/MODULE.bazel has the same problem, and no CI job builds it, so CI cannot catch this.
Suggestion: declare the host Kokkos the same way as the GPU @kokkos below, as an http_archive with a build_file. A repository rule also works when openroad is not the root module. The patch only adds MODULE.bazel and BUILD.bazel, so you can move its BUILD.bazel to something like //bazel/kokkos:BUILD.kokkos_host.bazel and remove the MODULE.bazel part. The root module already depends on bazel_skylib, rules_cc, and openmp, so the loads in that file still resolve, and the @kokkos_host//:backend / @kokkos_host//:openmp labels stay the same.
http_archive(
name = "kokkos_host",
build_file = "//bazel/kokkos:BUILD.kokkos_host.bazel",
sha256 = "d6557aaef39302282a15f9c770433d1fcdf4e961dfd6d9ed726b9d0d0f546b9f",
strip_prefix = "kokkos-5.2.2",
urls = ["https://github.com/kokkos/kokkos/releases/download/5.2.2/kokkos-5.2.2.tar.gz"],
)This example uses the same release asset and sha256 as the GPU @kokkos, so both paths fetch one tarball. Please check that the overlay builds against it; it only adds files, so I expect no difference.
There was a problem hiding this comment.
I modified it to use http_archive and verified that test/downstream/MODULE.bazel no longer fails on kokkos
| local cmake_bin=cmake | ||
| if [[ -n "${PREFIX}" && -x "${PREFIX}/bin/cmake" ]]; then | ||
| cmake_bin="${PREFIX}/bin/cmake" | ||
| fi |
There was a problem hiding this comment.
When PREFIX is not set, this runs cmake from PATH. The other installers in this script use ${PREFIX:-/usr/local}/bin/cmake, which is where _install_cmake installs it just before this step.
The RHEL base package list does not install cmake. With sudo, RHEL's default secure_path usually does not include /usr/local/bin. So sudo ./etc/DependencyInstaller.sh -common can fail at "Configuring Kokkos" with cmake: command not found. Docker CI runs as root with /usr/local/bin in PATH, so CI does not show this.
Suggestion, matching the Yosys step:
| local cmake_bin=cmake | |
| if [[ -n "${PREFIX}" && -x "${PREFIX}/bin/cmake" ]]; then | |
| cmake_bin="${PREFIX}/bin/cmake" | |
| fi | |
| local cmake_bin=${PREFIX:-/usr/local}/bin/cmake | |
| if [[ ! -f "${cmake_bin}" ]]; then | |
| cmake_bin="cmake" | |
| fi |
There was a problem hiding this comment.
applied the suggestion
Signed-off-by: Jakub Michalski <jmichalski@antmicro.com>
Signed-off-by: Jakub Michalski <jmichalski@antmicro.com>
Using more than 1 thread was giving different results due to different bin ordering and loss of precision on int64_t += float operations Signed-off-by: Jakub Michalski <jmichalski@antmicro.com>
Signed-off-by: Jakub Michalski <jmichalski@antmicro.com>
Signed-off-by: Jakub Michalski <jmichalski@antmicro.com>
Replace archive_override with http_archive to fix module resolution in downstream roots Signed-off-by: Jakub Michalski <jmichalski@antmicro.com>
Signed-off-by: Jakub Michalski <jmichalski@antmicro.com>
|
@calewis I modified it to reuse execution space. Its created in two places ( |
Signed-off-by: Jakub Michalski <jmichalski@antmicro.com>
Summary
Enable Kokkos by default in Bazel and add CMake dependency installer support. Replace native GPL OpenMP loops with Kokkos parallelism, preserving configurable host backends and GPU support. Upgrade Kokkos to 5.2.2 and fix CUDA build compatibility.
Impact
I ran a benchmark on i7-7820X. SMT & boost disabled, governor set to performance, pinned cores. Run order was random, with 5 second cooldown between attempts. 2 warmup runs, then 10 measured runs, results are medians of those 10 measured runs. Process time plus cooldowns: 21.78 hours.
Benchmark results:
My conclusion from the results is that it adds a bit of overhead which causes slowdowns on runs that are not taking long overall and when running on a single thread, whereas for larger designs and many threads there is relevant speedup.
Verification
./etc/Build.sh).