Merge branch 'master' into or_update_sta_latest_0629

Author: dsengupta0628

.bazelrc

@@ -108,6 +108,11 @@ build:asan --copt=-O1
 build:asan --copt=-g
 build:asan --copt=-fno-omit-frame-pointer
 build:asan --linkopt=-fsanitize=address
+# TCMalloc is incompatible with ASan: __asan_init's interceptor setup calls
+# dlsym -> malloc before TCMalloc is initialized, segfaulting at startup.
+# Override the per-target malloc attribute to use the system allocator, which
+# ASan intercepts cleanly.
+build:asan --custom_malloc=@bazel_tools//tools/cpp:malloc
 
 # Flags with enough debug symbols to get useful outputs with Linux `perf`
 build:profile --strip=never

AUTHORS

@@ -6,6 +6,7 @@
 # source control.
 
 Andrew Kennings
+Antmicro
 Federal University of Rio Grande do Sul (UFRGS)
 Google LLC
 Iowa State University

CMakeLists.txt

@@ -42,6 +42,13 @@ option(USE_SYSTEM_ABC "Use system shared ABC library" OFF)
 # Allow disabling tests
 option(ENABLE_TESTS "Enable OpenROAD tests" ON)
 
+# Opt-in GPU acceleration via Kokkos. The actual compute backend (CUDA, HIP,
+# SYCL, or host-only OpenMP/Threads) is determined by the installed Kokkos
+# package; OpenROAD inspects Kokkos_ENABLE_* and turns on the matching CMake
+# language and dependencies automatically. See the per-module CMakeLists for
+# how individual subsystems wire their GPU sources.
+option(ENABLE_GPU "Enable GPU acceleration via Kokkos" OFF)
+
 # Allow enabling address sanitizer
 option(ASAN "Enable Address Sanitizer" OFF)
 
@@ -88,6 +95,13 @@ if(NOT CMAKE_BUILD_TYPE)
   set(CMAKE_BUILD_TYPE RELEASE)
 endif()
 
+# GPU backend wiring (opt-in). All Kokkos / CUDA / HIP / SYCL detection,
+# compiler probing, and language enablement live in cmake/KokkosBackend.cmake
+# and are loaded only when the user opts in via ENABLE_GPU=ON.
+if(ENABLE_GPU)
+  include(KokkosBackend)
+endif()
+
 if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
   if(CMAKE_CXX_COMPILER_VERSION VERSION_LESS "8.3.0")
     message(FATAL_ERROR "Insufficient gcc version. Found ${CMAKE_CXX_COMPILER_VERSION}, but require  >= 8.3.0.")

MODULE.bazel

@@ -78,7 +78,7 @@ bazel_dep(name = "eigen", version = "3.4.0.bcr.3")
 bazel_dep(name = "fmt", version = "11.2.0.bcr.1")
 bazel_dep(name = "git", version = "2.54.0")
 bazel_dep(name = "googletest", version = "1.17.0.bcr.2")
-bazel_dep(name = "openmp", version = "21.1.5.bcr.1")
+bazel_dep(name = "openmp", version = "21.1.5.bcr.2")
 bazel_dep(name = "or-tools", version = "9.15")
 bazel_dep(name = "spdlog", version = "1.15.1")
 bazel_dep(name = "sv-lang", version = "10.0.1-20260316-f04e8156")

MODULE.bazel.lock

@@ -23,7 +23,10 @@ def _tcl_encode_or_impl(ctx):
         outputs = [output_file],
         inputs = ctx.files.srcs,
         arguments = [args],
-        tools = [ctx.executable._encode_script],
+        tools = depset(
+            direct = [ctx.executable._encode_script],
+            transitive = [ctx.toolchains["@rules_python//python:toolchain_type"].py3_runtime.files] if ctx.toolchains["@rules_python//python:toolchain_type"].py3_runtime and ctx.toolchains["@rules_python//python:toolchain_type"].py3_runtime.files else [],
+        ),
         executable = ctx.toolchains["@rules_python//python:toolchain_type"].py3_runtime.interpreter,
     )
     return [DefaultInfo(files = depset([output_file]))]

