[CUDA][Windows] Final fix for bug 38811 (Step 3 of 3)

[emankov]

Summary

Last fix for the clang Bug 38811 "Clang fails to compile with CUDA-9.x on Windows".

[IMPORTANT]
With that last fix, CUDA has just started being compiling by clang on Windows after nearly a year and two clang’s major releases (7 and 8).
As long as the last LLVM release, in which clang was compiling CUDA on Windows successfully, was 6.0.1, this fix and two previous have to be included into upcoming 7.1.0 and 8.0.1 releases.

[How to repro]

clang++.exe -x cuda "c:\ProgramData\NVIDIA Corporation\CUDA Samples\v9.0\0_Simple\simplePrintf\simplePrintf.cu" -I"c:\ProgramData\NVIDIA Corporation\CUDA Samples\v9.0\common\inc" --cuda-gpu-arch=sm_50 --cuda-path="C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0" -L"c:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\lib\x64" -lcudart.lib  -v

[Output]

In file included from C:\GIT\LLVM\trunk-for-submits\llvm-64-release-vs2017-15.9.5\dist\lib\clang\9.0.0\include\__clang_cuda_runtime_wrapper.h:327:
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:390:11: error: no matching function for call to '__isinfl'
  return (__isinfl(a) != 0);
          ^~~~~~~~
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:2662:14: note: candidate function not viable: call to __host__ function from __device__ function
__func__(int __isinfl(long double a))
             ^
In file included from <built-in>:1:
In file included from C:\GIT\LLVM\trunk-for-submits\llvm-64-release-vs2017-15.9.5\dist\lib\clang\9.0.0\include\__clang_cuda_runtime_wrapper.h:327:
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:438:11: error: no matching function for call to '__isnanl'
  return (__isnanl(a) != 0);
          ^~~~~~~~
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:2672:14: note: candidate function not viable: call to __host__ function from __device__ function
__func__(int __isnanl(long double a))
             ^
In file included from <built-in>:1:
In file included from C:\GIT\LLVM\trunk-for-submits\llvm-64-release-vs2017-15.9.5\dist\lib\clang\9.0.0\include\__clang_cuda_runtime_wrapper.h:327:
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:486:11: error: no matching function for call to '__finitel'
  return (__finitel(a) != 0);
          ^~~~~~~~~
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:2652:14: note: candidate function not viable: call to __host__ function from __device__ function
__func__(int __finitel(long double a))
             ^
3 errors generated when compiling for sm_50.

[Solution]
Add missing device functions' declarations and definitions.

[tra]

clang/lib/Headers/__clang_cuda_cmath.h

81–90

CUDA does not support long double on device side, so I would rather see a compilation error when someone attempts to use a long double, as opposed to providing inconsistent support for it.

I wonder if we could just provide declarations for these functions. This should allow math_functions.hpp to compile, but would still prevent any use of long double on device side.

[emankov]

Provide only declarations for missing long double device functions to prevent any use of long double on the device side, because CUDA does not support long double on the device side.

[Testing]
{Windows 10, Ubuntu 16.04.5}/{Visual C++ 2017 15.9.9, gcc+ 5.4.0}/CUDA {8.0, 9.0, 9.1, 9.2, 10.0, 10.1}

[tra]

Thank you for fixing this!

[tra]

One more thing -- perhaps the long double declarations should be put under #ifndef _MSC_VER in all the files to make the change unobservable on non-windows platforms.

Adding a comment why we only have declarations for these functions would also be helpful.

[emankov]

Oooh, sorry, but I've just pushed the fix. But with the following words: "Add missing long double device functions' declarations. Provide only declarations to prevent any use of long double on the device side, because CUDA does not support long double on the device side."

[tra]

In D60220#1456430, @emankov wrote:

Oooh, sorry, but I've just pushed the fix. But with the following words: "Add missing long double device functions' declarations. Provide only declarations to prevent any use of long double on the device side, because CUDA does not support long double on the device side."

It's not a big deal at the moment -- there are no long double users in CUDA on linux yet. You can clean up in another commit.
BTW, you may want to make commit description somewhat more concise than rC357779. Including all details of reproduction, etc is way too much info which is better suited for the bug report or review.

[emankov]

In D60220#1456449, @tra wrote:

It's not a big deal at the moment -- there are no long double users in CUDA on linux yet. You can clean up in another commit.
BTW, you may want to make commit description somewhat more concise than rC357779. Including all details of reproduction, etc is way too much info which is better suited for the bug report or review.

Ok, thanks, will do it.

[6yearold]

I'm seeing quite similar errors on FreeBSD with Clang 8 and 9:

In file included from <built-in>:1:
In file included from /usr/local/llvm90/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/local/llvm90/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); }
               ^
In file included from <built-in>:1:
In file included from /usr/local/llvm90/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:2934:7: error: no matching function for call to '__isnan'
  if (__isnan(a)) {
      ^~~~~~~
/usr/local/llvm90/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); }
               ^
In file included from <built-in>:1:
In file included from /usr/local/llvm90/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:2960:7: error: no matching function for call to '__isnan'
  if (__isnan(a)) {
      ^~~~~~~
/usr/local/llvm90/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); }
               ^
