Parameterize `thrust::*_vector` Resizing
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
We've had a request to change the resizing strategy for `thrust::*_vector`. Currently, when we need to reallocate, we double the size of storage. Some users are asking for us to switch to a smaller factor, like 1.5.
I doubt that we can change this globally for everyone, so I'm inclined to suggest that we make this parameterizable. We'd do this by making it a property of the allocator that the vector is using; for example, if the allocator has an `allocator::recommend_new_size` method, we'd call that.
Some background, from emails:
> As part of getting feedback from different DOE projects, I noticed this
> on a list from one of the AMReX/LBL developers:
>
> > A complaint [about] thrust vector. The growth factor when resizing is 2. This results in a lot of waste of memory. Sometimes the size of vector is 1GB before the resize. Adding one number to the container ends up with a vector with a capacity of 2GB. Given the limited amount of GPU memory, this is very annoying. We wish the growth factor is adjustable or smaller like Facebook's vector.
>
> I have not validated the behavior, nor do I know if they've already said
> something to you directly, but just in case, I'm passing this along.
> can you re-articulate the issues you've encountered with thrust's vector resizing within AMReX on Volta and perhaps create a small example of how memory usage grows/shrinks in practice?
>
> Finally, are you wanting ...
> - exponential 2x growth to 1GB followed by linear growth (e.g. 0.5GB increments)
> - (slower) exponential (1.4x) growth
> The main use of thrust vector in AMReX is for particles. These
> vectors are usually very big because of the number of particles and
> the number of attributes particles have. For the WarpX code using
> AMReX, most of the memory is used by particles. Particles move from
> one process to another. So the sizes of particle containers change.
> Memory allocation is very expensive for gpu. So AMReX has a number of
> memory arenas and they are used by the particle containers. When we
> run the WarpX code on summit, we find that a lot of memory is wasted
> because of the memory arena and the way how vector resize works.
> Suppose at some point a vector's size and capacity are 1GB and a new
> element is pushed to the back. Then suddenly the capacity becomes
> 2GB, whereas the size of 1+epsilon GB. This gets even worse when we
> have multiple vectors and use a memory arena. Suppose we have 4
> vectors that all have the capacity and size of 1GB. We add one
> element to each of the vectors. For each vector, it is resized and
> its capacity becomes 2GB. Then its original 1GB is released back to
> the memory arena. There will be 4 chunks of 1GB memory released back
> to the arena. If we could not coalesce them into a big chunk, we will
> end up with 12GB of memory allocated from the system although the data
> only require a tiny bit more than 4GB. We have found that, for CUDA
> memory (at least for unified memory), memory blocks allocated from
> different cudaMallocs cannot be coalesced even if their addresses are
> contiguous. Because of that, we have modified AMReX's arena. Now we
> allocate a large chunk of memory say 12GB in one cudaMalloc upfront at
> the beginning and put it in our memory arena. This helps a lot. But
> it is still the case the memory used by a vector could suddenly
> increase be factor of 2 when adding only one more element.
>
> As for the two strategies you mentioned, I don't know which way is
> better. Both gcc, llvm and thrust simply use 2. Facebook vector uses
> 1.5 as growth rate and they argues 2 is a very bad choice.
>
> https://github.com/facebook/folly/blob/master/folly/docs/FBVector.md
>
> I don't know what kind of examples you like to see. I am cc'ing
> Andrew. He is our main developer on particle codes.
Contributor guide
Assessment
This issue has not been assessed yet.