Skip to content

Commit a9147cf

Browse files
malibertymaliberty
authored
Merge pull request #10511 from ApeachM/gpl-gpu-fft
gpl: opt-in GPU acceleration via Kokkos
2 parents 1e21a07 + 214fa85 commit a9147cf

54 files changed

Lines changed: 8639 additions & 195 deletions

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎AUTHORS‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66
# source control.
77

88
Andrew Kennings
9+
Antmicro
910
Federal University of Rio Grande do Sul (UFRGS)
1011
Google LLC
1112
Iowa State University

‎CMakeLists.txt‎

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,13 @@ option(USE_SYSTEM_ABC "Use system shared ABC library" OFF)
4242
# Allow disabling tests
4343
option(ENABLE_TESTS "Enable OpenROAD tests" ON)
4444

45+
# Opt-in GPU acceleration via Kokkos. The actual compute backend (CUDA, HIP,
46+
# SYCL, or host-only OpenMP/Threads) is determined by the installed Kokkos
47+
# package; OpenROAD inspects Kokkos_ENABLE_* and turns on the matching CMake
48+
# language and dependencies automatically. See the per-module CMakeLists for
49+
# how individual subsystems wire their GPU sources.
50+
option(ENABLE_GPU "Enable GPU acceleration via Kokkos" OFF)
51+
4552
# Allow enabling address sanitizer
4653
option(ASAN "Enable Address Sanitizer" OFF)
4754

@@ -88,6 +95,13 @@ if(NOT CMAKE_BUILD_TYPE)
8895
set(CMAKE_BUILD_TYPE RELEASE)
8996
endif()
9097

98+
# GPU backend wiring (opt-in). All Kokkos / CUDA / HIP / SYCL detection,
99+
# compiler probing, and language enablement live in cmake/KokkosBackend.cmake
100+
# and are loaded only when the user opts in via ENABLE_GPU=ON.
101+
if(ENABLE_GPU)
102+
include(KokkosBackend)
103+
endif()
104+
91105
if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
92106
if(CMAKE_CXX_COMPILER_VERSION VERSION_LESS "8.3.0")
93107
message(FATAL_ERROR "Insufficient gcc version. Found ${CMAKE_CXX_COMPILER_VERSION}, but require >= 8.3.0.")

‎cmake/KokkosBackend.cmake‎

