Replace thrust::device_vector with torch::Tensor #17
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
thrust::device_vector owns the underlying CUDA memory, and it allocates/deallocates for every invocation of `solve_cuda_batch`. Since this code is written for Torch, it is more efficient to rely on Torch's CUDA caching allocator, which virtually eliminates the memory allocation/deallocation in TLA code.
mthrok
commented
Jan 15, 2025
| @@ -265,6 +271,7 @@ std::vector<torch::Tensor> batch_linear_assignment_cuda(torch::Tensor cost) { | |||
|
|
|||
| AT_DISPATCH_FLOATING_TYPES(cost.scalar_type(), "solve_cuda_batch", [&] { | |||
| solve_cuda_batch<scalar_t>( | |||
| cost.scalar_type(), | |||
There was a problem hiding this comment.
note: if we initialize u, v, shortestPathCosts before invoking solve_cuda_batch, then we don't have to pass scalar_type as argument. I took this approach because the diff is easier to see.
|
Thank you for your excellent work profiling the code and identifying the bottleneck! I will make a new release at Pypi soon. |
|
The updated version 0.0.3 is on PyPi. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
thrust::device_vectorowns the underlying CUDA memory, and it allocates/deallocates for every invocation ofsolve_cuda_batch.Since this code is written for Torch, it is more efficient to rely on Torch's CUDA caching allocator, which virtually eliminates the memory allocation/deallocation in TLA code.
Before
batch_linear_assignmenttook 7ms (3ms is device query which is addressed in Query device only once #16)The majority of time is
cudaFree. (2.6ms)After
batch_linear_assignmenttakes 0.37 ms.