add gputimer

cpp_to_py/gputimer/CMakeLists.txt

@@ -0,0 +1,38 @@
+# Files
+file(GLOB_RECURSE SRC_FILES_GT ${CMAKE_CURRENT_SOURCE_DIR}/*.cpp
+                                ${CMAKE_CURRENT_SOURCE_DIR}/*.hpp)
+file(GLOB_RECURSE SRC_FILES_GT_CUDA ${CMAKE_CURRENT_SOURCE_DIR}/*.cu)
+
+# OpenMP
+find_package(OpenMP REQUIRED)
+
+# Flute MP
+set(FLUTE_MP_INCLUDE_DIR ${PATH_THIRDPARTY_ROOT}/flute_mp)
+
+# CUDA/CPP GGR Kernel
+add_library(gt STATIC ${CMAKE_CURRENT_SOURCE_DIR}/../io_parser/gp/GPDatabase.cpp
+                       ${SRC_FILES_GT}
+                       ${SRC_FILES_GT_CUDA}
+                       ${FLUTE_MP_INCLUDE_DIR}/flute.cpp
+                       )
+
+set_target_properties(gt PROPERTIES
+                       CUDA_RESOLVE_DEVICE_SYMBOLS ON
+                       POSITION_INDEPENDENT_CODE ON)
+
+target_include_directories(gt PRIVATE ${PROJECT_SOURCE_DIR}/cpp_to_py ${TORCH_INCLUDE_DIRS} ${FLUTE_MP_INCLUDE_DIR})
+target_link_libraries(gt PRIVATE torch ${TORCH_PYTHON_LIBRARY} xplace_common io_parser OpenMP::OpenMP_CXX stdc++fs pthread)
+target_compile_options(gt PRIVATE -fPIC)
+target_compile_options(gt PRIVATE "$<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>")
+
+pybind11_add_module(gputimer MODULE PyBindCppMain.cpp)
+
+target_include_directories(gputimer PRIVATE ${TORCH_INCLUDE_DIRS} ${PROJECT_SOURCE_DIR}/cpp_to_py ${FLUTE_MP_INCLUDE_DIR})
+target_link_libraries(gputimer PRIVATE torch ${TORCH_PYTHON_LIBRARY} xplace_common io_parser gt)
+target_compile_definitions(gputimer PRIVATE 
+                            TORCH_EXTENSION_NAME=gputimer
+                            TORCH_VERSION_MAJOR=${TORCH_VERSION_MAJOR}
+                            TORCH_VERSION_MINOR=${TORCH_VERSION_MINOR}
+                            ENABLE_CUDA=${TORCH_ENABLE_CUDA})
+
+install(TARGETS gputimer DESTINATION ${XPLACE_LIB_DIR})

cpp_to_py/gputimer/PyBindCppMain.cpp

@@ -0,0 +1,124 @@
+#include "common/common.h"
+
+#include "common/db/Database.h"
+#include "io_parser/gp/GPDatabase.h"
+#include "gputimer/db/GTDatabase.h"
+#include "gputimer/core/GPUTimer.h"
+
+#include <flute.hpp>
+using namespace Flute;
+
+namespace Xplace {
+
+std::shared_ptr<gt::GPUTimer> create_gputimer(const py::dict& kwargs,
+                                              std::shared_ptr<db::Database> rawdb,
+                                              std::shared_ptr<gp::GPDatabase> gpdb,
+                                              std::shared_ptr<gt::TimingTorchRawDB> timing_raw_db) {
+    std::shared_ptr<gt::GTDatabase> gtdb = std::make_shared<gt::GTDatabase>(rawdb, gpdb, timing_raw_db);
+    auto sdc = std::make_shared<gt::sdc::SDC>();
+
+    try {
+        if (kwargs.contains("sdc")) sdc->read(kwargs["sdc"].cast<std::string>());
+    } catch (std::exception& e) {
+        logger.error("%s\n", e.what());
+    }
+
+    gtdb->ExtractTimingGraph();
+    gtdb->readSdc(*sdc);
+
+    std::shared_ptr<gt::GPUTimer> gputimer = std::make_shared<gt::GPUTimer>(gtdb, timing_raw_db);
+
+    readLUT("thirdparty/flute_mp/lut.ICCAD2015/POWV9.dat", "thirdparty/flute_mp/lut.ICCAD2015/POST9.dat");
+    
+    return gputimer;
+}
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+    pybind11::class_<gt::GPUTimer, std::shared_ptr<gt::GPUTimer>>(m, "GPUTimer")
+        .def(pybind11::init<std::shared_ptr<gt::GTDatabase>, std::shared_ptr<gt::TimingTorchRawDB>>())
+        .def("time_unit", &gt::GPUTimer::time_unit)
+        .def("read_spef", &gt::GPUTimer::read_spef)
+        .def("init", &gt::GPUTimer::initialize)
+        .def("levelize", &gt::GPUTimer::levelize)
+        .def("update_rc", &gt::GPUTimer::update_rc_timing)
+        .def("update_rc_flute", &gt::GPUTimer::update_rc_timing_flute)
+        .def("update_rc_spef", &gt::GPUTimer::update_rc_timing_spef)
+        .def("update_states", &gt::GPUTimer::update_states)
+        .def("update_timing", &gt::GPUTimer::update_timing)
+        .def("update_endpoints", &gt::GPUTimer::update_endpoints)
+        .def("report_wns", &gt::GPUTimer::report_wns)
+        .def("report_tns_elw", &gt::GPUTimer::report_tns_elw)
+        .def("report_wns_and_tns", &gt::GPUTimer::report_wns_and_tns)
+        .def("report_pin_slack", &gt::GPUTimer::report_pin_slack, py::return_value_policy::move)
+        .def("report_pin_at", &gt::GPUTimer::report_pin_at, py::return_value_policy::move)
+        .def("report_pin_rat", &gt::GPUTimer::report_pin_rat, py::return_value_policy::move)
+        .def("report_pin_slew", &gt::GPUTimer::report_pin_slew, py::return_value_policy::move)
+        .def("report_pin_load", &gt::GPUTimer::report_pin_load, py::return_value_policy::move)
+        .def("report_endpoint_slack", &gt::GPUTimer::report_endpoint_slack, py::return_value_policy::move)
+        .def("endpoints_index", &gt::GPUTimer::endpoints_index, py::return_value_policy::copy)
+        .def("report_path", &gt::GPUTimer::report_path, py::return_value_policy::copy)
+        .def("report_K_path", &gt::GPUTimer::report_K_path, py::return_value_policy::copy)
+        .def("report_criticality", &gt::GPUTimer::report_criticality, py::return_value_policy::copy)
+        .def("report_criticality_threshold", &gt::GPUTimer::report_criticality_threshold, py::return_value_policy::copy)
+        ;
+    pybind11::class_<gt::TimingTorchRawDB, std::shared_ptr<gt::TimingTorchRawDB>>(m, "TimingTorchRawDB")
+        .def(pybind11::init<torch::Tensor,
+                            torch::Tensor,
+                            torch::Tensor,
+                            torch::Tensor,
+                            torch::Tensor,
+                            torch::Tensor,
+                            torch::Tensor,
+                            torch::Tensor,
+                            torch::Tensor,
+                            torch::Tensor,
+                            int,
+                            float,
+                            int,
+                            float,
+                            float>())
+        .def("commit_from", &gt::TimingTorchRawDB::commit_from)
+        .def("get_curr_cposx", &gt::TimingTorchRawDB::get_curr_cposx, py::return_value_policy::move)
+        .def("get_curr_cposy", &gt::TimingTorchRawDB::get_curr_cposy, py::return_value_policy::move)
+        .def("get_curr_lposx", &gt::TimingTorchRawDB::get_curr_lposx, py::return_value_policy::move)
+        .def("get_curr_lposy", &gt::TimingTorchRawDB::get_curr_lposy, py::return_value_policy::move);
+
+    pybind11::class_<gt::GTDatabase, std::shared_ptr<gt::GTDatabase>>(m, "GTDatabase")
+        .def(pybind11::init<std::shared_ptr<db::Database>, std::shared_ptr<gp::GPDatabase>, std::shared_ptr<gt::TimingTorchRawDB>>());
+
+    m.def("create_gputimer", &create_gputimer, "Create gputimer object");
+    m.def("create_timing_rawdb",
+          [](torch::Tensor node_lpos_init_,
+             torch::Tensor node_size_,
+             torch::Tensor pin_rel_lpos_,
+             torch::Tensor pin_id2node_id_,
+             torch::Tensor pin_id2net_id_,
+             torch::Tensor node2pin_list_,
+             torch::Tensor node2pin_list_end_,
+             torch::Tensor hyperedge_list_,
+             torch::Tensor hyperedge_list_end_,
+             torch::Tensor net_mask_,
+             int num_movable_nodes_,
+             float scale_factor_,
+             int microns_,
+             float wire_resistance_per_micron_,
+             float wire_capacitance_per_micron_) {
+              return std::make_shared<gt::TimingTorchRawDB>(node_lpos_init_,
+                                                            node_size_,
+                                                            pin_rel_lpos_,
+                                                            pin_id2node_id_,
+                                                            pin_id2net_id_,
+                                                            node2pin_list_,
+                                                            node2pin_list_end_,
+                                                            hyperedge_list_,
+                                                            hyperedge_list_end_,
+                                                            net_mask_,
+                                                            num_movable_nodes_,
+                                                            scale_factor_,
+                                                            microns_,
+                                                            wire_resistance_per_micron_,
+                                                            wire_capacitance_per_micron_);
+          });
+}
+
+}  // namespace Xplace

cpp_to_py/gputimer/base.h

@@ -0,0 +1,27 @@
+
+
+#pragma once
+
+namespace gt {
+
+// #define BLOCK_SIZE 512
+#define BLOCK_SIZE 512
+#define BLOCK_NUMBER(n) (((n) + (BLOCK_SIZE)-1) / BLOCK_SIZE)
+// el rf rf: e0 l1 r0 f1
+
+#define NUM_ATTR 4
+
+// using index_type = int64_t;
+using index_type = int;
+
+
+// Overloadded.
+template <typename... Ts>
+struct Functors : Ts... { 
+  using Ts::operator()... ;
+};
+
+template <typename... Ts>
+Functors(Ts...) -> Functors<Ts...>;
+
+}

cpp_to_py/gputimer/core/GPUTimer.cpp

@@ -0,0 +1,138 @@
+
+
+
+#include "common/common.h"
+#include "common/db/Database.h"
+#include "io_parser/gp/GPDatabase.h"
+#include "gputimer/db/GTDatabase.h"
+#include "GPUTimer.h"
+
+namespace gt {
+
+GPUTimer::GPUTimer(std::shared_ptr<GTDatabase> gtdb_, shared_ptr<TimingTorchRawDB> timing_raw_db_)
+    : gtdb(*gtdb_),
+      timing_raw_db(*timing_raw_db_),
+      x(timing_raw_db.x.data_ptr<float>()),
+      y(timing_raw_db.y.data_ptr<float>()),
+      init_x(timing_raw_db.init_x.data_ptr<float>()),
+      init_y(timing_raw_db.init_y.data_ptr<float>()),
+      node_size_x(timing_raw_db.node_size_x.data_ptr<float>()),
+      node_size_y(timing_raw_db.node_size_y.data_ptr<float>()),
+      pin_offset_x(timing_raw_db.pin_offset_x.data_ptr<float>()),
+      pin_offset_y(timing_raw_db.pin_offset_y.data_ptr<float>()),
+      // GPU pin attributes array
+      pinSlew(timing_raw_db.pinSlew.data_ptr<float>()),
+      pinLoad(timing_raw_db.pinLoad.data_ptr<float>()),
+      pinRAT(timing_raw_db.pinRAT.data_ptr<float>()),
+      pinAT(timing_raw_db.pinAT.data_ptr<float>()),
+      pinImpulse(timing_raw_db.pinImpulse.data_ptr<float>()),
+      pinRootDelay(timing_raw_db.pinRootDelay.data_ptr<float>()),
+      arcDelay(timing_raw_db.arcDelay.data_ptr<float>()),
+      // Critical path prefix info
+      at_prefix_pin(timing_raw_db.at_prefix_pin.data_ptr<index_type>()),
+      at_prefix_arc(timing_raw_db.at_prefix_arc.data_ptr<index_type>()),
+      at_prefix_attr(timing_raw_db.at_prefix_attr.data_ptr<index_type>()),
+      // Timing graph topology
+      pin_forward_arc_list(timing_raw_db.pin_forward_arc_list.data_ptr<index_type>()),
+      pin_forward_arc_list_end(timing_raw_db.pin_forward_arc_list_end.data_ptr<index_type>()),
+      pin_backward_arc_list(timing_raw_db.pin_backward_arc_list.data_ptr<index_type>()),
+      pin_backward_arc_list_end(timing_raw_db.pin_backward_arc_list_end.data_ptr<index_type>()),
+      timing_arc_from_pin_id(timing_raw_db.timing_arc_from_pin_id.data_ptr<index_type>()),
+      timing_arc_to_pin_id(timing_raw_db.timing_arc_to_pin_id.data_ptr<index_type>()),
+      pin_num_fanin(timing_raw_db.pin_num_fanin.data_ptr<int>()),
+      pin_fanout_list(timing_raw_db.pin_fanout_list.data_ptr<index_type>()),
+      pin_fanout_list_end(timing_raw_db.pin_fanout_list_end.data_ptr<index_type>()),
+      // Timer timing liberty variables
+      timing_arc_id_map(timing_raw_db.timing_arc_id_map.data_ptr<int>()),
+      arc_types(timing_raw_db.arc_types.data_ptr<int>()),
+      arc_id2test_id(timing_raw_db.arc_id2test_id.data_ptr<int>()),
+      test_id2_arc_id(timing_raw_db.test_id2_arc_id.data_ptr<int>()),
+      // Circuit info
+      flat_node2pin_start_map(timing_raw_db.flat_node2pin_start_map.data_ptr<int>()),
+      flat_node2pin_map(timing_raw_db.flat_node2pin_map.data_ptr<int>()),
+      pin2node_map(timing_raw_db.pin2node_map.data_ptr<int>()),
+      flat_net2pin_start_map(timing_raw_db.flat_net2pin_start_map.data_ptr<int>()),
+      flat_net2pin_map(timing_raw_db.flat_net2pin_map.data_ptr<int>()),
+      pin2net_map(timing_raw_db.pin2net_map.data_ptr<int>()),
+      net_mask(timing_raw_db.net_mask.data_ptr<bool>()),
+      num_threads(timing_raw_db.num_threads),
+      num_nodes(timing_raw_db.num_nodes),
+      num_movable_nodes(timing_raw_db.num_movable_nodes),
+      num_nets(timing_raw_db.num_nets),
+      num_pins(timing_raw_db.num_pins),
+      scale_factor(timing_raw_db.scale_factor) {
+    num_arcs = gtdb.num_arcs;
+    num_timings = gtdb.num_timings;
+    total_num_fanouts = gtdb.total_num_fanouts;
+    num_tests = gtdb.num_tests;
+    num_POs = gtdb.num_POs;
+    wire_resistance_per_micron = timing_raw_db.wire_resistance_per_micron;
+    wire_capacitance_per_micron = timing_raw_db.wire_capacitance_per_micron;
+    microns = timing_raw_db.microns;
+    res_unit = gtdb.res_unit;
+    cap_unit = gtdb.cap_unit;
+    if (gtdb.clocks.empty())
+        clock_period = 0;
+    else
+        clock_period = gtdb.clocks.begin()->second.period();
+    gtdb_holder = gtdb_;
+    timing_raw_db_holder = timing_raw_db_;
+}
+
+torch::Tensor GPUTimer::report_pin_at() {return timing_raw_db.pinAT;}
+torch::Tensor GPUTimer::report_pin_rat() { return timing_raw_db.pinRAT; }
+torch::Tensor GPUTimer::report_pin_slew() { return timing_raw_db.pinSlew; }
+torch::Tensor GPUTimer::report_pin_load() { return timing_raw_db.pinLoad; }
+torch::Tensor GPUTimer::report_endpoint_slack() { return endpoint_slacks; }
+torch::Tensor GPUTimer::endpoints_index(){ return timing_raw_db.endpoints_id;}
+float GPUTimer::time_unit() const { return gtdb.time_unit; }
+
+float GPUTimer::report_wns(int el) {
+    auto ep_slacks = torch::nan_to_num(endpoint_slacks, FLT_MAX);
+    return torch::min(ep_slacks.index({"...", torch::indexing::Slice(2 * el, 2 * (el + 1))})).item<float>();
+}
+
+float GPUTimer::report_tns_elw(int el) {
+    auto ep_slacks = torch::nan_to_num(endpoint_slacks);
+    auto [slack_elw, order] = torch::min(ep_slacks.index({"...", torch::indexing::Slice(2 * el, 2 * (el + 1))}), 1);
+    slack_elw.clamp_max_(0);
+
+    return torch::sum(slack_elw, 0).item<float>();
+}
+
+tuple<torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor> GPUTimer::report_wns_and_tns() {
+    auto ep_slacks = torch::nan_to_num(endpoint_slacks, FLT_MAX);
+    auto [slack_e, order_e] = torch::min(ep_slacks.index({"...", torch::indexing::Slice(0, 2)}), 1);
+    slack_e.clamp_max_(0);
+
+    auto [slack_l, order_l] = torch::min(ep_slacks.index({"...", torch::indexing::Slice(2, 4)}), 1);
+    slack_l.clamp_max_(0);
+
+    return {torch::min(ep_slacks.index({"...", torch::indexing::Slice(0, 2)})),
+            torch::sum(slack_e, 0),
+            torch::min(ep_slacks.index({"...", torch::indexing::Slice(2, 4)})),
+            torch::sum(slack_l, 0)};
+}
+
+torch::Tensor GPUTimer::report_pin_slack() {
+    pin_slacks = torch::zeros_like(timing_raw_db.pinAT, torch::dtype(torch::kFloat32).device(timing_raw_db.pinAT.device()));
+    auto s1 = timing_raw_db.pinAT - timing_raw_db.pinRAT;
+    auto s2 = timing_raw_db.pinRAT - timing_raw_db.pinAT;
+    pin_slacks.index({"...", torch::indexing::Slice(0, 2)}).data().copy_(s1.index({"...", torch::indexing::Slice(0, 2)}));
+    pin_slacks.index({"...", torch::indexing::Slice(2, 4)}).data().copy_(s2.index({"...", torch::indexing::Slice(2, 4)}));
+
+    return pin_slacks.contiguous();
+}
+
+// void GPUTimer::update_endpoints() {
+//     pin_slacks = torch::zeros_like(timing_raw_db.pinAT, torch::dtype(torch::kFloat32).device(timing_raw_db.pinAT.device()));
+//     auto s1 = timing_raw_db.pinAT - timing_raw_db.pinRAT;
+//     auto s2 = timing_raw_db.pinRAT - timing_raw_db.pinAT;
+//     pin_slacks.index({"...", torch::indexing::Slice(0, 2)}).data().copy_(s1.index({"...", torch::indexing::Slice(0, 2)}));
+//     pin_slacks.index({"...", torch::indexing::Slice(2, 4)}).data().copy_(s2.index({"...", torch::indexing::Slice(2, 4)}));
+    
+//     auto [endpoints_id, tmp1] = torch::_unique(timing_raw_db.endpoints_id);
+//     endpoint_slacks = torch::nan_to_num(pin_slacks.index_select(0, endpoints_id));
+// }
+
+}  // namespace gt

cpp_to_py/gputimer/core/GPUTimer.cu

@@ -0,0 +1,142 @@
+
+
+#include "GPUTimer.h"
+#include "gputimer/db/GTDatabase.h"
+#include "gputiming.h"
+#include "utils.cuh"
+
+namespace gt {
+
+void GPUTimer::initialize() {
+    cudaMalloc(&pinCap, num_pins * (NUM_ATTR + 2) * sizeof(float));
+    cudaMalloc(&pinWireCap, num_pins * NUM_ATTR * sizeof(float));
+    cudaMalloc(&testRelatedAT, num_tests * NUM_ATTR * sizeof(float));
+    cudaMalloc(&testRAT, num_tests * NUM_ATTR * sizeof(float));
+    cudaMalloc(&testConstraint, num_tests * NUM_ATTR * sizeof(float));
+    cudaMalloc(&pinRootRes, num_pins * NUM_ATTR * sizeof(float));
+    cudaMalloc(&arcSlew, num_arcs * 2 * NUM_ATTR * sizeof(float));
+
+    cudaMalloc(&net_is_clock, num_nets * sizeof(int));
+    cudaMalloc(&level_list, num_pins * sizeof(int));
+    cudaMalloc(&primary_outputs, num_POs * sizeof(index_type));
+
+    cudaMemcpy(pinCap, gtdb.pin_capacitance.data(), num_pins * (NUM_ATTR + 2) * sizeof(float), cudaMemcpyHostToDevice);
+    cudaMemcpy(net_is_clock, gtdb.net_is_clock.data(), num_nets * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(primary_outputs, gtdb.primary_outputs.data(), gtdb.primary_outputs.size() * sizeof(index_type), cudaMemcpyHostToDevice);
+
+
+    allocator = new GPULutAllocator();
+    allocator->AllocateBatch(gtdb.liberty_timing_arcs);
+    allocator->CopyToGPU();
+    cudaMalloc((void **)&d_allocator, sizeof(GPULutAllocator));
+    cudaMemcpy(d_allocator, allocator, sizeof(GPULutAllocator), cudaMemcpyHostToDevice);
+    allocator->CopyToGPU(d_allocator);
+
+    logger.info("GPUTimer initialized");
+
+    cudaMalloc(&__pinSlew__, num_pins * NUM_ATTR * sizeof(float));
+    cudaMalloc(&__pinLoad__, num_pins * NUM_ATTR * sizeof(float));
+    cudaMalloc(&__pinRAT__, num_pins * NUM_ATTR * sizeof(float));
+    cudaMalloc(&__pinAT__, num_pins * NUM_ATTR * sizeof(float));
+
+    device_copy_batch<float><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(pinSlew, __pinSlew__, num_pins * NUM_ATTR);
+    device_copy_batch<float><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(pinLoad, __pinLoad__, num_pins * NUM_ATTR);
+    device_copy_batch<float><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(pinRAT, __pinRAT__, num_pins * NUM_ATTR);
+    device_copy_batch<float><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(pinAT, __pinAT__, num_pins * NUM_ATTR);
+}
+
+GPUTimer::~GPUTimer() {
+    logger.info("destruct GPUTimer");
+
+    cudaFree(pinCap);
+    cudaFree(pinWireCap);
+    cudaFree(testRelatedAT);
+    cudaFree(testRAT);
+    cudaFree(testConstraint);
+    cudaFree(pinRootRes);
+    cudaFree(arcSlew);
+
+    cudaFree(net_is_clock);
+    cudaFree(level_list);
+    cudaFree(primary_outputs);
+
+    cudaFree(__pinSlew__);
+    cudaFree(__pinLoad__);
+    cudaFree(__pinRAT__);
+    cudaFree(__pinAT__);
+
+    allocator->~GPULutAllocator();
+    cudaFree(d_allocator);
+}
+
+void GPUTimer::update_states() {
+    cudaMemset(pinImpulse, 0, num_pins * NUM_ATTR * sizeof(float));
+    cudaMemset(pinRootRes, 0, num_pins * NUM_ATTR * sizeof(float));
+    cudaMemset(pinRootDelay, 0, num_pins * NUM_ATTR * sizeof(float));
+    cudaMemset(pinWireCap, 0, num_pins * NUM_ATTR * sizeof(float));
+
+    reset_val<float><<<BLOCK_NUMBER(2 * num_arcs * NUM_ATTR), BLOCK_SIZE>>>(arcDelay, 2 * num_arcs * NUM_ATTR);
+    reset_val<float><<<BLOCK_NUMBER(2 * num_arcs * NUM_ATTR), BLOCK_SIZE>>>(arcSlew, 2 * num_arcs * NUM_ATTR);
+    reset_val<float><<<BLOCK_NUMBER(num_tests * NUM_ATTR), BLOCK_SIZE>>>(testRelatedAT, num_tests * NUM_ATTR);
+    reset_val<float><<<BLOCK_NUMBER(num_tests * NUM_ATTR), BLOCK_SIZE>>>(testRAT, num_tests * NUM_ATTR);
+    reset_val<float><<<BLOCK_NUMBER(num_tests * NUM_ATTR), BLOCK_SIZE>>>(testConstraint, num_tests * NUM_ATTR);
+
+    reset_val<index_type><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(at_prefix_pin, num_pins * NUM_ATTR);
+    reset_val<index_type><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(at_prefix_arc, num_pins * NUM_ATTR);
+    reset_val<index_type><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(at_prefix_attr, num_pins * NUM_ATTR);
+
+    device_copy_batch<float><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(__pinSlew__, pinSlew, num_pins * NUM_ATTR);
+    device_copy_batch<float><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(__pinLoad__, pinLoad, num_pins * NUM_ATTR);
+    device_copy_batch<float><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(__pinRAT__, pinRAT, num_pins * NUM_ATTR);
+    device_copy_batch<float><<<BLOCK_NUMBER(num_pins * NUM_ATTR), BLOCK_SIZE>>>(__pinAT__, pinAT, num_pins * NUM_ATTR);
+    cudaDeviceSynchronize();
+}
+
+__global__ void update_endpoints_kernel0(float *pinAT, float *testRAT, int *test_id2_arc_id, index_type *timing_arc_from_pin_id, index_type *timing_arc_to_pin_id, float *endpoints0, int num_tests) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int test_idx = idx >> 2;
+    const int i = idx & 0b11;
+    const int el = i >> 1;
+    const int rf = i & 1;
+    if (test_idx < num_tests) {
+        const int arc_id = test_id2_arc_id[test_idx];
+        const int from_pin_id = timing_arc_from_pin_id[arc_id];
+        const int to_pin_id = timing_arc_to_pin_id[arc_id];
+        if (isnan(pinAT[to_pin_id * NUM_ATTR + i]) || isnan(testRAT[test_idx * NUM_ATTR + i])) return;
+        if (el == 0) {
+            endpoints0[test_idx * NUM_ATTR + i] = pinAT[to_pin_id * NUM_ATTR + i] - testRAT[test_idx * NUM_ATTR + i];
+        } else {
+            endpoints0[test_idx * NUM_ATTR + i] = testRAT[test_idx * NUM_ATTR + i] - pinAT[to_pin_id * NUM_ATTR + i];
+        }
+    }
+}
+
+__global__ void update_endpoints_kernel1(float *pinAT, float *pinRAT, index_type *primary_outputs, float *endpoints1, int num_POs) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int po_idx = idx >> 2;
+    const int i = idx & 0b11;
+    const int el = i >> 1;
+    if (po_idx < num_POs) {
+        const int pin_idx = primary_outputs[po_idx];
+        if (isnan(pinAT[pin_idx * NUM_ATTR + i]) || isnan(pinRAT[pin_idx * NUM_ATTR + i])) return;
+        if (el == 0) {
+            endpoints1[po_idx * NUM_ATTR + i] = pinAT[pin_idx * NUM_ATTR + i] - pinRAT[pin_idx * NUM_ATTR + i];
+        } else {
+            endpoints1[po_idx * NUM_ATTR + i] = pinRAT[pin_idx * NUM_ATTR + i] - pinAT[pin_idx * NUM_ATTR + i];
+        }
+    }
+}
+
+void GPUTimer::update_endpoints() {
+    torch::Tensor endpoints0 = torch::zeros({num_tests, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::kCUDA)).contiguous();
+    torch::Tensor endpoints1 = torch::zeros({num_POs, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::kCUDA)).contiguous();
+    torch::fill_(endpoints0, nanf(""));
+    torch::fill_(endpoints1, nanf(""));
+
+    update_endpoints_kernel0<<<BLOCK_NUMBER(num_tests * NUM_ATTR), BLOCK_SIZE>>>(pinAT, testRAT, test_id2_arc_id, timing_arc_from_pin_id, timing_arc_to_pin_id, endpoints0.data_ptr<float>(), num_tests);
+    update_endpoints_kernel1<<<BLOCK_NUMBER(num_POs * NUM_ATTR), BLOCK_SIZE>>>(pinAT, pinRAT, primary_outputs, endpoints1.data_ptr<float>(), num_POs);
+
+    endpoint_slacks = torch::cat({endpoints0, endpoints1}, 0).contiguous();
+}
+
+}  // namespace gt

cpp_to_py/gputimer/core/GPUTimer.h

@@ -0,0 +1,130 @@
+#pragma once
+
+#include <torch/extension.h>
+#include "common/common.h"
+#include "common/lib/spef/parser-spef.hpp"
+#include "gputimer/base.h"
+
+using std::tuple;
+using std::vector;
+using std::shared_ptr;
+
+namespace gt {
+
+class TimingArc;
+class TimingTorchRawDB;
+class GTDatabase;
+class GPULutAllocator;
+
+class GPUTimer {
+public:
+    GPULutAllocator *allocator;
+    GPULutAllocator *d_allocator;
+    GTDatabase& gtdb;
+    TimingTorchRawDB& timing_raw_db;
+    shared_ptr<GTDatabase> gtdb_holder;
+    shared_ptr<TimingTorchRawDB> timing_raw_db_holder;
+    GPUTimer(shared_ptr<GTDatabase> gtdb_, shared_ptr<TimingTorchRawDB> timing_raw_db_);
+    ~GPUTimer();
+    spef::Spef spef;
+    void read_spef(const std::string& file);
+
+    // === functions ===
+    void initialize();
+    void levelize();
+    void update_rc_timing(torch::Tensor node_lpos, bool record = false, bool load = false, bool conpensation = false);
+    void update_rc_timing_flute(torch::Tensor node_lpos, bool record = false);
+    void update_rc_timing_spef();
+    void update_states();
+    void update_timing();
+    void update_endpoints();
+    
+    float report_wns(int el);
+    float report_tns_elw(int el);
+    tuple<torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor> report_wns_and_tns();
+    torch::Tensor report_pin_slack();
+    torch::Tensor endpoints_index();
+    torch::Tensor report_endpoint_slack();
+    torch::Tensor report_pin_at();
+    torch::Tensor report_pin_rat();
+    torch::Tensor report_pin_slew();
+    torch::Tensor report_pin_load();
+
+    tuple<vector<int64_t>, vector<float>, vector<float>> report_path(int ep_idx = -1, int el = -1, bool verbose = false);
+    vector<vector<int64_t>> report_K_path(int K, bool verbose = false);
+    tuple<torch::Tensor, torch::Tensor> report_criticality(int K, bool verbose = false, bool deterministic = true);
+    tuple<torch::Tensor, torch::Tensor> report_criticality_threshold(float thrs, bool verbose = false, bool deterministic = true);
+
+public:
+    float time_unit() const;
+
+public:
+    int num_pins, num_arcs, num_timings, num_tests, num_POs, total_num_fanouts;
+    float *pinSlew, *pinLoad, *pinRAT, *pinAT;
+    float *pinImpulse, *pinRootDelay, *pinRootRes;
+    float *arcDelay, *arcSlew;
+    float *pinCap, *pinWireCap;
+    float *testRelatedAT, *testConstraint, *testRAT;
+
+    float *__pinSlew__, *__pinLoad__, *__pinRAT__, *__pinAT__;
+    float *pinImpulse_ref, *pinLoad_ref, *pinRootDelay_ref;
+    float *pinLoad_ratio, *pinRootDelay_ratio;
+
+    int* pin_num_fanin;
+    index_type *pin_fanout_list_end, *pin_fanout_list;
+    index_type *pin_forward_arc_list_end, *pin_forward_arc_list;
+    index_type *pin_backward_arc_list_end, *pin_backward_arc_list;
+    index_type *timing_arc_from_pin_id, *timing_arc_to_pin_id;
+    int *arc_types, *timing_arc_id_map, *arc_id2test_id;
+    int* test_id2_arc_id;
+
+    index_type* primary_outputs;
+    TimingArc* liberty_timing_arcs;
+    index_type *level_list_end, *level_list;
+    vector<int> level_list_end_cpu;
+    int* net_is_clock;
+
+    float clock_period;
+
+public:
+    float* x;
+    float* y;
+    const float* init_x;
+    const float* init_y;
+    const float* node_size_x;
+    const float* node_size_y;
+
+    const float* pin_offset_x;
+    const float* pin_offset_y;
+    index_type *at_prefix_pin;
+    index_type *at_prefix_arc;
+    index_type *at_prefix_attr;
+
+    const int* flat_node2pin_start_map;
+    const int* flat_node2pin_map;
+    const int* pin2node_map;
+
+    const int* flat_net2pin_start_map;
+    const int* flat_net2pin_map;
+    const int* pin2net_map;
+    const bool* net_mask;
+
+    /* row info */
+    int num_nets;
+    int num_movable_nodes;
+    int num_nodes;
+
+    int num_threads;
+
+    float wire_resistance_per_micron;
+    float wire_capacitance_per_micron;
+    int microns;
+    float scale_factor;
+    float res_unit;
+    float cap_unit;
+
+    torch::Tensor pin_slacks;
+    torch::Tensor endpoint_slacks;
+};
+
+}  // namespace gt

