Merge remote-tracking branch 'private/master' into dpl-dpo-load-macro-flipping

Author: gudeh

.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.")

cmake/KokkosBackend.cmake

@@ -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()

src/dpl/src/NegotiationLegalizer.cpp

@@ -10,6 +10,7 @@
 #include <cstddef>
 #include <functional>
 #include <limits>
+#include <map>
 #include <queue>
 #include <tuple>
 #include <unordered_map>
@@ -523,13 +524,8 @@ bool NegotiationLegalizer::initFromDb()
   die_xhi_ = core_area.xMax();
   die_yhi_ = core_area.yMax();
 
-  // Site width from the DPL grid; row height from the first DB row.
   site_width_ = dpl_grid->getSiteWidth().v;
-  for (auto* row : block->getRows()) {
-    row_height_ = row->getSite()->getHeight();
-    break;
-  }
-  assert(site_width_ > 0 && row_height_ > 0);
+  assert(site_width_ > 0);
 
   // Grid dimensions from the DPL grid (accounts for actual DB rows).
   grid_w_ = dpl_grid->getRowSiteCount().v;
@@ -578,10 +574,7 @@ bool NegotiationLegalizer::initFromDb()
         1,
         static_cast<int>(
             std::round(static_cast<double>(master->getWidth()) / site_width_)));
-    cell.height = std::max(
-        1,
-        static_cast<int>(std::round(static_cast<double>(master->getHeight())
-                                    / row_height_)));
+    cell.height = dpl_grid->gridHeight(master).v;
 
     // Clamp to valid grid range – gridRoundY can return grid_h_ when the
     // instance is near the top edge.  Use (grid_w_ - width) so the full
@@ -647,9 +640,9 @@ bool NegotiationLegalizer::initFromDb()
           for (auto* box : odb_region->getBoundaries()) {
             RegionRectInline r;
             r.xlo = (box->xMin() - die_xlo_) / site_width_;
-            r.ylo = (box->yMin() - die_ylo_) / row_height_;
+            r.ylo = dpl_grid->gridEndY(DbuY{box->yMin() - die_ylo_}).v;
             r.xhi = (box->xMax() - die_xlo_) / site_width_;
-            r.yhi = (box->yMax() - die_ylo_) / row_height_;
+            r.yhi = dpl_grid->gridSnapDownY(DbuY{box->yMax() - die_ylo_}).v;
             rects.push_back(r);
           }
           it = region_rect_cache.emplace(odb_region, std::move(rects)).first;
@@ -684,11 +677,12 @@ bool NegotiationLegalizer::initFromDb()
                    db_x,
                    db_y,
                    die_xlo_ + cell.init_x * site_width_,
-                   die_ylo_ + cell.init_y * row_height_);
+                   die_ylo_ + dpl_grid->gridYToDbu(GridY{cell.init_y}).v);
 
         // Priority queue keyed on physical Manhattan distance (DBU) so the
         // search expands in true physical proximity, not grid-unit proximity.
-        // One step in X = site_width_ DBU; one step in Y = row_height_ DBU.
+        // One step in X = site_width_ DBU; Y distance is computed via the
+        // DPL grid since pixel rows may have non-uniform heights.
         using PQEntry = std::tuple<int, int, int>;  // physDist, gx, gy
         std::
             priority_queue<PQEntry, std::vector<PQEntry>, std::greater<PQEntry>>
@@ -704,8 +698,10 @@ bool NegotiationLegalizer::initFromDb()
           if (!visited.insert(gy * grid_w_ + gx).second) {
             return;
           }
-          const int dist = std::abs(gx - cell.init_x) * site_width_
-                           + std::abs(gy - cell.init_y) * row_height_;
+          const int dy_dbu = dpl_grid->gridYToDbu(GridY{gy}).v
+                             - dpl_grid->gridYToDbu(GridY{cell.init_y}).v;
+          const int dist
+              = std::abs(gx - cell.init_x) * site_width_ + std::abs(dy_dbu);
           pq.emplace(dist, gx, gy);
         };
 
@@ -745,6 +741,19 @@ bool NegotiationLegalizer::initFromDb()
     cells_.push_back(cell);
   }
 
+  std::map<int, int> neg_height_counts;
+  for (const NegCell& c : cells_) {
+    neg_height_counts[c.height]++;
+  }
+  logger_->info(
+      utl::DPL,
+      392,
+      "Negotiation cell height distribution ({} unique row-count(s)):",
+      neg_height_counts.size());
+  for (const auto& [height, count] : neg_height_counts) {
+    logger_->info(utl::DPL, 393, "  height {} row(s): {} cells", height, count);
+  }
+
   return true;
 }
 
