Fix assignments to uninitialized memory

[hwinkler]

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):

#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.