cpp_to_py/gputimer/core/gputiming.h

@@ -0,0 +1,370 @@
+#pragma once
+
+#include <vector>
+#include "common/lib/Lut.h"
+#include "common/lib/Timing.h"
+
+using std::vector;
+
+namespace gt {
+
+template <typename T>
+__device__ int lower_bound(T *arr, int size, T val) {
+    int l = 0, r = size - 1;
+    while (l < r) {
+        int m = (l + r) / 2;
+        if (arr[m] < val)
+            l = m + 1;
+        else
+            r = m;
+    }
+    return l;
+}
+
+template <typename T>
+__device__ float interpolate(T x1, T x2, T y1, T y2, T x) {
+    if (x1 == x2) return y1;
+    return y1 + (y2 - y1) * (x - x1) / (x2 - x1);
+}
+
+class GPULutAllocator {
+public:
+    // LUT attributes
+    int num_luts_in_timing = 6;
+    int num_luts;
+    int x_size = 0, y_size = 0, table_size = 0;
+
+    int *num_x, *num_y, *num_table;
+    float *x_array, *y_array, *table_array;
+    size_t *x_offset, *y_offset, *table_offset;
+    bool *allocated;
+
+    int *d_num_x, *d_num_y, *d_num_table;
+    float *d_x_array, *d_y_array, *d_table_array;
+    size_t *d_x_offset, *d_y_offset, *d_table_offset;
+    bool *d_allocated;
+
+    // Timing attributes
+    int num_timings;
+    int *timing_sense;
+    int *lut_template_var;
+    bool *is_rising_edge_triggered, *is_falling_edge_triggered, *is_constraint;
+
+    int *d_timing_sense;
+    int *d_lut_template_var;
+    bool *d_is_rising_edge_triggered, *d_is_falling_edge_triggered, *d_is_constraint;
+
+public:
+    GPULutAllocator() = default;
+    __host__ __forceinline__ void AllocateBatch(vector<TimingArc *> timings) {
+        auto check_lut = [&](Lut *lut) {
+            if (!lut) return;
+            if (lut->set_) {
+                x_size += lut->indices1.size();
+                y_size += lut->indices2.size();
+                table_size += lut->table.size();
+            }
+        };
+        num_timings = timings.size();
+        is_rising_edge_triggered = new bool[num_timings];
+        is_falling_edge_triggered = new bool[num_timings];
+        is_constraint = new bool[num_timings];
+        timing_sense = new int[num_timings];
+        num_timings = 0;
+        for (auto timing_ptr : timings) {
+            auto &timing = *timing_ptr;
+            check_lut(timing.cell_delay_[0]);
+            check_lut(timing.cell_delay_[1]);
+            check_lut(timing.transition_[0]);
+            check_lut(timing.transition_[1]);
+            check_lut(timing.constraint_[0]);
+            check_lut(timing.constraint_[1]);
+            is_rising_edge_triggered[num_timings] = timing.is_rising_edge_triggered();
+            is_falling_edge_triggered[num_timings] = timing.is_falling_edge_triggered();
+            is_constraint[num_timings] = timing.is_constraint();
+            if (timing.timing_sense_ != TimingSense::unknown)
+                timing_sense[num_timings] = static_cast<int>(timing.timing_sense_);
+            else
+                timing_sense[num_timings] = -1;
+            num_timings++;
+        }
+
+        num_luts = num_luts_in_timing * timings.size();  // delay/transitions/constraints * rise/fall
+        x_array = new float[x_size];
+        y_array = new float[y_size];
+        table_array = new float[table_size];
+        num_x = new int[num_luts];
+        num_y = new int[num_luts];
+        num_table = new int[num_luts];
+        x_offset = new size_t[num_luts + 1];
+        y_offset = new size_t[num_luts + 1];
+        table_offset = new size_t[num_luts + 1];
+        lut_template_var = new int[num_luts * 2];  // 0:capacitance/1:transition/2:constraint_transition/3:related_transition/4:input_transition
+        allocated = new bool[num_luts];
+        x_offset[0] = 0;
+        y_offset[0] = 0;
+        table_offset[0] = 0;
+
+        num_luts = 0;
+        auto insert_lut = [&](Lut *lut) {
+            if (lut->set_) {
+                num_x[num_luts] = lut->indices1.size();
+                num_y[num_luts] = lut->indices2.size();
+                num_table[num_luts] = lut->table.size();
+                x_offset[num_luts + 1] = x_offset[num_luts] + num_x[num_luts];
+                y_offset[num_luts + 1] = y_offset[num_luts] + num_y[num_luts];
+                table_offset[num_luts + 1] = table_offset[num_luts] + num_table[num_luts];
+                memcpy(x_array + x_offset[num_luts], lut->indices1.data(), lut->indices1.size() * sizeof(float));
+                memcpy(y_array + y_offset[num_luts], lut->indices2.data(), lut->indices2.size() * sizeof(float));
+                memcpy(table_array + table_offset[num_luts], lut->table.data(), lut->table.size() * sizeof(float));
+
+                if (lut->lut_template) {
+                    if (lut->lut_template->variable1)
+                        lut_template_var[num_luts * 2] = static_cast<int>(lut->lut_template->variable1.value());
+                    else
+                        lut_template_var[num_luts * 2] = -1;
+                    if (lut->lut_template->variable2)
+                        lut_template_var[num_luts * 2 + 1] = static_cast<int>(lut->lut_template->variable2.value());
+                    else
+                        lut_template_var[num_luts * 2 + 1] = -1;
+                } else {
+                    lut_template_var[num_luts * 2] = -1;
+                    lut_template_var[num_luts * 2 + 1] = -1;
+                }
+                allocated[num_luts] = true;
+            } else {
+                num_x[num_luts] = 0;
+                num_y[num_luts] = 0;
+                num_table[num_luts] = 0;
+                x_offset[num_luts + 1] = x_offset[num_luts];
+                y_offset[num_luts + 1] = y_offset[num_luts];
+                table_offset[num_luts + 1] = table_offset[num_luts];
+                lut_template_var[num_luts * 2] = -1;      // var1
+                lut_template_var[num_luts * 2 + 1] = -1;  // var2
+                allocated[num_luts] = false;
+            }
+            num_luts++;
+        };
+        for (auto timing_ptr : timings) {
+            auto &timing = *timing_ptr;
+            insert_lut(timing.cell_delay_[0]);
+            insert_lut(timing.cell_delay_[1]);
+            insert_lut(timing.transition_[0]);
+            insert_lut(timing.transition_[1]);
+            insert_lut(timing.constraint_[0]);
+            insert_lut(timing.constraint_[1]);
+        }
+    }
+    __host__ __forceinline__ void CopyToGPU() {
+        cudaMalloc(&d_x_array, x_size * sizeof(float));
+        cudaMalloc(&d_y_array, y_size * sizeof(float));
+        cudaMalloc(&d_table_array, table_size * sizeof(float));
+        cudaMalloc(&d_num_x, num_luts * sizeof(int));
+        cudaMalloc(&d_num_y, num_luts * sizeof(int));
+        cudaMalloc(&d_num_table, num_luts * sizeof(int));
+        cudaMalloc(&d_x_offset, (num_luts + 1) * sizeof(size_t));
+        cudaMalloc(&d_y_offset, (num_luts + 1) * sizeof(size_t));
+        cudaMalloc(&d_table_offset, (num_luts + 1) * sizeof(size_t));
+        cudaMalloc(&d_allocated, num_luts * sizeof(bool));
+        cudaMalloc(&d_is_rising_edge_triggered, num_timings * sizeof(bool));
+        cudaMalloc(&d_is_falling_edge_triggered, num_timings * sizeof(bool));
+        cudaMalloc(&d_is_constraint, num_timings * sizeof(bool));
+        cudaMalloc(&d_timing_sense, num_timings * sizeof(int));
+        cudaMalloc(&d_lut_template_var, 2 * num_luts * sizeof(int));
+
+        cudaMemcpy(d_x_array, x_array, x_size * sizeof(float), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_y_array, y_array, y_size * sizeof(float), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_table_array, table_array, table_size * sizeof(float), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_num_x, num_x, num_luts * sizeof(int), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_num_y, num_y, num_luts * sizeof(int), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_num_table, num_table, num_luts * sizeof(int), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_x_offset, x_offset, (num_luts + 1) * sizeof(size_t), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_y_offset, y_offset, (num_luts + 1) * sizeof(size_t), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_table_offset, table_offset, (num_luts + 1) * sizeof(size_t), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_allocated, allocated, num_luts * sizeof(bool), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_is_rising_edge_triggered, is_rising_edge_triggered, num_timings * sizeof(bool), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_is_falling_edge_triggered, is_falling_edge_triggered, num_timings * sizeof(bool), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_is_constraint, is_constraint, num_timings * sizeof(bool), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_timing_sense, timing_sense, num_timings * sizeof(int), cudaMemcpyHostToDevice);
+        cudaMemcpy(d_lut_template_var, lut_template_var, 2 * num_luts * sizeof(int), cudaMemcpyHostToDevice);
+    }
+    __host__ __forceinline__ void CopyToGPU(GPULutAllocator *d_gpuluts) {
+        cudaMemcpy(&(d_gpuluts->d_num_x), &d_num_x, sizeof(int *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_num_y), &d_num_y, sizeof(int *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_num_table), &d_num_table, sizeof(int *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_x_array), &d_x_array, sizeof(float *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_y_array), &d_y_array, sizeof(float *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_table_array), &d_table_array, sizeof(float *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_x_offset), &d_x_offset, sizeof(size_t *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_y_offset), &d_y_offset, sizeof(size_t *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_table_offset), &d_table_offset, sizeof(size_t *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_allocated), &d_allocated, sizeof(bool *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_is_rising_edge_triggered), &d_is_rising_edge_triggered, sizeof(bool *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_is_falling_edge_triggered), &d_is_falling_edge_triggered, sizeof(bool *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_timing_sense), &d_timing_sense, sizeof(int *), cudaMemcpyHostToDevice);
+        cudaMemcpy(&(d_gpuluts->d_lut_template_var), &d_lut_template_var, sizeof(int *), cudaMemcpyHostToDevice);
+    }
+
+    __device__ __forceinline__ bool is_input_transition_defined(int timing_id, int irf) {
+        if (d_is_rising_edge_triggered[timing_id] && irf != 0) return false;
+        if (d_is_falling_edge_triggered[timing_id] && irf != 1) return false;
+        return true;
+    }
+
+    __device__ __forceinline__ bool is_transition_defined(int timing_id, int irf, int orf) {
+        if (!is_input_transition_defined(timing_id, irf)) return false;
+        int sense = d_timing_sense[timing_id];
+
+        if (sense != -1) {
+            switch (sense) {
+                case 1:
+                    if (irf != orf) return false;
+                    break;
+                case 2:
+                    if (irf == orf) return false;
+                    break;
+                default:
+                    break;
+            }
+        }
+
+        return true;
+    }
+
+    __device__ __forceinline__ float lut(int in_timing_lut, float x, float y) {
+        if (d_num_x[in_timing_lut] < 1 || d_num_y[in_timing_lut] < 1) {
+            // return std::nullopt;
+            return nanf("");
+        }
+
+        if (d_num_table[in_timing_lut] == 1) {
+            // return std::nullopt;
+            return d_table_array[d_table_offset[in_timing_lut]];
+        }
+        int x_idx[2], y_idx[2];
+
+        x_idx[1] = lower_bound<float>(d_x_array + d_x_offset[in_timing_lut], d_num_x[in_timing_lut], x);
+        y_idx[1] = lower_bound<float>(d_y_array + d_y_offset[in_timing_lut], d_num_y[in_timing_lut], y);
+
+        x_idx[1] = max(1, min(d_num_x[in_timing_lut] - 1, x_idx[1]));
+        y_idx[1] = max(1, min(d_num_y[in_timing_lut] - 1, y_idx[1]));
+        x_idx[0] = x_idx[1] - 1;
+        y_idx[0] = y_idx[1] - 1;
+        if (d_num_x[in_timing_lut] == 1) x_idx[1] = 0;
+        if (d_num_y[in_timing_lut] == 1) y_idx[1] = 0;
+
+        // interpolation
+        float numeric[2];
+        numeric[0] = interpolate<float>(d_x_array[d_x_offset[in_timing_lut] + x_idx[0]],
+                                        d_x_array[d_x_offset[in_timing_lut] + x_idx[1]],
+                                        d_table_array[d_table_offset[in_timing_lut] + x_idx[0] * d_num_y[in_timing_lut] + y_idx[0]],
+                                        d_table_array[d_table_offset[in_timing_lut] + x_idx[1] * d_num_y[in_timing_lut] + y_idx[0]],
+                                        x);
+        numeric[1] = interpolate<float>(d_x_array[d_x_offset[in_timing_lut] + x_idx[0]],
+                                        d_x_array[d_x_offset[in_timing_lut] + x_idx[1]],
+                                        d_table_array[d_table_offset[in_timing_lut] + x_idx[0] * d_num_y[in_timing_lut] + y_idx[1]],
+                                        d_table_array[d_table_offset[in_timing_lut] + x_idx[1] * d_num_y[in_timing_lut] + y_idx[1]],
+                                        x);
+
+        return interpolate<float>(d_y_array[d_y_offset[in_timing_lut] + y_idx[0]], d_y_array[d_y_offset[in_timing_lut] + y_idx[1]], numeric[0], numeric[1], y);
+    }
+
+    __device__ __forceinline__ float query(int timing_id, int irf, int orf, float slew_or_related, float load_or_constraint, int type) {  // 0:cell/1:trans/3:constraint
+        if (!is_transition_defined(timing_id, irf, orf)) {
+            // return std::nullopt;
+            return nanf("");
+        }
+
+        int in_timing_lut = num_luts_in_timing * timing_id + orf + type * 2;
+        in_timing_lut = d_allocated[in_timing_lut] ? in_timing_lut : -1;
+
+        if (in_timing_lut == -1) {
+            // return std::nullopt;
+            return nanf("");
+        }
+
+        float val1{0.0f}, val2{0.0f};
+
+        if (type == 0 || type == 1) {
+            switch (d_lut_template_var[in_timing_lut * 2]) {
+                case 0:  // LutVar::TOTAL_OUTPUT_NET_CAPACITANCE
+                    if (d_lut_template_var[in_timing_lut * 2 + 1] != -1) {
+                        assert(d_lut_template_var[in_timing_lut * 2 + 1] == 1);  // LutVar::INPUT_NET_TRANSITION
+                    }
+                    val1 = load_or_constraint;
+                    val2 = slew_or_related;
+                    break;
+                case 1:  // LutVar::INPUT_NET_TRANSITION
+                    if (d_lut_template_var[in_timing_lut * 2 + 1] != -1) {
+                        assert(d_lut_template_var[in_timing_lut * 2 + 1] == 0);  // LutVar::TOTAL_OUTPUT_NET_CAPACITANCE
+                    }
+                    val1 = slew_or_related;
+                    val2 = load_or_constraint;
+                    break;
+                default:
+                    // printf("Invalid lut template variable\n");
+                    break;
+            }
+        } else if (type == 2) {
+            switch (d_lut_template_var[in_timing_lut * 2]) {
+                case 2:  // LutVar::CONSTRAINED_PIN_TRANSITION
+                    if (d_lut_template_var[in_timing_lut * 2 + 1] != -1) {
+                        assert(d_lut_template_var[in_timing_lut * 2 + 1] == 3);  // LutVar::RELATED_PIN_TRANSITION
+                    }
+                    val1 = load_or_constraint;
+                    val2 = slew_or_related;
+                    break;
+                case 3:  // LutVar::RELATED_PIN_TRANSITION
+                    if (d_lut_template_var[in_timing_lut * 2 + 1] != -1) {
+                        assert(d_lut_template_var[in_timing_lut * 2 + 1] == 2);  // LutVar::CONSTRAINED_PIN_TRANSITION
+                    }
+                    val1 = slew_or_related;
+                    val2 = load_or_constraint;
+                    break;
+                default:
+                    // printf("Invalid lut template variable\n");
+                    break;
+            }
+        }
+
+        return lut(in_timing_lut, val1, val2);
+    }
+
+    __host__ __forceinline__ void freeMem() {
+        if (allocated) {
+            logger.info("destruct gputiming");
+            delete[] num_x;
+            delete[] num_y;
+            delete[] num_table;
+            delete[] x_array;
+            delete[] y_array;
+            delete[] table_array;
+            delete[] x_offset;
+            delete[] y_offset;
+            delete[] table_offset;
+            delete[] allocated;
+            delete[] is_rising_edge_triggered;
+            delete[] is_falling_edge_triggered;
+            delete[] is_constraint;
+            delete[] timing_sense;
+            cudaFree(d_num_x);
+            cudaFree(d_num_y);
+            cudaFree(d_num_table);
+            cudaFree(d_x_array);
+            cudaFree(d_y_array);
+            cudaFree(d_table_array);
+            cudaFree(d_x_offset);
+            cudaFree(d_y_offset);
+            cudaFree(d_table_offset);
+            cudaFree(d_allocated);
+            cudaFree(d_is_rising_edge_triggered);
+            cudaFree(d_is_falling_edge_triggered);
+            cudaFree(d_is_constraint);
+            cudaFree(d_timing_sense);
+        }
+    }
+
+    __host__ ~GPULutAllocator() { freeMem(); }
+};
+
+}  // namespace gt

