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

[Justin Lebar via cfe-commits]

jlebar added a comment.

In https://reviews.llvm.org/D25403#573666, @mclow.lists wrote:

> Yesterday and today is the first time in a while that clang has been seriously broken for more than an hour or so.
>  I'm not inclined to worry about it yet.


Oh, awesome.  That sounds good to me.

The question about build configurations still stands, though -- should that be documented somewhere?  In particular I could not successfully run check-cxx by following the directions I found by googling.  I am volunteering to write the patch if we can agree on what the workflow should be.

>>> I don't see how this can possibly be constexpr.
>>>  it calls isfinite(), which is hoisted from ::isfinite(), which comes from the C library.
>>>  Since C knows nothing about constexpr, we're stuck.
>> 
>> Functions that call non-constexpr things can be marked constexpr; you just can't evaluate them in a constexpr context (as you demonstrated).
>> 
>> All I need is for the function to be marked as constexpr; I do not need the function be constexpr-evaluatable.
> 
> Then what's the point? How can you test if it is correct?

https://reviews.llvm.org/D24971 tests this.  I suppose we could write a test inside libcxx, but it would be kind of ugly.

> 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>.

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