@@ -813,12 +822,13 @@ void NegotiationLegalizer::initFenceRegions()
     FenceRegion fr;
     fr.id = region->getId();
 
+    const Grid* dpl_grid = opendp_->grid_.get();
     for (auto* box : region->getBoundaries()) {
       FenceRect r;
       r.xlo = (box->xMin() - die_xlo_) / site_width_;
-      r.ylo = (box->yMin() - die_ylo_) / row_height_;
+      r.ylo = dpl_grid->gridEndY(DbuY{box->yMin() - die_ylo_}).v;
       r.xhi = (box->xMax() - die_xlo_) / site_width_;
-      r.yhi = (box->yMax() - die_ylo_) / row_height_;
+      r.yhi = dpl_grid->gridSnapDownY(DbuY{box->yMax() - die_ylo_}).v;
       fr.rects.push_back(r);
     }
 

src/dpl/src/NegotiationLegalizer.h

@@ -233,7 +233,6 @@ class NegotiationLegalizer
   Network* network_{nullptr};
 
   int site_width_{0};
-  int row_height_{0};
   int die_xlo_{0};
   int die_ylo_{0};
   int die_xhi_{0};

src/dpl/src/NegotiationLegalizerPass.cpp

@@ -418,9 +418,10 @@ void NegotiationLegalizer::place(int cell_idx, int x, int y)
       pushNegotiationPixels();
       const NegCell& c = cells_[cell_idx];
       const int orig_x_dbu = die_xlo_ + c.init_x * site_width_;
-      const int orig_y_dbu = die_ylo_ + c.init_y * row_height_;
+      const int orig_y_dbu
+          = die_ylo_ + opendp_->grid_->gridYToDbu(GridY{c.init_y}).v;
       const int tgt_x_dbu = die_xlo_ + c.x * site_width_;
-      const int tgt_y_dbu = die_ylo_ + c.y * row_height_;
+      const int tgt_y_dbu = die_ylo_ + opendp_->grid_->gridYToDbu(GridY{c.y}).v;
       logger_->report(
           "Pause at placing of cell {}. orig=({},{}) target=({},{}) dbu. "
           "rowidx={}.",

src/dpl/src/dbToOpendp.cpp

@@ -187,6 +187,22 @@ void Opendp::importDb()
   createNetwork();
   createArchitecture();
   setUpPlacementGroups();
+
+  if (logger_->debugCheck(utl::DPL, "hybrid", 1)) {
+    std::map<int, int> height_counts;
+    for (const auto& node : network_->getNodes()) {
+      if (node->getType() != Node::CELL) {
+        continue;
+      }
+      height_counts[node->getHeight().v]++;
+    }
+    logger_->report("Cell height distribution ({} unique micron height(s)):",
+                    height_counts.size());
+    for (const auto& [height, count] : height_counts) {
+      logger_->report(
+          "  height {:.3f} um: {} cells", block_->dbuToMicrons(height), count);
+    }
+  }
 }
 
 void Opendp::importClear()

src/dpl/test/negotiation01.ok

@@ -6,6 +6,8 @@
 [INFO DPL-0006] Core area: 14266.91 um^2, Instances area: 637.60 um^2, Utilization: 4.5%
 [INFO DPL-0005] Diamond search max displacement: +/- 500 sites horizontally, +/- 100 rows vertically.
 [INFO DPL-1102] Legalizing using negotiation legalizer.
+[INFO DPL-0392] Negotiation cell height distribution (1 unique row-count(s)):
+[INFO DPL-0393]   height 1 row(s): 549 cells
 Negotiation iteration 0: total overflow 12.
 Negotiation iteration 1: total overflow 0.
 Placement Analysis

src/dpl/test/negotiation_one_site_gap.ok

@@ -6,6 +6,8 @@
 [INFO DPL-0006] Core area: 16.49 um^2, Instances area: 3.72 um^2, Utilization: 22.6%
 [INFO DPL-0005] Diamond search max displacement: +/- 500 sites horizontally, +/- 100 rows vertically.
 [INFO DPL-1102] Legalizing using negotiation legalizer.
+[INFO DPL-0392] Negotiation cell height distribution (1 unique row-count(s)):
+[INFO DPL-0393]   height 1 row(s): 4 cells
 Negotiation iteration 0: total overflow 0.
 Placement Analysis
 ---------------------------------