cpp_to_py/gputimer/core/levelize.cu

@@ -0,0 +1,82 @@
+
+#include "GPUTimer.h"
+#include "utils.cuh"
+#include "gputimer/db/GTDatabase.h"
+
+namespace gt {
+
+__global__ void advanceLevel(index_type *frontiers,
+                             index_type *next_frontiers,
+                             index_type *level_list,
+                             index_type *pin_fanout_list_end,
+                             index_type *pin_fanout_list,
+                             int *pin_num_fanin,
+                             int num_frontiers,
+                             int *next_num_frontiers,
+                             int *last_idx) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx < num_frontiers) {
+        index_type pin_id = frontiers[idx];
+        index_type ptr = atomicAdd(last_idx, 1);
+        level_list[ptr] = pin_id;
+        for (index_type i = pin_fanout_list_end[pin_id]; i < pin_fanout_list_end[pin_id + 1]; i++) {
+            index_type fo_pin_id = pin_fanout_list[i];
+            int prev_num = atomicAdd(&pin_num_fanin[fo_pin_id], -1);
+            if (prev_num == 1) {
+                index_type end = atomicAdd(next_num_frontiers, 1);
+                next_frontiers[end] = fo_pin_id;
+            }
+        }
+    }
+}
+
+void checkTimingGraph(index_type* level_list_cpu, int num_pins, vector<std::string> pin_names) {
+    // check which pins are not in timing graph
+    std::set<index_type> pins;
+    for (int i = 0; i < num_pins; i++) {
+        pins.insert(i);
+    }
+    for (int i = 0; i < num_pins; i++) {
+        index_type pin_id = level_list_cpu[i];
+        pins.erase(pin_id);
+    }
+    for (auto pin_id : pins) {
+        printf("Unconnected pin_id: %d, name: %s\n", pin_id, pin_names[pin_id].c_str());
+    }
+}
+
+
+void GPUTimer::levelize() {
+    index_type *frontiers, *next_frontiers;
+    int *next_num_frontiers, *last_idx;
+    int num_frontiers = gtdb.pin_frontiers.size();
+    cudaMalloc(&frontiers, num_pins * sizeof(index_type));
+    cudaMalloc(&next_frontiers, num_pins * sizeof(index_type));
+    cudaMalloc(&next_num_frontiers, sizeof(int));
+    cudaMalloc(&last_idx, sizeof(int));
+    cudaMemset(next_num_frontiers, 0, sizeof(int));
+    cudaMemset(last_idx, 0, sizeof(int));
+    cudaMemcpy(frontiers, gtdb.pin_frontiers.data(), num_frontiers * sizeof(index_type), cudaMemcpyHostToDevice);
+
+    level_list_end_cpu.clear();
+    level_list_end_cpu.push_back(0);
+    int total_num_frontiers = 0;
+    while (num_frontiers) {
+        total_num_frontiers += num_frontiers;
+        level_list_end_cpu.push_back(total_num_frontiers);
+        advanceLevel<<<BLOCK_NUMBER(num_pins), BLOCK_SIZE>>>(
+            frontiers, next_frontiers, level_list, pin_fanout_list_end, pin_fanout_list, pin_num_fanin, num_frontiers, next_num_frontiers, last_idx);
+        cudaMemcpy(&num_frontiers, next_num_frontiers, sizeof(int), cudaMemcpyDeviceToHost);
+        device_copy<index_type><<<1, 1>>>(next_frontiers, frontiers, num_frontiers);
+        cudaMemset(next_num_frontiers, 0, sizeof(int));
+        // debugPrint<int><<<1, 1>>>(next_num_frontiers, 1);
+    }
+    cudaMalloc(&level_list_end, level_list_end_cpu.size() * sizeof(index_type));
+    cudaMemcpy(level_list_end, level_list_end_cpu.data(), level_list_end_cpu.size() * sizeof(index_type), cudaMemcpyHostToDevice);
+    index_type *level_list_cpu = new index_type[total_num_frontiers];
+    cudaMemcpy(level_list_cpu, level_list, total_num_frontiers * sizeof(index_type), cudaMemcpyDeviceToHost);
+
+    // checkTimingGraph(level_list_cpu, num_pins, gtdb.pin_names);
+}
+
+} // namespace gt

cpp_to_py/gputimer/core/path.cpp

@@ -0,0 +1,169 @@
+
+#include "GPUTimer.h"
+#include "gputimer/db/GTDatabase.h"
+
+using std::setw;
+using std::string;
+using std::endl;
+using std::cout;
+using std::tuple;
+
+namespace gt {
+
+// ------------------------------------------------------------------------------------------------------------------------
+// Report timing paths
+//
+tuple<vector<int64_t>, vector<float>, vector<float>> GPUTimer::report_path(int ep_idx, int el, bool verbose) {
+    if (!pin_slacks.numel()) {
+        pin_slacks = torch::zeros_like(timing_raw_db.pinAT, torch::dtype(torch::kFloat32).device(timing_raw_db.pinAT.device()));
+        auto s1 = timing_raw_db.pinAT - timing_raw_db.pinRAT;
+        auto s2 = timing_raw_db.pinRAT - timing_raw_db.pinAT;
+        pin_slacks.index({"...", torch::indexing::Slice(0, 2)}).data().copy_(s1.index({"...", torch::indexing::Slice(0, 2)}));
+        pin_slacks.index({"...", torch::indexing::Slice(2, 4)}).data().copy_(s2.index({"...", torch::indexing::Slice(2, 4)}));
+    }
+    int worst_ep_idx;
+    int worst_ep_i;
+    if (ep_idx == -1) {
+        auto ep_slacks = torch::nan_to_num(pin_slacks.index_select(0, timing_raw_db.endpoints_id), FLT_MAX);
+        if (el != -1) ep_slacks = ep_slacks.index({"...", torch::indexing::Slice(2 * el, 2 * (el + 1))});
+
+        auto [slack_elw, order] = torch::min(ep_slacks, 1);
+        auto worst_ep = torch::argmin(slack_elw).item<int>();
+        worst_ep_idx = timing_raw_db.endpoints_id[worst_ep].item<int>();
+        worst_ep_i = torch::argmin(ep_slacks[worst_ep]).item<int>();
+    } else {
+        worst_ep_idx = ep_idx;
+        worst_ep_i = torch::argmin(pin_slacks[worst_ep_idx]).item<int>();
+    }
+    worst_ep_i = el == -1 ? worst_ep_i : worst_ep_i + 2 * el;
+
+    int cur = worst_ep_idx;
+    int to_i = worst_ep_i;
+    vector<long> path;
+    vector<float> path_at;
+    vector<float> path_rat;
+    vector<float> to_delay;
+    vector<float> path_slack;
+
+    while (cur != -1) {
+        int prev = timing_raw_db.at_prefix_pin[cur][to_i].item<int>();
+        int arc_id = timing_raw_db.at_prefix_arc[cur][to_i].item<int>();
+        int from_i = timing_raw_db.at_prefix_attr[cur][to_i].item<int>();
+        int from_el = from_i >> 1;
+        int to_el = to_i >> 1;
+
+        int arc_i = (from_i << 1) + (to_i & 0b1);
+        float at = 0;
+        float rat = 0;
+        float delay = 0;
+        float slack = 0;
+        if (prev != -1) {
+            at = timing_raw_db.pinAT[cur][to_i].item<float>();
+            rat = timing_raw_db.pinRAT[cur][to_i].item<float>();
+            delay = timing_raw_db.arcDelay[arc_id][arc_i].item<float>();
+            slack = pin_slacks[cur][to_i].item<float>();
+        }
+        path.push_back(cur);
+        path_at.push_back(at);
+        path_rat.push_back(rat);
+        to_delay.push_back(delay);
+        path_slack.push_back(slack);
+
+        cur = prev;
+        to_i = from_i;
+    }
+
+    if (verbose) {
+        cout << std::fixed << std::setprecision(3);
+        cout << '\n' << std::setw(10) << "Type" << std::setw(10) << "Delay" << std::setw(10) << "AT" << std::setw(10) << "RAT" << std::setw(10) << "Slack" << std::setw(10) << "Pin" << '\n';
+        
+        for (int i = path.size() - 1; i >= 0; i--) {
+            int cur = path[i];
+            cout << setw(10) << "pin " << setw(10) << to_delay[i] << setw(10) << path_at[i] << setw(10) << path_rat[i] << setw(10) << path_slack[i];
+            std::fill_n(std::ostream_iterator<char>(cout), 3, ' ');
+            cout << gtdb.pin_names[cur] << '\n';
+        }
+    }
+
+    return {path, path_at, to_delay};
+}
+
+vector<vector<int64_t>> GPUTimer::report_K_path(int K, bool verbose) {
+    if (!pin_slacks.numel()) {
+        pin_slacks = torch::zeros_like(timing_raw_db.pinAT, torch::dtype(torch::kFloat32).device(timing_raw_db.pinAT.device()));
+        auto s1 = timing_raw_db.pinAT - timing_raw_db.pinRAT;
+        auto s2 = timing_raw_db.pinRAT - timing_raw_db.pinAT;
+        pin_slacks.index({"...", torch::indexing::Slice(0, 2)}).data().copy_(s1.index({"...", torch::indexing::Slice(0, 2)}));
+        pin_slacks.index({"...", torch::indexing::Slice(2, 4)}).data().copy_(s2.index({"...", torch::indexing::Slice(2, 4)}));
+    }
+    auto [endpoints_id, tmp1] = torch::_unique(timing_raw_db.endpoints_id);
+    auto ep_slacks = torch::nan_to_num(pin_slacks.index_select(0, endpoints_id));
+    auto [ep_slack_elw, ep_i_indices] = torch::min(ep_slacks, 1);
+    auto [ep_slack_elw_ordered, indices] = torch::sort(ep_slack_elw, false);
+
+    indices = indices.contiguous();
+    ep_i_indices = ep_i_indices.contiguous();
+    endpoints_id = endpoints_id.contiguous();
+
+    vector<vector<int64_t>> paths;
+
+    for (int i = 0; i < K; i++) {
+        auto [path, path_at, to_delay] = report_path(endpoints_id[indices[i]].item<int>(), false);
+        paths.push_back(path);
+    }
+    return paths;
+}
+
+// ------------------------------------------------------------------------------------------------------------------------
+// Report timing paths
+//
+tuple<torch::Tensor, torch::Tensor> explore_path(index_type* at_prefix_pin,
+                                                 index_type* at_prefix_arc,
+                                                 int* at_prefix_attr,
+                                                 float* pinAT,
+                                                 int* arc_types,
+                                                 float* arcDelay,
+                                                 torch::Tensor indices,
+                                                 torch::Tensor ep_i_indices,
+                                                 torch::Tensor endpoints_id,
+                                                 int num_pins,
+                                                 int K,
+                                                 bool deterministic);
+
+tuple<torch::Tensor, torch::Tensor> GPUTimer::report_criticality(int K, bool verbose, bool deterministic) {
+    auto [endpoints_id, tmp1] = torch::_unique(timing_raw_db.endpoints_id);
+    auto ep_slacks = torch::nan_to_num(pin_slacks.index_select(0, endpoints_id));
+    auto [ep_slack_elw, ep_i_indices] = torch::min(ep_slacks, 1);
+    auto [ep_slack_elw_ordered, indices] = torch::sort(ep_slack_elw, false);
+
+    indices = indices.contiguous();
+    ep_i_indices = ep_i_indices.contiguous();
+    endpoints_id = endpoints_id.contiguous();
+
+    auto [from_pin_delay, pin_visited] = explore_path(
+        at_prefix_pin, at_prefix_arc, at_prefix_attr, pinAT, arc_types, arcDelay, indices, ep_i_indices, endpoints_id, num_pins, K, deterministic);
+
+    return {from_pin_delay, pin_visited};
+}
+
+tuple<torch::Tensor, torch::Tensor> GPUTimer::report_criticality_threshold(float thrs, bool verbose, bool deterministic) {
+    auto [endpoints_id, tmp1] = torch::_unique(timing_raw_db.endpoints_id);
+    auto ep_slacks = torch::nan_to_num(pin_slacks.index_select(0, endpoints_id));
+    auto [ep_slack_elw, ep_i_indices] = torch::min(ep_slacks, 1);
+    auto [ep_slack_elw_ordered, indices] = torch::sort(ep_slack_elw, false);
+
+    auto threshold = thrs * ep_slack_elw_ordered[0];
+    int K = (ep_slack_elw_ordered < threshold).sum().item<int>();
+
+    indices = indices.contiguous();
+    ep_i_indices = ep_i_indices.contiguous();
+    endpoints_id = endpoints_id.contiguous();
+
+    auto [from_pin_delay, pin_visited] = explore_path(
+        at_prefix_pin, at_prefix_arc, at_prefix_attr, pinAT, arc_types, arcDelay, indices, ep_i_indices, endpoints_id, num_pins, K, deterministic);
+
+    return {from_pin_delay, pin_visited};
+}
+
+
+}

cpp_to_py/gputimer/core/path.cu

@@ -0,0 +1,185 @@
+#include <ATen/cuda/CUDAContext.h>
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <torch/extension.h>
+
+#include "utils.cuh"
+
+namespace gt {
+
+__global__ void explore_path_kernel(index_type* at_prefix_pin,
+                                    index_type* at_prefix_arc,
+                                    int* at_prefix_attr,
+                                    float* pinAT,
+                                    int* arc_types,
+                                    float* arcDelay,
+                                    torch::PackedTensorAccessor32<int64_t, 1, torch::RestrictPtrTraits> indices,
+                                    torch::PackedTensorAccessor32<int64_t, 1, torch::RestrictPtrTraits> ep_i_indices,
+                                    torch::PackedTensorAccessor32<int, 1, torch::RestrictPtrTraits> endpoints_index,
+                                    float* from_pin_delay,
+                                    torch::PackedTensorAccessor32<int, 1, torch::RestrictPtrTraits> pin_visited,
+                                    int K) {
+    const int k = blockIdx.x * blockDim.x + threadIdx.x;
+    if (k < K) {
+        index_type cur_id = endpoints_index[indices[k]];
+        int to_i = ep_i_indices[indices[k]];
+        while (cur_id != -1) {
+            atomicAdd(&pin_visited[cur_id], 1);
+            int prev_id = at_prefix_pin[cur_id * NUM_ATTR + to_i];
+            int arc_id = at_prefix_arc[cur_id * NUM_ATTR + to_i];
+            int from_i = at_prefix_attr[cur_id * NUM_ATTR + to_i];
+            int arc_i = (from_i << 1) + (to_i & 0b1);
+            float at = 0;
+            float delay = 0;
+            if (prev_id != -1) {
+                at = pinAT[cur_id * NUM_ATTR + to_i];
+                delay = arcDelay[arc_id * 2 * NUM_ATTR + arc_i] / pow(1 + k, 2);
+
+                if (arc_types[arc_id] == 0) {
+                    atomicAdd(&from_pin_delay[cur_id], delay);
+                } else {
+                    atomicAdd(&from_pin_delay[prev_id], delay);
+                }
+            }
+            cur_id = prev_id;
+            to_i = from_i;
+        }
+    }
+}
+
+__global__ void explore_path_deterministic_kernel(index_type* at_prefix_pin,
+                                                  index_type* at_prefix_arc,
+                                                  int* at_prefix_attr,
+                                                  float* pinAT,
+                                                  int* arc_types,
+                                                  float* arcDelay,
+                                                  torch::PackedTensorAccessor32<int64_t, 1, torch::RestrictPtrTraits> indices,
+                                                  torch::PackedTensorAccessor32<int64_t, 1, torch::RestrictPtrTraits> ep_i_indices,
+                                                  torch::PackedTensorAccessor32<int, 1, torch::RestrictPtrTraits> endpoints_index,
+                                                  unsigned long long* from_pin_delay,
+                                                  torch::PackedTensorAccessor32<int, 1, torch::RestrictPtrTraits> pin_visited,
+                                                  int K,
+                                                  unsigned long long scalar) {
+    const int k = blockIdx.x * blockDim.x + threadIdx.x;
+    if (k < K) {
+        index_type cur_id = endpoints_index[indices[k]];
+        int to_i = ep_i_indices[indices[k]];
+        while (cur_id != -1) {
+            atomicAdd(&pin_visited[cur_id], 1);
+            int prev_id = at_prefix_pin[cur_id * NUM_ATTR + to_i];
+            int arc_id = at_prefix_arc[cur_id * NUM_ATTR + to_i];
+            int from_i = at_prefix_attr[cur_id * NUM_ATTR + to_i];
+            int arc_i = (from_i << 1) + (to_i & 0b1);
+            float at = 0;
+            float delay = 0;
+            if (prev_id != -1) {
+                at = pinAT[cur_id * NUM_ATTR + to_i];
+                delay = arcDelay[arc_id * 2 * NUM_ATTR + arc_i] / pow(1 + k, 2);
+
+                if (arc_types[arc_id] == 0) {
+                    atomicAdd(&from_pin_delay[cur_id], static_cast<unsigned long long>(delay * scalar));
+                } else {
+                    atomicAdd(&from_pin_delay[prev_id], static_cast<unsigned long long>(delay * scalar));
+                }
+            }
+            cur_id = prev_id;
+            to_i = from_i;
+        }
+    }
+}
+
+__global__ void copyFromFloatAuxArray(
+    unsigned long long* aux_array_uint64, float* aux_array, unsigned long long scalar, float inv_scalar, int num_pins) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < num_pins) {
+        aux_array_uint64[i] = static_cast<unsigned long long>(aux_array[i] * scalar);
+    }
+}
+
+__global__ void copyToFloatAuxArray(
+    unsigned long long* aux_array_uint64, float* aux_array, unsigned long long scalar, float inv_scalar, int num_pins) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < num_pins) {
+        aux_array[i] = static_cast<float>(inv_scalar * aux_array_uint64[i]);
+    }
+}
+
+std::tuple<torch::Tensor, torch::Tensor> explore_path(index_type* at_prefix_pin,
+                                                      index_type* at_prefix_arc,
+                                                      int* at_prefix_attr,
+                                                      float* pinAT,
+                                                      int* arc_types,
+                                                      float* arcDelay,
+                                                      torch::Tensor indices,
+                                                      torch::Tensor ep_i_indices,
+                                                      torch::Tensor endpoints_index,
+                                                      int num_pins,
+                                                      int K,
+                                                      bool deterministic) {
+    cudaSetDevice(indices.get_device());
+    auto stream = at::cuda::getCurrentCUDAStream();
+    torch::Tensor from_pin_delay = torch::zeros({num_pins}, torch::dtype(torch::kFloat32).device(indices.device())).contiguous();
+    torch::Tensor pin_visited = torch::zeros({num_pins}, torch::dtype(torch::kInt32).device(indices.device())).contiguous();
+
+    int numThreads = 512;
+    int numBlocks = (K + numThreads - 1) / numThreads;
+
+    if (deterministic) {
+        // use cache to save runtime
+        int cp_threads = 512;
+        int cp_blocks = (num_pins + cp_threads - 1) / cp_threads;
+
+        // TODO: use a better way to determine the scalar
+        int max_value_bits = 32;
+        int scalar_bits = max(64 - max_value_bits, 0);
+        unsigned long long scalar = (1UL << scalar_bits);
+        float inv_scalar = 1.0 / static_cast<float>(scalar);
+        static unsigned long long* aux_array_uint64_ptr = nullptr;
+        static int aux_array_uint64_size = -1;
+        if (aux_array_uint64_ptr == nullptr) {
+            aux_array_uint64_size = num_pins;
+            cudaMalloc(&aux_array_uint64_ptr, num_pins * sizeof(unsigned long long));
+        } else if (num_pins != aux_array_uint64_size) {
+            cudaFree(aux_array_uint64_ptr);
+            aux_array_uint64_ptr = nullptr;
+            aux_array_uint64_size = num_pins;
+            cudaMalloc(&aux_array_uint64_ptr, aux_array_uint64_size * sizeof(unsigned long long));
+        }
+        copyFromFloatAuxArray<<<cp_blocks, cp_threads, 0, stream>>>(
+            aux_array_uint64_ptr, from_pin_delay.data_ptr<float>(), scalar, inv_scalar, num_pins);
+        explore_path_deterministic_kernel<<<numBlocks, numThreads, 0, stream>>>(
+            at_prefix_pin,
+            at_prefix_arc,
+            at_prefix_attr,
+            pinAT,
+            arc_types,
+            arcDelay,
+            indices.packed_accessor32<int64_t, 1, torch::RestrictPtrTraits>(),
+            ep_i_indices.packed_accessor32<int64_t, 1, torch::RestrictPtrTraits>(),
+            endpoints_index.packed_accessor32<int, 1, torch::RestrictPtrTraits>(),
+            aux_array_uint64_ptr,
+            pin_visited.packed_accessor32<int, 1, torch::RestrictPtrTraits>(),
+            K,
+            scalar);
+        copyToFloatAuxArray<<<cp_blocks, cp_threads, 0, stream>>>(
+            aux_array_uint64_ptr, from_pin_delay.data_ptr<float>(), scalar, inv_scalar, num_pins);
+
+    } else {
+        explore_path_kernel<<<numBlocks, numThreads, 0, stream>>>(at_prefix_pin,
+                                                                  at_prefix_arc,
+                                                                  at_prefix_attr,
+                                                                  pinAT,
+                                                                  arc_types,
+                                                                  arcDelay,
+                                                                  indices.packed_accessor32<int64_t, 1, torch::RestrictPtrTraits>(),
+                                                                  ep_i_indices.packed_accessor32<int64_t, 1, torch::RestrictPtrTraits>(),
+                                                                  endpoints_index.packed_accessor32<int, 1, torch::RestrictPtrTraits>(),
+                                                                  from_pin_delay.data_ptr<float>(),
+                                                                  pin_visited.packed_accessor32<int, 1, torch::RestrictPtrTraits>(),
+                                                                  K);
+    }
+
+    return {from_pin_delay, pin_visited};
+}
+
+}  // namespace gt

cpp_to_py/gputimer/core/propagate.cpp

@@ -0,0 +1,67 @@
+
+#include "GPUTimer.h"
+
+namespace gt {
+
+void update_timing_cuda(index_type *level_list,
+                        vector<int> level_list_end_cpu,
+                        index_type *pin_forward_arc_list_end,
+                        index_type *pin_forward_arc_list,
+                        index_type *timing_arc_to_pin_id,
+                        index_type *pin_backward_arc_list_end,
+                        index_type *pin_backward_arc_list,
+                        index_type *timing_arc_from_pin_id,
+                        int *arc_types,
+                        int *arc_id2test_id,
+                        float *pinSlew,
+                        float *pinLoad,
+                        float *pinImpulse,
+                        float *pinRootDelay,
+                        float *pinAT,
+                        float *pinRAT,
+                        float *testRelatedAT,
+                        float *testRAT,
+                        float *testConstraint,
+                        float *arcDelay,
+                        int *timing_arc_id_map,
+                        index_type *at_prefix_pin,
+                        index_type *at_prefix_arc,
+                        index_type *at_prefix_attr,
+                        float clock_period,
+                        GPULutAllocator *d_allocator,
+                        int num_pins,
+                        bool deterministic);
+
+
+void GPUTimer::update_timing() {
+    update_timing_cuda(level_list,
+                       level_list_end_cpu,
+                       pin_forward_arc_list_end,
+                       pin_forward_arc_list,
+                       timing_arc_to_pin_id,
+                       pin_backward_arc_list_end,
+                       pin_backward_arc_list,
+                       timing_arc_from_pin_id,
+                       arc_types,
+                       arc_id2test_id,
+                       pinSlew,
+                       pinLoad,
+                       pinImpulse,
+                       pinRootDelay,
+                       pinAT,
+                       pinRAT,
+                       testRelatedAT,
+                       testRAT,
+                       testConstraint,
+                       arcDelay,
+                       timing_arc_id_map,
+                       at_prefix_pin,
+                       at_prefix_arc,
+                       at_prefix_attr,
+                       clock_period,
+                       d_allocator,
+                       num_pins,
+                       true);
+}
+
+}  // namespace gt

cpp_to_py/gputimer/core/propagate.cu

