[PATCH] D25403: [CUDA] Mark __libcpp_{isnan, isinf, isfinite} as constexpr.

[Hal Finkel via cfe-commits]

hfinkel added a comment.

...

> 
> 
>> I guess I don't understand the motivation for this change.
> 
> To make <complex> work in CUDA, we do the following terrible, awful, horrible, no good thing: ...well, I can just show you the code.  https://github.com/llvm-project/llvm-project/blob/master/clang/lib/Headers/cuda_wrappers/complex
> 
> The significant part here is
> 
>   #pragma clang force_cuda_host_device begin
>   #include_next <complex>
>   #pragma clang force_cuda_host_device end
>    
> 
> This tells clang, everything between the two pragmas is something that we can run on the host (CPU) and device (GPU).  And that works fine for libstdc++.  But for libc++, marking everything inside <complex> as host+device is not enough -- we also need to mark these four functions, which are called from <complex>.

Is this because the functions are in <cmath> instead of in <complex> are you don't want to mark all of <cmath> as host/device?

> We could mark them as __host__ __device__ explicitly, but then we'd need checks for CUDA compilation inside of libc++, and I've been avoiding asking for that.  Instead, marking these functions as constexpr works, because constexpr functions are implicitly host+device in our CUDA implementation.




https://reviews.llvm.org/D25403