NVIDIA / NVIDIA/cccl

set_intersection improperly compares keys not in the input

Open
#822 6 comments 0 reactions 0 assignees View on GitHub
thrust
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

Bug in `set_intersection` using keys that are `struct`s, only for THRUST_DEVICE_SYSTEM_CUDA.
(It works fine in THRUST_DEVICE_SYSTEM_CPP).
The `assert` in Key::operator<() will fail.
That means one of the two Keys it is comparing has been constructed with the default constructor.
But this example ensures the input keys it passes to `set_intersection` were initialized
to a different value. `set_intersection` ought never compare against any default constructed
Keys.

(I think!)

```c++
#include
#include
#include

// Bug in `set_intersection` using keys that are `struct`s, only for THRUST_DEVICE_SYSTEM_CUDA.
// (It works fine in THRUST_DEVICE_SYSTEM_CPP).
// The `assert` in Key::operator<() will fail.
// That means one of the two Keys it is comparing has been constructed with the default constructor.
// But this example ensures the input keys it passes to `set_intersection` were initialized
// to a different value. `set_intersection` ought never compare against any default constructed
// Keys.

struct Key
{
int some_member;

__host__ __device__ Key(); // some_member <- 0
__host__ __device__ Key(int s); // some_member <- 1

bool __host__ __device__ operator<(const Key& b) const;
};

__host__ __device__ Key::Key()
: some_member(0)
{
}

__host__ __device__ Key::Key(int)
: some_member(1)
{
}

bool __host__ __device__ Key::operator<(const Key& b) const
{
if (this->some_member != b.some_member) {
assert(this->some_member == b.some_member);
printf("some_member mismatch %d %d\n", this->some_member, b.some_member);
}
// all keys are equal for our purposes:
return false;
}

auto doSetIntersectionByKey()
{
// This Key will have `some_member` initialized to 1:
Key hostKey(1);

Key hostKeysA[] = { hostKey, hostKey, hostKey, hostKey };
Key hostKeysB[] = { hostKey, hostKey, hostKey, hostKey };

thrust::device_vector aKey(hostKeysA, hostKeysA + 4);
thrust::device_vector bKey(hostKeysB, hostKeysB + 4);

thrust::device_vector outKeys(4);

// The two input vectors have Keys all of whose `some_member` is 1.
// Yet when compiled for CUDA backend, the assertion in
// operator<() fails. When compiled for CPP, the assertion passes.
// Under CUDA, set_intersection must be comparing to Key instances constructed
// using the default constructor (or else to uninitialized Key instances).
return thrust::set_intersection(
aKey.begin(), aKey.end(), bKey.begin(), bKey.end(), outKeys.begin());
}

int main()
{
// Dereferencing just to force optimized builds not to eliminate code.
Key x = *doSetIntersectionByKey();
return (x.some_member) ? 0 : 1;
}
```

to build:
```sh
echo Building repro-cuda-debug
/usr/local/cuda-11.7/bin/nvcc \
-DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA \
-g -G -O0 \
-std=c++17 \
-arch=sm_70 -gencode=arch=compute_86,code=sm_86 \
-x cu \
-o "repro-cuda-debug" \
"set_intersection_repro.cu"
```

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.