use overloading with clang to avoid taking a reference to a __global_…

[andrewcorrigan]

…_ function in device code

"So the canonical way of solving this in clang, I think, is to write a host and a device overload of the function in question." #840 (comment)

Since Clang allows overloading, how about making use of that? @jaredhoberock @egaburov @jlebar

[3gx]

You beat me on the idea :) Having came to same idea independently probably tells it is not such a bad one after all.

I'd suggest one slight improvement though. It would be good to fix macro in https://github.com/thrust/thrust/blob/master/thrust/detail/config/compiler.h#L77 to set THRUST_DEVICE_COMPILER to THRUST_DEVICE_COMPILER_CLANG

It appears we use

#if (defined(__clang__) && defined(__CUDA__))

in many places in Thrust, and it would be a nightmare to change if the signature of clang device compiler will change.

@andrewcorrigan can you please adjust https://github.com/thrust/thrust/blob/master/thrust/detail/config/compiler.h#L77 correspondingly, so that we can just type

#if THRUST_DEVICE_COMPILER  == THRUST_DEVICE_COPMILER_CLANG
<clang code>
#else
<default code>
#endif

Thanks

[andrewcorrigan]

working on it.

[3gx]

The change looks good to me. Unit tests pass with nvcc in both separable and whole program compilation mode.

[jlebar]

it would be a nightmare to change if the signature of clang device compiler will change.

We would never dream of such a thing. :)

I'm happy we were able to settle on something we're all happy with.

[3gx]

Can we add a comment saying that clang doesn't support dynamic parallelism, and add assert(0) or something like this to terminate execution should this code path be ever taken at runtime?

[andrewcorrigan]

Is assert(0) ok in __device__ code? How about bulk::detail::terminate() instead? If you want analogous changes in thrust/system/cuda/detail/detail/launch_closure.inl, can that call directly into bulk::detail::terminate() too?

[3gx]

assert(0) works in nvcc, don't know about clang. bulk::detail::terminate() is better I think.

[jlebar]

assert works in clang too, fwiw.

[andrewcorrigan]

I used bulk::detail::terminate() in bulk and assert(0) otherwise. Anything else?

[jaredhoberock]

This looks good to me. There's also terminate_with_message(). If you wanted, you could add a message indicating why the program had to be terminated.

[Artem-B]

#if defined __CUDACC__ on line 66 would still set compiler to THRUST_DEVICE_COMPILER_NVCC and we'll never make it here.

[Artem-B]

Furthermore, if THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_CLANG, then it will trigger static asserts in number of places that expect compiler for CUDA code to be NVCC.
For instance there's a static assert in thrust/system/cuda/detail/for_each.inl:109