@@ -0,0 +1,420 @@
+
+#include "gputiming.h"
+#include "utils.cuh"
+
+namespace gt {
+
+__device__ void propagateSlew(index_type arc_id,
+                              index_type from_pin_id,
+                              index_type to_pin_id,
+                              float *pinSlew,
+                              float *pinLoad,
+                              float *pinImpulse,
+                              float *pinRootDelay,
+                              float *arcDelay,
+                              int arc_type,
+                              int *timing_arc_id_map,
+                              GPULutAllocator *d_allocator) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int i = idx & 0b111;
+    if ((arc_type == 0) && (i < NUM_ATTR)) {
+        float si = pinSlew[from_pin_id * NUM_ATTR + i];
+        if (isnan(si)) return;
+        float imp = pinImpulse[to_pin_id * NUM_ATTR + i];
+        float so = si < 0.0 ? -sqrt(si * si + imp * imp) : sqrt(si * si + imp * imp);
+        pinSlew[to_pin_id * NUM_ATTR + i] = so;
+    } else if (arc_type == 1) {
+        int el = i >> 2;
+        int fel_rf = i >> 1;
+        int tel_rf = ((i & 0b100) >> 1) + (i & 1);
+        int irf = fel_rf & 1;
+        int orf = tel_rf & 1;
+        if ((timing_arc_id_map[arc_id * 2 + el] == -1) || isnan(pinSlew[from_pin_id * NUM_ATTR + fel_rf])) return;
+        float si = pinSlew[from_pin_id * NUM_ATTR + fel_rf];
+        float lc = pinLoad[to_pin_id * NUM_ATTR + tel_rf];
+        int timing_id = timing_arc_id_map[arc_id * 2 + el];
+        float so = d_allocator->query(timing_id, irf, orf, si, lc, 1);
+        if (isnan(so)) return;
+        if (isnan(pinSlew[to_pin_id * NUM_ATTR + tel_rf]) || ((pinSlew[to_pin_id * NUM_ATTR + tel_rf] > so) ^ el)) {
+            atomicExch(&pinSlew[to_pin_id * NUM_ATTR + tel_rf], so);
+        }
+    }
+}
+
+__device__ void propagateDelay(index_type arc_id,
+                               index_type from_pin_id,
+                               index_type to_pin_id,
+                               float *pinSlew,
+                               float *pinLoad,
+                               float *pinImpulse,
+                               float *pinRootDelay,
+                               float *arcDelay,
+                               int arc_type,
+                               int *timing_arc_id_map,
+                               GPULutAllocator *d_allocator) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int i = idx & 0b111;
+    if ((arc_type == 0) && (i < NUM_ATTR)) {
+        float delay = pinRootDelay[to_pin_id * NUM_ATTR + i];
+        int el_rf_rf = (i << 1) + (i & 1);
+        arcDelay[arc_id * 2 * NUM_ATTR + el_rf_rf] = delay;
+    } else if (arc_type == 1) {
+        int el = i >> 2;
+        int fel_rf = i >> 1;
+        int tel_rf = ((i & 0b100) >> 1) + (i & 1);
+        int irf = fel_rf & 1;
+        int orf = tel_rf & 1;
+        if ((timing_arc_id_map[arc_id * 2 + el] == -1) || isnan(pinSlew[from_pin_id * NUM_ATTR + fel_rf])) return;
+        float si = pinSlew[from_pin_id * NUM_ATTR + fel_rf];
+        float lc = pinLoad[to_pin_id * NUM_ATTR + tel_rf];
+        int timing_id = timing_arc_id_map[arc_id * 2 + el];
+        float delay = d_allocator->query(timing_id, irf, orf, si, lc, 0);
+        if (isnan(delay)) return;
+        arcDelay[arc_id * 2 * NUM_ATTR + i] = delay;
+    }
+}
+
+__device__ void propagateAT(index_type arc_id,
+                            index_type from_pin_id,
+                            index_type to_pin_id,
+                            float *pinAt,
+                            float *arcDelay,
+                            index_type *at_prefix_pin,
+                            index_type *at_prefix_arc,
+                            index_type *at_prefix_attr) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int i = idx & 0b111;
+    int el = i >> 2;
+    int fel_rf = i >> 1;
+    int tel_rf = ((i & 0b100) >> 1) + (i & 1);
+    int irf = fel_rf & 1;
+    int orf = tel_rf & 1;
+    if (isnan(pinAt[from_pin_id * NUM_ATTR + fel_rf]) || isnan(arcDelay[arc_id * 2 * NUM_ATTR + i])) return;
+    float delay = arcDelay[arc_id * 2 * NUM_ATTR + i];
+    float at = pinAt[from_pin_id * NUM_ATTR + fel_rf] + delay;
+
+    // FIXME: conflict
+    if (isnan(pinAt[to_pin_id * NUM_ATTR + tel_rf]) || ((pinAt[to_pin_id * NUM_ATTR + tel_rf] > at) ^ el)) {
+        atomicExch(&pinAt[to_pin_id * NUM_ATTR + tel_rf], at);
+        at_prefix_pin[to_pin_id * NUM_ATTR + tel_rf] = from_pin_id;
+        at_prefix_arc[to_pin_id * NUM_ATTR + tel_rf] = arc_id;
+        at_prefix_attr[to_pin_id * NUM_ATTR + tel_rf] = fel_rf;
+    }
+}
+
+__device__ void propagateTest(index_type arc_id,
+                              index_type test_id,
+                              index_type from_pin_id,
+                              index_type to_pin_id,
+                              int *timing_arc_id_map,
+                              float *pinSlew,
+                              float *pinAt,
+                              float *pinRat,
+                              float *testRelatedAT,
+                              float *testRAT,
+                              float *testConstraint,
+                              float clock_period,
+                              GPULutAllocator *d_allocator) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int i = idx & 0b111;
+    if (i < NUM_ATTR) {
+        const int el = i >> 1;
+        const int rf = i & 1;
+        const int el_rf_rf = (i << 1) + (i & 1);
+        if ((timing_arc_id_map[arc_id * 2 + el] == -1) || (isnan(pinSlew[to_pin_id * NUM_ATTR + i]))) return;
+        int fel = el ^ 1;
+        int timing_id = timing_arc_id_map[arc_id * 2 + el];
+        int frf = d_allocator->d_is_rising_edge_triggered[timing_id] ? 0 : 1;
+        if (frf && !d_allocator->d_is_falling_edge_triggered[timing_id]) {
+            return;
+        }
+        const int fel_rf = (fel << 1) + frf;
+        if (isnan(pinAt[from_pin_id * NUM_ATTR + fel_rf]) || isnan(pinSlew[from_pin_id * NUM_ATTR + fel_rf])) return;
+
+        if (el == 0) {
+            testRelatedAT[test_id * NUM_ATTR + i] = pinAt[from_pin_id * NUM_ATTR + fel_rf];
+        } else {
+            testRelatedAT[test_id * NUM_ATTR + i] = pinAt[from_pin_id * NUM_ATTR + fel_rf] + clock_period;
+        }
+
+        float sr = pinSlew[from_pin_id * NUM_ATTR + fel_rf];
+        float sc = pinSlew[to_pin_id * NUM_ATTR + i];
+        testConstraint[test_id * NUM_ATTR + i] = d_allocator->query(timing_id, frf, rf, sr, sc, 2);
+
+        if (!isnan(testConstraint[test_id * NUM_ATTR + i]) && !isnan(testRelatedAT[test_id * NUM_ATTR + i])) {
+            if (el == 0) {
+                pinRat[to_pin_id * NUM_ATTR + i] = testRelatedAT[test_id * NUM_ATTR + i] + testConstraint[test_id * NUM_ATTR + i];
+            } else {
+                pinRat[to_pin_id * NUM_ATTR + i] = testRelatedAT[test_id * NUM_ATTR + i] - testConstraint[test_id * NUM_ATTR + i];
+            }
+            testRAT[test_id * NUM_ATTR + i] = pinRat[to_pin_id * NUM_ATTR + i];
+        }
+    }
+}
+
+__global__ void propagatePin(index_type *level_list,
+                             index_type *pin_backward_arc_list_end,
+                             index_type *pin_backward_arc_list,
+                             index_type *timing_arc_from_pin_id,
+                             int *arc_types,
+                             int *arc_id2test_id,
+                             float *pinSlew,
+                             float *pinLoad,
+                             float *pinImpulse,
+                             float *pinRootDelay,
+                             float *pinAt,
+                             float *pinRat,
+                             float *testRelatedAT,
+                             float *testRAT,
+                             float *testConstraint,
+                             float *arcDelay,
+                             int *timing_arc_id_map,
+                             index_type *at_prefix_pin,
+                             index_type *at_prefix_arc,
+                             index_type *at_prefix_attr,
+                             index_type level_start_offset,
+                             int num_pins_level,
+                             float clock_period,
+                             GPULutAllocator *d_allocator) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int pin_idx = idx >> 3;
+    if (pin_idx < num_pins_level) {
+        index_type to_pin_id = level_list[level_start_offset + pin_idx];
+        for (index_type i = pin_backward_arc_list_end[to_pin_id]; i < pin_backward_arc_list_end[to_pin_id + 1]; i++) {
+            index_type arc_id = pin_backward_arc_list[i];
+            index_type from_pin_id = timing_arc_from_pin_id[arc_id];
+            int arc_type = arc_types[arc_id];
+            propagateSlew(arc_id, from_pin_id, to_pin_id, pinSlew, pinLoad, pinImpulse, pinRootDelay, arcDelay, arc_type, timing_arc_id_map, d_allocator);
+            propagateDelay(arc_id, from_pin_id, to_pin_id, pinSlew, pinLoad, pinImpulse, pinRootDelay, arcDelay, arc_type, timing_arc_id_map, d_allocator);
+            propagateAT(arc_id, from_pin_id, to_pin_id, pinAt, arcDelay, at_prefix_pin, at_prefix_arc, at_prefix_attr);
+            int test_id = arc_id2test_id[arc_id];
+            if (clock_period > 0 && test_id != -1) {
+                propagateTest(arc_id, test_id, from_pin_id, to_pin_id, timing_arc_id_map, pinSlew, pinAt, pinRat, testRelatedAT, testRAT, testConstraint, clock_period, d_allocator);
+            }
+        }
+    }
+}
+
+__device__ void propagateRAT(index_type arc_id,
+                             int arc_type,
+                             index_type from_pin_id,
+                             index_type to_pin_id,
+                             float *pinAt,
+                             float *pinRat,
+                             float *arcDelay,
+                             int *timing_arc_id_map,
+                             float *from_rats,
+                             GPULutAllocator *d_allocator) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int i = idx & 0b111;
+    if ((arc_type == 0) && (i < NUM_ATTR)) {
+        const int el_rf_rf = (i << 1) + (i & 1);
+        const int el = i >> 1;
+        if (isnan(pinRat[to_pin_id * NUM_ATTR + i]) || isnan(arcDelay[arc_id * 2 * NUM_ATTR + el_rf_rf])) return;
+        float delay = arcDelay[arc_id * 2 * NUM_ATTR + el_rf_rf];
+        float rat = pinRat[to_pin_id * NUM_ATTR + i] - delay;
+        if (isnan(pinRat[from_pin_id * NUM_ATTR + i]) || ((pinRat[from_pin_id * NUM_ATTR + i] < rat) ^ el)) {
+            atomicExch(&pinRat[from_pin_id * NUM_ATTR + i], rat);
+        }
+    } else if (arc_type == 1) {
+        int el = i >> 2;
+        int fel_rf = i >> 1;
+        int tel_rf = ((i & 0b100) >> 1) + (i & 1);
+        int irf = fel_rf & 1;
+        int orf = tel_rf & 1;
+        if (timing_arc_id_map[arc_id * 2 + el] == -1) return;
+        int timing_id = timing_arc_id_map[arc_id * 2 + el];
+        if (!d_allocator->d_is_constraint[timing_id]) {
+            if (isnan(pinRat[to_pin_id * NUM_ATTR + tel_rf]) || isnan(arcDelay[arc_id * 2 * NUM_ATTR + i])) return;
+            float delay = arcDelay[arc_id * 2 * NUM_ATTR + i];
+            float rat = pinRat[to_pin_id * NUM_ATTR + tel_rf] - delay;
+            from_rats[threadIdx.x] = rat;
+        } else {
+            if (!d_allocator->is_transition_defined(timing_id, irf, orf)) return;
+            if (el == 0) {
+                const int fel_rf = 2 + irf;
+                const int tel_rf = orf;
+                float at = pinAt[from_pin_id * NUM_ATTR + fel_rf];
+                if (isnan(pinRat[to_pin_id * NUM_ATTR + tel_rf]) || isnan(pinAt[to_pin_id * NUM_ATTR + tel_rf]) || isnan(at)) return;
+                float slack = (pinRat[to_pin_id * NUM_ATTR + tel_rf] - pinAt[to_pin_id * NUM_ATTR + tel_rf]) * -1;
+                float rat = at + slack;
+                from_rats[threadIdx.x] = rat;
+            } else {
+                const int fel_rf = irf;
+                const int tel_rf = 2 + orf;
+                float at = pinAt[from_pin_id * NUM_ATTR + fel_rf];
+                if (isnan(pinRat[to_pin_id * NUM_ATTR + tel_rf]) || isnan(pinAt[to_pin_id * NUM_ATTR + tel_rf]) || isnan(at)) return;
+                float slack = (pinRat[to_pin_id * NUM_ATTR + tel_rf] - pinAt[to_pin_id * NUM_ATTR + tel_rf]);
+                float rat = at - slack;
+                from_rats[threadIdx.x] = rat;
+            }
+        }
+    }
+}
+
+__global__ void propagatePinBack(index_type *level_list,
+                                 index_type *pin_forward_arc_list_end,
+                                 index_type *pin_forward_arc_list,
+                                 index_type *timing_arc_to_pin_id,
+                                 int *arc_types,
+                                 int *arc_id2test_id,
+                                 float *pinSlew,
+                                 float *pinLoad,
+                                 float *pinImpulse,
+                                 float *pinRootDelay,
+                                 float *pinAt,
+                                 float *pinRat,
+                                 float *testRelatedAT,
+                                 float *testConstraint,
+                                 float *arcDelay,
+                                 int *timing_arc_id_map,
+                                 index_type level_start_offset,
+                                 int num_pins_level,
+                                 float clock_period,
+                                 GPULutAllocator *d_allocator) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int pin_idx = idx >> 3;
+    extern __shared__ float from_rats[];
+
+    if (pin_idx < num_pins_level) {
+        index_type from_pin_id = level_list[level_start_offset + pin_idx];
+        for (index_type i = pin_forward_arc_list_end[from_pin_id]; i < pin_forward_arc_list_end[from_pin_id + 1]; i++) {
+            index_type arc_id = pin_forward_arc_list[i];
+            index_type to_pin_id = timing_arc_to_pin_id[arc_id];
+            int arc_type = arc_types[arc_id];
+            if ((threadIdx.x % (2 * NUM_ATTR)) == 0) {
+                for (int i = threadIdx.x; i < threadIdx.x + 2 * NUM_ATTR; i++) from_rats[i] = nanf("");
+            }
+            __syncthreads();
+
+            propagateRAT(arc_id, arc_type, from_pin_id, to_pin_id, pinAt, pinRat, arcDelay, timing_arc_id_map, from_rats, d_allocator);
+
+            __syncthreads();
+            if ((threadIdx.x % (2 * NUM_ATTR)) == 0) {
+                for (int ti = threadIdx.x; ti < threadIdx.x + 2 * NUM_ATTR; ti++) {
+                    const int i = ti & 0b111;
+                    if (isnan(from_rats[ti])) continue;
+                    int el = i >> 2;
+                    int fel_rf = i >> 1;
+                    int tel_rf = ((i & 0b100) >> 1) + (i & 1);
+                    int irf = fel_rf & 1;
+                    int orf = tel_rf & 1;
+                    int timing_id = timing_arc_id_map[arc_id * 2 + el];
+                    float rat = from_rats[ti];
+                    if (!d_allocator->d_is_constraint[timing_id]) {
+                        if (isnan(pinRat[from_pin_id * NUM_ATTR + fel_rf]) || ((pinRat[from_pin_id * NUM_ATTR + fel_rf] < rat) ^ el)) {
+                            atomicExch(&pinRat[from_pin_id * NUM_ATTR + fel_rf], rat);
+                        }
+                    } else {
+                        if (el == 0) {
+                            const int fel_rf = 2 + irf;
+                            const int tel_rf = orf;
+                            if (isnan(pinRat[from_pin_id * NUM_ATTR + fel_rf]) || (pinRat[from_pin_id * NUM_ATTR + fel_rf] > rat)) {
+                                atomicExch(&pinRat[from_pin_id * NUM_ATTR + fel_rf], rat);
+                            }
+                        } else {
+                            const int fel_rf = irf;
+                            const int tel_rf = 2 + orf;
+                            if (isnan(pinRat[from_pin_id * NUM_ATTR + fel_rf]) || (pinRat[from_pin_id * NUM_ATTR + fel_rf] < rat)) {
+                                atomicExch(&pinRat[from_pin_id * NUM_ATTR + fel_rf], rat);
+                            }
+                        }
+                    }
+                }
+            }
+        }
+    }
+}
+
+void update_timing_cuda(index_type *level_list,
+                        vector<int> level_list_end_cpu,
+                        index_type *pin_forward_arc_list_end,
+                        index_type *pin_forward_arc_list,
+                        index_type *timing_arc_to_pin_id,
+                        index_type *pin_backward_arc_list_end,
+                        index_type *pin_backward_arc_list,
+                        index_type *timing_arc_from_pin_id,
+                        int *arc_types,
+                        int *arc_id2test_id,
+                        float *pinSlew,
+                        float *pinLoad,
+                        float *pinImpulse,
+                        float *pinRootDelay,
+                        float *pinAt,
+                        float *pinRat,
+                        float *testRelatedAT,
+                        float *testRAT,
+                        float *testConstraint,
+                        float *arcDelay,
+                        int *timing_arc_id_map,
+                        index_type *at_prefix_pin,
+                        index_type *at_prefix_arc,
+                        index_type *at_prefix_attr,
+                        float clock_period,
+                        GPULutAllocator *d_allocator,
+                        int num_pins,
+                        bool deterministic) {
+    for (int i = 1; i < level_list_end_cpu.size() - 1; i++) {
+        int num_pins_level = level_list_end_cpu[i + 1] - level_list_end_cpu[i];
+        index_type level_start_offset = level_list_end_cpu[i];
+        // printf("==== level %d ======= %d \n", i, num_pins_level);
+        propagatePin<<<BLOCK_NUMBER(num_pins_level * 2 * NUM_ATTR), BLOCK_SIZE>>>(level_list,
+                                                                                  pin_backward_arc_list_end,
+                                                                                  pin_backward_arc_list,
+                                                                                  timing_arc_from_pin_id,
+                                                                                  arc_types,
+                                                                                  arc_id2test_id,
+                                                                                  pinSlew,
+                                                                                  pinLoad,
+                                                                                  pinImpulse,
+                                                                                  pinRootDelay,
+                                                                                  pinAt,
+                                                                                  pinRat,
+                                                                                  testRelatedAT,
+                                                                                  testRAT,
+                                                                                  testConstraint,
+                                                                                  arcDelay,
+                                                                                  timing_arc_id_map,
+                                                                                  at_prefix_pin,
+                                                                                  at_prefix_arc,
+                                                                                  at_prefix_attr,
+                                                                                  level_start_offset,
+                                                                                  num_pins_level,
+                                                                                  clock_period,
+                                                                                  d_allocator);
+
+        cudaDeviceSynchronize();
+    }
+    cudaDeviceSynchronize();
+
+    for (int i = level_list_end_cpu.size() - 3; i >= 0; i--) {
+        int num_pins_level = level_list_end_cpu[i + 1] - level_list_end_cpu[i];
+        index_type level_start_offset = level_list_end_cpu[i];
+        // printf("==== level %d ======= %d \n", i, num_pins_level);
+        propagatePinBack<<<BLOCK_NUMBER(num_pins_level * 2 * NUM_ATTR), BLOCK_SIZE, BLOCK_SIZE * sizeof(float)>>>(level_list,
+                                                                                                                  pin_forward_arc_list_end,
+                                                                                                                  pin_forward_arc_list,
+                                                                                                                  timing_arc_to_pin_id,
+                                                                                                                  arc_types,
+                                                                                                                  arc_id2test_id,
+                                                                                                                  pinSlew,
+                                                                                                                  pinLoad,
+                                                                                                                  pinImpulse,
+                                                                                                                  pinRootDelay,
+                                                                                                                  pinAt,
+                                                                                                                  pinRat,
+                                                                                                                  testRelatedAT,
+                                                                                                                  testConstraint,
+                                                                                                                  arcDelay,
+                                                                                                                  timing_arc_id_map,
+                                                                                                                  level_start_offset,
+                                                                                                                  num_pins_level,
+                                                                                                                  clock_period,
+                                                                                                                  d_allocator);
+
+        cudaDeviceSynchronize();
+    }
+    cudaDeviceSynchronize();
+}
+
+}  // namespace gt

cpp_to_py/gputimer/core/rctree.cpp

