inclusive_scan passing invalid data to AssociativeOperator
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
The `thrust::inclusive_scan` passes random data not included in the input vector/iterators to the AssociateOperator functor parameter. The following code illustrates the issue.
```
#include
#include
#include
#include
#include
int main(int argc, const char **argv)
{
int count = argc > 1 ? std::atoi(argv[1]) : 1537; // 1536 and below work ok
thrust::device_vector input(count);
thrust::fill(thrust::device, input.begin(), input.end(), 8);
thrust::device_vector output(count);
thrust::inclusive_scan(thrust::device, input.begin(), input.end(), output.begin(),
[] __device__(auto const &lhs, auto const &rhs) {
if (lhs != 8 || rhs != 8)
printf("(%d,%d)\n", lhs, rhs);
return lhs < rhs ? lhs : rhs;
});
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 1537 to see the issue.
The `input` vector is initialized to all 8s. When `inclusive_scan` is called, the `printf` will only fire if a passed in value is not `8`. The result from this call does not appear to be used since the scan result is correct as verified by the `count_if` checking for only 8s in the result.
For a more complicated data element, the AssociativeOperator may crash if invalid data is passed.
For example, the following structure contains a pointer to device memory. The pointers are all valid in the input vector but the AssociativeOperator functor is sometimes passed random data.
```
struct mydata
{
int const *a{};
int b{};
mutable int c{};
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' may crash here
return b < rhs.b;
}
};
int main(int argc, const char **argv)
{
int count = argc > 1 ? std::atoi(argv[1]) : 385; // 384 and below are ok
thrust::device_vector adata(1, 7);
auto d_adata = adata.data().get();
thrust::device_vector input(count);
thrust::fill(thrust::device, input.begin(), input.end(), mydata{d_adata, 8});
thrust::device_vector output(count);
thrust::inclusive_scan(thrust::device, input.begin(), input.end(), output.begin(),
[d_adata] __device__(mydata const &lhs, mydata const &rhs) {
if (lhs.a != d_adata || lhs.b != 8 || rhs.a != d_adata || rhs.b != 8)
printf("(%p,%d),(%p,%d)\n", lhs.a, lhs.b, rhs.a, rhs.b);
return lhs < rhs ? lhs : rhs;
});
auto result = thrust::count_if(thrust::device, output.begin(), output.end(),
[d_adata] __device__(mydata const val) { return val.a != d_adata || val.b != 8; });
cudaStreamSynchronize(0);
printf("result = %d\n", (int)result);
return 0;
}
```
The `count` value here must be at least 385 to see the issue.
The printf results for the above example show up `nullptrs` being passed but in our larger [RAPIDS/libcudf](https://github.com/rapidsai/cudf) application, some objects passed actually contain invalid, non-null pointers.
The command line compile options are `nvcc -std=c++17 --expt-extended-lambda inclscan.cu -o inclscan`
This issue was reproduced on nvcc 11.0.221, gcc 9.3.0, ubuntu-1804, with the builtin thrust as well as the latest from the main branch downloaded today. I've also reproduced this on nvcc 11.1.105 and nvcc 11.2.142
Contributor guide
Assessment
This issue has not been assessed yet.