Lines changed: 158 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,158 @@
1+
# SPDX-License-Identifier: BSD-3-Clause
2+
# Copyright (c) 2026, The OpenROAD Authors
3+
4+
# Kokkos GPU backend wiring for OpenROAD. Included from the root
5+
# CMakeLists.txt only when ENABLE_GPU=ON; not loaded otherwise.
6+
#
7+
# Discovers the user's Kokkos install, inherits its compute backend, turns
8+
# on the matching CMake language so downstream targets can mark kernel
9+
# sources with set_source_files_properties(... LANGUAGE CUDA|HIP), and
10+
# applies the small set of nvcc / fmt / host-compiler workarounds that the
11+
# CUDA backend currently needs in modern Linux toolchains. Per-module
12+
# CMakeLists (e.g. src/gpl) key off ENABLE_GPU and Kokkos_ENABLE_*; they
13+
# do not need to call find_package(Kokkos) or enable_language() themselves.
14+
15+
find_package(Kokkos QUIET)
16+
if(NOT Kokkos_FOUND)
17+
message(FATAL_ERROR
18+
"OpenROAD: ENABLE_GPU=ON requires the Kokkos package to be "
19+
"installed and discoverable by CMake, but Kokkos was not found.\n"
20+
" - If Kokkos is already installed: pass "
21+
"-DKokkos_ROOT=/path/to/kokkos (or extend CMAKE_PREFIX_PATH).\n"
22+
" - If not: build and install Kokkos from "
23+
"https://github.com/kokkos/kokkos with the desired backend "
24+
"(CUDA / HIP / SYCL / OpenMP) and a target architecture that "
25+
"matches the host GPU.\n"
26+
" - A future etc/DependencyInstaller.sh -gpu option will "
27+
"automate this step.")
28+
endif()
29+
30+
# KokkosFFT — required by the gpl GPU FFT backend (src/gpl/src/gpu/dct.cpp).
31+
# A separate package from Kokkos core.
32+
find_package(KokkosFFT QUIET)
33+
if(NOT KokkosFFT_FOUND)
34+
message(FATAL_ERROR
35+
"ENABLE_GPU=ON requires KokkosFFT, which was not found.\n"
36+
" - Install KokkosFFT (https://github.com/kokkos/kokkos-fft) against\n"
37+
" your Kokkos build, then re-configure with -DKokkosFFT_ROOT=<prefix>.\n"
38+
" - A future etc/DependencyInstaller.sh -gpu will install Kokkos and\n"
39+
" KokkosFFT together.")
40+
endif()
41+
42+
message(STATUS "OpenROAD: GPU acceleration enabled (Kokkos ${Kokkos_VERSION})")
43+
44+
if(Kokkos_ENABLE_CUDA)
45+
# Auto-discover nvcc when the user has CUDA installed at a standard
46+
# location but their environment does not expose it on PATH (common
47+
# with IDE-launched configures: the bundled CMake does not inherit
48+
# the shell PATH). enable_language(CUDA) below would otherwise abort
49+
# with "No CMAKE_CUDA_COMPILER could be found" even though Kokkos's
50+
# find_package already located the toolkit.
51+
if(NOT DEFINED CMAKE_CUDA_COMPILER AND NOT DEFINED ENV{CUDACXX})
52+
find_program(_OPENROAD_NVCC nvcc
53+
HINTS ENV CUDA_HOME ENV CUDA_PATH ENV CUDA_ROOT
54+
/usr/local/cuda/bin
55+
/usr/local/cuda-13.0/bin
56+
/usr/local/cuda-12.8/bin /usr/local/cuda-12.0/bin
57+
/opt/cuda/bin
58+
)
59+
if(_OPENROAD_NVCC)
60+
set(CMAKE_CUDA_COMPILER "${_OPENROAD_NVCC}" CACHE FILEPATH "")
61+
message(STATUS "OpenROAD: auto-discovered nvcc at ${_OPENROAD_NVCC}")
62+
endif()
63+
endif()
64+
# nvcc < 13 cannot parse glibc 2.38+'s _Float128 type that ships with
65+
# gcc 13+'s C++ standard library headers (math.h template specialization
66+
# for __iseqsig_type<_Float128>). When a known-broken pairing is detected,
67+
# pin a compatible older g++ as the CUDA host compiler (the system C++
68+
# compiler stays unchanged for non-CUDA TUs). Override is always
69+
# available via -DCMAKE_CUDA_HOST_COMPILER or CUDAHOSTCXX.
70+
if(NOT DEFINED CMAKE_CUDA_HOST_COMPILER AND NOT DEFINED ENV{CUDAHOSTCXX}
71+
AND CMAKE_CXX_COMPILER_ID STREQUAL "GNU"
72+
AND CMAKE_CXX_COMPILER_VERSION VERSION_GREATER_EQUAL "13.0"
73+
AND _OPENROAD_NVCC)
74+
execute_process(
75+
COMMAND "${_OPENROAD_NVCC}" --version
76+
OUTPUT_VARIABLE _OPENROAD_NVCC_VERSION_OUTPUT
77+
ERROR_QUIET
78+
OUTPUT_STRIP_TRAILING_WHITESPACE)
79+
if(_OPENROAD_NVCC_VERSION_OUTPUT MATCHES "release ([0-9]+)")
80+
set(_OPENROAD_NVCC_MAJOR "${CMAKE_MATCH_1}")
81+
if(_OPENROAD_NVCC_MAJOR LESS 13)
82+
foreach(_OPENROAD_GXX_VER 12 11)
83+
find_program(_OPENROAD_CUDAHOST g++-${_OPENROAD_GXX_VER}
84+
HINTS /usr/bin /usr/local/bin)
85+
if(_OPENROAD_CUDAHOST)
86+
set(CMAKE_CUDA_HOST_COMPILER "${_OPENROAD_CUDAHOST}"
87+
CACHE FILEPATH "")
88+
message(STATUS
89+
"OpenROAD: pinning CUDA host compiler to "
90+
"${_OPENROAD_CUDAHOST} (nvcc ${_OPENROAD_NVCC_MAJOR}.x + "
91+
"glibc/gcc 13+ _Float128 compat)")
92+
break()
93+
endif()
94+
unset(_OPENROAD_CUDAHOST CACHE)
95+
endforeach()
96+
if(NOT DEFINED CMAKE_CUDA_HOST_COMPILER)
97+
message(FATAL_ERROR
98+
"OpenROAD: nvcc ${_OPENROAD_NVCC_MAJOR}.x cannot parse "
99+
"_Float128 declarations in glibc 2.38+ system headers used "
100+
"by gcc ${CMAKE_CXX_COMPILER_VERSION}, and no compatible "
101+
"g++-12 / g++-11 was found in /usr/bin or /usr/local/bin. "
102+
"Install one (e.g. apt install g++-12) or set "
103+
"-DCMAKE_CUDA_HOST_COMPILER=/path/to/older-g++ explicitly.")
104+
endif()
105+
endif()
106+
endif()
107+
endif()
108+
if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES OR "${CMAKE_CUDA_ARCHITECTURES}" STREQUAL "")
109+
if(DEFINED Kokkos_CUDA_ARCHITECTURES
110+
AND NOT "${Kokkos_CUDA_ARCHITECTURES}" STREQUAL "")
111+
set(CMAKE_CUDA_ARCHITECTURES "${Kokkos_CUDA_ARCHITECTURES}")
112+
else()
113+
message(FATAL_ERROR
114+
"OpenROAD: ENABLE_GPU=ON with Kokkos CUDA backend, but the "
115+
"Kokkos package does not advertise Kokkos_CUDA_ARCHITECTURES "
116+
"and CMAKE_CUDA_ARCHITECTURES was not provided. Set "
117+
"-DCMAKE_CUDA_ARCHITECTURES=<arch> explicitly (e.g. 89 for "
118+
"RTX 4070, 120 for RTX 5090) or rebuild Kokkos with the "
119+
"target architecture baked in.")
120+
endif()
121+
endif()
122+
enable_language(CUDA)
123+
find_package(CUDAToolkit REQUIRED)
124+
message(STATUS "OpenROAD: CUDA backend (arch=${CMAKE_CUDA_ARCHITECTURES})")
125+
# A GPU driver (the kernel module exposing libcuda.so.1) is needed only to
126+
# *run* CUDA code, never to build it -- nvcc cross-compiles device code on a
127+
# host with no GPU. Note its absence so the resulting libcuda.so.1 load
128+
# errors on this host (e.g. ctest, or running openroad) read as expected
129+
# rather than as a misconfiguration. This is informational only: a GPU build
130+
# on a driverless host is a supported cross-compile workflow, not an error.
131+
if(NOT EXISTS "/proc/driver/nvidia")
132+
message(STATUS
133+
"OpenROAD: no NVIDIA driver on this host -- GPU code is being "
134+
"cross-compiled. Run the GPU binaries and tests on a GPU machine.")
135+
endif()
136+
# nvcc 12.8 cannot parse fmt 11's nontype-template-parameter user-defined
137+
# literals (fmt/bundled/format.h: operator""_a with fixed_string). The
138+
# legacy literal fallback is still available; opt into it for CUDA TUs
139+
# only. Project-wide CXX compilation is unaffected.
140+
add_compile_definitions(
141+
$<$<COMPILE_LANGUAGE:CUDA>:FMT_USE_NONTYPE_TEMPLATE_ARGS=0>)
142+
# On aarch64, Boost's unordered_flat_map detects __ARM_NEON and includes
143+
# <arm_neon.h> for SIMD-accelerated hashing. nvcc cannot parse gcc's
144+
# arm_neon.h (it contains gcc-specific intrinsics), so disable the NEON
145+
# path for CUDA TUs. The CPU TUs (compiled by g++) are unaffected.
146+
if(CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64|ARM64")
147+
add_compile_definitions(
148+
$<$<COMPILE_LANGUAGE:CUDA>:BOOST_UNORDERED_DISABLE_NEON>)
149+
endif()
150+
elseif(Kokkos_ENABLE_HIP)
151+
enable_language(HIP)
152+
message(STATUS "OpenROAD: HIP backend")
153+
elseif(Kokkos_ENABLE_SYCL)
154+
message(STATUS "OpenROAD: SYCL backend (driven by Kokkos host compiler)")
155+
else()
156+
message(STATUS
157+
"OpenROAD: host-only Kokkos backend (Serial / OpenMP / Threads)")
158+
endif()

‎src/gpl/BUILD‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,13 +38,22 @@ cc_library(
3838
name = "gpl",
3939
srcs = [
4040
"src/AbstractGraphics.cpp",
41+
"src/backendContext.h",
4142
"src/clockBase.cpp",
4243
"src/clockBase.h",
44+
"src/densityGradient.cpp",
45+
"src/densityGradientBackend.h",
4346
"src/fft.cpp",
4447
"src/fft.h",
48+
"src/fftBackend.h",
4549
"src/fftsg.cpp",
4650
"src/fftsg2d.cpp",
51+
"src/gpu/deviceState.h",
52+
"src/gpu/nesterovDeviceContext.h",
53+
"src/gpu/regionDensityField.h",
4754
"src/graphicsNone.cpp",
55+
"src/hpwl.cpp",
56+
"src/hpwlBackend.h",
4857
"src/initialPlace.cpp",
4958
"src/initialPlace.h",
5059
"src/mbff.cpp",
@@ -57,6 +66,8 @@ cc_library(
5766
"src/solver.h",
5867
"src/timingBase.cpp",
5968
"src/timingBase.h",
69+
"src/wirelengthGradient.cpp",
70+
"src/wirelengthGradientBackend.h",
6071
],
6172
hdrs = [
6273
"include/gpl/Replace.h",

‎src/gpl/CMakeLists.txt‎

Lines changed: 88 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -34,6 +34,9 @@ add_library(gpl_lib
3434
src/fft.cpp
3535
src/fftsg.cpp
3636
src/fftsg2d.cpp
37+
src/hpwl.cpp
38+
src/wirelengthGradient.cpp
39+
src/densityGradient.cpp
3740
src/routeBase.cpp
3841
src/timingBase.cpp
3942
src/clockBase.cpp
@@ -42,6 +45,84 @@ add_library(gpl_lib
4245
src/mbff.cpp
4346
)
4447

48+
# --- HPWL & FFT backends: runtime switch (Strategy + Factory) ---
49+
# The CPU backends (CpuHpwlBackend in src/hpwl.cpp, CpuFftBackend in
50+
# src/fft.cpp, + the Ooura src/fftsg*.cpp) are always compiled. When
51+
# ENABLE_GPU=ON the Kokkos GPU backends in src/gpu/ are also compiled in;
52+
# makeHpwlBackend() / makeFftBackend() pick the backend per process at run
53+
# time (gpl::gpuEnabled(), driven by the ENABLE_GPU env var). ENABLE_GPU is a
54+
# compile definition gating the #ifdef in those two factories; the consumer
55+
# headers (nesterovBase.h, fft.h) stay preprocessor-free. gpu/ is a
56+
# file-layout subdirectory only (no nested CMakeLists.txt) so kernel build
57+
# settings stay in this module's CMakeLists with the rest of gpl_lib.
58+
if(ENABLE_GPU)
59+
target_sources(gpl_lib PRIVATE
60+
src/gpu/gpuHpwlBackend.cpp
61+
src/gpu/gpuRuntime.cpp
62+
src/gpu/gpuFftBackend.cpp
63+
src/gpu/poissonSolver.cpp
64+
src/gpu/dct.cpp
65+
src/gpu/deviceState.cpp
66+
src/gpu/regionDensityField.cpp
67+
src/gpu/gpuWirelengthGradientBackend.cpp
68+
src/gpu/wirelengthOp.cpp
69+
src/gpu/gpuDensityGradientBackend.cpp
70+
src/gpu/densityOp.cpp
71+
src/gpu/nesterovOp.cpp
72+
src/gpu/nesterovDeviceContext.cpp)
73+
target_compile_definitions(gpl_lib PRIVATE ENABLE_GPU)
74+
# nesterovBase.h and other private gpl headers live in src/; sources
75+
# under src/gpu/ need that on the include path explicitly because
76+
# the compiler's default same-dir lookup points into src/gpu/ instead.
77+
target_include_directories(gpl_lib PRIVATE src)
78+
# The src/gpu/ TUs are device kernels. gpu/gpuRuntime.cpp carries no device
79+
# code itself, but it includes <Kokkos_Core.hpp> for the lazy Kokkos
80+
# initialize()/finalize(): when Kokkos is built with the CUDA (or HIP)
81+
# backend, that header bakes KOKKOS_ENABLE_CUDA into its config and refuses
82+
# to compile under a plain host compiler (it requires __CUDACC__). The same
83+
# applies to src/fft.cpp, whose makeFftBackend() factory includes
84+
# gpu/gpuFftBackend.h (Kokkos-dependent) to construct a GpuFftBackend. All
85+
# such TUs are flagged with the device language to match the Kokkos backend.
86+
# src/hpwl.cpp stays a plain CXX TU — gpu/gpuHpwlBackend.h is Kokkos-free, so
87+
# its makeHpwlBackend() factory needs no device language.
88+
# src/fftsg.cpp / src/fftsg2d.cpp are pure C++ Ooura code — left as CXX.
89+
if(Kokkos_ENABLE_CUDA)
90+
set_source_files_properties(
91+
src/gpu/gpuHpwlBackend.cpp src/gpu/gpuRuntime.cpp src/gpu/gpuFftBackend.cpp
92+
src/gpu/poissonSolver.cpp src/gpu/dct.cpp src/gpu/deviceState.cpp
93+
src/gpu/regionDensityField.cpp
94+
src/gpu/gpuWirelengthGradientBackend.cpp src/gpu/wirelengthOp.cpp
95+
src/gpu/gpuDensityGradientBackend.cpp src/gpu/densityOp.cpp
96+
src/gpu/nesterovOp.cpp src/gpu/nesterovDeviceContext.cpp
97+
src/fft.cpp
98+
PROPERTIES LANGUAGE CUDA)
99+
elseif(Kokkos_ENABLE_HIP)
100+
set_source_files_properties(
101+
src/gpu/gpuHpwlBackend.cpp src/gpu/gpuRuntime.cpp src/gpu/gpuFftBackend.cpp
102+
src/gpu/poissonSolver.cpp src/gpu/dct.cpp src/gpu/deviceState.cpp
103+
src/gpu/regionDensityField.cpp
104+
src/gpu/gpuWirelengthGradientBackend.cpp src/gpu/wirelengthOp.cpp
105+
src/gpu/gpuDensityGradientBackend.cpp src/gpu/densityOp.cpp
106+
src/gpu/nesterovOp.cpp src/gpu/nesterovDeviceContext.cpp
107+
src/fft.cpp
108+
PROPERTIES LANGUAGE HIP)
109+
endif()
110+
# Disable FP contraction for kernels that share gpl_lib's compile
111+
# context so they stay bit-stable across compilers. Scoped to gpl_lib
112+
# but the CXX flag is also harmless on the existing CPU TUs.
113+
target_compile_options(gpl_lib PRIVATE
114+
$<$<COMPILE_LANGUAGE:CXX>:-ffp-contract=off>
115+
$<$<COMPILE_LANGUAGE:CUDA>:--fmad=false>
116+
$<$<COMPILE_LANGUAGE:HIP>:-ffp-contract=off>
117+
)
118+
target_link_libraries(gpl_lib Kokkos::kokkos KokkosFFT::fft)
119+
if(Kokkos_ENABLE_CUDA)
120+
# cuda runtime symbols are referenced from the CUDA TU; expose cudart
121+
# so that gpl_lib (and the openroad binary) link against libcudart.
122+
target_link_libraries(gpl_lib CUDA::cudart)
123+
endif()
124+
endif()
125+
45126
target_sources(gpl
46127
PRIVATE
47128
src/MakeReplace.cpp
@@ -60,6 +141,13 @@ target_include_directories(gpl_lib
60141
PUBLIC
61142
include
62143
${LEMON_INCLUDE_DIRS}
144+
PRIVATE
145+
# The PIMPL headers under src/gpu/ (deviceState.h, nesterovDeviceContext.h)
146+
# are included from src/nesterovBase.cpp on both ENABLE_GPU=ON and OFF
147+
# paths, and they need to find sibling headers like src/point.h. Add the
148+
# src/ directory to the private include path unconditionally; previously
149+
# it was only added inside the if(ENABLE_GPU) block.
150+
src
63151
)
64152

65153
target_link_libraries(gpl_lib

‎src/gpl/src/backendContext.h‎

Lines changed: 47 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,47 @@
1+
// SPDX-License-Identifier: BSD-3-Clause
2+
// Copyright (c) 2026, The OpenROAD Authors
3+
4+
// BackendContext — a single bundle of construction parameters passed to each
5+
// of the gpl Strategy backend factories (makeHpwlBackend,
6+
// makeWirelengthGradientBackend, makeDensityGradientBackend, makeFftBackend).
7+
//
8+
// Each factory consumes the subset of fields it needs and ignores the rest;
9+
// callers build one context per construction site and reuse it across the
10+
// four factory calls. Plain C++ — Kokkos types are forward-declared elsewhere
11+
// and pointers (DeviceState*, NesterovBase*, NesterovBaseCommon*,
12+
// RegionDensityField*) are only dereferenced inside backend translation units.
13+
14+
#pragma once
15+
16+
namespace gpl {
17+
18+
class DeviceState;
19+
class NesterovBase;
20+
class NesterovBaseCommon;
21+
class RegionDensityField;
22+
23+
struct BackendContext
24+
{
25+
// Owning / context pointers. nbc is required by the wirelength gradient
26+
// backend; nb is required by the density gradient backend; device_state is
27+
// borrowed by the GPU HPWL, wirelength-gradient, and density-gradient
28+
// backends (ignored by the CPU backends and by the GPU FFT backend, which
29+
// uses region_field instead). region_field carries the per-region FFT field
30+
// Views; borrowed by the GPU FFT and density-gradient backends (one per
31+
// placement region), ignored by the CPU and wirelength/HPWL backends.
32+
NesterovBaseCommon* nbc = nullptr;
33+
NesterovBase* nb = nullptr;
34+
DeviceState* device_state = nullptr;
35+
RegionDensityField* region_field = nullptr;
36+
37+
// OpenMP fan-out for the CPU backends.
38+
int num_threads = 1;
39+
40+
// FFT-only grid geometry. Required by makeFftBackend; ignored elsewhere.
41+
int bin_cnt_x = 0;
42+
int bin_cnt_y = 0;
43+
float bin_size_x = 0;
44+
float bin_size_y = 0;
45+
};
46+
47+
} // namespace gpl

0 commit comments

Comments
 (0)