@@ -0,0 +1,554 @@
+
+#include "GPUTimer.h"
+#include "common/utils/utils.h"
+#include "common/db/Database.h"
+#include "gputimer/db/GTDatabase.h"
+#include <flute.hpp>
+using namespace Flute;
+
+namespace gt {
+
+void update_rc_timing_cuda(float* x,
+                           float* y,
+                           const float* pin_offset_x,
+                           const float* pin_offset_y,
+                           const int* pin2node_map,
+                           const int* flat_net2pin_start_map,
+                           const int* flat_net2pin_map,
+                           float* pinLoad,
+                           float* pinImpulse,
+                           float* pinCap,
+                           float* pinWireCap,
+                           float* pinRootDelay,
+                           float* pinRootRes,
+                           int num_nets,
+                           int num_pins,
+                           float unit_to_micron,
+                           int* net_is_clock,
+                           float cf,
+                           float rf);
+
+void GPUTimer::update_rc_timing(torch::Tensor node_lpos, bool record, bool load, bool conpensation) {
+    timing_raw_db.commit_from(node_lpos.index({"...", 0}).contiguous(), node_lpos.index({"...", 1}).contiguous());
+    float unit_to_micron = scale_factor * microns;
+    float rf = wire_resistance_per_micron / res_unit;
+    float cf = wire_capacitance_per_micron / cap_unit;
+    update_rc_timing_cuda(x,
+                          y,
+                          pin_offset_x,
+                          pin_offset_y,
+                          pin2node_map,
+                          flat_net2pin_start_map,
+                          flat_net2pin_map,
+                          pinLoad,
+                          pinImpulse,
+                          pinCap,
+                          pinWireCap,
+                          pinRootDelay,
+                          pinRootRes,
+                          num_nets,
+                          num_pins,
+                          unit_to_micron,
+                          net_is_clock,
+                          cf,
+                          rf);
+    if (record) {
+        auto ratio_load = torch::nan_to_num(timing_raw_db.pinLoad_ref / timing_raw_db.pinLoad, 1.0);
+        timing_raw_db.pinLoad_ratio.data().copy_(ratio_load.contiguous().data());
+        auto ratio_delay = torch::sqrt(torch::nan_to_num(timing_raw_db.pinRootDelay_ref / timing_raw_db.pinRootDelay, 1.0));
+        timing_raw_db.pinRootDelay_ratio.data().copy_(ratio_delay.contiguous().data());
+
+        auto delay_comp = (torch::nan_to_num(timing_raw_db.pinRootDelay_ref - timing_raw_db.pinRootDelay, 0)).clamp(0.0);
+        timing_raw_db.pinRootDelay_compensation.data().copy_(delay_comp.contiguous().data());
+    }
+    if (load) {
+        timing_raw_db.pinImpulse.data().copy_(timing_raw_db.pinImpulse_ref.data());
+        timing_raw_db.pinLoad *= timing_raw_db.pinLoad_ratio;
+        if (conpensation)
+            timing_raw_db.pinRootDelay += timing_raw_db.pinRootDelay_compensation;
+        else
+            timing_raw_db.pinRootDelay *= timing_raw_db.pinRootDelay_ratio;
+    }
+}
+
+// ------------------------------------------------------------------------------------------------------------------------
+//
+
+auto& retrieve_pins_from_pos(std::map<utils::PointT<int>, std::set<int>>& pos2pins_map, const utils::PointT<int>& point, int& index) {
+    if (pos2pins_map.find(point) != pos2pins_map.end()) return pos2pins_map[point];
+    pos2pins_map.emplace(point, std::set<int>{index++});
+    return pos2pins_map[point];
+}
+
+tuple<vector<int>, vector<int>, vector<float>, vector<int>, vector<int>, vector<int>, int, int> FluteRCTree(TimingTorchRawDB& timing_raw_db,
+                                                                                                            float rf,
+                                                                                                            float cf) {
+    torch::Tensor flat_node2pin_start_map_at = timing_raw_db.flat_node2pin_start_map.clone().cpu().contiguous();
+    torch::Tensor flat_node2pin_map_at = timing_raw_db.flat_node2pin_map.clone().cpu().contiguous();
+    torch::Tensor pin2node_map_at = timing_raw_db.pin2node_map.clone().cpu().contiguous();
+    torch::Tensor flat_net2pin_start_map_at = timing_raw_db.flat_net2pin_start_map.clone().cpu().contiguous();
+    torch::Tensor flat_net2pin_map_at = timing_raw_db.flat_net2pin_map.clone().cpu().contiguous();
+    torch::Tensor pin2net_map_at = timing_raw_db.pin2net_map.clone().cpu().contiguous();
+    torch::Tensor x_at = timing_raw_db.x.clone().cpu().contiguous();
+    torch::Tensor y_at = timing_raw_db.y.clone().cpu().contiguous();
+    torch::Tensor pin_offset_x_at = timing_raw_db.pin_offset_x.clone().cpu().contiguous();
+    torch::Tensor pin_offset_y_at = timing_raw_db.pin_offset_y.clone().cpu().contiguous();
+
+    const int* flat_node2pin_start_map = flat_node2pin_start_map_at.data_ptr<int>();
+    const int* flat_node2pin_map = flat_node2pin_map_at.data_ptr<int>();
+    const int* pin2node_map = pin2node_map_at.data_ptr<int>();
+    const int* flat_net2pin_start_map = flat_net2pin_start_map_at.data_ptr<int>();
+    const int* flat_net2pin_map = flat_net2pin_map_at.data_ptr<int>();
+    const int* pin2net_map = pin2net_map_at.data_ptr<int>();
+    const float* x = x_at.data_ptr<float>();
+    const float* y = y_at.data_ptr<float>();
+    const float* pin_offset_x = pin_offset_x_at.data_ptr<float>();
+    const float* pin_offset_y = pin_offset_y_at.data_ptr<float>();
+    int& num_nets = timing_raw_db.num_nets;
+
+    constexpr const int scale = 1000;  // flute only supports integers.
+    using Point2i = utils::PointT<int>;
+
+    vector<int> edge_from;
+    vector<int> edge_to;
+    vector<float> edge_wl;
+    vector<int> flat_net2node_start_map;
+    vector<int> flat_net2edge_start_map;
+    vector<int> node2pin_map;
+    int node_count = 0;
+    int edge_count = 0;
+    flat_net2node_start_map.push_back(0);
+    flat_net2edge_start_map.push_back(0);
+
+    vector<vector<int>> net_id2edge_from(num_nets);
+    vector<vector<int>> net_id2edge_to(num_nets);
+    vector<vector<float>> net_id2edge_wl(num_nets);
+    vector<vector<int>> net_id2node2pin_map(num_nets);
+
+    omp_lock_t lock;
+    omp_init_lock(&lock);
+#pragma omp parallel for
+    for (int i = 0; i < num_nets; ++i) {
+        const int degree = flat_net2pin_start_map[i + 1] - flat_net2pin_start_map[i];
+        const int root = flat_net2pin_map[flat_net2pin_start_map[i]];
+        std::map<Point2i, std::set<int>> pos2pins_map;
+        std::vector<int> vx, vy;
+        vx.reserve(degree);
+        vy.reserve(degree);
+
+        std::map<int, int> global2inner_map;
+
+        for (int j = 0; j < degree; ++j) {
+            int pin = flat_net2pin_map[j + flat_net2pin_start_map[i]];
+            int node = pin2node_map[pin];
+            float offset_x = pin_offset_x[pin], offset_y = pin_offset_y[pin];
+            // Find the correct pin locations given cell locations.
+            auto x_ = static_cast<int>((x[node] + offset_x) * scale);
+            auto y_ = static_cast<int>((y[node] + offset_y) * scale);
+            global2inner_map[pin] = j;
+
+            if (pos2pins_map.find(Point2i(x_, y_)) != pos2pins_map.end())
+                pos2pins_map[Point2i(x_, y_)].insert(j);
+            else {
+                pos2pins_map.emplace(Point2i(x_, y_), std::set<int>{j});
+                vx.emplace_back(x_);
+                vy.emplace_back(y_);
+            }
+        }
+        const int valid_size = static_cast<int>(vx.size());
+        int num_pins = degree;
+        std::set<Point2i> multipin_pos;
+        std::map<Point2i, Point2i> pos2neighbor_map;
+
+        if (valid_size > 1) {
+            Tree flutetree = flute(valid_size, vx.data(), vy.data(), 8);
+
+            for (int bid = 0; bid < 2 * valid_size - 2; ++bid) {
+                Branch& branch1 = flutetree.branch[bid];
+                Branch& branch2 = flutetree.branch[branch1.n];
+
+                Point2i p1(branch1.x, branch1.y), p2(branch2.x, branch2.y);
+
+                if (p1 == p2) continue;
+
+                pos2neighbor_map.emplace(p2, p1);
+                auto& id1 = retrieve_pins_from_pos(pos2pins_map, p1, num_pins);
+                auto& id2 = retrieve_pins_from_pos(pos2pins_map, p2, num_pins);
+
+                auto distance = Dist(p1, p2);
+                float wl = static_cast<float>(distance * 1.0) / scale;
+
+                if (!id1.empty() && !id2.empty()) {
+                    auto base1 = id1.begin(), base2 = id2.begin();
+                    if (*base1 != *base2) {
+                        net_id2edge_from[i].emplace_back(*base1);
+                        net_id2edge_to[i].emplace_back(*base2);
+                        net_id2edge_wl[i].emplace_back(wl);
+                    }
+                    if (id1.size() > 1) multipin_pos.insert(p1);
+                    if (id2.size() > 1) multipin_pos.insert(p2);
+                }
+            }
+            free(flutetree.branch);
+        } else if (valid_size == 1 && degree > 1) {
+            multipin_pos.emplace(vx[0], vy[0]);
+        }
+        for (const auto& pos : multipin_pos) {
+            const auto& pins = pos2pins_map[pos];
+            int adj_pin = global2inner_map[root];
+            const auto& _ppos = pos2neighbor_map[pos];
+            if (auto itr = pos2pins_map.find(_ppos); itr != pos2pins_map.end()) {
+                adj_pin = *itr->second.cbegin();
+            }
+            auto distance = Dist(pos, _ppos);
+            float wl = static_cast<float>(distance * 1.0) / scale;
+            for (auto it = std::next(pins.cbegin()); it != pins.cend(); ++it) {
+                net_id2edge_from[i].emplace_back(adj_pin);
+                net_id2edge_to[i].emplace_back(*it);
+                net_id2edge_wl[i].emplace_back(0);
+            }
+        }
+
+        for (int j = 0; j < num_pins; ++j) {
+            if (j < degree)
+                net_id2node2pin_map[i].push_back(flat_net2pin_map[j + flat_net2pin_start_map[i]]);
+            else
+                net_id2node2pin_map[i].push_back(-1);
+        }
+    }
+    omp_destroy_lock(&lock);
+
+    for (int i = 0; i < num_nets; ++i) {
+        for (int j = 0; j < net_id2edge_from[i].size(); ++j) {
+            edge_from.push_back(node_count + net_id2edge_from[i][j]);
+            edge_to.push_back(node_count + net_id2edge_to[i][j]);
+            edge_wl.push_back(net_id2edge_wl[i][j]);
+            edge_count++;
+        }
+        node_count += net_id2node2pin_map[i].size();
+        for (int j = 0; j < net_id2node2pin_map[i].size(); ++j) {
+            node2pin_map.push_back(net_id2node2pin_map[i][j]);
+        }
+        flat_net2node_start_map.push_back(node_count);
+        flat_net2edge_start_map.push_back(edge_count);
+    }
+
+    return {edge_from, edge_to, edge_wl, flat_net2node_start_map, flat_net2edge_start_map, node2pin_map, node_count, edge_count};
+}
+
+void flatten_rc_tree(std::vector<int> host_edge_from,
+                     std::vector<int> host_edge_to,
+                     float* edge_res,
+                     float* node_cap,
+                     std::vector<int> host_flat_net2node_start_map,
+                     std::vector<int> host_flat_net2edge_start_map,
+                     std::vector<int> host_node2pin_map,
+                     int* node_order,
+                     int* edge_order,
+                     int* parent_node,
+                     float* res_parent,
+                     float* pinLoad,
+                     float* pinImpulse,
+                     float* pinCap,
+                     float* pinWireCap,
+                     float* pinRootDelay,
+                     float* pinRootRes,
+                     int num_nets,
+                     int num_pins,
+                     int num_nodes,
+                     int num_edges);
+
+void propagate_rc_tree(std::vector<int> host_edge_from,
+                       std::vector<int> host_edge_to,
+                       float* edge_res,
+                       float* node_cap,
+                       std::vector<int> host_flat_net2node_start_map,
+                       std::vector<int> host_flat_net2edge_start_map,
+                       std::vector<int> host_node2pin_map,
+                       int* node_order,
+                       int* parent_node,
+                       float* res_parent,
+                       float* pinLoad,
+                       float* pinImpulse,
+                       float* pinCap,
+                       float* pinWireCap,
+                       float* pinRootDelay,
+                       float* pinRootRes,
+                       int num_nets,
+                       int num_pins,
+                       int num_nodes,
+                       int num_edges);
+
+
+void calc_res_cap(std::vector<int> host_edge_from,
+                  std::vector<int> host_edge_to,
+                  int* edge_order,
+                  float* edge_res,
+                  float* node_cap,
+                  std::vector<int> host_flat_net2node_start_map,
+                  std::vector<int> host_flat_net2edge_start_map,
+                  std::vector<int> host_node2pin_map,
+                  std::vector<float> host_edge_wl,
+                  int num_nets,
+                  int num_edges,
+                  int num_nodes,
+                  int* net_is_clock,
+                  float unit_to_micron,
+                  float rf,
+                  float cf);
+
+void GPUTimer::update_rc_timing_flute(torch::Tensor node_lpos, bool record) {
+    timing_raw_db.commit_from(node_lpos.index({"...", 0}).contiguous(), node_lpos.index({"...", 1}).contiguous());
+
+    float unit_to_micron = scale_factor * microns;
+    float rf = wire_resistance_per_micron / res_unit;
+    float cf = wire_capacitance_per_micron / cap_unit;
+
+    auto [edge_from, edge_to, edge_wl, flat_net2node_start_map, flat_net2edge_start_map, node2pin_map, num_nodes, num_edges] =
+        FluteRCTree(timing_raw_db, rf, cf);
+    auto device = timing_raw_db.node_size.device();
+    torch::Tensor node_order = torch::zeros({num_nodes}, torch::dtype(torch::kInt32).device(device)).contiguous();
+    torch::Tensor edge_order = torch::zeros({num_edges}, torch::dtype(torch::kInt32).device(device)).contiguous();
+    torch::Tensor parent_node = -torch::ones({num_nodes}, torch::dtype(torch::kInt32).device(device)).contiguous();
+    torch::Tensor res_parent = torch::zeros({num_nodes * NUM_ATTR}, torch::dtype(torch::kFloat32).device(device)).contiguous();
+    torch::Tensor node_cap = torch::zeros({num_nodes * NUM_ATTR}, torch::dtype(torch::kFloat32).device(device)).contiguous();
+    torch::Tensor edge_res = torch::zeros({num_edges}, torch::dtype(torch::kFloat32).device(device)).contiguous();
+
+    calc_res_cap(edge_from,
+                 edge_to,
+                 edge_order.data_ptr<int>(),
+                 edge_res.data_ptr<float>(),
+                 node_cap.data_ptr<float>(),
+                 flat_net2node_start_map,
+                 flat_net2edge_start_map,
+                 node2pin_map,
+                 edge_wl,
+                 num_nets,
+                 num_edges,
+                 num_nodes,
+                 net_is_clock,
+                 unit_to_micron,
+                 rf,
+                 cf);
+
+    flatten_rc_tree(edge_from,
+                    edge_to,
+                    edge_res.data_ptr<float>(),
+                    node_cap.data_ptr<float>(),
+                    flat_net2node_start_map,
+                    flat_net2edge_start_map,
+                    node2pin_map,
+                    node_order.data_ptr<int>(),
+                    edge_order.data_ptr<int>(),
+                    parent_node.data_ptr<int>(),
+                    res_parent.data_ptr<float>(),
+                    pinLoad,
+                    pinImpulse,
+                    pinCap,
+                    pinWireCap,
+                    pinRootDelay,
+                    pinRootRes,
+                    num_nets,
+                    num_pins,
+                    num_nodes,
+                    num_edges);
+
+    propagate_rc_tree(edge_from,
+                      edge_to,
+                      edge_res.data_ptr<float>(),
+                      node_cap.data_ptr<float>(),
+                      flat_net2node_start_map,
+                      flat_net2edge_start_map,
+                      node2pin_map,
+                      node_order.data_ptr<int>(),
+                      parent_node.data_ptr<int>(),
+                      res_parent.data_ptr<float>(),
+                      pinLoad,
+                      pinImpulse,
+                      pinCap,
+                      pinWireCap,
+                      pinRootDelay,
+                      pinRootRes,
+                      num_nets,
+                      num_pins,
+                      num_nodes,
+                      num_edges);
+
+    if (record) {
+        timing_raw_db.pinImpulse_ref.data().copy_(timing_raw_db.pinImpulse.data());
+        timing_raw_db.pinLoad_ref.data().copy_(timing_raw_db.pinLoad.data());
+        timing_raw_db.pinRootDelay_ref.data().copy_(timing_raw_db.pinRootDelay.data());
+    }
+}
+
+void GPUTimer::update_rc_timing_spef() {
+
+    torch::Tensor flat_net2pin_start_map_at = timing_raw_db.flat_net2pin_start_map.clone().cpu().contiguous();
+    torch::Tensor flat_net2pin_map_at = timing_raw_db.flat_net2pin_map.clone().cpu().contiguous();
+    const int* flat_net2pin_start_map = flat_net2pin_start_map_at.data_ptr<int>();
+    const int* flat_net2pin_map = flat_net2pin_map_at.data_ptr<int>();
+
+    vector<vector<int>> net_id2edge_from(num_nets);
+    vector<vector<int>> net_id2edge_to(num_nets);
+    vector<vector<float>> net_id2node_cap(num_nets);
+    vector<vector<float>> net_id2edge_res(num_nets);
+    vector<vector<int>> net_id2node2pin_map(num_nets);
+    vector<std::unordered_map<std::string, int>> node_name2node_id_map(num_nets);
+
+    printf("num_nets: %d\n", num_nets);
+    printf("num_nets in spef file: %d\n", spef.nets.size());
+    float spef_res_ratio = *gtdb.spef_res_unit / gtdb.res_unit;
+    float spef_cap_ratio = *gtdb.spef_cap_unit / gtdb.cap_unit;
+    float spef_time_ratio = *gtdb.spef_time_unit / gtdb.time_unit;
+    logger.info("spef lib ratios: res %.5E cap %.5E time %.5E", spef_res_ratio, spef_cap_ratio, spef_time_ratio);
+
+    auto add_node_cap = [&](const std::string& node_name, float cap, int net_idx) {
+        if (auto itr = std::find(gtdb.pin_names.begin(), gtdb.pin_names.end(), node_name); itr != gtdb.pin_names.end()) {
+            int in_net_idx = flat_net2pin_map[std::distance(gtdb.pin_names.begin(), itr)] - flat_net2pin_start_map[net_idx];
+            net_id2node_cap[net_idx][in_net_idx] = cap * spef_cap_ratio;
+            node_name2node_id_map[net_idx][node_name] = in_net_idx;
+        } else {
+            net_id2node_cap[net_idx].push_back(cap * spef_cap_ratio);
+            node_name2node_id_map[net_idx][node_name] = net_id2node_cap[net_idx].size() - 1;
+            net_id2node2pin_map[net_idx].push_back(-1);
+        }
+    };
+    for (const auto& n : spef.nets) {
+        string net_name = n.name;
+        net_name = validate_token(net_name);
+        if (auto itr = std::find(gtdb.net_names.begin(), gtdb.net_names.end(), net_name); itr == gtdb.net_names.end()) {
+            continue;
+        } else {
+            int net_idx = std::distance(gtdb.net_names.begin(), itr);
+
+            // Put pin nodes in the front
+            for (int j = 0; j < flat_net2pin_start_map[net_idx + 1] - flat_net2pin_start_map[net_idx]; ++j) {
+                net_id2node2pin_map[net_idx].push_back(flat_net2pin_map[j + flat_net2pin_start_map[net_idx]]); // Pins in the front
+            }
+            net_id2node_cap[net_idx].resize(net_id2node2pin_map[net_idx].size(), 0);
+
+            // Add ground-node capacitance
+            for (const auto& [node1, node2, cap] : n.caps) {
+                if (node2.empty()) {
+                    add_node_cap(node1, cap, net_idx);
+                }
+            }
+
+            // Add node-node resistance
+            for (const auto& [node1, node2, value] : n.ress) {
+                if (node_name2node_id_map[net_idx].find(node1) == node_name2node_id_map[net_idx].end()) {
+                    add_node_cap(node1, 0, net_idx);
+                }
+                if (node_name2node_id_map[net_idx].find(node2) == node_name2node_id_map[net_idx].end()) {
+                    add_node_cap(node2, 0, net_idx);
+                }
+                int from = node_name2node_id_map[net_idx][node1];
+                int to = node_name2node_id_map[net_idx][node2];
+                net_id2edge_from[net_idx].push_back(from);
+                net_id2edge_to[net_idx].push_back(to);
+                net_id2edge_res[net_idx].push_back(value * spef_res_ratio);
+            }
+        }
+    }
+    for (int i = 0; i < num_nets; ++i) {
+        if (net_id2node2pin_map[i].empty()) {
+            logger.warning("net %s has no spef rc, assign 0", gtdb.net_names[i].c_str());
+            for (int j = 0; j < flat_net2pin_start_map[i + 1] - flat_net2pin_start_map[i]; ++j) {
+                net_id2node2pin_map[i].push_back(flat_net2pin_map[j + flat_net2pin_start_map[i]]); // Pins in the front
+                net_id2node_cap[i].push_back(0);
+                if (j != 0) {
+                    net_id2edge_from[i].push_back(0);
+                    net_id2edge_to[i].push_back(j);
+                    net_id2edge_res[i].push_back(0);
+                }
+            }
+        }
+        
+    }
+    
+
+    // data vectors
+    vector<int> edge_from;
+    vector<int> edge_to;
+    vector<float> node_cap_vec;
+    vector<float> edge_res_vec;
+    vector<int> flat_net2node_start_map;
+    vector<int> flat_net2edge_start_map;
+    vector<int> node2pin_map;
+    int num_nodes = 0;
+    int num_edges = 0;
+    flat_net2node_start_map.push_back(0);
+    flat_net2edge_start_map.push_back(0);
+    for (int i = 0; i < num_nets; ++i) {
+        for (int j = 0; j < net_id2edge_from[i].size(); ++j) {
+            edge_from.push_back(num_nodes + net_id2edge_from[i][j]);
+            edge_to.push_back(num_nodes + net_id2edge_to[i][j]);
+            float res = gtdb.net_is_clock[i] == 1 ? 0 : net_id2edge_res[i][j];
+            edge_res_vec.push_back(res);
+            num_edges++;
+        }
+        num_nodes += net_id2node2pin_map[i].size();
+        for (int j = 0; j < net_id2node_cap[i].size(); ++j) {
+            node2pin_map.push_back(net_id2node2pin_map[i][j]);
+            float cap = gtdb.net_is_clock[i] == 1 ? 0 : net_id2node_cap[i][j];
+            for (int k = 0; k < NUM_ATTR; k++) {
+                node_cap_vec.push_back(cap);
+            }
+        }
+        flat_net2node_start_map.push_back(num_nodes);
+        flat_net2edge_start_map.push_back(num_edges);
+    }
+
+    auto device = timing_raw_db.node_size.device();
+    torch::Tensor edge_res = torch::from_blob(edge_res_vec.data(), {num_edges}, torch::dtype(torch::kFloat32)).contiguous().to(device);
+    torch::Tensor node_cap = torch::from_blob(node_cap_vec.data(), {num_nodes * NUM_ATTR}, torch::dtype(torch::kFloat32)).contiguous().to(device);
+    torch::Tensor node_order = torch::zeros({num_nodes}, torch::kInt32).contiguous().to(device);
+    torch::Tensor edge_order = torch::zeros({num_edges}, torch::kInt32).contiguous().to(device);
+    torch::Tensor parent_node = -torch::ones({num_nodes}, torch::dtype(torch::kInt32).device(device));
+    torch::Tensor res_parent = torch::zeros({num_nodes * NUM_ATTR}, torch::dtype(torch::kFloat32).device(device));
+
+    flatten_rc_tree(edge_from,
+                    edge_to,
+                    edge_res.data_ptr<float>(),
+                    node_cap.data_ptr<float>(),
+                    flat_net2node_start_map,
+                    flat_net2edge_start_map,
+                    node2pin_map,
+                    node_order.data_ptr<int>(),
+                    edge_order.data_ptr<int>(),
+                    parent_node.data_ptr<int>(),
+                    res_parent.data_ptr<float>(),
+                    pinLoad,
+                    pinImpulse,
+                    pinCap,
+                    pinWireCap,
+                    pinRootDelay,
+                    pinRootRes,
+                    num_nets,
+                    num_pins,
+                    num_nodes,
+                    num_edges);
+
+    propagate_rc_tree(edge_from,
+                      edge_to,
+                      edge_res.data_ptr<float>(),
+                      node_cap.data_ptr<float>(),
+                      flat_net2node_start_map,
+                      flat_net2edge_start_map,
+                      node2pin_map,
+                      node_order.data_ptr<int>(),
+                      parent_node.data_ptr<int>(),
+                      res_parent.data_ptr<float>(),
+                      pinLoad,
+                      pinImpulse,
+                      pinCap,
+                      pinWireCap,
+                      pinRootDelay,
+                      pinRootRes,
+                      num_nets,
+                      num_pins,
+                      num_nodes,
+                      num_edges);
+}
+
+}  // namespace gt

cpp_to_py/gputimer/core/rctree.cu