In file included from <built-in>:1:
In file included from /usr/local/llvm90/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:3090:7: error: no matching function for call to '__isnan'
  if (__isnan(a)) {
      ^~~~~~~
/usr/local/llvm90/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); }
               ^
In file included from <built-in>:1:
In file included from /usr/local/llvm90/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:3196:7: error: no matching function for call to '__isnan'
  if (__isnan(a)) {
      ^~~~~~~
/usr/local/llvm90/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); }
               ^

Any idea how to fix this?

[emankov]

In D60220#1723350, @6yearold wrote:

I'm seeing quite similar errors on FreeBSD with Clang 8 and 9:
Any idea how to fix this?

It looks like CUDA doesn't support double argument for device function __isnan on FreeBSD.

1. I'd look at LLVM trunk (10.0.0svn).
2. If the issue were not eliminated in the trunk, I'd make a change for FreeBSD similar to https://reviews.llvm.org/rL358654 to provide declarations for that function, allowing math_functions.hpp to compile, but with preventing from any use of it on the device side.

[tra]

In D60220#1723458, @emankov wrote:

In D60220#1723350, @6yearold wrote:

I'm seeing quite similar errors on FreeBSD with Clang 8 and 9:
Any idea how to fix this?

It looks like CUDA doesn't support double argument for device function __isnan on FreeBSD.

It's actually the opposite -- 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.

FreeBSD does have host-side implementation of isnan, only it's apparently called __inline_isnan() https://github.com/freebsd/freebsd/blob/master/lib/msun/src/math.h#L197
I think the right thing to do here would probably be to define a wrapper __isnan(double) which would call it.

1. I'd look at LLVM trunk (10.0.0svn).
2. If the issue were not eliminated in the trunk, I'd make a change for FreeBSD similar to https://reviews.llvm.org/rL358654 to provide declarations for that function, allowing math_functions.hpp to compile, but with preventing from any use of it on the device side.

[emankov]

It looks like CUDA doesn't support double argument for device function __isnan on FreeBSD.

It's actually the opposite -- 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.

Sure, you right, quite opposite: call to __device__ function from __host__ function. I agree: the wrapper for the existing __inline_isnan() is the solution.

[dim]

Hm, I would really say that __isnan and the other __ prefixed functions are Linuxisms, or more accurately, glibc-isms. They also don't exist on e.g. macOS:

$ cat check-isnan.cpp
#include <math.h>

int check_isnan(double d)
{
  return ::__isnan(d);
}

$ clang -v
Apple clang version 11.0.0 (clang-1100.0.33.8)
Target: x86_64-apple-darwin18.7.0
Thread model: posix
InstalledDir: /Applications/Xcode.app/Contents/Developer/Toolchains/XcodeDefault.xctoolchain/usr/bin

$ clang -c check-isnan.cpp
check-isnan.cpp:5:12: error: no member named '__isnan' in the global namespace; did you mean 'isnan'?
  return ::__isnan(d);
         ~~^~~~~~~
           isnan
/Applications/Xcode.app/Contents/Developer/Toolchains/XcodeDefault.xctoolchain/usr/bin/../include/c++/v1/math.h:519:1: note: 'isnan' declared here
isnan(_A1 __lcpp_x) _NOEXCEPT
^
1 error generated.

Why can't the regular isnan be used instead? Or is this a CUDA-specific requirement? (Apologies, but I know next to nothing about CUDA :) )

[tra]

In D60220#1725633, @dim wrote:

Hm, I would really say that __isnan and the other __ prefixed functions are Linuxisms, or more accurately, glibc-isms. They also don't exist on e.g. macOS:

Why can't the regular isnan be used instead? Or is this a CUDA-specific requirement? (Apologies, but I know next to nothing about CUDA :) )

Well, that's what *CUDA* headers use, as written by NVIDIA. We have no control over them. :-(

[emankov]

In D60220#1725633, @dim wrote:

$ cat check-isnan.cpp
#include <math.h>

int check_isnan(double d)
{
  return ::__isnan(d);
}
$ clang -c check-isnan.cpp
Why can't the regular `isnan` be used instead?  Or is this a CUDA-specific requirement?  (Apologies, but I know next to nothing about CUDA :) )

1. #include "cuda_runtime.h"
2. as long as __isnan is a devide function, it should be called from __devide__ or __global__ function
3. clang -c check-isnan.cpp -x cuda

clang/lib/Headers/__clang_cuda_cmath.h

__DEVICE__ float frexp(float __arg, int *__exp) {
   return ::frexpf(__arg, __exp);
 }
 
 // For inscrutable reasons, the CUDA headers define these functions for us on
 // Windows.
 #ifndef _MSC_VER
 __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
 __DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
+__DEVICE__ bool isinf(long double __x) { return ::__isinfl(__x); }
 __DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); }
 // For inscrutable reasons, __finite(), the double-precision version of
 // __finitef, does not exist when compiling for MacOS.  __isfinited is available
 // everywhere and is just as good.
 __DEVICE__ bool isfinite(double __x) { return ::__isfinited(__x); }
