Reduce binary size random select primitive#5574
Conversation
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
ChuckHastings
added
improvement
|
I am still thinking about the best approach, but I am thinking more in the line of splitting I guess the edge sampling part varies a lot for different sampling methods while the follow up transform and gather part is largely same (IIRC, this is the main motivation for creating Let me think more about this. |
seunghwak
left a comment
seunghwak
left a comment
There was a problem hiding this comment.
I am still reviewing this and thinking more about the best approach but few quick early comments.
| void compute_heterogeneous_biased_sampling_index_without_replacement( | ||
| raft::handle_t const& handle, | ||
| std::optional<raft::device_span<size_t const>> | ||
| input_frontier_indices, // input_per_tyep_degree_offsets & input_biases are already packed if |
| } | ||
| CUGRAPH_EXPECTS(num_invalid_biases == 0, | ||
| "invalid_input_argument: bias_e_op return values should be non-negative and " | ||
| "should not exceed std::numeirc_limits<bias_t>::max()."); |
| assert(graph_view.number_of_local_edge_partitions() == minor_comm_size); | ||
|
|
||
| if (do_expensive_check) { | ||
| // FIXME: better re-factor this check function? |
There was a problem hiding this comment.
Yeah... and this is an easy target, not sure how big will be the gain, but we have
count_invalid_vertex_pairs
https://github.com/rapidsai/cugraph/blob/main/cpp/include/cugraph/utilities/error_check_utils.cuh#L59
We may create
count_invalid_vertices as well.
And explicitly instantiate these two functions, and use them everywhere we check for invalid vertices or invalid vertex pairs.
| raft::host_span<size_t const>(local_key_list_sizes.data(), local_key_list_sizes.size()), | ||
| raft::host_span<size_t const>(local_key_list_offsets.data(), | ||
| local_key_list_offsets.size() - 1), | ||
| handle.get_stream()); |
There was a problem hiding this comment.
And another thing to experiment is that how much we can save by explicitly instantiating device_comm functions (device_allgatherv, device_bcast, and so on).
Maybe not worth the effort as most heavy lifting happens inside the NCCL functions, so the saving might be negligible but we might be able to quickly experiment (e.g. just delete few device comm calls and see whether we see any meaningful change in the resulting binary size).
There was a problem hiding this comment.
If this turns out to be non-negligible, we can do something similar to what we are doing for thrust algorithms.
| } | ||
| auto edge_partition_e_value_input = edge_partition_e_input_device_view_t(edge_value_input, i); | ||
|
|
||
| if (sample_key_indices) { |
There was a problem hiding this comment.
Need to double check but sample_key_indices.has_value() is true only when minor_comm_size > 1.
If we accept some overhead when GraphViewType::is_multi_gpu is true and minor_comm_size == 1, we may create std::conditional_t<GraphViewType::is_multi_gpu, rmm::device_uvector<size_t>, std::byte /* dummy */> sample_key_indices and here we can use if constexpr (GraphViewType::is_multi_gpu) instead of if. So, only one of the two transform calls will be instantiated. Try to delete the if pass or the else pass and see how much this reduces the binary size. If the difference is non-negligible, this might be another place to cut binary size.
There was a problem hiding this comment.
I didn't look extensively through the code, but I did add some constexpr blocks in a few places where we were just checking minor_comm_size so that in SG code we wouldn't even compile the block. I'm sure there are other examples like this, I can take a look.
|
A data point to consider in this work.
Via random_walks_impl.cuh
Via sample_edges.cuh
Merged these files to four, SG/v32/e32, SG/v64/e64, MG/v32/e32, MG/v64/e64. After fixing a minor build error ( -rwxr-xr-x 1 seunghwak domain-users 21814664 Jul 8 11:16 libcugraph_common.so We may re-structure the primitive sub-functions little more to further facilitate a function re-use within a TU (e.g. create functions with only edge_t template type). I think we might be able to get similar 10% cut in this way without sacrificing primitives' capability to support arbitrary type combinations. |
|
I thought more about splitting The main motivation for this was
But it seems like the build overhead of the follow-up transform part is not very significant and removing redundancy requires additional plumbing code largely offsetting the code size reduction while also introducing some performance overhead. More details here.
This works in the following steps.
1 & 2 can go to In step 3,
They will be compiled only once if template parameters are same (mainly In step 4, we are mainly using cugraph thrust wrappers (except for one thrust::transform and one thrust::tabulate) or Steps 3 & 4 look cheap to build, so splitting may not be very helpful in build time by increasing parallelism (this is most effective when steps 1&2 and 3&4 took comparable amount of time, but it seems like steps 1&2 will be much more expensive than 3&4). Complication part. We are not really gathering edges here but gathering key frontier (major vertex can be tagged) and their sampled outgoing edges. The key distinction here is that the edge major can be tagged. This is not a simple gathering of edges. This can cause several complications. Another source of complication is that So, I am leaning towards not splitting the primitive for now and focusing more on building in a single TU, use if constexpr if it helps cutting binary size/build time, refactor some sub-routines if it can prevent multiple explicit instantiations within a single TU. |
3 participants
The random select primitive is our most complex primitive, and many portions of it do not vary based on the graph, the bias operator or the edge operator. This PR splits some of these common pieces out into separate helper functions that can be compiled once for each variation that occurs rather than being compiled redundantly with the combinatorial explosion of template parameters of the primitive.
On CUDA 12 (the CUDA version driving our work), here is the savings:
libcugraph_common.solibcugraph.solibcugraph_mg.so