@@ -0,0 +1,621 @@
+#include <ATen/cuda/CUDAContext.h>
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <torch/extension.h>
+
+#include "utils.cuh"
+
+namespace gt {
+
+__global__ void RCTreeNet(float *x,
+                          float *y,
+                          const float *pin_offset_x,
+                          const float *pin_offset_y,
+                          const int *pin2node_map,
+                          const int *flat_net2pin_start_map,
+                          const int *flat_net2pin_map,
+                          float *pinLoad,
+                          float *pinImpulse,
+                          float *pinCap,
+                          float *pinWireCap,
+                          float *pinRootDelay,
+                          float *pinRootRes,
+                          int num_nets,
+                          float unit_to_micron,
+                          int *net_is_clock,
+                          float cf,
+                          float rf) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx < num_nets) {
+        int start_idx = flat_net2pin_start_map[idx];
+        int end_idx = flat_net2pin_start_map[idx + 1];
+        int root = flat_net2pin_map[start_idx];
+        float x_root = x[pin2node_map[root]] + pin_offset_x[root];
+        float y_root = y[pin2node_map[root]] + pin_offset_y[root];
+        float root_cap = 0;
+
+        // Load
+        for (int i = start_idx + 1; i < end_idx; i++) {
+            int pin_id = flat_net2pin_map[i];
+            float x_pin = x[pin2node_map[pin_id]] + pin_offset_x[pin_id];
+            float y_pin = y[pin2node_map[pin_id]] + pin_offset_y[pin_id];
+            float dist = abs(x_pin - x_root) + abs(y_pin - y_root);
+            float wl = dist / unit_to_micron;
+            if (net_is_clock[idx]) wl = 0;
+            float pin_cap = cf * wl * 0.5;
+            float pin_res = rf * wl;
+            root_cap += pin_cap;
+
+            for (int j = 0; j < NUM_ATTR; j++) {
+                float pin_cap_lib =
+                    isnan(pinCap[pin_id * (NUM_ATTR + 2) + j]) ? pinCap[pin_id * (NUM_ATTR + 2) + 4 + (j >> 1)] : pinCap[pin_id * (NUM_ATTR + 2) + j];
+                float load = pinLoad[pin_id * NUM_ATTR + j];
+
+                pinLoad[pin_id * NUM_ATTR + j] = isnan(load) ? pin_cap + pin_cap_lib : load + pin_cap + pin_cap_lib;
+                pinRootRes[pin_id * NUM_ATTR + j] = pin_res;
+                pinLoad[root * NUM_ATTR + j] = isnan(pinLoad[root * NUM_ATTR + j]) ? pinLoad[pin_id * NUM_ATTR + j]
+                                                                                   : pinLoad[root * NUM_ATTR + j] + pinLoad[pin_id * NUM_ATTR + j];
+            }
+        }
+        // Root
+        for (int j = 0; j < NUM_ATTR; j++) {
+            float pin_cap_lib =
+                isnan(pinCap[root * (NUM_ATTR + 2) + j]) ? pinCap[root * (NUM_ATTR + 2) + 4 + (j >> 1)] : pinCap[root * (NUM_ATTR + 2) + j];
+            float load = pinLoad[root * NUM_ATTR + j];
+            pinLoad[root * NUM_ATTR + j] = isnan(load) ? root_cap + pin_cap_lib : load + root_cap + pin_cap_lib;
+        }
+        // Delay
+        for (int i = start_idx + 1; i < end_idx; i++) {
+            int pin_id = flat_net2pin_map[i];
+            for (int j = 0; j < NUM_ATTR; j++) {
+                pinRootDelay[pin_id * NUM_ATTR + j] = pinRootRes[pin_id * NUM_ATTR + j] * pinLoad[pin_id * NUM_ATTR + j];
+                pinImpulse[pin_id * NUM_ATTR + j] = 0;
+            }
+        }
+        // Impulse
+        for (int i = start_idx + 1; i < end_idx; i++) {
+            int pin_id = flat_net2pin_map[i];
+            for (int j = 0; j < NUM_ATTR; j++) {
+                float pin_cap_lib =
+                    isnan(pinCap[pin_id * (NUM_ATTR + 2) + j]) ? pinCap[pin_id * (NUM_ATTR + 2) + 4 + (j >> 1)] : pinCap[pin_id * (NUM_ATTR + 2) + j];
+                float res = pinRootRes[pin_id * NUM_ATTR + j];
+                float cap = pinLoad[pin_id * NUM_ATTR + j];
+                float delay = pinRootDelay[pin_id * NUM_ATTR + j];
+                pinImpulse[pin_id * NUM_ATTR + j] = sqrt(2 * res * cap * delay - delay * delay);
+            }
+        }
+        if (end_idx - start_idx == 1) {
+            for (int j = 0; j < NUM_ATTR; j++) {
+                pinRootDelay[root * NUM_ATTR + j] = 0;
+                pinImpulse[root * NUM_ATTR + j] = 0;
+            }
+        }
+    }
+}
+
+void update_rc_timing_cuda(float *x,
+                           float *y,
+                           const float *pin_offset_x,
+                           const float *pin_offset_y,
+                           const int *pin2node_map,
+                           const int *flat_net2pin_start_map,
+                           const int *flat_net2pin_map,
+                           float *pinLoad,
+                           float *pinImpulse,
+                           float *pinCap,
+                           float *pinWireCap,
+                           float *pinRootDelay,
+                           float *pinRootRes,
+                           int num_nets,
+                           int num_pins,
+                           float unit_to_micron,
+                           int *net_is_clock,
+                           float cf,
+                           float rf) {
+    RCTreeNet<<<BLOCK_NUMBER(num_nets), BLOCK_SIZE>>>(x,
+                                                      y,
+                                                      pin_offset_x,
+                                                      pin_offset_y,
+                                                      pin2node_map,
+                                                      flat_net2pin_start_map,
+                                                      flat_net2pin_map,
+                                                      pinLoad,
+                                                      pinImpulse,
+                                                      pinCap,
+                                                      pinWireCap,
+                                                      pinRootDelay,
+                                                      pinRootRes,
+                                                      num_nets,
+                                                      unit_to_micron,
+                                                      net_is_clock,
+                                                      cf,
+                                                      rf);
+}
+
+__global__ void flatten_rc_kernel(const int *edge_from,
+                                  const int *edge_to,
+                                  const int *flat_net2node_start_map,
+                                  const int *flat_net2edge_start_map,
+                                  float *edge_res,
+                                  float *res_parent,
+                                  int *parent_node,
+                                  int *root_dist,
+                                  int *cnts,
+                                  int *edge_cnts,
+                                  int *node_order,
+                                  int *edge_order,
+                                  int num_nets,
+                                  int num_nodes,
+                                  int num_edges) {
+    const int idx = blockIdx.x;
+    if (idx < num_nets) {
+        int nst = flat_net2node_start_map[idx];
+        int nend = flat_net2node_start_map[idx + 1];
+        int root = nst;
+
+        int est = flat_net2edge_start_map[idx];
+        int eend = flat_net2edge_start_map[idx + 1];
+
+        if (threadIdx.x == 0) {
+            parent_node[root] = -1;
+            root_dist[root] = 0;
+        }
+
+        __syncthreads();
+
+        for (int d = 0; d < nend - nst; d++) {
+            for (int i = est + threadIdx.x; i < eend; i += blockDim.x) {
+                int from = edge_from[i];
+                int to = edge_to[i];
+                float res = edge_res[i];
+                if ((root_dist[from] == d) && (root_dist[to] == -1)) {
+                    parent_node[to] = from;
+                    root_dist[to] = d + 1;
+                    atomicAdd(&cnts[d + nst], 1);
+                    for (int j = 0; j < NUM_ATTR; j++) {
+                        atomicAdd(&res_parent[to * NUM_ATTR + j], res);
+                    }
+                } else if ((root_dist[to] == d) && (root_dist[from] == -1)) {
+                    parent_node[from] = to;
+                    root_dist[from] = d + 1;
+                    atomicAdd(&cnts[d + nst], 1);
+                    for (int j = 0; j < NUM_ATTR; j++) {
+                        atomicAdd(&res_parent[from * NUM_ATTR + j], res);
+                    }
+                }
+            }
+            __syncthreads();
+            if (cnts[d + nst] == 0) break;
+        }
+
+        if (threadIdx.x == 0) {
+            const int num_edges_local = eend - est;
+
+            // calculate accumulation
+            int edge_count = 0;
+            for (int i = 0; i < num_edges_local; i++) {
+                edge_count += cnts[i + nst];  // FIXME:
+                cnts[i + nst] = edge_count;
+            }
+
+            // calculate order
+            for (int i = 0; i < num_edges_local; i++) {
+                int from = edge_from[i + est];
+                int to = edge_to[i + est];
+                int min_d = min(root_dist[from], root_dist[to]);
+
+                int start = min_d == 0 ? 0 : cnts[min_d - 1 + nst];
+                edge_order[est + start + edge_cnts[min_d + est]] = i + est;
+                atomicAdd(&edge_cnts[min_d + est], 1);
+            }
+        }
+
+        __syncthreads();
+
+        // sort according to dist and cnts
+        extern __shared__ int offset;
+        if (threadIdx.x == 0) {
+            offset = 0;
+        }
+        __syncthreads();
+        for (int d = 0; d < nend - nst; d++) {
+            // if (threadIdx.x == 0) {
+            //     offset += cnts[d + nst];
+            // }
+            __syncthreads();
+            for (int i = nst + threadIdx.x; i < nend; i += blockDim.x) {
+                if (root_dist[i] == d) {
+                    int pos = atomicAdd(&cnts[d + nst], -1);
+                    int off = atomicAdd(&offset, 1);
+                    // order[nst + offset - pos] = i;
+                    node_order[nst + off] = i;
+                }
+            }
+            __syncthreads();
+        }
+    }
+}
+
+__global__ void propagate_rc_kernel(const int *flat_net2node_start_map,
+                                    const int *parent_node,
+                                    const int *node_order,
+                                    const int *node2pin_map,
+                                    const float *res_parent,
+                                    const float *pinCap,
+                                    const float *pinLoad,
+                                    const float *node_cap,
+                                    float *node_load,
+                                    float *node_delay,
+                                    float *node_ldelay,
+                                    float *node_impulse,
+                                    float *node_beta,
+                                    int num_nets,
+                                    int num_nodes) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int cond = threadIdx.y;
+    if (idx < num_nets) {
+        int nst = flat_net2node_start_map[idx];
+        int nend = flat_net2node_start_map[idx + 1];
+
+        for (int i = nend - 1; i >= nst; i--) {
+            int node = node_order[i];
+            int pnode = parent_node[node];
+            int pin = node2pin_map[node];
+            float wire_cap = node_cap[node * NUM_ATTR + cond];
+            if (pin != -1) {
+                float pin_cap_lib =
+                    isnan(pinCap[pin * (NUM_ATTR + 2) + cond]) ? pinCap[pin * (NUM_ATTR + 2) + 4 + (cond >> 1)] : pinCap[pin * (NUM_ATTR + 2) + cond];
+                float pin_load = pinLoad[pin * NUM_ATTR + cond];
+                wire_cap = wire_cap + pin_cap_lib + pin_load;
+            }
+            atomicAdd(&node_load[node * NUM_ATTR + cond], wire_cap);
+            if (pnode != -1) atomicAdd(&node_load[pnode * NUM_ATTR + cond], node_load[node * NUM_ATTR + cond]);
+        }
+        for (int i = nst + 1; i < nend; i++) {
+            int node = node_order[i];
+            int pnode = parent_node[node];
+            int pin = node2pin_map[node];
+            float t = node_load[node * NUM_ATTR + cond] * res_parent[node * NUM_ATTR + cond];
+            node_delay[node * NUM_ATTR + cond] = node_delay[pnode * NUM_ATTR + cond] + t;
+        }
+        for (int i = nend - 1; i >= nst; i--) {
+            int node = node_order[i];
+            int pnode = parent_node[node];
+            int pin = node2pin_map[node];
+            float wire_cap = node_cap[node * NUM_ATTR + cond];
+            if (pin != -1) {
+                float pin_cap_lib =
+                    isnan(pinCap[pin * (NUM_ATTR + 2) + cond]) ? pinCap[pin * (NUM_ATTR + 2) + 4 + (cond >> 1)] : pinCap[pin * (NUM_ATTR + 2) + cond];
+                float pin_load = pinLoad[pin * NUM_ATTR + cond];
+                wire_cap = wire_cap + pin_cap_lib + pin_load;
+            }
+            float l = wire_cap * node_delay[node * NUM_ATTR + cond];
+            atomicAdd(&node_ldelay[node * NUM_ATTR + cond], l);
+            if (pnode != -1) atomicAdd(&node_ldelay[pnode * NUM_ATTR + cond], node_ldelay[node * NUM_ATTR + cond]);
+        }
+        for (int i = nst + 1; i < nend; i++) {
+            int node = node_order[i];
+            int pnode = parent_node[node];
+            int pin = node2pin_map[node];
+            float t = node_ldelay[node * NUM_ATTR + cond] * res_parent[node * NUM_ATTR + cond];
+            node_beta[node * NUM_ATTR + cond] = node_beta[pnode * NUM_ATTR + cond] + t;
+            node_impulse[node * NUM_ATTR + cond] =
+                sqrt(2 * node_beta[node * NUM_ATTR + cond] - node_delay[node * NUM_ATTR + cond] * node_delay[node * NUM_ATTR + cond]);
+        }
+    }
+}
+
+__global__ void move_to_timing_graph(const int *flat_net2node_start_map,
+                                     const int *node2pin_map,
+                                     const float *node_load,
+                                     const float *node_delay,
+                                     const float *node_impulse,
+                                     float *pinLoad,
+                                     float *pinImpulse,
+                                     float *pinRootDelay,
+                                     int num_nets) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    const int cond = threadIdx.y;
+    if (idx < num_nets) {
+        int nst = flat_net2node_start_map[idx];
+        int nend = flat_net2node_start_map[idx + 1];
+        for (int i = nst; i < nend; i++) {
+            int node = i;
+            int pin = node2pin_map[node];
+            if (pin != -1) {
+                pinLoad[pin * NUM_ATTR + cond] = node_load[node * NUM_ATTR + cond];
+                pinRootDelay[pin * NUM_ATTR + cond] = node_delay[node * NUM_ATTR + cond];
+                pinImpulse[pin * NUM_ATTR + cond] = node_impulse[node * NUM_ATTR + cond];
+            }
+        }
+    }
+}
+
+void flatten_rc_tree(std::vector<int> host_edge_from,
+                     std::vector<int> host_edge_to,
+                     float *edge_res,
+                     float *node_cap,
+                     std::vector<int> host_flat_net2node_start_map,
+                     std::vector<int> host_flat_net2edge_start_map,
+                     std::vector<int> host_node2pin_map,
+                     int *node_order,
+                     int *edge_order,
+                     int *parent_node,
+                     float *res_parent,
+                     float *pinLoad,
+                     float *pinImpulse,
+                     float *pinCap,
+                     float *pinWireCap,
+                     float *pinRootDelay,
+                     float *pinRootRes,
+                     int num_nets,
+                     int num_pins,
+                     int num_nodes,
+                     int num_edges) {
+    int *edge_from, *edge_to, *flat_net2node_start_map, *flat_net2edge_start_map, *node2pin_map;
+
+    cudaMalloc(&edge_from, host_edge_from.size() * sizeof(int));
+    cudaMalloc(&edge_to, host_edge_to.size() * sizeof(int));
+    cudaMalloc(&flat_net2node_start_map, host_flat_net2node_start_map.size() * sizeof(int));
+    cudaMalloc(&flat_net2edge_start_map, host_flat_net2edge_start_map.size() * sizeof(int));
+    cudaMalloc(&node2pin_map, host_node2pin_map.size() * sizeof(int));
+
+    cudaMemcpy(edge_from, host_edge_from.data(), host_edge_from.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(edge_to, host_edge_to.data(), host_edge_to.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(flat_net2node_start_map, host_flat_net2node_start_map.data(), host_flat_net2node_start_map.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(flat_net2edge_start_map, host_flat_net2edge_start_map.data(), host_flat_net2edge_start_map.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(node2pin_map, host_node2pin_map.data(), host_node2pin_map.size() * sizeof(int), cudaMemcpyHostToDevice);
+
+    int *root_dist, *cnts;
+    int *edge_cnts;
+    cudaMalloc(&root_dist, num_nodes * sizeof(int));
+    cudaMalloc(&cnts, num_nodes * sizeof(int));
+    cudaMalloc(&edge_cnts, num_edges * sizeof(int));
+
+    reset_val<int><<<BLOCK_NUMBER(num_nodes), BLOCK_SIZE>>>(root_dist, num_nodes);
+    cudaMemset(cnts, 0, num_nodes * sizeof(int));
+    cudaMemset(edge_cnts, 0, num_edges * sizeof(int));
+
+    int thread_count = 64;
+    int numBlocks = num_nets;
+    flatten_rc_kernel<<<numBlocks, thread_count>>>(edge_from,
+                                                   edge_to,
+                                                   flat_net2node_start_map,
+                                                   flat_net2edge_start_map,
+                                                   edge_res,
+                                                   res_parent,
+                                                   parent_node,
+                                                   root_dist,
+                                                   cnts,
+                                                   edge_cnts,
+                                                   node_order,
+                                                   edge_order,
+                                                   num_nets,
+                                                   num_nodes,
+                                                   num_edges);
+    cudaFree(edge_from);
+    cudaFree(edge_to);
+    cudaFree(flat_net2node_start_map);
+    cudaFree(flat_net2edge_start_map);
+    cudaFree(node2pin_map);
+    cudaFree(root_dist);
+    cudaFree(cnts);
+    cudaFree(edge_cnts);
+
+    // device sync
+    cudaDeviceSynchronize();
+}
+
+void propagate_rc_tree(std::vector<int> host_edge_from,
+                       std::vector<int> host_edge_to,
+                       float *edge_res,
+                       float *node_cap,
+                       std::vector<int> host_flat_net2node_start_map,
+                       std::vector<int> host_flat_net2edge_start_map,
+                       std::vector<int> host_node2pin_map,
+                       int *node_order,
+                       int *parent_node,
+                       float *res_parent,
+                       float *pinLoad,
+                       float *pinImpulse,
+                       float *pinCap,
+                       float *pinWireCap,
+                       float *pinRootDelay,
+                       float *pinRootRes,
+                       int num_nets,
+                       int num_pins,
+                       int num_nodes,
+                       int num_edges) {
+    int *edge_from, *edge_to, *flat_net2node_start_map, *flat_net2edge_start_map, *node2pin_map;
+    cudaMalloc(&edge_from, host_edge_from.size() * sizeof(int));
+    cudaMalloc(&edge_to, host_edge_to.size() * sizeof(int));
+    cudaMalloc(&flat_net2node_start_map, host_flat_net2node_start_map.size() * sizeof(int));
+    cudaMalloc(&flat_net2edge_start_map, host_flat_net2edge_start_map.size() * sizeof(int));
+    cudaMalloc(&node2pin_map, host_node2pin_map.size() * sizeof(int));
+    cudaMemcpy(edge_from, host_edge_from.data(), host_edge_from.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(edge_to, host_edge_to.data(), host_edge_to.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(flat_net2node_start_map, host_flat_net2node_start_map.data(), host_flat_net2node_start_map.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(flat_net2edge_start_map, host_flat_net2edge_start_map.data(), host_flat_net2edge_start_map.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(node2pin_map, host_node2pin_map.data(), host_node2pin_map.size() * sizeof(int), cudaMemcpyHostToDevice);
+
+    float *node_load, *node_delay, *node_ldelay, *node_impulse, *node_beta;
+
+    cudaMalloc(&node_load, num_nodes * NUM_ATTR * sizeof(float));
+    cudaMalloc(&node_delay, num_nodes * NUM_ATTR * sizeof(float));
+    cudaMalloc(&node_ldelay, num_nodes * NUM_ATTR * sizeof(float));
+    cudaMalloc(&node_impulse, num_nodes * NUM_ATTR * sizeof(float));
+    cudaMalloc(&node_beta, num_nodes * NUM_ATTR * sizeof(float));
+
+    cudaMemset(node_load, 0, num_nodes * NUM_ATTR * sizeof(float));
+    cudaMemset(node_delay, 0, num_nodes * NUM_ATTR * sizeof(float));
+    cudaMemset(node_ldelay, 0, num_nodes * NUM_ATTR * sizeof(float));
+    cudaMemset(node_impulse, 0, num_nodes * NUM_ATTR * sizeof(float));
+    cudaMemset(node_beta, 0, num_nodes * NUM_ATTR * sizeof(float));
+
+    int thread_count2 = 64;
+    dim3 block_size(thread_count2, NUM_ATTR);
+    int numBlocks2 = num_nets - 1 + thread_count2 / thread_count2;
+
+    propagate_rc_kernel<<<numBlocks2, block_size>>>(flat_net2node_start_map,
+                                                    parent_node,
+                                                    node_order,
+                                                    node2pin_map,
+                                                    res_parent,
+                                                    pinCap,
+                                                    pinLoad,
+                                                    node_cap,
+                                                    node_load,
+                                                    node_delay,
+                                                    node_ldelay,
+                                                    node_impulse,
+                                                    node_beta,
+                                                    num_nets,
+                                                    num_nodes);
+
+    move_to_timing_graph<<<numBlocks2, block_size>>>(flat_net2node_start_map, node2pin_map, node_load, node_delay, node_impulse, pinLoad, pinImpulse, pinRootDelay, num_nets);
+
+    cudaFree(edge_from);
+    cudaFree(edge_to);
+    cudaFree(flat_net2node_start_map);
+    cudaFree(flat_net2edge_start_map);
+    cudaFree(node2pin_map);
+
+    cudaFree(node_load);
+    cudaFree(node_delay);
+    cudaFree(node_ldelay);
+    cudaFree(node_impulse);
+    cudaFree(node_beta);
+    cudaDeviceSynchronize();
+}
+
+__global__ void calc_rc_kernel(const int *edge_from,
+                               const int *edge_to,
+                               const int *flat_net2node_start_map,
+                               const int *flat_net2edge_start_map,
+                               int *root_dist,
+                               int *cnts,
+                               const int *edge_order,
+                               const float *edge_wl,
+                               float *node_cap,
+                               float *edge_res,
+                               int num_nets,
+                               int num_edges,
+                               int *net_is_clock,
+                               float unit_to_micron,
+                               float rf,
+                               float cf) {
+    const int idx = blockIdx.x;
+    if (idx < num_nets) {
+        int nst = flat_net2node_start_map[idx];
+        int nend = flat_net2node_start_map[idx + 1];
+        int root = nst;
+
+        int est = flat_net2edge_start_map[idx];
+        int eend = flat_net2edge_start_map[idx + 1];
+
+        if (threadIdx.x == 0) {
+            root_dist[root] = 0;
+        }
+        __syncthreads();
+
+        for (int d = 0; d < nend - nst; d++) {
+            for (int i = est + threadIdx.x; i < eend; i += blockDim.x) {
+                int from = edge_from[i];
+                int to = edge_to[i];
+                float wl = edge_wl[i];
+                if (net_is_clock[idx] == 1) wl = 0;
+                float cap = wl * cf * 0.5 / unit_to_micron;
+                float res = wl * rf / unit_to_micron;
+                if ((root_dist[from] == d) && (root_dist[to] == -1)) {
+                    root_dist[to] = d + 1;
+                    atomicAdd(&cnts[d + nst], 1);
+                    atomicAdd(&edge_res[i], res);
+                    for (int j = 0; j < NUM_ATTR; j++) {
+                        atomicAdd(&node_cap[to * NUM_ATTR + j], cap);
+                        atomicAdd(&node_cap[from * NUM_ATTR + j], cap);
+                    }
+                } else if ((root_dist[to] == d) && (root_dist[from] == -1)) {
+                    root_dist[from] = d + 1;
+                    atomicAdd(&cnts[d + nst], 1);
+                    atomicAdd(&edge_res[i], res);
+                    for (int j = 0; j < NUM_ATTR; j++) {
+                        atomicAdd(&node_cap[to * NUM_ATTR + j], cap);
+                        atomicAdd(&node_cap[from * NUM_ATTR + j], cap);
+                    }
+                }
+            }
+            __syncthreads();
+            if (cnts[d + nst] == 0) break;
+        }
+    }
+}
+
+void calc_res_cap(std::vector<int> host_edge_from,
+                  std::vector<int> host_edge_to,
+                  int *edge_order,
+                  float *edge_res,
+                  float *node_cap,
+                  std::vector<int> host_flat_net2node_start_map,
+                  std::vector<int> host_flat_net2edge_start_map,
+                  std::vector<int> host_node2pin_map,
+                  std::vector<float> host_edge_wl,
+                  int num_nets,
+                  int num_edges,
+                  int num_nodes,
+                  int *net_is_clock,
+                  float unit_to_micron,
+                  float rf,
+                  float cf) {
+    int *edge_from, *edge_to, *flat_net2node_start_map, *flat_net2edge_start_map, *node2pin_map;
+    float *edge_wl;
+    cudaMalloc(&edge_from, host_edge_from.size() * sizeof(int));
+    cudaMalloc(&edge_to, host_edge_to.size() * sizeof(int));
+    cudaMalloc(&flat_net2node_start_map, host_flat_net2node_start_map.size() * sizeof(int));
+    cudaMalloc(&flat_net2edge_start_map, host_flat_net2edge_start_map.size() * sizeof(int));
+    cudaMalloc(&node2pin_map, host_node2pin_map.size() * sizeof(int));
+    cudaMalloc(&edge_wl, host_edge_wl.size() * sizeof(float));
+
+    cudaMemcpy(edge_from, host_edge_from.data(), host_edge_from.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(edge_to, host_edge_to.data(), host_edge_to.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(flat_net2node_start_map, host_flat_net2node_start_map.data(), host_flat_net2node_start_map.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(flat_net2edge_start_map, host_flat_net2edge_start_map.data(), host_flat_net2edge_start_map.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(node2pin_map, host_node2pin_map.data(), host_node2pin_map.size() * sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(edge_wl, host_edge_wl.data(), host_edge_wl.size() * sizeof(float), cudaMemcpyHostToDevice);
+
+    int *root_dist, *cnts;
+    cudaMalloc(&root_dist, num_nodes * sizeof(int));
+    cudaMalloc(&cnts, num_nodes * sizeof(int));
+
+    reset_val<int><<<BLOCK_NUMBER(num_nodes), BLOCK_SIZE>>>(root_dist, num_nodes);
+    cudaMemset(cnts, 0, num_nodes * sizeof(int));
+
+    int thread_count = 64;
+    int numBlocks = num_nets;
+    calc_rc_kernel<<<numBlocks, thread_count>>>(edge_from,
+                                                edge_to,
+                                                flat_net2node_start_map,
+                                                flat_net2edge_start_map,
+                                                root_dist,
+                                                cnts,
+                                                edge_order,
+                                                edge_wl,
+                                                node_cap,
+                                                edge_res,
+                                                num_nets,
+                                                num_edges,
+                                                net_is_clock,
+                                                unit_to_micron,
+                                                rf,
+                                                cf);
+
+    cudaFree(edge_from);
+    cudaFree(edge_to);
+    cudaFree(flat_net2node_start_map);
+    cudaFree(flat_net2edge_start_map);
+    cudaFree(node2pin_map);
+    cudaFree(root_dist);
+    cudaFree(cnts);
+    cudaFree(edge_wl);
+}
+
+}  // namespace gt

cpp_to_py/gputimer/core/spef.cpp

@@ -0,0 +1,49 @@
+#include "GPUTimer.h"
+#include "gputimer/db/GTDatabase.h"
+
+using std::ofstream;
+using std::string;
+using std::cerr;
+using std::endl;
+using std::stringstream;
+
+namespace gt {
+
+void GPUTimer::read_spef(const std::string& file) {
+    logger.info("reading spef: %s", file.c_str());
+    if (not std::filesystem::exists(file)) {
+        std::cerr << "can't find " << file << '\n';
+        std::exit(EXIT_FAILURE);
+    }
+
+    // Invoke the read function and check the return value
+    if (not spef.read(file)) {
+        std::cerr << *spef.error;
+        std::exit(EXIT_FAILURE);
+    }
+
+
+    if (spef.time_unit == "1 PS") gtdb.spef_time_unit = 1e-12;
+    if (spef.time_unit == "1 NS") gtdb.spef_time_unit = 1e-9;
+    if (spef.time_unit == "1 US") gtdb.spef_time_unit = 1e-6;
+    if (spef.time_unit == "1 MS") gtdb.spef_time_unit = 1e-3;
+    if (spef.time_unit == "1 S") gtdb.spef_time_unit = 1.0; ;
+
+    if (spef.capacitance_unit == "1 FF") gtdb.spef_cap_unit = 1e-15;
+    if (spef.capacitance_unit == "1 PF") gtdb.spef_cap_unit = 1e-12;
+    if (spef.capacitance_unit == "1 NF") gtdb.spef_cap_unit = 1e-9;
+    if (spef.capacitance_unit == "1 UF") gtdb.spef_cap_unit = 1e-6;
+    if (spef.capacitance_unit == "1 F") gtdb.spef_cap_unit = 1.0;
+
+    if (spef.resistance_unit == "1 OHM") gtdb.spef_res_unit = 1.0;
+    if (spef.resistance_unit == "1 KOHM") gtdb.spef_res_unit = 1e3;
+    if (spef.resistance_unit == "1 MOHM") gtdb.spef_res_unit = 1e6;
+
+    logger.info("spef time_unit: %.5E s", *gtdb.spef_time_unit);
+    logger.info("spef capacitance_unit: %.5E F", *gtdb.spef_cap_unit);
+    logger.info("spef resistance_unit: %.5E Ohm", *gtdb.spef_res_unit);
+
+    spef.expand_name();
+}
+
+}  // namespace gt

cpp_to_py/gputimer/core/utils.cuh

@@ -0,0 +1,81 @@
+#pragma once
+
+#include "gputimer/base.h"
+namespace gt {
+
+template <typename T>
+__global__ void debugPrint(T *arr, int size) {
+    for (int i = 0; i < size; i++) {
+        if constexpr (std::is_same_v<T, int>) {
+            printf("%d %d\n", i, arr[i]);
+        } else if constexpr (std::is_same_v<T, float>) {
+            printf("%d %f\n", i, arr[i]);
+        } else if constexpr (std::is_same_v<T, index_type>) {
+            printf("%d %d\n", i, arr[i]);
+        }
+    }
+    printf("\n");
+}
+template <typename T>
+__global__ void debugPrint1(T *arr, int size, int m) {
+    for (int i = 0; i < size; i++) {
+        printf("%d", i);
+        for (int j = 0; j < m; j++) {
+            if constexpr (std::is_same_v<T, int>) {
+                printf(" %d ", arr[m * i + j]);
+            } else if constexpr (std::is_same_v<T, float>) {
+                printf(" %f ", arr[m * i + j]);
+            } else if constexpr (std::is_same_v<T, index_type>) {
+                printf(" %d ", arr[m * i + j]);
+            }
+        }
+        printf("\n");
+    }
+    printf("\n");
+}
+template <typename T>
+__global__ void debugPrintIdx(int idx, T *arr) {
+    printf("idx: %d, value: %d\n", idx, arr[idx]);
+}
+
+template <typename T>
+__global__ void reset(float *array, int size) {
+    for (int i = 0; i < size; i++) {
+        array[i] = nanf("");
+    }
+}
+
+template <typename T>
+__global__ void reset_batch(float *array, int size) {
+    const int index = blockIdx.x * blockDim.x + threadIdx.x;
+    if (index < size) array[index] = nanf("");
+}
+
+template <typename T>
+__global__ void reset_val(T *array, int size) {
+    const int index = blockIdx.x * blockDim.x + threadIdx.x;
+    if (index < size) {
+        if constexpr (std::is_same_v<T, int>) {
+            array[index] = -1;
+        }
+        if constexpr (std::is_same_v<T, float>) {
+            array[index] = nanf("");
+        }
+    }
+}
+
+template <typename T>
+__global__ void device_copy(T *src, T *dst, int size) {
+    for (int i = 0; i < size; i++) {
+        dst[i] = src[i];
+    }
+}
+
+template <typename T>
+__global__ void device_copy_batch(T *src, T *dst, int size) {
+    const int index = blockIdx.x * blockDim.x + threadIdx.x;
+    if (index < size) dst[index] = src[index];
+}
+
+
+}  // namespace gt

cpp_to_py/gputimer/db/GTDatabase.cpp

@@ -0,0 +1,594 @@
+
+
+#include "GTDatabase.h"
+
+#include "common/common.h"
+#include "common/db/Cell.h"
+#include "common/db/Database.h"
+#include "common/db/Pin.h"
+#include "common/lib/Liberty.h"
+#include "common/lib/Timing.h"
+#include "common/lib/sdc/sdc.h"
+#include "io_parser/gp/GPDatabase.h"
+
+namespace gt {
+
+bool GTDatabase::is_redundant_timing(const TimingArc* timing_arc, Split el) {
+    if (timing_arc->from_port_->name == timing_arc->to_port_->name) return true;
+    if (timing_arc->related_port_name_.empty()) return true;
+    if (timing_arc->timing_type_ == TimingType::non_seq_setup_rising || timing_arc->timing_type_ == TimingType::non_seq_setup_falling || timing_arc->timing_type_ == TimingType::non_seq_hold_rising ||
+        timing_arc->timing_type_ == TimingType::non_seq_hold_falling)
+        return true;
+    switch (el) {
+        case MIN:
+            if (timing_arc->is_max_constraint()) {
+                return true;
+            }
+            break;
+        case MAX:
+            if (timing_arc->is_min_constraint()) {
+                return true;
+            }
+            break;
+    }
+    return false;
+}
+
+GTDatabase::GTDatabase(shared_ptr<db::Database> rawdb_, shared_ptr<gp::GPDatabase> gpdb_, shared_ptr<TimingTorchRawDB> timing_raw_db_) : rawdb(*rawdb_), gpdb(*gpdb_), timing_raw_db(*timing_raw_db_) {
+    cell_libs_[MIN] = rawdb.cell_libs_[MIN];
+    cell_libs_[MAX] = rawdb.cell_libs_[MAX];
+}
+
+
+void GTDatabase::ExtractTimingGraph() {
+    res_unit = cell_libs_[MIN]->resistance_unit_->value();
+    cap_unit = cell_libs_[MIN]->capacitance_unit_->value();
+    time_unit = cell_libs_[MIN]->time_unit_->value();
+    pin_names = gpdb.getPinNames();
+    net_names = gpdb.getNetNames();
+    
+    //  Flatten Liberty Cell Timing
+    for (db::CellType* cell_type : rawdb.celltypes) {
+        string cell_type_name = cell_type->name;
+        array<LibertyCell*, 2> liberty_cell_view = {cell_libs_[MIN]->get_cell(cell_type_name), cell_libs_[MAX]->get_cell(cell_type_name)};
+        if (!liberty_cell_view[MIN] || !liberty_cell_view[MAX]) {
+            liberty_cell_type2port_list_end.push_back(liberty_cell_type2port_list_end.back());
+            continue;
+        }
+        liberty_cell_type2port_list_end.push_back(liberty_cell_type2port_list_end.back() + liberty_cell_view[MIN]->ports_.size());
+        for (int i = 0; i < liberty_cell_view[MIN]->ports_.size(); i++) {
+            array<LibertyPort*, 2> liberty_port_view = {liberty_cell_view[MIN]->ports_[i], liberty_cell_view[MAX]->ports_[i]};
+            for_each_el(el) {
+                liberty_port_capacitance.push_back(liberty_port_view[el]->port_capacitance_[0].value_or(nanf("")));
+                liberty_port_capacitance.push_back(liberty_port_view[el]->port_capacitance_[1].value_or(nanf("")));
+                liberty_port_capacitance.push_back(liberty_port_view[el]->port_capacitance_[2].value_or(0.0f));
+            }
+
+            for_each_el(el) {
+                liberty_port2timing_list_end.push_back(liberty_port2timing_list_end.back() + liberty_port_view[el]->timing_arcs_non_cond_non_bundle_.size());
+                for (int j = 0; j < liberty_port_view[el]->timing_arcs_non_cond_non_bundle_.size(); j++) {
+                    liberty_timing_arcs.push_back(liberty_port_view[el]->timing_arcs_non_cond_non_bundle_[j]);
+                }
+            }
+        }
+    }
+
+    //  Traverse Circuit Pins
+    //
+    num_pins = gpdb.getPins().size();
+    pin_names = gpdb.getPinNames();
+    net_names = gpdb.getNetNames();
+    pin_id2cell_type_id.resize(num_pins);
+    pin_id2port_offset_id.resize(num_pins);
+    STA_pins.resize(num_pins, nullptr);
+    pin_capacitance.resize(2 * 3 * num_pins, 0.0f);
+    for (auto& gppin : gpdb.getPins()) {
+        int pin_id = gppin.getId();
+        string pin_name = gppin.getName();
+        string pin_macro_name = gppin.getMacroName();
+        STA_pins[pin_id] = new STAPin();
+        auto [ori_node_id, ori_node_pin_id, ori_net_id] = gppin.getOriDBInfo();
+        if (ori_node_pin_id == -1) {
+            auto dbiopin = rawdb.iopins[ori_node_id];
+            pin_id2cell_type_id[pin_id] = -1;
+            if (dbiopin->type->direction() == 'i') {
+                primary_outputs.push_back(pin_id);
+                endpoints_id.push_back(pin_id);
+                primary_output2pin_id[pin_name] = pin_id;
+            } else if (dbiopin->type->direction() == 'o') {
+                primary_inputs.push_back(pin_id);
+                primary_input2pin_id[pin_name] = pin_id;
+            }
+        } else {
+            auto& dbcell = rawdb.cells[ori_node_id];
+            LibertyCell* liberty_cell = dbcell->ctype()->liberty_cell;
+            pin_id2cell_type_id[pin_id] = dbcell->ctype()->libcell();
+            pin_id2port_offset_id[pin_id] = liberty_cell->ports_map_[pin_macro_name];
+
+            int liberty_port_id = liberty_cell_type2port_list_end[pin_id2cell_type_id[pin_id]] + pin_id2port_offset_id[pin_id];
+
+            for_each_el(el) {
+                pin_capacitance[6 * pin_id + el * 2 + 0] = liberty_port_capacitance[6 * liberty_port_id + el * 3 + 0];
+                pin_capacitance[6 * pin_id + el * 2 + 1] = liberty_port_capacitance[6 * liberty_port_id + el * 3 + 1];
+                pin_capacitance[6 * pin_id + 4 + el] = liberty_port_capacitance[6 * liberty_port_id + el * 3 + 2];
+            }
+        }
+    }
+    num_POs = primary_outputs.size();
+
+
+    //  Map Pin to Liberty Timing
+    //
+    auto connect_from_to_pin = [&](int from_pin_id, int to_pin_id) -> pair<STAPin*, STAPin*> {
+        STAPin* from_pin = STA_pins[from_pin_id];
+        STAPin* to_pin = STA_pins[to_pin_id];
+        from_pin->fanout_pin_ids.insert(to_pin_id);
+        to_pin->fanin_pin_ids.insert(from_pin_id);
+        timing_arc_from_pin_id.push_back(from_pin_id);
+        timing_arc_to_pin_id.push_back(to_pin_id);
+        from_pin->timing_arc_out.push_back(num_arcs);
+        to_pin->timing_arc_in.push_back(num_arcs);
+        return {from_pin, to_pin};
+    };
+
+    for (auto& gpnet : gpdb.getNets()) {
+        int driver_pin_id = gpnet.pins()[0];
+        for (index_type i = 1; i < static_cast<index_type>(gpnet.pins().size()); i++) {
+            int sink_pin_id = gpnet.pins()[i];
+            auto [from_pin, to_pin] = connect_from_to_pin(driver_pin_id, sink_pin_id);
+            timing_arc_id_map.push_back(-1);
+            timing_arc_id_map.push_back(-1);
+            arc_types.push_back(0);
+            arc_id2test_id.push_back(-1);
+            num_arcs++;
+        }
+    }
+
+    cell_node_type_map.resize(gpdb.getNodes().size(), -1);
+    for (auto& dbcell : rawdb.cells) {
+        int gpdb_id = dbcell->gpdb_id;
+        int libcell_id = dbcell->ctype()->libcell();
+        cell_node_type_map[gpdb_id] = libcell_id;
+        for_each_el(el) {
+            for (int pin_id : gpdb.getNodes()[gpdb_id].pins()) {
+                int pin_id2port_start = liberty_cell_type2port_list_end[libcell_id];
+                int pin_id2port_offset = pin_id2port_offset_id[pin_id];
+                int port_id = pin_id2port_start + pin_id2port_offset;
+                int start = liberty_port2timing_list_end[2 * port_id + el];
+                int end = liberty_port2timing_list_end[2 * port_id + el + 1];
+                for (int i = start; i < end; i++) {
+                    TimingArc* timing_arc = liberty_timing_arcs[i];
+                    if (is_redundant_timing(timing_arc, el)) {
+                        continue;
+                    }
+                    array<int, 2> timing_view = {-1, -1};
+                    timing_view[el] = i;
+
+                    int from_pin_id = gpdb.getNodes()[gpdb_id].getPinbyPortName(timing_arc->from_port_->name);;
+                    int to_pin_id = gpdb.getNodes()[gpdb_id].getPinbyPortName(timing_arc->to_port_->name);
+                    auto [from_pin, to_pin] = connect_from_to_pin(from_pin_id, to_pin_id);
+                    timing_arc_id_map.push_back(timing_view[MIN]);
+                    timing_arc_id_map.push_back(timing_view[MAX]);
+                    arc_types.push_back(1);
+                    num_arcs++;
+
+                    if (timing_arc->is_constraint()) {
+                        arc_id2test_id.push_back(num_tests++);
+                        test_id2_arc_id.push_back(num_arcs - 1);
+                        endpoints_id.push_back(to_pin_id);
+                    } else {
+                        arc_id2test_id.push_back(-1);
+                    }
+                }
+            }
+        }
+    }
+
+    //  Construct Connectivity Graph
+    //
+    for (int i = 0; i < num_pins; i++) total_num_fanouts += STA_pins[i]->fanout_pin_ids.size();
+
+    pin_fanout_list_end.resize(num_pins + 1);
+    pin_fanout_list_end[0] = 0;
+    pin_num_fanin.resize(num_pins);
+    pin_fanout_list.resize(total_num_fanouts);
+
+    index_type ptr = 0;
+    index_type last_idx = 0;
+    for (index_type i = 0; i < static_cast<index_type>(num_pins); i++) {
+        for (auto fanout_pin_id : STA_pins[i]->fanout_pin_ids) pin_fanout_list[ptr++] = fanout_pin_id;
+        last_idx += STA_pins[i]->fanout_pin_ids.size();
+        pin_fanout_list_end[i + 1] = last_idx;
+        pin_num_fanin[i] = STA_pins[i]->fanin_pin_ids.size();
+    }
+    for (int i = 0; i < num_pins; i++) {
+        if (pin_num_fanin[i] == 0) pin_frontiers.push_back(i);
+    }
+
+    pin_forward_arc_list_end.push_back(0);
+    pin_backward_arc_list_end.push_back(0);
+    for (index_type i = 0; i < static_cast<index_type>(num_pins); i++) {
+        for (auto fanout_arc : STA_pins[i]->timing_arc_out) {
+            pin_forward_arc_list.push_back(fanout_arc);
+        }
+        pin_forward_arc_list_end.push_back(pin_forward_arc_list.size());
+        for (auto fanin_arc : STA_pins[i]->timing_arc_in) {
+            pin_backward_arc_list.push_back(fanin_arc);
+        }
+        pin_backward_arc_list_end.push_back(pin_backward_arc_list.size());
+    }
+
+    // gputimer arrays
+    auto device = timing_raw_db.node_size.device();
+    auto options = torch::TensorOptions().dtype(torch::kInt32);
+    // Timer graph topology variables
+    timing_raw_db.pin_forward_arc_list = torch::from_blob(pin_forward_arc_list.data(), {static_cast<index_type>(pin_forward_arc_list.size())}, options).contiguous().to(device);
+    timing_raw_db.pin_forward_arc_list_end = torch::from_blob(pin_forward_arc_list_end.data(), {static_cast<index_type>(pin_forward_arc_list_end.size())}, options).contiguous().to(device);
+    timing_raw_db.pin_backward_arc_list = torch::from_blob(pin_backward_arc_list.data(), {static_cast<index_type>(pin_backward_arc_list.size())}, options).contiguous().to(device);
+    timing_raw_db.pin_backward_arc_list_end = torch::from_blob(pin_backward_arc_list_end.data(), {static_cast<index_type>(pin_backward_arc_list_end.size())}, options).contiguous().to(device);
+    timing_raw_db.timing_arc_from_pin_id = torch::from_blob(timing_arc_from_pin_id.data(), {static_cast<index_type>(timing_arc_from_pin_id.size())}, options).contiguous().to(device);
+    timing_raw_db.timing_arc_to_pin_id = torch::from_blob(timing_arc_to_pin_id.data(), {static_cast<index_type>(timing_arc_to_pin_id.size())}, options).contiguous().to(device);
+    timing_raw_db.pin_num_fanin = torch::from_blob(pin_num_fanin.data(), {static_cast<index_type>(pin_num_fanin.size())}, options).contiguous().to(device);
+    timing_raw_db.pin_fanout_list = torch::from_blob(pin_fanout_list.data(), {static_cast<index_type>(pin_fanout_list.size())}, options).contiguous().to(device);
+    timing_raw_db.pin_fanout_list_end = torch::from_blob(pin_fanout_list_end.data(), {static_cast<index_type>(pin_fanout_list_end.size())}, options).contiguous().to(device);
+
+    // Timer timing liberty variables
+    timing_raw_db.arc_types = torch::from_blob(arc_types.data(), {static_cast<int>(arc_types.size())}, options).contiguous().to(device);
+    timing_raw_db.timing_arc_id_map = torch::from_blob(timing_arc_id_map.data(), {static_cast<int>(timing_arc_id_map.size())}, options).contiguous().to(device);
+    timing_raw_db.arc_id2test_id = torch::from_blob(arc_id2test_id.data(), {static_cast<int>(arc_id2test_id.size())}, options).contiguous().to(device);
+    timing_raw_db.test_id2_arc_id = torch::from_blob(test_id2_arc_id.data(), {static_cast<int>(test_id2_arc_id.size())}, options).contiguous().to(device);
+    timing_raw_db.endpoints_id = torch::from_blob(endpoints_id.data(), {static_cast<index_type>(endpoints_id.size())}, options).contiguous().to(device);
+
+    timing_raw_db.pinSlew = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinLoad = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinRAT = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinAT = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinImpulse = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinRootDelay = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    torch::fill_(timing_raw_db.pinSlew, nanf(""));
+    torch::fill_(timing_raw_db.pinRAT, nanf(""));
+    torch::fill_(timing_raw_db.pinAT, nanf(""));
+    torch::fill_(timing_raw_db.pinImpulse, nanf(""));
+    torch::fill_(timing_raw_db.pinRootDelay, nanf(""));
+
+    timing_raw_db.arcDelay = torch::zeros({num_arcs, 2 * NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinImpulse_ref = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinLoad_ref = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinLoad_ratio = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinRootDelay_ref = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinRootDelay_ratio = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+    timing_raw_db.pinRootDelay_compensation = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(device))).contiguous();
+
+    logger.info("Design info: %d pins, %d arcs, %d tests", num_pins, num_arcs, num_tests);
+}
+
+void GTDatabase::readSdc(sdc::SDC& sdc) {
+    for (auto& command : sdc.commands) {
+        std::visit(Functors{[this](auto&& cmd) { _read_sdc(cmd); }}, command);
+    }
+
+    // string clock_name = clocks.begin()->second.source_name();
+    string clock_name = gpdb.getPins()[clocks.begin()->second.source_id()].getName();
+    float period = clocks.begin()->second.period();
+    logger.info("clock: %s, period: %.2f", clock_name.c_str(), period);
+
+    net_is_clock.resize(gpdb.getNets().size(), 0);
+    for (auto& gpnet : gpdb.getNets()) {
+        if (gpnet.getName() == clock_name) {
+            net_is_clock[gpnet.getId()] = 1;
+        }
+    }
+
+    // set nan slew of PIs to half period
+    for (auto& pi : primary_inputs) {
+        if (torch::isnan(timing_raw_db.pinSlew[pi][0]).item<bool>()) timing_raw_db.pinSlew[pi][0] = 0.0f;
+        if (torch::isnan(timing_raw_db.pinSlew[pi][1]).item<bool>()) timing_raw_db.pinSlew[pi][1] = 0.0f;
+        if (torch::isnan(timing_raw_db.pinSlew[pi][2]).item<bool>()) timing_raw_db.pinSlew[pi][2] = 0.0f;
+        if (torch::isnan(timing_raw_db.pinSlew[pi][3]).item<bool>()) timing_raw_db.pinSlew[pi][3] = 0.0f;
+        // if (torch::isnan(pinAT[pi][0]).item<bool>()) pinAT[pi][0] = 0.0f;
+        // if (torch::isnan(pinAT[pi][1]).item<bool>()) pinAT[pi][1] = period / 2.0;
+        // if (torch::isnan(pinAT[pi][2]).item<bool>()) pinAT[pi][2] = 0.0f;
+        // if (torch::isnan(pinAT[pi][3]).item<bool>()) pinAT[pi][3] = period / 2.0;
+    }
+
+    if (clocks.begin()->second.source_id() != -1) {
+        int clock_pin_id = clocks.begin()->second.source_id();
+        if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][0]).item<bool>()) timing_raw_db.pinAT[clock_pin_id][0] = 0.0f;
+        if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][1]).item<bool>()) timing_raw_db.pinAT[clock_pin_id][1] = 0.0f;
+        if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][2]).item<bool>()) timing_raw_db.pinAT[clock_pin_id][2] = 0.0f;
+        if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][3]).item<bool>()) timing_raw_db.pinAT[clock_pin_id][3] = 0.0f;
+        // if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][0]).item<bool>()) timing_raw_db.pinAT[clock_pin_id][0] = 0.0f;
+        // if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][1]).item<bool>()) timing_raw_db.pinAT[clock_pin_id][1] = period / 2.0;
+        // if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][2]).item<bool>()) timing_raw_db.pinAT[clock_pin_id][2] = 0.0f;
+        // if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][3]).item<bool>()) timing_raw_db.pinAT[clock_pin_id][3] = period / 2.0;
+    }
+}
+
+// Sets input delay on pins or input ports relative to a clock signal.
+void GTDatabase::_read_sdc(sdc::SetUnits& obj) {
+    if (obj.time.has_value()) {
+        auto s = *obj.time;
+        if (s == "ps") sdc_time_unit = 1e-12;
+        if (s == "ns") sdc_time_unit = 1e-9;
+        if (s == "us") sdc_time_unit = 1e-6;
+        if (s == "ms") sdc_time_unit = 1e-3;
+        if (s == "s") sdc_time_unit = 1.0;
+    }
+    if (obj.capacitance.has_value()) {
+        auto s = *obj.capacitance;
+        if (s == "fF") sdc_cap_unit = 1e-15;
+        if (s == "pF") sdc_cap_unit = 1e-12;
+        if (s == "nF") sdc_cap_unit = 1e-9;
+        if (s == "uF") sdc_cap_unit = 1e-6;
+        if (s == "F") sdc_cap_unit = 1.0;
+    }
+    if (obj.resistance.has_value()) {
+        auto s = *obj.resistance;
+        if (s == "Ohm") sdc_res_unit = 1.0;
+        if (s == "kOhm") sdc_res_unit = 1e3;
+        if (s == "MOhm") sdc_res_unit = 1e6;
+    }
+    if (sdc_time_unit.has_value()) printf("sdc time unit: %.2E\n", *sdc_time_unit);
+    if (sdc_cap_unit.has_value()) printf("sdc capacitance unit: %.2E\n", *sdc_cap_unit);
+    if (sdc_res_unit.has_value()) printf("sdc resistance unit: %.2E\n", *sdc_res_unit);
+}
+
+// Sets input delay on pins or input ports relative to a clock signal.
+void GTDatabase::_read_sdc(sdc::SetInputDelay& obj) {
+    assert(obj.delay_value && obj.port_pin_list);
+
+    auto mask = sdc::TimingMask(obj.min, obj.max, obj.rise, obj.fall);
+
+    std::visit(Functors{[&](sdc::AllInputs&) {
+                            for (auto& pi : primary_inputs) {
+                                for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                    float delay = *obj.delay_value;
+                                    if (sdc_time_unit.has_value()) delay = delay * *sdc_time_unit / time_unit;
+                                    timing_raw_db.pinAT[pi][(el << 1) + rf] = delay;
+                                }
+                            }
+                        },
+                        [&](sdc::GetPorts& get_ports) {
+                            for (auto& port : get_ports.ports) {
+                                if (auto itr = primary_input2pin_id.find(port); itr != primary_input2pin_id.end()) {
+                                    for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                        float delay = *obj.delay_value;
+                                        if (sdc_time_unit.has_value()) delay = delay * *sdc_time_unit / time_unit;
+                                        timing_raw_db.pinAT[itr->second][(el << 1) + rf] = delay;
+                                    }
+                                } else {
+                                    printf(obj.command, ": port ", std::quoted(port), " not found");
+                                }
+                            }
+                        },
+                        [](auto&&) { assert(false); }},
+               *obj.port_pin_list);
+}
+
+// Sets input transition on pins or input ports relative to a clock signal.
+void GTDatabase::_read_sdc(sdc::SetInputTransition& obj) {
+    assert(obj.transition && obj.port_list);
+
+    auto mask = sdc::TimingMask(obj.min, obj.max, obj.rise, obj.fall);
+
+    std::visit(Functors{[&](sdc::AllInputs&) {
+                            for (auto& pi : primary_inputs) {
+                                for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                    float transition = *obj.transition;
+                                    if (sdc_time_unit.has_value()) transition = transition * *sdc_time_unit / time_unit;
+                                    timing_raw_db.pinSlew[pi][(el << 1) + rf] = transition;
+                                }
+                            }
+                        },
+                        [&](sdc::GetPorts& get_ports) {
+                            for (auto& port : get_ports.ports) {
+                                if (auto itr = primary_input2pin_id.find(port); itr != primary_input2pin_id.end()) {
+                                    for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                        float transition = *obj.transition;
+                                        if (sdc_time_unit.has_value()) transition = transition * *sdc_time_unit / time_unit;
+                                        timing_raw_db.pinSlew[itr->second][(el << 1) + rf] = transition;
+                                    }
+                                } else {
+                                    printf(obj.command, ": port ", std::quoted(port), " not found");
+                                }
+                            }
+                        },
+                        [](auto&&) { assert(false); }},
+               *obj.port_list);
+}
+
+// Sets input transition on pins or input ports relative to a clock signal.
+void GTDatabase::_read_sdc(sdc::SetDrivingCell& obj) {
+    assert((obj.transitions[0] || obj.transitions[1]) && obj.port_list);
+
+    auto mask = sdc::TimingMask(obj.min, obj.max, obj.rise, obj.fall);
+
+    std::visit(Functors{[&](sdc::AllInputs&) {
+                            for (auto& pi : primary_inputs) {
+                                for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                    float transition = *obj.transitions[el];
+                                    if (sdc_time_unit.has_value()) transition = transition * *sdc_time_unit / time_unit;
+                                    timing_raw_db.pinSlew[pi][(el << 1) + rf] = transition;
+                                }
+                            }
+                        },
+                        [&](sdc::GetPorts& get_ports) {
+                            for (auto& port : get_ports.ports) {
+                                if (auto itr = primary_input2pin_id.find(port); itr != primary_input2pin_id.end()) {
+                                    for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                        float transition = *obj.transitions[el];
+                                        if (sdc_time_unit.has_value()) transition = transition * *sdc_time_unit / time_unit;
+                                        timing_raw_db.pinSlew[itr->second][(el << 1) + rf] = transition;
+                                    }
+                                } else {
+                                    printf(obj.command, ": port ", std::quoted(port), " not found");
+                                }
+                            }
+                        },
+                        [](auto&&) { assert(false); }},
+               *obj.port_list);
+}
+
+// Sets output delay on pins or input ports relative to a clock signal.
+void GTDatabase::_read_sdc(sdc::SetOutputDelay& obj) {
+    assert(obj.delay_value && obj.port_pin_list);
+
+    if (clocks.find(obj.clock) == clocks.end()) {
+        printf(obj.command, ": clock ", std::quoted(obj.clock), " not found");
+        return;
+    }
+
+    auto& clock = clocks.at(obj.clock);
+
+    auto mask = sdc::TimingMask(obj.min, obj.max, obj.rise, obj.fall);
+
+    std::visit(Functors{[&](sdc::AllOutputs&) {
+                            for (auto& po : primary_outputs) {
+                                for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                    float delay = *obj.delay_value;
+                                    if (sdc_time_unit.has_value()) delay = delay * *sdc_time_unit / time_unit;
+                                    timing_raw_db.pinRAT[po][(el << 1) + rf] = el == MIN ? -delay : clock.period() - delay;
+                                }
+                            }
+                        },
+                        [&](sdc::GetPorts& get_ports) {
+                            for (auto& port : get_ports.ports) {
+                                if (auto itr = primary_output2pin_id.find(port); itr != primary_output2pin_id.end()) {
+                                    for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                        float delay = *obj.delay_value;
+                                        if (sdc_time_unit.has_value()) delay = delay * *sdc_time_unit / time_unit;
+                                        timing_raw_db.pinRAT[itr->second][(el << 1) + rf] = el == MIN ? -delay : clock.period() - delay;
+                                    }
+                                } else {
+                                    printf(obj.command, ": port ", std::quoted(port), " not found");
+                                }
+                            }
+                        },
+                        [](auto&&) { assert(false); }},
+               *obj.port_pin_list);
+}
+
+// Sets the load attribute to a specified value on specified ports and nets.
+void GTDatabase::_read_sdc(sdc::SetLoad& obj) {
+    assert(obj.value && obj.objects);
+
+    auto mask = sdc::TimingMask(obj.min, obj.max, std::nullopt, std::nullopt);
+
+    std::visit(Functors{[&](sdc::AllOutputs&) {
+                            for (auto& po : primary_outputs) {
+                                for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                    float load = *obj.value;
+                                    if (sdc_res_unit.has_value()) load = load * *sdc_res_unit / res_unit;
+                                    timing_raw_db.pinLoad[po][(el << 1) + rf] = load;
+                                }
+                            }
+                        },
+                        [&](sdc::GetPorts& get_ports) {
+                            for (auto& port : get_ports.ports) {
+                                if (auto itr = primary_output2pin_id.find(port); itr != primary_output2pin_id.end()) {
+                                    for_each_el_rf_if(el, rf, (mask | el) && (mask | rf)) {
+                                        float load = *obj.value;
+                                        if (sdc_res_unit.has_value()) load = load * *sdc_res_unit / res_unit;
+                                        timing_raw_db.pinLoad[itr->second][(el << 1) + rf] = load;
+                                    }
+                                } else {
+                                    printf(obj.command, ": port ", std::quoted(port), " not found");
+                                }
+                            }
+                        },
+                        [](auto&&) { assert(false); }},
+               *obj.objects);
+}
+
+void GTDatabase::_read_sdc(sdc::CreateClock& obj) {
+    assert(obj.period && !obj.name.empty());
+
+    // create clock from given sources
+    if (obj.port_pin_list) {
+        std::visit(Functors{[&](sdc::GetPorts& get_ports) {
+                                auto& ports = get_ports.ports;
+                                assert(ports.size() == 1);
+                                if (auto itr = primary_input2pin_id.find(ports.front()); itr != primary_input2pin_id.end()) {
+                                    clocks.try_emplace(obj.name, obj.name, itr->second, *obj.period);
+                                } else {
+                                    printf(obj.command, ": port ", std::quoted(ports.front()), " not found");
+                                }
+                            },
+                            [](auto&&) { assert(false); }},
+                   *obj.port_pin_list);
+    }
+    // create virtual clock
+    else {
+        clocks.try_emplace(obj.name, obj.name, *obj.period);
+    }
+}
+
+TimingTorchRawDB::TimingTorchRawDB(torch::Tensor node_lpos_init_,
+                                   torch::Tensor node_size_,
+                                   torch::Tensor pin_rel_lpos_,
+                                   torch::Tensor pin_id2node_id_,
+                                   torch::Tensor pin_id2net_id_,
+                                   torch::Tensor node2pin_list_,
+                                   torch::Tensor node2pin_list_end_,
+                                   torch::Tensor hyperedge_list_,
+                                   torch::Tensor hyperedge_list_end_,
+                                   torch::Tensor net_mask_,
+                                   int num_movable_nodes_,
+                                   float scale_factor_,
+                                   int microns_,
+                                   float wire_resistance_per_micron_,
+                                   float wire_capacitance_per_micron_) {
+    node_lpos_init = node_lpos_init_;
+    node_size = node_size_;
+    pin_rel_lpos = pin_rel_lpos_;
+
+    node_size_x = node_size.index({"...", 0}).clone().contiguous();
+    node_size_y = node_size.index({"...", 1}).clone().contiguous();
+    init_x = node_lpos_init.index({"...", 0}).clone().contiguous();
+    init_y = node_lpos_init.index({"...", 1}).clone().contiguous();
+    pin_offset_x = pin_rel_lpos.index({"...", 0}).clone().contiguous();
+    pin_offset_y = pin_rel_lpos.index({"...", 1}).clone().contiguous();
+    x = init_x.clone().contiguous();
+    y = init_y.clone().contiguous();
+
+    num_nodes = node_size.size(0);
+    num_pins = pin_id2node_id_.size(0);
+    num_nets = hyperedge_list_end_.size(0);
+    num_movable_nodes = num_movable_nodes_;
+    net_mask = net_mask_;
+
+    pinAT = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(node_size.device()))).contiguous();
+    pinRAT = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kFloat32).device(torch::Device(node_size.device()))).contiguous();
+    at_prefix_pin = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kInt32).device(torch::Device(node_size.device()))).contiguous();
+    at_prefix_arc = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kInt32).device(torch::Device(node_size.device()))).contiguous();
+    at_prefix_attr = torch::zeros({num_pins, NUM_ATTR}, torch::dtype(torch::kInt32).device(torch::Device(node_size.device()))).contiguous();
+
+    flat_node2pin_start_map = torch::cat({torch::zeros({1}, torch::dtype(torch::kInt32).device(torch::Device(node_size.device()))), node2pin_list_end_}, 0).to(torch::kInt32).contiguous();
+    flat_node2pin_map = node2pin_list_.to(torch::kInt32);
+    pin2node_map = pin_id2node_id_.to(torch::kInt32);
+
+    flat_net2pin_start_map = torch::cat({torch::zeros({1}, torch::dtype(torch::kInt32).device(torch::Device(node_size.device()))), hyperedge_list_end_}, 0).to(torch::kInt32).contiguous();
+    flat_net2pin_map = hyperedge_list_.to(torch::kInt32);
+    pin2net_map = pin_id2net_id_.to(torch::kInt32);
+
+    num_threads = std::max(6, 1);
+    scale_factor = scale_factor_;
+    microns = microns_;
+    wire_resistance_per_micron = wire_resistance_per_micron_;
+    wire_capacitance_per_micron = wire_capacitance_per_micron_;
+}
+
+void TimingTorchRawDB::commit_from(torch::Tensor x_, torch::Tensor y_) {
+    // commit external pos to original pos
+    init_x.index({torch::indexing::Slice(0, num_movable_nodes)}).data().copy_(x_.index({torch::indexing::Slice(0, num_movable_nodes)}));
+    init_y.index({torch::indexing::Slice(0, num_movable_nodes)}).data().copy_(y_.index({torch::indexing::Slice(0, num_movable_nodes)}));
+    x.index({torch::indexing::Slice(0, num_movable_nodes)}).data().copy_(x_.index({torch::indexing::Slice(0, num_movable_nodes)}));
+    y.index({torch::indexing::Slice(0, num_movable_nodes)}).data().copy_(y_.index({torch::indexing::Slice(0, num_movable_nodes)}));
+}
+
+torch::Tensor TimingTorchRawDB::get_curr_cposx() { return x + node_size_x / 2; }
+torch::Tensor TimingTorchRawDB::get_curr_cposy() { return y + node_size_y / 2; }
+torch::Tensor TimingTorchRawDB::get_curr_lposx() { return x; }
+torch::Tensor TimingTorchRawDB::get_curr_lposy() { return y; }
+
+}  // namespace gt

