diff --git a/cpp_to_py/gpudp/PyBindCppMain.cpp b/cpp_to_py/gpudp/PyBindCppMain.cpp index a7f5b35..bf10fb2 100644 --- a/cpp_to_py/gpudp/PyBindCppMain.cpp +++ b/cpp_to_py/gpudp/PyBindCppMain.cpp @@ -110,8 +110,8 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { }); m.def( "globalSwap", - [](std::shared_ptr at_db_ptr, int num_bins_x, int num_bins_y, int batch_size, int max_iters) { - return dp::globalSwap(*at_db_ptr, num_bins_x, num_bins_y, batch_size, max_iters); + [](std::shared_ptr 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, displacement_region_ratio); }); m.def("independentSetMatching", [](std::shared_ptr at_db_ptr, diff --git a/cpp_to_py/gpudp/db/dp_torch.h b/cpp_to_py/gpudp/db/dp_torch.h index f9d5df3..1a59e0d 100644 --- a/cpp_to_py/gpudp/db/dp_torch.h +++ b/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); diff --git a/cpp_to_py/gpudp/dp/detail_placement.cpp b/cpp_to_py/gpudp/dp/detail_placement.cpp index bc01ac5..6c91707 100644 --- a/cpp_to_py/gpudp/dp/detail_placement.cpp +++ b/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) { - globalSwapCUDA(at_db, num_bins_x, num_bins_y, batch_size, 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, displacement_region_ratio); } void independentSetMatching( diff --git a/cpp_to_py/gpudp/dp/global_swap_cuda.cu b/cpp_to_py/gpudp/dp/global_swap_cuda.cu index bb36382..39d29c8 100644 --- a/cpp_to_py/gpudp/dp/global_swap_cuda.cu +++ b/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 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; diff --git a/src/detail_placement.py b/src/detail_placement.py index b6ac062..87941fa 100644 --- a/src/detail_placement.py +++ b/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()