dp global_swap with displacement constraint

cpp_to_py/gpudp/PyBindCppMain.cpp

@@ -110,8 +110,8 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
           });
     m.def(
         "globalSwap",
-        [](std::shared_ptr<dp::DPTorchRawDB> at_db_ptr, int num_bins_x, int num_bins_y, int batch_size, int max_iters) {
+        [](std::shared_ptr<dp::DPTorchRawDB> at_db_ptr, int num_bins_x, int num_bins_y, int batch_size, int max_iters, float displacement_region_ratio) {
-            return dp::globalSwap(*at_db_ptr, num_bins_x, num_bins_y, batch_size, max_iters);
+            return dp::globalSwap(*at_db_ptr, num_bins_x, num_bins_y, batch_size, max_iters, displacement_region_ratio);
         });
     m.def("independentSetMatching",
           [](std::shared_ptr<dp::DPTorchRawDB> at_db_ptr,

cpp_to_py/gpudp/db/dp_torch.h

@@ -114,7 +114,7 @@ void fillerLegalization(DPTorchRawDB& at_db);
 
 // Detailed Placement
 void kReorder(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int K, int max_iters);
-void globalSwap(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int max_iters);
+void globalSwap(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int max_iters, float displacement_region_ratio);
 void independentSetMatching(
     DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int set_size, int max_iters);
 

cpp_to_py/gpudp/dp/detail_placement.cpp

@@ -4,7 +4,7 @@
 namespace dp {
 
 void kReorderCUDA(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int K, int max_iters);
-void globalSwapCUDA(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int max_iters);
+void globalSwapCUDA(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int max_iters, float displacement_region_ratio);
 void independentSetMatchingCUDA(
     DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int set_size, int max_iters);
 
@@ -12,8 +12,8 @@ void kReorder(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int K, int ma
     kReorderCUDA(at_db, num_bins_x, num_bins_y, K, max_iters);
 }
 
-void globalSwap(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int max_iters) {
+void globalSwap(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int max_iters, float displacement_region_ratio) {
-    globalSwapCUDA(at_db, num_bins_x, num_bins_y, batch_size, max_iters);
+    globalSwapCUDA(at_db, num_bins_x, num_bins_y, batch_size, max_iters, displacement_region_ratio);
 }
 
 void independentSetMatching(

cpp_to_py/gpudp/dp/global_swap_cuda.cu

@@ -71,6 +71,8 @@ struct SwapState {
     int max_num_candidates;
     int max_num_candidates_all;
 
+    float displacement_region_ratio = 0;
+
     int pair_hpwl_computing_strategy;  ///< 0: for the original node2pin_map and
                                        ///< net2pin_map; 1: for node2net_map and
                                        ///< net2node_map, which requires
@@ -351,6 +353,19 @@ __global__ void compute_search_bins(DetailedPlaceData db, SwapState<float> state
                                                                       db.y[node_id],
                                                                       db.x[node_id] + db.node_size_x[node_id],
                                                                       db.y[node_id] + db.node_size_y[node_id]);
+
+        if (state.search_bin_strategy and state.displacement_region_ratio != 0) {
+            float dxh = opt_box.xh - (db.x[node_id] + db.node_size_x[node_id]);
+            float dxl = db.x[node_id] - opt_box.xl;
+            float dyh = opt_box.yh - (db.y[node_id] + db.node_size_y[node_id]);
+            float dyl = db.y[node_id] - opt_box.yl;
+            if (dxh > state.displacement_region_ratio * (db.xh - db.xl)) {
+                opt_box.xh = db.x[node_id] + db.node_size_x[node_id] + state.displacement_region_ratio * (db.xh - db.xl);
+            }
+            if (dxl > state.displacement_region_ratio * (db.xh - db.xl)) {
+                opt_box.xl = db.x[node_id] - state.displacement_region_ratio * (db.xh - db.xl);
+            }
+        }
         int cx = db.pos2bin_x(opt_box.center_x());
         int cy = db.pos2bin_y(opt_box.center_y());
         state.search_bins[node_id] = cx * db.num_bins_y + cy;
@@ -834,7 +849,7 @@ int compute_max_num_nodes_per_bin(const DetailedPlaceData& db) {
     return max_num_nodes_per_bin;
 }
 
-void globalSwapCUDA(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int max_iters) {
+void globalSwapCUDA(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int batch_size, int max_iters, float displacement_region_ratio) {
     cudaSetDevice(at_db.node_size_x.get_device());
     DetailedPlaceData db(at_db);
     db.set_num_bins(num_bins_x, num_bins_y);
@@ -852,6 +867,7 @@ void globalSwapCUDA(DPTorchRawDB& at_db, int num_bins_x, int num_bins_y, int bat
 
     const float stop_threshold = 0.1 / 100;
     state.batch_size = batch_size;
+    state.displacement_region_ratio = displacement_region_ratio;
     int max_num_nodes_per_bin = compute_max_num_nodes_per_bin(db);
     state.max_num_candidates = max_num_nodes_per_bin * 5;
     state.max_num_candidates_all = state.batch_size * state.max_num_candidates;

src/detail_placement.py

@@ -498,7 +498,7 @@ def run_dp(node_pos: torch.Tensor, data: PlaceData, args, logger):
     
     dp_handler(gpudp.kReorder, "K-Reorder 1", num_bins_x, num_bins_y, kr_K, kr_iter)
     dp_handler(gpudp.independentSetMatching, "Independent Set Match", num_bins_x, num_bins_y, ism_bs, ism_set, ism_iter)
-    dp_handler(gpudp.globalSwap, "Global Swap", num_bins_x // 2, num_bins_y // 2, gs_bs, gs_iter)
+    dp_handler(gpudp.globalSwap, "Global Swap", num_bins_x // 2, num_bins_y // 2, gs_bs, gs_iter, 0)
     dp_handler(gpudp.kReorder, "K-Reorder 2", num_bins_x, num_bins_y, kr_K, kr_iter)
 
     del dp_rawdb
@@ -506,6 +506,55 @@ def run_dp(node_pos: torch.Tensor, data: PlaceData, args, logger):
     return node_pos
 
 
+def run_dp_local(node_pos: torch.Tensor, data: PlaceData, args, logger, displacement_ratio = 0.1):
+    # GPU Detailed Placement
+    dp_rawdb = setup_detailed_rawdb(node_pos, False, data, args, logger)
+
+    num_bins_x = data.num_bin_x
+    num_bins_y = data.num_bin_y
+    kr_K = 4
+    kr_iter = 2
+    gs_bs = 256
+    gs_iter = 2
+
+    # use integer coordinate systems in DP for better quality
+    scalar = compute_scalar(get_ori_scale_factor(data))
+    # scalar = 1.0
+
+    def dp_handler(dp_func, func_name, *func_args):
+        logger.info("Start running %s..." % func_name)
+        start_time = time.time()
+        if scalar != 1.0:
+            # NOTE: we assume site_width is integer, so 1 / scalar should be an integer
+            logger.info("scale dp_rawdb by %g" % round(1.0 / scalar))
+            dp_rawdb.scale(round(1.0 / scalar), True)
+        dp_func(dp_rawdb, *func_args)
+        if scalar != 1.0:
+            logger.info("scale dp_rawdb back by %g" % scalar)
+            dp_rawdb.scale(scalar, False)
+        # commit lpos for legality check
+        torch.cuda.synchronize(node_pos.device)
+        logger.info("Start checking...")
+        if not dp_rawdb.check(get_ori_scale_factor(data)):
+            dp_rawdb.rollback()
+            logger.error("Check failed in %s. Rollback to previous DP iteration." % func_name)
+            return
+        logger.info("Check Pass. Commit solution...")
+        # update dp_rawdb for next step dp_func and update the final solution
+        dp_rawdb.commit()
+        commit_to_node_pos(node_pos, data, dp_rawdb)
+        logger.info("***** Finish %s, HPWL: %.6E Time: %.4f *****" % (
+            func_name, get_obj_hpwl(node_pos, data, args).item(), time.time() - start_time
+        ))
+    
+    dp_handler(gpudp.kReorder, "K-Reorder 1", num_bins_x, num_bins_y, kr_K, kr_iter)
+    dp_handler(gpudp.globalSwap, "Global Swap", num_bins_x // 2, num_bins_y // 2, gs_bs, gs_iter, displacement_ratio)
+    dp_handler(gpudp.kReorder, "K-Reorder 2", num_bins_x, num_bins_y, kr_K, kr_iter)
+
+    del dp_rawdb
+
+    return node_pos
+
 def run_dp_route_opt(node_pos: torch.Tensor, gpdb, rawdb, ps, data: PlaceData, args, logger):
     # NOTE: we suppose M1's prefer routing direction is 0 (horizontal)
     if ps.enable_route and gpdb.m1direction() == 0:
@@ -703,8 +752,10 @@ def default_detail_placement(node_pos, gpdb, rawdb, ps, data: PlaceData, args, l
         draw_fig_with_cairo_cpp(node_pos, data.node_size, data, info, args)
         torch.cuda.synchronize(node_pos.device)
     dp_start_time = time.time()
-    if args.detail_placement:
+    if args.detail_placement and not args.timing_opt:
         node_pos = run_dp(node_pos, data, args, logger)
+    elif args.detail_placement and args.timing_opt:
+        node_pos = run_dp_local(node_pos, data, args, logger)
     torch.cuda.synchronize(node_pos.device)
     node_pos = run_dp_route_opt(node_pos, gpdb, rawdb, ps, data, args, logger)
     dp_end_time = time.time()