cpp_to_py/gputimer/db/GTDatabase.h

@@ -0,0 +1,294 @@
+
+
+#pragma once
+#include <torch/extension.h>
+
+#include <memory>
+
+#include "common/common.h"
+#include "common/lib/sdc/sdc.h"
+#include "gputimer/base.h"
+
+using std::set;
+using std::shared_ptr;
+using std::string;
+using std::unordered_map;
+using std::vector;
+using std::array;
+
+namespace gp {
+class GPDatabase;
+class GPPin;
+class GPNet;
+};  // namespace gp
+
+namespace db {
+class Database;
+};  // namespace db
+
+namespace gt {
+class CellLib;
+class LibertyCell;
+class LibertyPort;
+class TimingArc;
+class Lut;
+};  // namespace gt
+
+namespace gt {
+class TimingTorchRawDB;
+class STAPin;
+class clock;
+class Net;
+class Arc;
+class cellpin;
+
+class Clock {
+private:
+    std::string _name;
+    float _period = .0f;
+    int _source_id = -1;
+public:
+    Clock(const std::string& name, float period) : _name(name), _source_id(-1), _period(period) {};
+    Clock(const std::string& name, int source_id, float period) : _name(name), _source_id(source_id), _period(period) {};
+    inline const std::string& name() const;
+    inline float period() const { return _period; }
+    inline int source_id() { return _source_id; }
+};
+
+class STAPin {
+public:
+    vector<index_type> timing_arc_in;
+    vector<index_type> timing_arc_out;
+    set<index_type> fanin_pin_ids;
+    set<index_type> fanout_pin_ids;
+};
+
+class GTDatabase {
+public:
+    db::Database& rawdb;
+    gp::GPDatabase& gpdb;
+    TimingTorchRawDB& timing_raw_db;
+    std::array<std::shared_ptr<gt::CellLib>, MAX_SPLIT> cell_libs_;
+
+    GTDatabase(shared_ptr<db::Database> rawdb_, shared_ptr<gp::GPDatabase> gpdb_, shared_ptr<TimingTorchRawDB> timing_raw_db_);
+    ~GTDatabase() { logger.info("destruct gtdb"); }
+
+public:
+    void ExtractTimingGraph();
+    void readSpef(const std::string& file);
+    void readSdc(sdc::SDC& sdc);
+    void _read_sdc(sdc::SetInputDelay&);
+    void _read_sdc(sdc::SetDrivingCell&);
+    void _read_sdc(sdc::SetInputTransition&);
+    void _read_sdc(sdc::SetOutputDelay&);
+    void _read_sdc(sdc::SetLoad&);
+    void _read_sdc(sdc::CreateClock&);
+    void _read_sdc(sdc::SetUnits&);
+    bool is_redundant_timing(const TimingArc* timing_arc, Split el);
+
+    // Units
+    float res_unit;
+    float cap_unit;
+    float time_unit;
+
+    std::optional<float> sdc_res_unit;
+    std::optional<float> sdc_cap_unit;
+    std::optional<float> sdc_time_unit;
+
+    std::optional<float> spef_res_unit;
+    std::optional<float> spef_cap_unit;
+    std::optional<float> spef_time_unit;
+
+public:
+    vector<string> pin_names;
+    vector<string> net_names;
+    unordered_map<std::string, Clock> clocks;
+    unordered_map<string, index_type> primary_input2pin_id;
+    unordered_map<string, index_type> primary_output2pin_id;
+
+    vector<STAPin*> STA_pins;
+    vector<int> endpoints_id;
+
+    vector<int> liberty_cell_type2port_list_end = {0};
+    vector<int> liberty_port2timing_list_end = {0};
+    vector<float> liberty_port_capacitance;
+    vector<TimingArc*> liberty_timing_arcs;
+
+    vector<int> pin_id2cell_type_id;
+    vector<int> pin_id2port_offset_id;
+    vector<int> cell_node_type_map;
+    vector<float> pin_capacitance;
+
+    int num_arcs = 0;
+    int num_tests = 0;
+    int num_timings = 0;
+    int total_num_fanouts = 0;
+    int num_pins;
+    int num_POs;
+
+    vector<int> timing_arc_from_pin_id, timing_arc_to_pin_id;
+    vector<int> timing_arc_id_map;
+    vector<int> arc_types, arc_id2test_id;
+    vector<int> test_id2_arc_id;
+    vector<int> net_is_clock;
+
+    // Timing Graph
+    /// @param primary_inputs                 primary input pins
+    /// @param primary_outputs                primary output pins
+    /// @param pin_frontiers                  frontier pins
+    /// @param pin_fanout_list_end            pin_fanout start and end index
+    /// @param pin_fanout_list                pin_fanout list of index
+    /// @param pin_num_fanin                  number of fanin pins
+    /// @param pin_forward_arc_list_end       pin_forward_arc start and end index
+    /// @param pin_forward_arc_list           pin_forward_arc list of index
+    /// @param pin_backward_arc_list_end      pin_backward_arc start and end index
+    /// @param pin_backward_arc_list          pin_backward_arc list of index
+    vector<index_type> primary_inputs, primary_outputs;
+    vector<index_type> pin_frontiers;
+    vector<index_type> pin_fanout_list_end, pin_fanout_list;
+    vector<int> pin_num_fanin;
+    vector<index_type> pin_forward_arc_list_end, pin_forward_arc_list;
+    vector<index_type> pin_backward_arc_list_end, pin_backward_arc_list;
+
+};
+
+class TimingTorchRawDB {
+public:
+    TimingTorchRawDB(torch::Tensor node_lpos_init_,
+                     torch::Tensor node_size_,
+                     torch::Tensor pin_rel_lpos_,
+                     torch::Tensor pin_id2node_id_,
+                     torch::Tensor pin_id2net_id_,
+                     torch::Tensor node2pin_list_,
+                     torch::Tensor node2pin_list_end_,
+                     torch::Tensor hyperedge_list_,
+                     torch::Tensor hyperedge_list_end_,
+                     torch::Tensor net_mask_,
+                     int num_movable_nodes_,
+                     float scale_factor_,
+                     int microns_,
+                     float wire_resistance_per_micron_,
+                     float wire_capacitance_per_micron_);
+
+    void commit_from(torch::Tensor x_, torch::Tensor y_);
+    torch::Tensor get_curr_cposx();
+    torch::Tensor get_curr_cposy();
+    torch::Tensor get_curr_lposx();
+    torch::Tensor get_curr_lposy();
+
+public:
+    /* node info */
+    // for backup
+    torch::Tensor node_lpos_init;
+    torch::Tensor node_size;
+    torch::Tensor pin_rel_lpos;
+
+    torch::Tensor init_x;  // original pos (keep it const except committing)
+    torch::Tensor init_y;  // original pos (keep it const except committing)
+    torch::Tensor x;       // mutable/cached pos (current)
+    torch::Tensor y;       // mutable/cached pos (current)
+    torch::Tensor node_size_x;
+    torch::Tensor node_size_y;
+
+    /* pin info */
+    torch::Tensor pin_offset_x;
+    torch::Tensor pin_offset_y;
+
+    // gputimer api tensors
+    torch::Tensor at_prefix_pin;
+    torch::Tensor at_prefix_arc;
+    torch::Tensor at_prefix_attr;
+
+    torch::Tensor flat_node2pin_start_map;
+    torch::Tensor flat_node2pin_map;
+    torch::Tensor pin2node_map;
+
+    /* net info */
+    torch::Tensor flat_net2pin_start_map;
+    torch::Tensor flat_net2pin_map;
+    torch::Tensor pin2net_map;
+    torch::Tensor net_mask;
+
+    /* chip info */
+    int num_pins;
+    int num_nets;
+    int num_nodes;
+    int num_movable_nodes;
+
+    int num_threads;
+
+public:
+    float scale_factor;
+    float wire_resistance_per_micron;
+    float wire_capacitance_per_micron;
+    int microns;
+
+    // Timer model variables
+    /// @param pinSlew                  Slew value on a pin
+    /// @param pinLoad                  Load value on a pin
+    /// @param pinRAT                   Required arrival time on a pin
+    /// @param pinAT                    Arrival time on a pin
+    /// @param pinImpulse               Impulse value on a pin
+    /// @param pinRootDelay             Root delay value on a sink pin
+public:
+    torch::Tensor pinSlew;
+    torch::Tensor pinLoad;
+    torch::Tensor pinRAT;
+    torch::Tensor pinAT;
+    torch::Tensor pinImpulse;
+    torch::Tensor pinRootDelay;
+
+    // Timer RC Tree variables
+    /// @param endpoints_id             Index of the endpoints
+    /// @param arcDelay                 Delay value of an arc
+    /// @param pinImpulse_ref           Reference impulse value of a accurate Timer
+    /// @param pinLoad_ref              Reference load value of a accurate Timer
+    /// @param pinLoad_ratio            Load ratio value of a accurate Timer
+    /// @param pinRootDelay_ref         Reference root delay value of a accurate Timer
+    /// @param pinRootDelay_ratio       Root delay ratio value of a accurate Timer
+    /// @param pinRootDelay_compensation Root delay compensation value compared to a accurate Timer
+public:
+    // vector<float> arcDelay;
+    torch::Tensor endpoints_id;
+    torch::Tensor arcDelay;
+    torch::Tensor pinImpulse_ref;
+    torch::Tensor pinLoad_ref;
+    torch::Tensor pinLoad_ratio;
+    torch::Tensor pinRootDelay_ref;
+    torch::Tensor pinRootDelay_ratio;
+    torch::Tensor pinRootDelay_compensation;
+
+    // Timer graph topology variables
+    /// @param pin_forward_arc_list             List of forward arcs of a pin
+    /// @param pin_forward_arc_list_end         Star & End index of the forward arcs lists
+    /// @param pin_backward_arc_list            List of backward arcs of a pin
+    /// @param pin_backward_arc_list_end        Star & End index of the backward arcs lists
+    /// @param timing_arc_from_pin_id           From pin index of an arc
+    /// @param timing_arc_to_pin_id             To pin index of an arc
+    /// @param pin_num_fanin                    Number of fanin pins of a pin
+    /// @param pin_fanout_list                  List of fanout pins of a pin
+    /// @param pin_fanout_list_end              Star & End index of the fanout pins lists
+public:
+    torch::Tensor pin_forward_arc_list;
+    torch::Tensor pin_forward_arc_list_end;
+    torch::Tensor pin_backward_arc_list;
+    torch::Tensor pin_backward_arc_list_end;
+    torch::Tensor timing_arc_from_pin_id;
+    torch::Tensor timing_arc_to_pin_id;
+    torch::Tensor pin_num_fanin;
+    torch::Tensor pin_fanout_list;
+    torch::Tensor pin_fanout_list_end;
+
+    // Timer timing liberty variables
+    /// @param arc_types                        Types of an arc: 0/1
+    /// @param timing_arc_id_map                Timing liberty index of an arc
+    /// @param arc_id2test_id                   Timing test index of an arc: -1 for non-test arcs
+    /// @param test_id2_arc_id                  Timing arc index of a test
+public:
+    torch::Tensor arc_types;
+    torch::Tensor timing_arc_id_map;
+    torch::Tensor arc_id2test_id;
+    torch::Tensor test_id2_arc_id;
+};
+
+}  // namespace gt

