Bug 241550 - Base Clang can't compile trivial CUDA programs: error: no matching function for call to '__isnan'
Summary: Base Clang can't compile trivial CUDA programs: error: no matching function f...
Status: Closed FIXED
Alias: None
Product: Base System
Classification: Unclassified
Component: bin (show other bugs)
Version: CURRENT
Hardware: Any Any
: --- Affects Many People
Assignee: Dimitry Andric
URL: https://reviews.llvm.org/D60220#1723350
Keywords:
Depends on: 241611
Blocks:
  Show dependency treegraph
 
Reported: 2019-10-28 17:52 UTC by Gleb Popov
Modified: 2019-11-10 19:08 UTC (History)
6 users (show)

See Also:
koobs: mfc-stable11?
koobs: mfc-stable12?


Attachments

Note You need to log in before you can comment on or make changes to this bug.
Description Gleb Popov freebsd_committer 2019-10-28 17:52:33 UTC
I've obtain CUDA headers from NVIDIA distribution [1] and tried using Clang to compile simple program. However, I got errors from CUDA's internal headers saying:

In file included from <built-in>:1:
In file included from /usr/lib/clang/9.0.0/include/__clang_cuda_runtime_wrapper.h:204:
/usr/home/arr/cuda/var/cuda-repo-10-0-local-10.0.130-410.48/usr/local/cuda-10.0//include/crt/math_functions.hpp:2910:7: error: no matching function for call to '__isnan'
  if (__isnan(a)) {
      ^~~~~~~
/usr/lib/clang/9.0.0/include/__clang_cuda_device_functions.h:460:16: note: candidate function not viable: call to __device__ function from __host__ function
__DEVICE__ int __isnan(double __a) { return __nv_isnand(__a); }
               ^

I searched the net a bit and found following upstream commit [2]. Asking there, I was told that it due, citing,

>  FreeBSD does not provide *host*-side __isnan(double) -- the error complains that it's the host code that tried to use __isnan and failed when overload resolution produced a device variant.

Can anything be done about this? Sorry if this PR is not relevant for FreeBSD.

[1] http://developer.download.nvidia.com/compute/cuda/10.1/Prod/local_installers/cuda-repo-rhel7-10-1-local-10.1.243-418.87.00-1.0-1.x86_64.rpm
[2] https://reviews.llvm.org/D60220#1723350
Comment 1 Kubilay Kocak freebsd_committer freebsd_triage 2019-10-29 00:58:13 UTC
From upstream review [1]

"I think the right thing to do here would probably be to define a wrapper __isnan(double) which would call it."

[1] https://reviews.llvm.org/D60220#1723350
Comment 2 Dimitry Andric freebsd_committer 2019-10-29 18:05:48 UTC
It's unfortunate that the cuda headers use these __isnan functions, which seem to be glibc specific.

That said, I see the error is not in the clang-supplied cuda wrapper headers, but in /usr/home/arr/cuda/var/cuda-repo-10-0-local-10.0.130-410.48/usr/local/cuda-10.0//include/crt/math_functions.hpp.  So this is some specific thing that is (maybe) being used all of the place in CUDA code?

We do indeed have an __inline_isnan, but it is currently only used for the __fp_type_select() macro in math.h.  The name originates in base r253215 by David, which says "Cleanups to math.h that prevent namespace conflicts with C++".

It looks like the name used to be __isnan before that commit, but why it was changed is unclear.  There is no __isnan in the C++ standards, as far as I can see.
Comment 3 Dimitry Andric freebsd_committer 2019-10-29 18:40:07 UTC
Ah, I see we also have a different implementation of __isnan and __isnanf (but not __isnanl) in lib/libc/gen/isnan.c.  That one actually tests IEEE bits, instead of just doing "x != x".

So maybe that is the reason we can't expose __isnan, for backwards compatibility?
Comment 4 Dimitry Andric freebsd_committer 2019-10-29 19:27:02 UTC
Hm even looking at the tree in the state of just before base r253215, I see only three instances of __isnan:

lib/libc/gen/isnan.c: __weak_reference(__isnan, isnan);
lib/libc/gen/isnan.c: __isnan(double d)
lib/libc/gen/Symbol.map:        __isnan;

I'm now testing a universe build with this diff:

Index: lib/msun/src/math.h
===================================================================
--- lib/msun/src/math.h (revision 354146)
+++ lib/msun/src/math.h (working copy)
@@ -111,7 +111,7 @@ extern const union __nan_un {
 #define        isfinite(x) __fp_type_select(x, __isfinitef, __isfinite, __isfinitel)
 #define        isinf(x) __fp_type_select(x, __isinff, __isinf, __isinfl)
 #define        isnan(x) \
-       __fp_type_select(x, __inline_isnanf, __inline_isnan, __inline_isnanl)
+       __fp_type_select(x, __isnanf, __isnan, __isnanl)
 #define        isnormal(x) __fp_type_select(x, __isnormalf, __isnormal, __isnormall)

 #ifdef __MATH_BUILTIN_RELOPS
@@ -194,7 +194,7 @@ int __signbitf(float) __pure2;
 int    __signbitl(long double) __pure2;

 static __inline int
-__inline_isnan(__const double __x)
+__isnan(__const double __x)
 {

        return (__x != __x);
@@ -201,7 +201,7 @@ static __inline int
 }

 static __inline int
-__inline_isnanf(__const float __x)
+__isnanf(__const float __x)
 {

        return (__x != __x);
@@ -208,7 +208,7 @@ static __inline int
 }

 static __inline int
-__inline_isnanl(__const long double __x)
+__isnanl(__const long double __x)
 {

        return (__x != __x);


If that goes well, maybe it's best to do an exp-run too?  It's more likely that some ports fall over due to this change, than anything in our base system...
Comment 5 Dimitry Andric freebsd_committer 2019-10-29 21:54:02 UTC
(In reply to Dimitry Andric from comment #4)

Hmm, and that obviously doesn't build:

--- isnan.pico ---
/home/dim/src/head/lib/libc/gen/isnan.c:50:1: error: redefinition of '__isnan'
__isnan(double d)
^
/home/dim/src/head/lib/msun/src/math.h:197:1: note: previous definition is here
__isnan(__const double __x)
^
/home/dim/src/head/lib/libc/gen/isnan.c:59:1: error: redefinition of '__isnanf'
__isnanf(float f)
^
/home/dim/src/head/lib/msun/src/math.h:204:1: note: previous definition is here
__isnanf(__const float __x)
^
2 errors generated.

I'm unsure what a good way is to work around it...

Maybe it's best to just add wrapper __isnan() and __isnanf() to the CUDA specific headers, since they are only used in that particular case.  But that would only work for our own copy of clang, not for the ports version.

Or use some special hack to not define __isnan() and __isnanf() in math.h if libc's isnan.c is compiled. :)
Comment 6 Dimitry Andric freebsd_committer 2019-10-31 07:07:50 UTC
I have put up a different but working patch in bug 241611, where I requested a ports exp-run with it, since past experiences have shown that changes to these prototypes and macros is a minefield. :)
Comment 7 commit-hook freebsd_committer 2019-11-02 17:00:38 UTC
A commit references this bug:

Author: dim
Date: Sat Nov  2 16:59:54 UTC 2019
New revision: 354255
URL: https://svnweb.freebsd.org/changeset/base/354255

Log:
  Add __isnan()/__isnanf() aliases for compatibility with glibc and CUDA

  Even though clang comes with a number of internal CUDA wrapper headers,
  compiling sample CUDA programs will result in errors similar to:

  In file included from <built-in>:1:
  In file included from /usr/lib/clang/9.0.0/include/__clang_cuda_runtime_wrapper.h:204:
  /usr/home/arr/cuda/var/cuda-repo-10-0-local-10.0.130-410.48/usr/local/cuda-10.0//include/crt/math_functions.hpp:2910:7: error: no matching function for call to '__isnan'
    if (__isnan(a)) {
        ^~~~~~~
  /usr/lib/clang/9.0.0/include/__clang_cuda_device_functions.h:460:16: note: candidate function not viable: call to __device__ function from __host__ function
  __DEVICE__ int __isnan(double __a) { return __nv_isnand(__a); }
                 ^

  CUDA expects __isnan() and __isnanf() declarations to be available,
  which are glibc specific extensions, equivalent to the regular isnan()
  and isnanf().

  To provide these, define __isnan() and __isnanf() as aliases of the
  already existing static inline functions __inline_isnan() and
  __inline_isnanf() from math.h.

  Reported by:	arrowd
  PR:		241550
  MFC after:	1 week

Changes:
  head/lib/libc/gen/isnan.c
  head/lib/msun/src/math.h
  head/lib/msun/src/s_isnan.c
Comment 8 Gleb Popov freebsd_committer 2019-11-05 11:59:39 UTC
This change made the problem go away, but compilation still fails due to other errors.
Comment 9 commit-hook freebsd_committer 2019-11-10 17:34:11 UTC
A commit references this bug:

Author: dim
Date: Sun Nov 10 17:33:12 UTC 2019
New revision: 354596
URL: https://svnweb.freebsd.org/changeset/base/354596

Log:
  MFC r354255:

  Add __isnan()/__isnanf() aliases for compatibility with glibc and CUDA

  Even though clang comes with a number of internal CUDA wrapper headers,
  compiling sample CUDA programs will result in errors similar to:

  In file included from <built-in>:1:
  In file included from /usr/lib/clang/9.0.0/include/__clang_cuda_runtime_wrapper.h:204:
  /usr/home/arr/cuda/var/cuda-repo-10-0-local-10.0.130-410.48/usr/local/cuda-10.0//include/crt/math_functions.hpp:2910:7: error: no matching function for call to '__isnan'
    if (__isnan(a)) {
        ^~~~~~~
  /usr/lib/clang/9.0.0/include/__clang_cuda_device_functions.h:460:16: note: candidate function not viable: call to __device__ function from __host__ function
  __DEVICE__ int __isnan(double __a) { return __nv_isnand(__a); }
                 ^

  CUDA expects __isnan() and __isnanf() declarations to be available,
  which are glibc specific extensions, equivalent to the regular isnan()
  and isnanf().

  To provide these, define __isnan() and __isnanf() as aliases of the
  already existing static inline functions __inline_isnan() and
  __inline_isnanf() from math.h.

  Reported by:	arrowd
  PR:		241550

Changes:
_U  stable/10/
  stable/10/lib/libc/gen/isnan.c
  stable/10/lib/msun/src/math.h
  stable/10/lib/msun/src/s_isnan.c
_U  stable/11/
  stable/11/lib/libc/gen/isnan.c
  stable/11/lib/msun/src/math.h
  stable/11/lib/msun/src/s_isnan.c
_U  stable/12/
  stable/12/lib/libc/gen/isnan.c
  stable/12/lib/msun/src/math.h
  stable/12/lib/msun/src/s_isnan.c
Comment 10 Dimitry Andric freebsd_committer 2019-11-10 18:43:44 UTC
Gleb, do you want to keep this bug open for attempting to get any form of CUDA working on FreeBSD?
Comment 11 Gleb Popov freebsd_committer 2019-11-10 18:45:58 UTC
(In reply to Dimitry Andric from comment #10)

I'm slowly working on that, but lets close this bug. I'll just create a new one, when I need something from toolchain team again.

Thanks for working on this!