NVIDIA / NVIDIA/cccl

Fix assignments to uninitialized memory

Open
#860 3 comments 0 reactions 1 assignee Claimed by @miscco View on GitHub
cub
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

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.