+__DEVICE__ bool isfinite(long double __x) { return ::__finitel(__x); }
 __DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); }
 __DEVICE__ bool isnan(double __x) { return ::__isnan(__x); }
+__DEVICE__ bool isnan(long double __x) { return ::__isnanl(__x); }
 #endif
 
 __DEVICE__ bool isgreater(float __x, float __y) {
   return __builtin_isgreater(__x, __y);
 }
 __DEVICE__ bool isgreater(double __x, double __y) {
   return __builtin_isgreater(__x, __y);
 }

clang/lib/Headers/__clang_cuda_device_functions.h

 }
 __DEVICE__ float __fdividef(float __a, float __b) {
   return __nv_fast_fdividef(__a, __b);
 }
 __DEVICE__ int __ffs(int __a) { return __nv_ffs(__a); }
 __DEVICE__ int __ffsll(long long __a) { return __nv_ffsll(__a); }
 __DEVICE__ int __finite(double __a) { return __nv_isfinited(__a); }
 __DEVICE__ int __finitef(float __a) { return __nv_finitef(__a); }
+__DEVICE__ int __finitel(long double __a) { return __finite((double)__a); }
 __DEVICE__ int __float2int_rd(float __a) { return __nv_float2int_rd(__a); }
 __DEVICE__ int __float2int_rn(float __a) { return __nv_float2int_rn(__a); }
 __DEVICE__ int __float2int_ru(float __a) { return __nv_float2int_ru(__a); }
 __DEVICE__ int __float2int_rz(float __a) { return __nv_float2int_rz(__a); }
 __DEVICE__ long long __float2ll_rd(float __a) { return __nv_float2ll_rd(__a); }
 __DEVICE__ long long __float2ll_rn(float __a) { return __nv_float2ll_rn(__a); }
 __DEVICE__ long long __float2ll_ru(float __a) { return __nv_float2ll_ru(__a); }
 __DEVICE__ long long __float2ll_rz(float __a) { return __nv_float2ll_rz(__a); }

 __DEVICE__ float __int2float_rd(int __a) { return __nv_int2float_rd(__a); }
 __DEVICE__ float __int2float_rn(int __a) { return __nv_int2float_rn(__a); }
 __DEVICE__ float __int2float_ru(int __a) { return __nv_int2float_ru(__a); }
 __DEVICE__ float __int2float_rz(int __a) { return __nv_int2float_rz(__a); }
 __DEVICE__ float __int_as_float(int __a) { return __nv_int_as_float(__a); }
 __DEVICE__ int __isfinited(double __a) { return __nv_isfinited(__a); }
 __DEVICE__ int __isinf(double __a) { return __nv_isinfd(__a); }
 __DEVICE__ int __isinff(float __a) { return __nv_isinff(__a); }
+__DEVICE__ int __isinfl(long double __a) { return __isinf((double)__a); }
 __DEVICE__ int __isnan(double __a) { return __nv_isnand(__a); }
 __DEVICE__ int __isnanf(float __a) { return __nv_isnanf(__a); }
+__DEVICE__ int __isnanl(long double __a) { return __isnan((double)__a); }
 __DEVICE__ double __ll2double_rd(long long __a) {
   return __nv_ll2double_rd(__a);
 }
 __DEVICE__ double __ll2double_rn(long long __a) {
   return __nv_ll2double_rn(__a);
 }
 __DEVICE__ double __ll2double_ru(long long __a) {
   return __nv_ll2double_ru(__a);

clang/lib/Headers/__clang_cuda_math_forward_declares.h

 __DEVICE__ int fpclassify(double);
 __DEVICE__ int fpclassify(float);
 __DEVICE__ double frexp(double, int *);
 __DEVICE__ float frexp(float, int *);
 __DEVICE__ double hypot(double, double);
 __DEVICE__ float hypot(float, float);
 __DEVICE__ int ilogb(double);
 __DEVICE__ int ilogb(float);
+__DEVICE__ bool isfinite(long double);
 __DEVICE__ bool isfinite(double);
 __DEVICE__ bool isfinite(float);
 __DEVICE__ bool isgreater(double, double);
 __DEVICE__ bool isgreaterequal(double, double);
 __DEVICE__ bool isgreaterequal(float, float);
 __DEVICE__ bool isgreater(float, float);
+__DEVICE__ bool isinf(long double);
 __DEVICE__ bool isinf(double);
 __DEVICE__ bool isinf(float);
 __DEVICE__ bool isless(double, double);
 __DEVICE__ bool islessequal(double, double);
 __DEVICE__ bool islessequal(float, float);
 __DEVICE__ bool isless(float, float);
 __DEVICE__ bool islessgreater(double, double);
 __DEVICE__ bool islessgreater(float, float);
+__DEVICE__ bool isnan(long double);
 __DEVICE__ bool isnan(double);
 __DEVICE__ bool isnan(float);
 __DEVICE__ bool isnormal(double);
 __DEVICE__ bool isnormal(float);
 __DEVICE__ bool isunordered(double, double);
 __DEVICE__ bool isunordered(float, float);
 __DEVICE__ long labs(long);
 __DEVICE__ double ldexp(double, int);