From 6aa1b0e2534850703ee4ae8af6841b648b3baf0c Mon Sep 17 00:00:00 2001 From: bunchgrape Date: Fri, 2 May 2025 15:27:40 +0800 Subject: [PATCH] add gputimer --- cpp_to_py/gputimer/CMakeLists.txt | 38 ++ cpp_to_py/gputimer/PyBindCppMain.cpp | 124 +++++ cpp_to_py/gputimer/base.h | 27 ++ cpp_to_py/gputimer/core/GPUTimer.cpp | 138 ++++++ cpp_to_py/gputimer/core/GPUTimer.cu | 142 ++++++ cpp_to_py/gputimer/core/GPUTimer.h | 130 ++++++ cpp_to_py/gputimer/core/gputiming.h | 370 +++++++++++++++ cpp_to_py/gputimer/core/levelize.cu | 82 ++++ cpp_to_py/gputimer/core/path.cpp | 169 +++++++ cpp_to_py/gputimer/core/path.cu | 185 ++++++++ cpp_to_py/gputimer/core/propagate.cpp | 67 +++ cpp_to_py/gputimer/core/propagate.cu | 420 +++++++++++++++++ cpp_to_py/gputimer/core/rctree.cpp | 554 +++++++++++++++++++++++ cpp_to_py/gputimer/core/rctree.cu | 621 ++++++++++++++++++++++++++ cpp_to_py/gputimer/core/spef.cpp | 49 ++ cpp_to_py/gputimer/core/utils.cuh | 81 ++++ cpp_to_py/gputimer/db/GTDatabase.cpp | 594 ++++++++++++++++++++++++ cpp_to_py/gputimer/db/GTDatabase.h | 294 ++++++++++++ src/run_placement_nesterov.py | 13 + 19 files changed, 4098 insertions(+) create mode 100755 cpp_to_py/gputimer/CMakeLists.txt create mode 100644 cpp_to_py/gputimer/PyBindCppMain.cpp create mode 100755 cpp_to_py/gputimer/base.h create mode 100644 cpp_to_py/gputimer/core/GPUTimer.cpp create mode 100644 cpp_to_py/gputimer/core/GPUTimer.cu create mode 100644 cpp_to_py/gputimer/core/GPUTimer.h create mode 100644 cpp_to_py/gputimer/core/gputiming.h create mode 100755 cpp_to_py/gputimer/core/levelize.cu create mode 100644 cpp_to_py/gputimer/core/path.cpp create mode 100644 cpp_to_py/gputimer/core/path.cu create mode 100644 cpp_to_py/gputimer/core/propagate.cpp create mode 100644 cpp_to_py/gputimer/core/propagate.cu create mode 100644 cpp_to_py/gputimer/core/rctree.cpp create mode 100644 cpp_to_py/gputimer/core/rctree.cu create mode 100644 cpp_to_py/gputimer/core/spef.cpp create mode 100755 cpp_to_py/gputimer/core/utils.cuh create mode 100644 cpp_to_py/gputimer/db/GTDatabase.cpp create mode 100644 cpp_to_py/gputimer/db/GTDatabase.h diff --git a/cpp_to_py/gputimer/CMakeLists.txt b/cpp_to_py/gputimer/CMakeLists.txt new file mode 100755 index 0000000..342bc2f --- /dev/null +++ b/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 "$<$:--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}) \ No newline at end of file diff --git a/cpp_to_py/gputimer/PyBindCppMain.cpp b/cpp_to_py/gputimer/PyBindCppMain.cpp new file mode 100644 index 0000000..c01ba62 --- /dev/null +++ b/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 +using namespace Flute; + +namespace Xplace { + +std::shared_ptr create_gputimer(const py::dict& kwargs, + std::shared_ptr rawdb, + std::shared_ptr gpdb, + std::shared_ptr timing_raw_db) { + std::shared_ptr gtdb = std::make_shared(rawdb, gpdb, timing_raw_db); + auto sdc = std::make_shared(); + + try { + if (kwargs.contains("sdc")) sdc->read(kwargs["sdc"].cast()); + } catch (std::exception& e) { + logger.error("%s\n", e.what()); + } + + gtdb->ExtractTimingGraph(); + gtdb->readSdc(*sdc); + + std::shared_ptr gputimer = std::make_shared(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_>(m, "GPUTimer") + .def(pybind11::init, std::shared_ptr>()) + .def("time_unit", >::GPUTimer::time_unit) + .def("read_spef", >::GPUTimer::read_spef) + .def("init", >::GPUTimer::initialize) + .def("levelize", >::GPUTimer::levelize) + .def("update_rc", >::GPUTimer::update_rc_timing) + .def("update_rc_flute", >::GPUTimer::update_rc_timing_flute) + .def("update_rc_spef", >::GPUTimer::update_rc_timing_spef) + .def("update_states", >::GPUTimer::update_states) + .def("update_timing", >::GPUTimer::update_timing) + .def("update_endpoints", >::GPUTimer::update_endpoints) + .def("report_wns", >::GPUTimer::report_wns) + .def("report_tns_elw", >::GPUTimer::report_tns_elw) + .def("report_wns_and_tns", >::GPUTimer::report_wns_and_tns) + .def("report_pin_slack", >::GPUTimer::report_pin_slack, py::return_value_policy::move) + .def("report_pin_at", >::GPUTimer::report_pin_at, py::return_value_policy::move) + .def("report_pin_rat", >::GPUTimer::report_pin_rat, py::return_value_policy::move) + .def("report_pin_slew", >::GPUTimer::report_pin_slew, py::return_value_policy::move) + .def("report_pin_load", >::GPUTimer::report_pin_load, py::return_value_policy::move) + .def("report_endpoint_slack", >::GPUTimer::report_endpoint_slack, py::return_value_policy::move) + .def("endpoints_index", >::GPUTimer::endpoints_index, py::return_value_policy::copy) + .def("report_path", >::GPUTimer::report_path, py::return_value_policy::copy) + .def("report_K_path", >::GPUTimer::report_K_path, py::return_value_policy::copy) + .def("report_criticality", >::GPUTimer::report_criticality, py::return_value_policy::copy) + .def("report_criticality_threshold", >::GPUTimer::report_criticality_threshold, py::return_value_policy::copy) + ; + pybind11::class_>(m, "TimingTorchRawDB") + .def(pybind11::init()) + .def("commit_from", >::TimingTorchRawDB::commit_from) + .def("get_curr_cposx", >::TimingTorchRawDB::get_curr_cposx, py::return_value_policy::move) + .def("get_curr_cposy", >::TimingTorchRawDB::get_curr_cposy, py::return_value_policy::move) + .def("get_curr_lposx", >::TimingTorchRawDB::get_curr_lposx, py::return_value_policy::move) + .def("get_curr_lposy", >::TimingTorchRawDB::get_curr_lposy, py::return_value_policy::move); + + pybind11::class_>(m, "GTDatabase") + .def(pybind11::init, std::shared_ptr, std::shared_ptr>()); + + 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(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 diff --git a/cpp_to_py/gputimer/base.h b/cpp_to_py/gputimer/base.h new file mode 100755 index 0000000..fdaec42 --- /dev/null +++ b/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 +struct Functors : Ts... { + using Ts::operator()... ; +}; + +template +Functors(Ts...) -> Functors; + +} \ No newline at end of file diff --git a/cpp_to_py/gputimer/core/GPUTimer.cpp b/cpp_to_py/gputimer/core/GPUTimer.cpp new file mode 100644 index 0000000..8d9d856 --- /dev/null +++ b/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 gtdb_, shared_ptr timing_raw_db_) + : gtdb(*gtdb_), + timing_raw_db(*timing_raw_db_), + x(timing_raw_db.x.data_ptr()), + y(timing_raw_db.y.data_ptr()), + init_x(timing_raw_db.init_x.data_ptr()), + init_y(timing_raw_db.init_y.data_ptr()), + node_size_x(timing_raw_db.node_size_x.data_ptr()), + node_size_y(timing_raw_db.node_size_y.data_ptr()), + pin_offset_x(timing_raw_db.pin_offset_x.data_ptr()), + pin_offset_y(timing_raw_db.pin_offset_y.data_ptr()), + // GPU pin attributes array + pinSlew(timing_raw_db.pinSlew.data_ptr()), + pinLoad(timing_raw_db.pinLoad.data_ptr()), + pinRAT(timing_raw_db.pinRAT.data_ptr()), + pinAT(timing_raw_db.pinAT.data_ptr()), + pinImpulse(timing_raw_db.pinImpulse.data_ptr()), + pinRootDelay(timing_raw_db.pinRootDelay.data_ptr()), + arcDelay(timing_raw_db.arcDelay.data_ptr()), + // Critical path prefix info + at_prefix_pin(timing_raw_db.at_prefix_pin.data_ptr()), + at_prefix_arc(timing_raw_db.at_prefix_arc.data_ptr()), + at_prefix_attr(timing_raw_db.at_prefix_attr.data_ptr()), + // Timing graph topology + pin_forward_arc_list(timing_raw_db.pin_forward_arc_list.data_ptr()), + pin_forward_arc_list_end(timing_raw_db.pin_forward_arc_list_end.data_ptr()), + pin_backward_arc_list(timing_raw_db.pin_backward_arc_list.data_ptr()), + pin_backward_arc_list_end(timing_raw_db.pin_backward_arc_list_end.data_ptr()), + timing_arc_from_pin_id(timing_raw_db.timing_arc_from_pin_id.data_ptr()), + timing_arc_to_pin_id(timing_raw_db.timing_arc_to_pin_id.data_ptr()), + pin_num_fanin(timing_raw_db.pin_num_fanin.data_ptr()), + pin_fanout_list(timing_raw_db.pin_fanout_list.data_ptr()), + pin_fanout_list_end(timing_raw_db.pin_fanout_list_end.data_ptr()), + // Timer timing liberty variables + timing_arc_id_map(timing_raw_db.timing_arc_id_map.data_ptr()), + arc_types(timing_raw_db.arc_types.data_ptr()), + arc_id2test_id(timing_raw_db.arc_id2test_id.data_ptr()), + test_id2_arc_id(timing_raw_db.test_id2_arc_id.data_ptr()), + // Circuit info + flat_node2pin_start_map(timing_raw_db.flat_node2pin_start_map.data_ptr()), + flat_node2pin_map(timing_raw_db.flat_node2pin_map.data_ptr()), + pin2node_map(timing_raw_db.pin2node_map.data_ptr()), + flat_net2pin_start_map(timing_raw_db.flat_net2pin_start_map.data_ptr()), + flat_net2pin_map(timing_raw_db.flat_net2pin_map.data_ptr()), + pin2net_map(timing_raw_db.pin2net_map.data_ptr()), + net_mask(timing_raw_db.net_mask.data_ptr()), + 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 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(); +} + +tuple 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 diff --git a/cpp_to_py/gputimer/core/GPUTimer.cu b/cpp_to_py/gputimer/core/GPUTimer.cu new file mode 100644 index 0000000..8ea78c1 --- /dev/null +++ b/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<<>>(pinSlew, __pinSlew__, num_pins * NUM_ATTR); + device_copy_batch<<>>(pinLoad, __pinLoad__, num_pins * NUM_ATTR); + device_copy_batch<<>>(pinRAT, __pinRAT__, num_pins * NUM_ATTR); + device_copy_batch<<>>(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<<>>(arcDelay, 2 * num_arcs * NUM_ATTR); + reset_val<<>>(arcSlew, 2 * num_arcs * NUM_ATTR); + reset_val<<>>(testRelatedAT, num_tests * NUM_ATTR); + reset_val<<>>(testRAT, num_tests * NUM_ATTR); + reset_val<<>>(testConstraint, num_tests * NUM_ATTR); + + reset_val<<>>(at_prefix_pin, num_pins * NUM_ATTR); + reset_val<<>>(at_prefix_arc, num_pins * NUM_ATTR); + reset_val<<>>(at_prefix_attr, num_pins * NUM_ATTR); + + device_copy_batch<<>>(__pinSlew__, pinSlew, num_pins * NUM_ATTR); + device_copy_batch<<>>(__pinLoad__, pinLoad, num_pins * NUM_ATTR); + device_copy_batch<<>>(__pinRAT__, pinRAT, num_pins * NUM_ATTR); + device_copy_batch<<>>(__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<<>>(pinAT, testRAT, test_id2_arc_id, timing_arc_from_pin_id, timing_arc_to_pin_id, endpoints0.data_ptr(), num_tests); + update_endpoints_kernel1<<>>(pinAT, pinRAT, primary_outputs, endpoints1.data_ptr(), num_POs); + + endpoint_slacks = torch::cat({endpoints0, endpoints1}, 0).contiguous(); +} + +} // namespace gt diff --git a/cpp_to_py/gputimer/core/GPUTimer.h b/cpp_to_py/gputimer/core/GPUTimer.h new file mode 100644 index 0000000..dbd404b --- /dev/null +++ b/cpp_to_py/gputimer/core/GPUTimer.h @@ -0,0 +1,130 @@ +#pragma once + +#include +#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 gtdb_holder; + shared_ptr timing_raw_db_holder; + GPUTimer(shared_ptr gtdb_, shared_ptr 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 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, vector> report_path(int ep_idx = -1, int el = -1, bool verbose = false); + vector> report_K_path(int K, bool verbose = false); + tuple report_criticality(int K, bool verbose = false, bool deterministic = true); + tuple 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 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 diff --git a/cpp_to_py/gputimer/core/gputiming.h b/cpp_to_py/gputimer/core/gputiming.h new file mode 100644 index 0000000..fd89889 --- /dev/null +++ b/cpp_to_py/gputimer/core/gputiming.h @@ -0,0 +1,370 @@ +#pragma once + +#include +#include "common/lib/Lut.h" +#include "common/lib/Timing.h" + +using std::vector; + +namespace gt { + +template +__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 +__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 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(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(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(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(d_x_array + d_x_offset[in_timing_lut], d_num_x[in_timing_lut], x); + y_idx[1] = lower_bound(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(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(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(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 \ No newline at end of file diff --git a/cpp_to_py/gputimer/core/levelize.cu b/cpp_to_py/gputimer/core/levelize.cu new file mode 100755 index 0000000..4d1441c --- /dev/null +++ b/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 pin_names) { + // check which pins are not in timing graph + std::set 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<<>>( + 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<<<1, 1>>>(next_frontiers, frontiers, num_frontiers); + cudaMemset(next_num_frontiers, 0, sizeof(int)); + // debugPrint<<<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 \ No newline at end of file diff --git a/cpp_to_py/gputimer/core/path.cpp b/cpp_to_py/gputimer/core/path.cpp new file mode 100644 index 0000000..d4bba37 --- /dev/null +++ b/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, vector> 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(); + worst_ep_idx = timing_raw_db.endpoints_id[worst_ep].item(); + worst_ep_i = torch::argmin(ep_slacks[worst_ep]).item(); + } else { + worst_ep_idx = ep_idx; + worst_ep_i = torch::argmin(pin_slacks[worst_ep_idx]).item(); + } + 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 path; + vector path_at; + vector path_rat; + vector to_delay; + vector path_slack; + + while (cur != -1) { + int prev = timing_raw_db.at_prefix_pin[cur][to_i].item(); + int arc_id = timing_raw_db.at_prefix_arc[cur][to_i].item(); + int from_i = timing_raw_db.at_prefix_attr[cur][to_i].item(); + 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(); + rat = timing_raw_db.pinRAT[cur][to_i].item(); + delay = timing_raw_db.arcDelay[arc_id][arc_i].item(); + slack = pin_slacks[cur][to_i].item(); + } + 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(cout), 3, ' '); + cout << gtdb.pin_names[cur] << '\n'; + } + } + + return {path, path_at, to_delay}; +} + +vector> 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> paths; + + for (int i = 0; i < K; i++) { + auto [path, path_at, to_delay] = report_path(endpoints_id[indices[i]].item(), false); + paths.push_back(path); + } + return paths; +} + +// ------------------------------------------------------------------------------------------------------------------------ +// Report timing paths +// +tuple 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 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 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(); + + 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}; +} + + +} \ No newline at end of file diff --git a/cpp_to_py/gputimer/core/path.cu b/cpp_to_py/gputimer/core/path.cu new file mode 100644 index 0000000..1e0c0da --- /dev/null +++ b/cpp_to_py/gputimer/core/path.cu @@ -0,0 +1,185 @@ +#include +#include +#include +#include + +#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 indices, + torch::PackedTensorAccessor32 ep_i_indices, + torch::PackedTensorAccessor32 endpoints_index, + float* from_pin_delay, + torch::PackedTensorAccessor32 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 indices, + torch::PackedTensorAccessor32 ep_i_indices, + torch::PackedTensorAccessor32 endpoints_index, + unsigned long long* from_pin_delay, + torch::PackedTensorAccessor32 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(delay * scalar)); + } else { + atomicAdd(&from_pin_delay[prev_id], static_cast(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(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(inv_scalar * aux_array_uint64[i]); + } +} + +std::tuple 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(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<<>>( + aux_array_uint64_ptr, from_pin_delay.data_ptr(), scalar, inv_scalar, num_pins); + explore_path_deterministic_kernel<<>>( + at_prefix_pin, + at_prefix_arc, + at_prefix_attr, + pinAT, + arc_types, + arcDelay, + indices.packed_accessor32(), + ep_i_indices.packed_accessor32(), + endpoints_index.packed_accessor32(), + aux_array_uint64_ptr, + pin_visited.packed_accessor32(), + K, + scalar); + copyToFloatAuxArray<<>>( + aux_array_uint64_ptr, from_pin_delay.data_ptr(), scalar, inv_scalar, num_pins); + + } else { + explore_path_kernel<<>>(at_prefix_pin, + at_prefix_arc, + at_prefix_attr, + pinAT, + arc_types, + arcDelay, + indices.packed_accessor32(), + ep_i_indices.packed_accessor32(), + endpoints_index.packed_accessor32(), + from_pin_delay.data_ptr(), + pin_visited.packed_accessor32(), + K); + } + + return {from_pin_delay, pin_visited}; +} + +} // namespace gt \ No newline at end of file diff --git a/cpp_to_py/gputimer/core/propagate.cpp b/cpp_to_py/gputimer/core/propagate.cpp new file mode 100644 index 0000000..286ee8c --- /dev/null +++ b/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 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 \ No newline at end of file diff --git a/cpp_to_py/gputimer/core/propagate.cu b/cpp_to_py/gputimer/core/propagate.cu new file mode 100644 index 0000000..6e38047 --- /dev/null +++ b/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 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<<>>(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<<>>(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 diff --git a/cpp_to_py/gputimer/core/rctree.cpp b/cpp_to_py/gputimer/core/rctree.cpp new file mode 100644 index 0000000..1916d34 --- /dev/null +++ b/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 +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, std::set>& pos2pins_map, const utils::PointT& point, int& index) { + if (pos2pins_map.find(point) != pos2pins_map.end()) return pos2pins_map[point]; + pos2pins_map.emplace(point, std::set{index++}); + return pos2pins_map[point]; +} + +tuple, vector, vector, vector, vector, vector, 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(); + const int* flat_node2pin_map = flat_node2pin_map_at.data_ptr(); + const int* pin2node_map = pin2node_map_at.data_ptr(); + const int* flat_net2pin_start_map = flat_net2pin_start_map_at.data_ptr(); + const int* flat_net2pin_map = flat_net2pin_map_at.data_ptr(); + const int* pin2net_map = pin2net_map_at.data_ptr(); + const float* x = x_at.data_ptr(); + const float* y = y_at.data_ptr(); + const float* pin_offset_x = pin_offset_x_at.data_ptr(); + const float* pin_offset_y = pin_offset_y_at.data_ptr(); + int& num_nets = timing_raw_db.num_nets; + + constexpr const int scale = 1000; // flute only supports integers. + using Point2i = utils::PointT; + + vector edge_from; + vector edge_to; + vector edge_wl; + vector flat_net2node_start_map; + vector flat_net2edge_start_map; + vector 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> net_id2edge_from(num_nets); + vector> net_id2edge_to(num_nets); + vector> net_id2edge_wl(num_nets); + vector> 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> pos2pins_map; + std::vector vx, vy; + vx.reserve(degree); + vy.reserve(degree); + + std::map 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((x[node] + offset_x) * scale); + auto y_ = static_cast((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{j}); + vx.emplace_back(x_); + vy.emplace_back(y_); + } + } + const int valid_size = static_cast(vx.size()); + int num_pins = degree; + std::set multipin_pos; + std::map 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(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(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 host_edge_from, + std::vector host_edge_to, + float* edge_res, + float* node_cap, + std::vector host_flat_net2node_start_map, + std::vector host_flat_net2edge_start_map, + std::vector 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 host_edge_from, + std::vector host_edge_to, + float* edge_res, + float* node_cap, + std::vector host_flat_net2node_start_map, + std::vector host_flat_net2edge_start_map, + std::vector 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 host_edge_from, + std::vector host_edge_to, + int* edge_order, + float* edge_res, + float* node_cap, + std::vector host_flat_net2node_start_map, + std::vector host_flat_net2edge_start_map, + std::vector host_node2pin_map, + std::vector 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(), + edge_res.data_ptr(), + node_cap.data_ptr(), + 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(), + node_cap.data_ptr(), + flat_net2node_start_map, + flat_net2edge_start_map, + node2pin_map, + node_order.data_ptr(), + edge_order.data_ptr(), + parent_node.data_ptr(), + res_parent.data_ptr(), + 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(), + node_cap.data_ptr(), + flat_net2node_start_map, + flat_net2edge_start_map, + node2pin_map, + node_order.data_ptr(), + parent_node.data_ptr(), + res_parent.data_ptr(), + 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(); + const int* flat_net2pin_map = flat_net2pin_map_at.data_ptr(); + + vector> net_id2edge_from(num_nets); + vector> net_id2edge_to(num_nets); + vector> net_id2node_cap(num_nets); + vector> net_id2edge_res(num_nets); + vector> net_id2node2pin_map(num_nets); + vector> 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 edge_from; + vector edge_to; + vector node_cap_vec; + vector edge_res_vec; + vector flat_net2node_start_map; + vector flat_net2edge_start_map; + vector 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(), + node_cap.data_ptr(), + flat_net2node_start_map, + flat_net2edge_start_map, + node2pin_map, + node_order.data_ptr(), + edge_order.data_ptr(), + parent_node.data_ptr(), + res_parent.data_ptr(), + 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(), + node_cap.data_ptr(), + flat_net2node_start_map, + flat_net2edge_start_map, + node2pin_map, + node_order.data_ptr(), + parent_node.data_ptr(), + res_parent.data_ptr(), + pinLoad, + pinImpulse, + pinCap, + pinWireCap, + pinRootDelay, + pinRootRes, + num_nets, + num_pins, + num_nodes, + num_edges); +} + +} // namespace gt \ No newline at end of file diff --git a/cpp_to_py/gputimer/core/rctree.cu b/cpp_to_py/gputimer/core/rctree.cu new file mode 100644 index 0000000..dadc47a --- /dev/null +++ b/cpp_to_py/gputimer/core/rctree.cu @@ -0,0 +1,621 @@ +#include +#include +#include +#include + +#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<<>>(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 host_edge_from, + std::vector host_edge_to, + float *edge_res, + float *node_cap, + std::vector host_flat_net2node_start_map, + std::vector host_flat_net2edge_start_map, + std::vector 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<<>>(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<<>>(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 host_edge_from, + std::vector host_edge_to, + float *edge_res, + float *node_cap, + std::vector host_flat_net2node_start_map, + std::vector host_flat_net2edge_start_map, + std::vector 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<<>>(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<<>>(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 host_edge_from, + std::vector host_edge_to, + int *edge_order, + float *edge_res, + float *node_cap, + std::vector host_flat_net2node_start_map, + std::vector host_flat_net2edge_start_map, + std::vector host_node2pin_map, + std::vector 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<<>>(root_dist, num_nodes); + cudaMemset(cnts, 0, num_nodes * sizeof(int)); + + int thread_count = 64; + int numBlocks = num_nets; + calc_rc_kernel<<>>(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 \ No newline at end of file diff --git a/cpp_to_py/gputimer/core/spef.cpp b/cpp_to_py/gputimer/core/spef.cpp new file mode 100644 index 0000000..ca3a569 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/cpp_to_py/gputimer/core/utils.cuh b/cpp_to_py/gputimer/core/utils.cuh new file mode 100755 index 0000000..cc405a1 --- /dev/null +++ b/cpp_to_py/gputimer/core/utils.cuh @@ -0,0 +1,81 @@ +#pragma once + +#include "gputimer/base.h" +namespace gt { + +template +__global__ void debugPrint(T *arr, int size) { + for (int i = 0; i < size; i++) { + if constexpr (std::is_same_v) { + printf("%d %d\n", i, arr[i]); + } else if constexpr (std::is_same_v) { + printf("%d %f\n", i, arr[i]); + } else if constexpr (std::is_same_v) { + printf("%d %d\n", i, arr[i]); + } + } + printf("\n"); +} +template +__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) { + printf(" %d ", arr[m * i + j]); + } else if constexpr (std::is_same_v) { + printf(" %f ", arr[m * i + j]); + } else if constexpr (std::is_same_v) { + printf(" %d ", arr[m * i + j]); + } + } + printf("\n"); + } + printf("\n"); +} +template +__global__ void debugPrintIdx(int idx, T *arr) { + printf("idx: %d, value: %d\n", idx, arr[idx]); +} + +template +__global__ void reset(float *array, int size) { + for (int i = 0; i < size; i++) { + array[i] = nanf(""); + } +} + +template +__global__ void reset_batch(float *array, int size) { + const int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index < size) array[index] = nanf(""); +} + +template +__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) { + array[index] = -1; + } + if constexpr (std::is_same_v) { + array[index] = nanf(""); + } + } +} + +template +__global__ void device_copy(T *src, T *dst, int size) { + for (int i = 0; i < size; i++) { + dst[i] = src[i]; + } +} + +template +__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 \ No newline at end of file diff --git a/cpp_to_py/gputimer/db/GTDatabase.cpp b/cpp_to_py/gputimer/db/GTDatabase.cpp new file mode 100644 index 0000000..f0e797c --- /dev/null +++ b/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 rawdb_, shared_ptr gpdb_, shared_ptr 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 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 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* 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(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 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(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(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(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(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(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(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(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(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(pin_num_fanin.size())}, options).contiguous().to(device); + timing_raw_db.pin_fanout_list = torch::from_blob(pin_fanout_list.data(), {static_cast(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(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(arc_types.size())}, options).contiguous().to(device); + timing_raw_db.timing_arc_id_map = torch::from_blob(timing_arc_id_map.data(), {static_cast(timing_arc_id_map.size())}, options).contiguous().to(device); + timing_raw_db.arc_id2test_id = torch::from_blob(arc_id2test_id.data(), {static_cast(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(test_id2_arc_id.size())}, options).contiguous().to(device); + timing_raw_db.endpoints_id = torch::from_blob(endpoints_id.data(), {static_cast(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()) timing_raw_db.pinSlew[pi][0] = 0.0f; + if (torch::isnan(timing_raw_db.pinSlew[pi][1]).item()) timing_raw_db.pinSlew[pi][1] = 0.0f; + if (torch::isnan(timing_raw_db.pinSlew[pi][2]).item()) timing_raw_db.pinSlew[pi][2] = 0.0f; + if (torch::isnan(timing_raw_db.pinSlew[pi][3]).item()) timing_raw_db.pinSlew[pi][3] = 0.0f; + // if (torch::isnan(pinAT[pi][0]).item()) pinAT[pi][0] = 0.0f; + // if (torch::isnan(pinAT[pi][1]).item()) pinAT[pi][1] = period / 2.0; + // if (torch::isnan(pinAT[pi][2]).item()) pinAT[pi][2] = 0.0f; + // if (torch::isnan(pinAT[pi][3]).item()) 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()) timing_raw_db.pinAT[clock_pin_id][0] = 0.0f; + if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][1]).item()) timing_raw_db.pinAT[clock_pin_id][1] = 0.0f; + if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][2]).item()) timing_raw_db.pinAT[clock_pin_id][2] = 0.0f; + if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][3]).item()) timing_raw_db.pinAT[clock_pin_id][3] = 0.0f; + // if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][0]).item()) timing_raw_db.pinAT[clock_pin_id][0] = 0.0f; + // if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][1]).item()) timing_raw_db.pinAT[clock_pin_id][1] = period / 2.0; + // if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][2]).item()) timing_raw_db.pinAT[clock_pin_id][2] = 0.0f; + // if (torch::isnan(timing_raw_db.pinAT[clock_pin_id][3]).item()) 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 \ No newline at end of file diff --git a/cpp_to_py/gputimer/db/GTDatabase.h b/cpp_to_py/gputimer/db/GTDatabase.h new file mode 100644 index 0000000..3a4c8d3 --- /dev/null +++ b/cpp_to_py/gputimer/db/GTDatabase.h @@ -0,0 +1,294 @@ + + +#pragma once +#include + +#include + +#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 timing_arc_in; + vector timing_arc_out; + set fanin_pin_ids; + set fanout_pin_ids; +}; + +class GTDatabase { +public: + db::Database& rawdb; + gp::GPDatabase& gpdb; + TimingTorchRawDB& timing_raw_db; + std::array, MAX_SPLIT> cell_libs_; + + GTDatabase(shared_ptr rawdb_, shared_ptr gpdb_, shared_ptr 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 sdc_res_unit; + std::optional sdc_cap_unit; + std::optional sdc_time_unit; + + std::optional spef_res_unit; + std::optional spef_cap_unit; + std::optional spef_time_unit; + +public: + vector pin_names; + vector net_names; + unordered_map clocks; + unordered_map primary_input2pin_id; + unordered_map primary_output2pin_id; + + vector STA_pins; + vector endpoints_id; + + vector liberty_cell_type2port_list_end = {0}; + vector liberty_port2timing_list_end = {0}; + vector liberty_port_capacitance; + vector liberty_timing_arcs; + + vector pin_id2cell_type_id; + vector pin_id2port_offset_id; + vector cell_node_type_map; + vector 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 timing_arc_from_pin_id, timing_arc_to_pin_id; + vector timing_arc_id_map; + vector arc_types, arc_id2test_id; + vector test_id2_arc_id; + vector 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 primary_inputs, primary_outputs; + vector pin_frontiers; + vector pin_fanout_list_end, pin_fanout_list; + vector pin_num_fanin; + vector pin_forward_arc_list_end, pin_forward_arc_list; + vector 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 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 \ No newline at end of file diff --git a/src/run_placement_nesterov.py b/src/run_placement_nesterov.py index df3fb6d..2c4ad05 100644 --- a/src/run_placement_nesterov.py +++ b/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" % (