[BUG] Reduction core-dump with transform iterator using device only lambda.

[aschaffer]

thrust::reduce() crashes with a transform_iterator defined in terms of a device-only lambda (i.e., qualified only by __device__). The exception is:

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  after reduction step 2: cudaErrorInvalidDeviceFunction: invalid device function
Aborted (core dumped)

Example:

auto dlambda_sq = [] __device__ (T x) {
    return x * x;
  };

  auto t0_iter_begin = thrust::make_transform_iterator(d_v.begin(), dlambda_sq);
  auto t0_iter_end = thrust::make_transform_iterator(d_v.end(), dlambda_sq);

  //crashes w/ device lambda:
  //
  auto res0 = thrust::reduce(t0_iter_begin, t0_iter_end);//crashes...

If dlambda_sq above is used in a thrust::transform() the call is successful; or, if the lambda is qualified by __host__ __device__, instead of just __device__, again the test passes.

[aschaffer]

Repro here (internal): https://nvidia.slack.com/files/U9TQTPHJR/F01979G9KHC/reduce_transform_it_cu.cpp

I've repro the bug with 10.2 and 11.0 CUDA toolkits.

[brycelelbach]

@jaredhoberock indicated this may be because of how result_of works with device lambdas. Jared, could you elaborate?

[harrism]

Is there any hope of a fix for this?

[alliepiper]

Probably not. This came up on #cdd-thrust about a month ago, and @jaredhoberock pointed out this list of extended lambda restrictions here.

The tl;dr is that __device__ lambdas are replaced by a placeholder type in host code, and the placeholder does not behave as expected when queried with metaprogramming techniques, which thrust relies on.

Using a __host__ __device__ lambda avoids some of the issues in some cases and is a reasonable workaround when it works, but for a pure __device__ function, a fully defined functor is usually necessary.

Closing as wontfix since this is a restriction of the language and beyond Thrust's ability to fix.

[alliepiper]

Actually, it looks like there's another discussion of this in #779 that proposes some ideas to triage the placeholder problem. I'll remove wontfix.

Follow that issue to track this, but for now use __host__ __device__ lambda or explicit functors to work around the issue.