Fix assignments to uninitialized memory
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
The C++ assignment operator presumes the left-hand side is in a valid state, but for many places in CUB code, the left-hand side is uninitialized memory. This causes failures for CUB algorithms that manipulate complex objects.
It's generally easy to replace assignment by placement `new` in the cases I've found, e.g. https://github.com/thrust/cub/pull/32. This replacement ought not degrade performance when the template parameter is a primitive.
I discovered the problem implementing a reference counting scheme, a common idiom. Within `operator=`, the object must call `if(this->someMember) this->someMember->release()`. In several CUB routines, `this` points to uninitialized memory and that line crashes.
I ran across this problem diagnosing `thrust::copy_if`, but it's a pervasive CUB issue. See:
- https://github.com/thrust/thrust/issues/1153
- https://github.com/thrust/cub/pull/32
- https://github.com/thrust/cub/pull/34
To repro the problem, replace `struct TestBar` in `test/test_util.h` by the following definition, which sets a `magic` member to a certain value and verifies, during assignment, that the `magic` member has been properly initialized (I'll submit a PR with this test code):
```c++
#define MAGIC 895245
struct TestBar
{
long long x;
int y;
int magic;
// Constructor
__host__ __device__ __forceinline__ TestBar() : x(0), y(0), magic(MAGIC)
{}
// Constructor
__host__ __device__ __forceinline__ TestBar(int b) : x(b), y(b), magic(MAGIC)
{}
// Constructor
__host__ __device__ __forceinline__ TestBar(long long x, int y) : x(x), y(y), magic(MAGIC)
{}
// Assignment operator
__host__ __device__ __forceinline__ TestBar& operator =(const TestBar& that)
{
assert (magic == MAGIC);
x = that.x;
y = that.y;
return *this;
}
// Assignment from int operator
__host__ __device__ __forceinline__ TestBar& operator =(int b)
{
assert (magic == MAGIC);
x = b;
y = b;
return *this;
}
///...elided....
}
```
# Deliverables
- [ ] `[1.13]` Add tests for each device and collective algorithm that use a custom type designed to detect writes to uninitialized memory.
- The above report provides the definition of a class that can be used for this purpose, though it should be extended to also detect when destructors are skipped. This can be done by changing the value of `magic` in the destructor, and then validating that the buffer contains the new value afterwards.
- [ ] `[1.13]` Fix all errors detected by the new tests.
- [ ] `[1.13]` Specifically validate that the reported cases above are fixed.
Contributor guide
Assessment
This issue has not been assessed yet.