Consider always discarding __device__ lambdas' results

[jaredhoberock]

Host functions can't rely on the result of std::result_of when it is instantiated with a __device__ lambda.

If we have the means to detect their presence at compile time, then one option is to entirely ban them from bulk_invoke and friends. A more relaxed policy might be to always interpret their result to be void. That would still allow them to be used, but it would not be possible to return results from them.

[brycelelbach]

I don't think there's a way to detect that it's a lambda.

I'm not sure I understand why result_of doesn't work in this instance?

Would decltype work?

[jaredhoberock]

The issue is that __device__ lambdas' always return int! Even if there is no return used inside at all. Ugh! That means utilities like result_of or invoke_result or decltype to can't reliably determine what __device__ lambdas return. :-(

Fortunately, we are able to distinguish __device__ lambdas from other types.

Agency uses this inside its implementation of result_of to treat all __device__ lambdas as if they return void.

[alliepiper]

I'm going to start using this issue to track the various "__device__ lambda aren't working" issues.

The issues Jared mentioned above are documented here in the current toolkit documentation.

Until we have a better system in place, users should stick with __host__ __device__ lambda or explicit __device__ functor objects instead.

[jrhemstad]

Just for completeness, the exact restriction is G.6.2p13:

As described above, the CUDA compiler replaces a device extended lambda defined in a host function with a placeholder type defined in namespace scope. This placeholder type does not define a operator() function equivalent to the original lambda declaration. An attempt to determine the return type or parameter types of the operator() function may therefore work incorrectly in host code, as the code processed by the host compiler will be semantically different than the input code processed by the CUDA compiler.

I've stopped using __device__ lambdas all together with Thrust after I learned about this restriction. Too great a chance of silent or non-obvious failures.

[harrism]

@jaredhoberock said "device lambdas always return int". That can't be true, I have used device lambdas that return non-int before. I guess what Jared means is that result_of on a device lambda always evaluates to int?

[alliepiper]

Correct. The device lambda object is replaced with a placeholder object in host code, and that placeholder's API does not accurately represent the device lambda's actual call signature.

[jrhemstad]

I guess what Jared means is that result_of on a device lambda always evaluates to int

https://godbolt.org/z/6ob7q5

Not necessarily, no. It's moreso that the return type of the callable (as determined from host code) cannot be relied upon to be correct.

[harrism]

I just wanted reassurance that if I call a device lambda that returns float from a kernel, it doesn't return int.

[pauleonix]

I'm going to start using this issue to track the various "__device__ lambda aren't working" issues.

#1650 (comment) has a reproducer for an issue with both __device__ and __host__ __device__ lambdas causing thrust::transform_output_iterator< ... >::operator= to be deleted. Further information in #1650 (comment)

[jrhemstad]

Thanks to NVIDIA/libcudacxx#284, we will automatically error at compile time in instances where we attempt to query the return type of an extended lambda in a context where that is not supported.