[BUG]: maybe add documentation about mixing thrust::device_vectors with different allocators
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Type of Bug
Silent Failure
### Component
Thrust
### Describe the bug
Consider this snippet:
```
thrust::device_vector> vec3(1024, allocator);
vec3 = thrust::device_vector();
```
My goal was to deallocate the memory of vec3.
In the assignment, the memory of vec3 is deallocated using the default allocator, not the custom allocator.
This led to a problem in a program where memory was not properly returned to a memory pool.
### How to Reproduce
https://godbolt.org/z/39qbMEM1z
```
#include
#include
#include
template
struct custom_allocator : public thrust::device_malloc_allocator {
public:
using Base = thrust::device_malloc_allocator;
using pointer = typename Base::pointer;
using size_type = typename Base::size_type;
template
struct rebind {
using other = custom_allocator;
};
custom_allocator() = default;
template
custom_allocator(custom_allocator other) {}
pointer allocate(size_type num){
std::cout << "allocate async " << num << "\n";
T* result = nullptr;
cudaError_t status = cudaMallocAsync(&result, sizeof(T) * num, stream);
if(status != cudaSuccess){
throw std::runtime_error("custom_allocator error allocate");
}
return thrust::device_pointer_cast(result);
}
void deallocate(pointer ptr, size_type num){
std::cout << "deallocate async " << num << "\n";
cudaError_t status = cudaFreeAsync(thrust::raw_pointer_cast(ptr), stream);
if(status != cudaSuccess){
throw std::runtime_error("custom_allocator error deallocate");
}
}
cudaStream_t stream = cudaStreamPerThread;
};
int main(){
cudaSetDevice(0);
custom_allocator allocator;
//default alloc <- default alloc
std::cout << "AAA\n";
thrust::device_vector vec1(1024);
vec1 = thrust::device_vector();
std::cout << "AAA\n";
//custom alloc <- custom alloc
std::cout << "BBB\n";
thrust::device_vector> vec2(1024, allocator);
vec2 = thrust::device_vector>(allocator);
std::cout << "BBB\n";
//custom alloc <- default alloc
std::cout << "CCC\n";
thrust::device_vector> vec3(1024, allocator);
vec3 = thrust::device_vector();
std::cout << "CCC\n";
//default alloc <- custom alloc
std::cout << "DDD\n";
thrust::device_vector vec4(1024);
vec4 = thrust::device_vector>(allocator);
std::cout << "DDD\n";
}
```
Will output
```
AAA
AAA
BBB
allocate async 1024
deallocate async 1024
BBB
CCC
allocate async 1024
CCC
DDD
DDD
deallocate async 1024
```
### Expected behavior
I expected that memory will always be deallocated by the allocator used for allocation.
Expected output
```
AAA
AAA
BBB
allocate async 1024
deallocate async 1024
BBB
CCC
allocate async 1024
deallocate async 1024
CCC
DDD
DDD
```
### Reproduction link
https://godbolt.org/z/39qbMEM1z
### Operating System
_No response_
### nvidia-smi output
_No response_
### NVCC version
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.