NVIDIA / NVIDIA/cccl

[BUG]: maybe add documentation about mixing thrust::device_vectors with different allocators

Open
#4,623 2 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.