NVIDIA / NVIDIA/cccl

reduce_by_key passing invalid data to BinaryFunction operator

Open
#789 0 comments 0 reactions 1 assignee Assigned to @elstehle View on GitHub
thrust
Dominant language
C++
Stars
2.5k
Forks
487
Avg merge
2d 7h
Merged PRs (30d)
296

Description

The `thrust::reduce_by_key` function passes random data not included in the input vector/iterators to the BinaryFunction parameter as defined here: https://thrust.github.io/doc/group__reductions_gaafd6f34b72f1ea30d58ba916d53ee754.html#gaafd6f34b72f1ea30d58ba916d53ee754

The result from the binary-operator when given the invalid data is not actually used by the `reduce_by_key` so most simple binary operators will not actually fail. The following code illustrates the issue:
```
#include
#include
#include
#include
#include
#include
#include

struct min_op
{
__device__ int operator()(int const &lhs, int const &rhs)
{
if (lhs != 8 || rhs != 8)
printf("(%d,%d)\n", lhs, rhs); // this should not happen
return lhs < rhs ? lhs : rhs;
}
};

int main(int argc, const char **argv)
{
int count = argc > 1 ? std::atoi(argv[1]) : 2305; // 2304 and below work ok
thrust::device_vector input(count);
thrust::fill(thrust::device, input.begin(), input.end(), 8);

thrust::device_vector output(count);
auto begin = thrust::make_constant_iterator(0);
thrust::reduce_by_key(thrust::device, begin, begin + count,
input.begin(), thrust::make_discard_iterator(),
output.begin(), thrust::equal_to{},
min_op{});

auto result = thrust::count_if(thrust::device, output.begin(), output.end(),
[] __device__(auto const &val) { return val != 8; });
cudaStreamSynchronize(0);
printf("result = %d\n", (int)result);
return 0;
}
```

The `count` value must be at least 2304 to see the issue.

The `input` vector used for the `values_first` parameter is initialized to all 8s. When `reduce_by_key` is called, the `printf` function is only called if a passed in value is not `8`. The `result` is always correct so that is evidence the value returned by `min_op::operator()` when an invalid (`!=8`) value is passed is not actually being used.

However, passing invalid data of a more complex type may cause the BinaryFunction to crash. This is illustrated in the code example below. The values data type `mydata` contains a pointer to device memory.
```
#include
#include
#include
#include
#include
#include
#include

struct mydata
{
int const *a{};
int b{};

mydata() = default;
~mydata() = default;
mydata(const mydata &) = default;
mydata(mydata &&) = default;
mydata &operator=(const mydata &) = default;
mydata &operator=(mydata &&) = default;

__host__ __device__ mydata(int *a, int b) : a(a), b(b) {}

__device__ bool operator<(const mydata &rhs) const
{
// accessing '*a' or '*(rhs.a)' may crash here
return b < rhs.b;
}
};

struct min_op
{
__device__ mydata operator()(mydata const &lhs, mydata const &rhs)
{
if (lhs.b != 8 || rhs.b != 8)
printf("(%p,%d),(%p,%d)\n", lhs.a, lhs.b, rhs.a, rhs.b); // this should not happen
return lhs < rhs ? lhs : rhs;
}
};

int main(int argc, const char **argv)
{
int count = argc > 1 ? std::atoi(argv[1]) : 1025; // 1024 and below work ok

thrust::device_vector adata(1, 8);
auto d_adata = adata.data().get();
// initialize a vector of identical mydata objects into device memory
thrust::device_vector input(count);
thrust::fill(thrust::device, input.begin(), input.end(), mydata{d_adata, 8});

thrust::device_vector output(count);

auto begin = thrust::make_constant_iterator(1);
thrust::reduce_by_key(thrust::device, begin, begin + count,
input.begin(), thrust::make_discard_iterator(),
output.begin(), thrust::equal_to{},
min_op{});

auto result = thrust::count_if(thrust::device, output.begin(), output.end(),
[d_adata] __device__(mydata const &v) { return v.a != d_adata || v.b != 8; });
cudaStreamSynchronize(0);
printf("result = %d\n", (int)result);
return 0;
}

```
The `count` value must be 1025 or greater to see the issue.

The `reduce_by_key` function arbitrarily builds a `mydata` instance with invalid data. Again, the `printf` line will only be called if the invalid data is detected. In this simple case, this is usually a `nullptr` for member `a` and `0` for member `b` but in our larger https://github.com/rapidsai/cudf application, some objects passed contained invalid, non-null pointers.

This issue was reproduced on nvcc 11.2 with thrust 1.12 and nvcc 11.4 with the latest thrust code from github at this writing.

If temporary objects are required for the internal logic and the result is ignored anyway, it would be great if the temporary objects were created from one of the valid input values.

This issue appears to be similar to NVIDIA/cccl#780 where invalid data is passed to the AssociativeOperator parameter in `thrust::inclusive_scan`

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.