bazel/tcl_encode_or.bzl

@@ -0,0 +1,158 @@
+# SPDX-License-Identifier: BSD-3-Clause
+# Copyright (c) 2026, The OpenROAD Authors
+
+# Kokkos GPU backend wiring for OpenROAD. Included from the root
+# CMakeLists.txt only when ENABLE_GPU=ON; not loaded otherwise.
+#
+# Discovers the user's Kokkos install, inherits its compute backend, turns
+# on the matching CMake language so downstream targets can mark kernel
+# sources with set_source_files_properties(... LANGUAGE CUDA|HIP), and
+# applies the small set of nvcc / fmt / host-compiler workarounds that the
+# CUDA backend currently needs in modern Linux toolchains. Per-module
+# CMakeLists (e.g. src/gpl) key off ENABLE_GPU and Kokkos_ENABLE_*; they
+# do not need to call find_package(Kokkos) or enable_language() themselves.
+
+find_package(Kokkos QUIET)
+if(NOT Kokkos_FOUND)
+  message(FATAL_ERROR
+    "OpenROAD: ENABLE_GPU=ON requires the Kokkos package to be "
+    "installed and discoverable by CMake, but Kokkos was not found.\n"
+    "  - If Kokkos is already installed: pass "
+    "-DKokkos_ROOT=/path/to/kokkos (or extend CMAKE_PREFIX_PATH).\n"
+    "  - If not: build and install Kokkos from "
+    "https://github.com/kokkos/kokkos with the desired backend "
+    "(CUDA / HIP / SYCL / OpenMP) and a target architecture that "
+    "matches the host GPU.\n"
+    "  - A future etc/DependencyInstaller.sh -gpu option will "
+    "automate this step.")
+endif()
+
+# KokkosFFT — required by the gpl GPU FFT backend (src/gpl/src/gpu/dct.cpp).
+# A separate package from Kokkos core.
+find_package(KokkosFFT QUIET)
+if(NOT KokkosFFT_FOUND)
+  message(FATAL_ERROR
+    "ENABLE_GPU=ON requires KokkosFFT, which was not found.\n"
+    "  - Install KokkosFFT (https://github.com/kokkos/kokkos-fft) against\n"
+    "    your Kokkos build, then re-configure with -DKokkosFFT_ROOT=<prefix>.\n"
+    "  - A future etc/DependencyInstaller.sh -gpu will install Kokkos and\n"
+    "    KokkosFFT together.")
+endif()
+
+message(STATUS "OpenROAD: GPU acceleration enabled (Kokkos ${Kokkos_VERSION})")
+
+if(Kokkos_ENABLE_CUDA)
+  # Auto-discover nvcc when the user has CUDA installed at a standard
+  # location but their environment does not expose it on PATH (common
+  # with IDE-launched configures: the bundled CMake does not inherit
+  # the shell PATH). enable_language(CUDA) below would otherwise abort
+  # with "No CMAKE_CUDA_COMPILER could be found" even though Kokkos's
+  # find_package already located the toolkit.
+  if(NOT DEFINED CMAKE_CUDA_COMPILER AND NOT DEFINED ENV{CUDACXX})
+    find_program(_OPENROAD_NVCC nvcc
+      HINTS ENV CUDA_HOME ENV CUDA_PATH ENV CUDA_ROOT
+            /usr/local/cuda/bin
+            /usr/local/cuda-13.0/bin
+            /usr/local/cuda-12.8/bin /usr/local/cuda-12.0/bin
+            /opt/cuda/bin
+    )
+    if(_OPENROAD_NVCC)
+      set(CMAKE_CUDA_COMPILER "${_OPENROAD_NVCC}" CACHE FILEPATH "")
+      message(STATUS "OpenROAD: auto-discovered nvcc at ${_OPENROAD_NVCC}")
+    endif()
+  endif()
+  # nvcc < 13 cannot parse glibc 2.38+'s _Float128 type that ships with
+  # gcc 13+'s C++ standard library headers (math.h template specialization
+  # for __iseqsig_type<_Float128>). When a known-broken pairing is detected,
+  # pin a compatible older g++ as the CUDA host compiler (the system C++
+  # compiler stays unchanged for non-CUDA TUs). Override is always
+  # available via -DCMAKE_CUDA_HOST_COMPILER or CUDAHOSTCXX.
+  if(NOT DEFINED CMAKE_CUDA_HOST_COMPILER AND NOT DEFINED ENV{CUDAHOSTCXX}
+     AND CMAKE_CXX_COMPILER_ID STREQUAL "GNU"
+     AND CMAKE_CXX_COMPILER_VERSION VERSION_GREATER_EQUAL "13.0"
+     AND _OPENROAD_NVCC)
+    execute_process(
+      COMMAND "${_OPENROAD_NVCC}" --version
+      OUTPUT_VARIABLE _OPENROAD_NVCC_VERSION_OUTPUT
+      ERROR_QUIET
+      OUTPUT_STRIP_TRAILING_WHITESPACE)
+    if(_OPENROAD_NVCC_VERSION_OUTPUT MATCHES "release ([0-9]+)")
+      set(_OPENROAD_NVCC_MAJOR "${CMAKE_MATCH_1}")
+      if(_OPENROAD_NVCC_MAJOR LESS 13)
+        foreach(_OPENROAD_GXX_VER 12 11)
+          find_program(_OPENROAD_CUDAHOST g++-${_OPENROAD_GXX_VER}
+            HINTS /usr/bin /usr/local/bin)
+          if(_OPENROAD_CUDAHOST)
+            set(CMAKE_CUDA_HOST_COMPILER "${_OPENROAD_CUDAHOST}"
+              CACHE FILEPATH "")
+            message(STATUS
+              "OpenROAD: pinning CUDA host compiler to "
+              "${_OPENROAD_CUDAHOST} (nvcc ${_OPENROAD_NVCC_MAJOR}.x + "
+              "glibc/gcc 13+ _Float128 compat)")
+            break()
+          endif()
+          unset(_OPENROAD_CUDAHOST CACHE)
+        endforeach()
+        if(NOT DEFINED CMAKE_CUDA_HOST_COMPILER)
+          message(FATAL_ERROR
+            "OpenROAD: nvcc ${_OPENROAD_NVCC_MAJOR}.x cannot parse "
+            "_Float128 declarations in glibc 2.38+ system headers used "
+            "by gcc ${CMAKE_CXX_COMPILER_VERSION}, and no compatible "
+            "g++-12 / g++-11 was found in /usr/bin or /usr/local/bin. "
+            "Install one (e.g. apt install g++-12) or set "
+            "-DCMAKE_CUDA_HOST_COMPILER=/path/to/older-g++ explicitly.")
+        endif()
+      endif()
+    endif()
+  endif()
+  if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES OR "${CMAKE_CUDA_ARCHITECTURES}" STREQUAL "")
+    if(DEFINED Kokkos_CUDA_ARCHITECTURES
+       AND NOT "${Kokkos_CUDA_ARCHITECTURES}" STREQUAL "")
+      set(CMAKE_CUDA_ARCHITECTURES "${Kokkos_CUDA_ARCHITECTURES}")
+    else()
+      message(FATAL_ERROR
+        "OpenROAD: ENABLE_GPU=ON with Kokkos CUDA backend, but the "
+        "Kokkos package does not advertise Kokkos_CUDA_ARCHITECTURES "
+        "and CMAKE_CUDA_ARCHITECTURES was not provided. Set "
+        "-DCMAKE_CUDA_ARCHITECTURES=<arch> explicitly (e.g. 89 for "
+        "RTX 4070, 120 for RTX 5090) or rebuild Kokkos with the "
+        "target architecture baked in.")
+    endif()
+  endif()
+  enable_language(CUDA)
+  find_package(CUDAToolkit REQUIRED)
+  message(STATUS "OpenROAD: CUDA backend (arch=${CMAKE_CUDA_ARCHITECTURES})")
+  # A GPU driver (the kernel module exposing libcuda.so.1) is needed only to
+  # *run* CUDA code, never to build it -- nvcc cross-compiles device code on a
+  # host with no GPU. Note its absence so the resulting libcuda.so.1 load
+  # errors on this host (e.g. ctest, or running openroad) read as expected
+  # rather than as a misconfiguration. This is informational only: a GPU build
+  # on a driverless host is a supported cross-compile workflow, not an error.
+  if(NOT EXISTS "/proc/driver/nvidia")
+    message(STATUS
+      "OpenROAD: no NVIDIA driver on this host -- GPU code is being "
+      "cross-compiled. Run the GPU binaries and tests on a GPU machine.")
+  endif()
+  # nvcc 12.8 cannot parse fmt 11's nontype-template-parameter user-defined
+  # literals (fmt/bundled/format.h: operator""_a with fixed_string). The
+  # legacy literal fallback is still available; opt into it for CUDA TUs
+  # only. Project-wide CXX compilation is unaffected.
+  add_compile_definitions(
+    $<$<COMPILE_LANGUAGE:CUDA>:FMT_USE_NONTYPE_TEMPLATE_ARGS=0>)
+  # On aarch64, Boost's unordered_flat_map detects __ARM_NEON and includes
+  # <arm_neon.h> for SIMD-accelerated hashing.  nvcc cannot parse gcc's
+  # arm_neon.h (it contains gcc-specific intrinsics), so disable the NEON
+  # path for CUDA TUs.  The CPU TUs (compiled by g++) are unaffected.
+  if(CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64|ARM64")
+    add_compile_definitions(
+      $<$<COMPILE_LANGUAGE:CUDA>:BOOST_UNORDERED_DISABLE_NEON>)
+  endif()
+elseif(Kokkos_ENABLE_HIP)
+  enable_language(HIP)
+  message(STATUS "OpenROAD: HIP backend")
+elseif(Kokkos_ENABLE_SYCL)
+  message(STATUS "OpenROAD: SYCL backend (driven by Kokkos host compiler)")
+else()
+  message(STATUS
+          "OpenROAD: host-only Kokkos backend (Serial / OpenMP / Threads)")
+endif()

cmake/KokkosBackend.cmake

@@ -883,9 +883,35 @@ _install_bazel() {
             _execute "Installing bazelisk..." mv bazelisk "${bazel_prefix}/bin/bazelisk"
         )
         if _command_exists "apt-get"; then
+            # Ubuntu 26.04 ships the libxml2 runtime with soname
+            # libxml2.so.16, but the prebuilt LLVM toolchain (lld) pulled in
+            # by the Bazel build is linked against the old libxml2.so.2.
+            # Pull in libxml2-dev there (and add a compatibility symlink
+            # below); older Ubuntu still provides .so.2 via libxml2.
+            local ubuntu_version=""
+            if [[ -f /etc/os-release ]]; then
+                ubuntu_version=$(awk -F= '/^VERSION_ID/{print $2}' /etc/os-release | sed 's/"//g')
+            fi
+            local libxml2_pkg="libxml2"
+            if [[ -n "${ubuntu_version}" ]] && _version_compare "${ubuntu_version}" -ge "26.04"; then
+                libxml2_pkg="libxml2-dev"
+            fi
             _execute "Installing bazel required libraries..." \
                 apt-get -y install --no-install-recommends \
-                libc6-dev libxml2 libtinfo6 zlib1g libstdc++6
+                libc6-dev "${libxml2_pkg}" libtinfo6 zlib1g libstdc++6
+            # lld only uses libxml2 for Windows COFF manifests, never during a
+            # Linux link, so the .so.16 -> .so.2 compatibility symlink is safe.
+            # Gated to 26.04+ only.
+            if [[ -n "${ubuntu_version}" ]] && _version_compare "${ubuntu_version}" -ge "26.04"; then
+                local libdir="/usr/lib/$(uname -m)-linux-gnu"
+                local libxml2_so
+                libxml2_so=$(ls "${libdir}"/libxml2.so.* 2>/dev/null \
+                    | grep -v 'libxml2.so.2$' | head -n1)
+                if [[ ! -e "${libdir}/libxml2.so.2" && -n "${libxml2_so}" ]]; then
+                    _execute "Adding libxml2.so.2 compatibility symlink for prebuilt LLVM lld..." \
+                        ln -sf "$(basename "${libxml2_so}")" "${libdir}/libxml2.so.2"
+                fi
+            fi
         elif _command_exists "yum"; then
             _execute "Installing bazel required libraries..." \
                 yum install -y \

etc/DependencyInstaller.sh

@@ -73,6 +73,9 @@ detailed_placement
     [-report_file_name filename]
     [-use_negotiation]
     [-abacus]
+    [-site_search_window sites]
+    [-row_search_window rows]
+    [-drc_penalty penalty]
 ```
 
 #### Options
@@ -86,6 +89,9 @@ detailed_placement
 | `-report_file_name` | File name for saving the report to (e.g. `report.json`.) |
 | `-use_negotiation` | Use the NegotiationLegalizer instead of the default diamond search engine. |
 | `-abacus` | Enable the Abacus pre-pass within the NegotiationLegalizer. Only effective when `-use_negotiation` is set. |
+| `-site_search_window` | NegotiationLegalizer: maximum number of sites a cell may be moved left or right of its initial position. Default `20`, `0` allowed (no horizontal movement). |
+| `-row_search_window` | NegotiationLegalizer: maximum number of rows a cell may be moved up or down from its initial position. Default `5`, `0` allowed (no row changes). |
+| `-drc_penalty` | NegotiationLegalizer: priority to DRC violations, ramped up each iteration to push DRC cleanup later in the run. Lower values tolerate DRC violations early on while overlaps are resolved. Default `5`, `0` allowed (disables the escalating per-candidate penalty, DRC-violating cells still accrue history cost separately). |
 
 ### Set Placement Padding
 

src/dpl/README.md

@@ -105,12 +105,18 @@ class Opendp
 
   // legalize/report
   // max_displacment is in sites. use zero for defaults.
+  // site_search_window/row_search_window/drc_penalty use a negative value to
+  // mean "unset" (use the negotiation legalizer's own default); 0 is a valid
+  // explicit value for all three.
   void detailedPlacement(int max_displacement_x,
                          int max_displacement_y,
                          const std::string& report_file_name = std::string(""),
                          bool incremental = false,
                          bool use_negotiation = false,
-                         bool run_abacus = false);
+                         bool run_abacus = false,
+                         int site_search_window = -1,
+                         int row_search_window = -1,
+                         double drc_penalty = -1.0);
   void reportLegalizationStats() const;
 
   void setPaddingGlobal(int left, int right);
@@ -120,6 +126,8 @@ class Opendp
   void setJumpMoves(int jump_moves);
   void setIterativePlacement(bool iterative);
   void setDeepIterativePlacement(bool deep_iterative);
+  void setNegotiationDebugInterval(int iterative_jump);
+  void setNegotiationDebugStart(int iterative_start);
 
   // Global padding.
   int padGlobalLeft() const;
@@ -398,6 +406,8 @@ class Opendp
   int move_count_ = 1;
   bool iterative_debug_ = false;
   bool deep_iterative_debug_ = false;
+  int negotiation_debug_interval_ = 1;
+  int negotiation_debug_start_ = 0;
   bool incremental_ = false;
   bool use_negotiation_ = false;