set_intersection improperly compares keys not in the input
- 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
Assessment
This issue has not been assessed yet.