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
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
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.
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?
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...
(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. :)
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. :)
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
This change made the problem go away, but compilation still fails due to other errors.
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
Gleb, do you want to keep this bug open for attempting to get any form of CUDA working on FreeBSD?
(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!
^Triage: Track merges to stable branches