src/run_placement_nesterov.py

@@ -411,6 +411,16 @@ def run_placement_main_nesterov(args, logger):
 
     ps = ParamScheduler(data, args, logger)
 
+    gputimer = None
+    if args.timing_opt:        
+        gputimer = GPUTimer(data, rawdb, gpdb, params, args)
+        data.gputimer = gputimer
+        def timing_eval_func(node_pos):
+            gputimer.update_timing_eval(node_pos)
+            wns_early, tns_early, wns_late, tns_late = gputimer.report_timing_slack()
+            logger.info("early WNS/TNS: %.4f/%.4f (ns) | late WNS/TNS: %.4f/%.4f (ns)" % (wns_early, tns_early, wns_late, tns_late))
+            return wns_early, tns_early, wns_late, tns_late
+            
     # global placement
     node_pos, iteration, gp_hpwl, overflow, gp_time, gp_per_iter = global_placement_main(
         gpdb, rawdb, ps, data, args, logger
@@ -419,6 +429,8 @@ def run_placement_main_nesterov(args, logger):
     node_pos, dp_hpwl, top5overflow, lg_time, dp_time = detail_placement_main(
         node_pos, gpdb, rawdb, ps, data, args, logger
     )
+    if args.timing_opt:
+        wns_early_dp, tns_early_dp, wns_late_dp, tns_late_dp = timing_eval_func(node_pos)
     iteration += 1
 
     route_metrics = None
@@ -435,6 +447,7 @@ def run_placement_main_nesterov(args, logger):
 
     if args.load_from_raw:
         del gpdb, rawdb
+        del gputimer
 
     place_time = time.time() - total_start
     logger.info("GP Time: %.4f LG Time: %.4f DP Time: %.4f Total Place Time: %.4f" % (