reduce_by_key passing invalid data to BinaryFunction operator

[davidwendt]

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 <thrust/device_vector.h>
#include <thrust/execution_policy.h>
#include <thrust/fill.h>
#include <thrust/reduce.h>
#include <thrust/count.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/constant_iterator.h>

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<int> input(count);
    thrust::fill(thrust::device, input.begin(), input.end(), 8);

    thrust::device_vector<int> output(count);
    auto begin = thrust::make_constant_iterator<int>(0);
    thrust::reduce_by_key(thrust::device, begin, begin + count,
                          input.begin(), thrust::make_discard_iterator(),
                          output.begin(), thrust::equal_to<int>{},
                          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 <thrust/device_vector.h>
#include <thrust/execution_policy.h>
#include <thrust/fill.h>
#include <thrust/reduce.h>
#include <thrust/count.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/constant_iterator.h>

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<int> adata(1, 8);
  auto d_adata = adata.data().get();
  // initialize a vector of identical mydata objects into device memory
  thrust::device_vector<mydata> input(count);
  thrust::fill(thrust::device, input.begin(), input.end(), mydata{d_adata, 8});

  thrust::device_vector<mydata> output(count);

  auto begin = thrust::make_constant_iterator<int>(1);
  thrust::reduce_by_key(thrust::device, begin, begin + count,
                        input.begin(), thrust::make_discard_iterator(),
                        output.begin(), thrust::equal_to<int>{},
                        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 #1479 where invalid data is passed to the AssociativeOperator parameter in thrust::inclusive_scan