This project is read-only.
1

Closed

Mixing of host code and device code in same file results in compliation issues

description

Attached is a simplified file demonstrating a problem when trying to verify code that mixes both host and device code. The attached code is successfully compiled by nvcc.

In this example, it is trivial to comment out the math.h include and get GPUVerify to run on the code, however, the actual code that I'm working on is far more complex.

Here are the (truncated) errors:
$ ~/ws/gpuverify/gpuverify/GPUVerify.py -I ../../include/ --cuda --blockDim=1 --gridDim=1 bug.cu

In file included from bug.cu:2:
In file included from /usr/include/math.h:71:
/usr/include/bits/mathcalls.h:55:13: error: exception specification in declaration does not match previous declaration
__MATHCALL (acos,, (_Mdouble_ __x));
            ^
/usr/include/math.h:55:25: note: expanded from macro '__MATHCALL'
  __MATHDECL (_Mdouble_,function,suffix, args)
                        ^
/usr/include/math.h:57:22: note: expanded from macro '__MATHDECL'
  __MATHDECL_1(type, function,suffix, args); \
                     ^
/usr/include/math.h:65:31: note: expanded from macro '__MATHDECL_1'
  extern type __MATH_PRECNAME(function,suffix) args __THROW
                              ^
/usr/include/math.h:68:42: note: expanded from macro '__MATH_PRECNAME'
#define __MATH_PRECNAME(name,r) __CONCAT(name,r)
                                         ^
/usr/include/sys/cdefs.h:80:23: note: expanded from macro '__CONCAT'
#define __CONCAT(x,y)   x ## y
                        ^
/h2/sreepai/ws/gpuverify/bugle/src/include-blang/cuda_math_functions.h:172:19: note: previous declaration is here
__device__ double acos(double x);
                  ^
In file included from bug.cu:2:
In file included from /usr/include/math.h:71:
/usr/include/bits/mathcalls.h:57:13: error: exception specification in declaration does not match previous declaration
__MATHCALL (asin,, (_Mdouble_ __x));
            ^
/usr/include/math.h:55:25: note: expanded from macro '__MATHCALL'
  __MATHDECL (_Mdouble_,function,suffix, args)
                        ^
/usr/include/math.h:57:22: note: expanded from macro '__MATHDECL'
  __MATHDECL_1(type, function,suffix, args); \
                     ^
/usr/include/math.h:65:31: note: expanded from macro '__MATHDECL_1'
  extern type __MATH_PRECNAME(function,suffix) args __THROW
                              ^
/usr/include/math.h:68:42: note: expanded from macro '__MATH_PRECNAME'
#define __MATH_PRECNAME(name,r) __CONCAT(name,r)
                                         ^
/usr/include/sys/cdefs.h:80:23: note: expanded from macro '__CONCAT'
#define __CONCAT(x,y)   x ## y
                        ^
/h2/sreepai/ws/gpuverify/bugle/src/include-blang/cuda_math_functions.h:171:19: note: previous declaration is here
__device__ double asin(double x);
                  ^
In file included from bug.cu:2:
In file included from /usr/include/math.h:71:
/usr/include/bits/mathcalls.h:59:13: error: exception specification in declaration does not match previous declaration
__MATHCALL (atan,, (_Mdouble_ __x));
            ^
/usr/include/math.h:55:25: note: expanded from macro '__MATHCALL'
  __MATHDECL (_Mdouble_,function,suffix, args)
                        ^
/usr/include/math.h:57:22: note: expanded from macro '__MATHDECL'
  __MATHDECL_1(type, function,suffix, args); \
                     ^
/usr/include/math.h:65:31: note: expanded from macro '__MATHDECL_1'
  extern type __MATH_PRECNAME(function,suffix) args __THROW
                              ^
/usr/include/math.h:68:42: note: expanded from macro '__MATH_PRECNAME'
#define __MATH_PRECNAME(name,r) __CONCAT(name,r)
                                         ^
/usr/include/sys/cdefs.h:80:23: note: expanded from macro '__CONCAT'
#define __CONCAT(x,y)   x ## y
                        ^
/h2/sreepai/ws/gpuverify/bugle/src/include-blang/cuda_math_functions.h:173:19: note: previous declaration is here
__device__ double atan(double x);
                  ^
In file included from bug.cu:2:
In file included from /usr/include/math.h:71:
/usr/include/bits/mathcalls.h:61:13: error: exception specification in declaration does not match previous declaration
__MATHCALL (atan2,, (_Mdouble_ __y, _Mdouble_ __x));
            ^
/usr/include/math.h:55:25: note: expanded from macro '__MATHCALL'
  __MATHDECL (_Mdouble_,function,suffix, args)
                        ^
/usr/include/math.h:57:22: note: expanded from macro '__MATHDECL'
  __MATHDECL_1(type, function,suffix, args); \
                     ^
/usr/include/math.h:65:31: note: expanded from macro '__MATHDECL_1'
  extern type __MATH_PRECNAME(function,suffix) args __THROW
                              ^
/usr/include/math.h:68:42: note: expanded from macro '__MATH_PRECNAME'
#define __MATH_PRECNAME(name,r) __CONCAT(name,r)
                                         ^
/usr/include/sys/cdefs.h:80:23: note: expanded from macro '__CONCAT'
#define __CONCAT(x,y)   x ## y
                        ^
/h2/sreepai/ws/gpuverify/bugle/src/include-blang/cuda_math_functions.h:174:19: note: previous declaration is here
__device__ double atan2(double y, double x);
                  ^
In file included from bug.cu:2:
In file included from /usr/include/math.h:71:
/usr/include/bits/mathcalls.h:64:13: error: exception specification in declaration does not match previous declaration
__MATHCALL (cos,, (_Mdouble_ __x));
            ^
/usr/include/math.h:55:25: note: expanded from macro '__MATHCALL'
  __MATHDECL (_Mdouble_,function,suffix, args)
                        ^
/usr/include/math.h:57:22: note: expanded from macro '__MATHDECL'
  __MATHDECL_1(type, function,suffix, args); \
                     ^
/usr/include/math.h:65:31: note: expanded from macro '__MATHDECL_1'
  extern type __MATH_PRECNAME(function,suffix) args __THROW
                              ^
/usr/include/math.h:68:42: note: expanded from macro '__MATH_PRECNAME'
#define __MATH_PRECNAME(name,r) __CONCAT(name,r)
                                         ^
/usr/include/sys/cdefs.h:80:23: note: expanded from macro '__CONCAT'
#define __CONCAT(x,y)   x ## y
                        ^
/h2/sreepai/ws/gpuverify/bugle/src/include-blang/cuda_math_functions.h:161:19: note: previous declaration is here
__device__ double cos(double x);
                  ^
In file included from bug.cu:2:
In file included from /usr/include/math.h:71:
/usr/include/bits/mathcalls.h:66:13: error: exception specification in declaration does not match previous declaration
__MATHCALL (sin,, (_Mdouble_ __x));
            ^
/usr/include/math.h:55:25: note: expanded from macro '__MATHCALL'
  __MATHDECL (_Mdouble_,function,suffix, args)
                        ^
/usr/include/math.h:57:22: note: expanded from macro '__MATHDECL'
  __MATHDECL_1(type, function,suffix, args); \
                     ^
/usr/include/math.h:65:31: note: expanded from macro '__MATHDECL_1'
  extern type __MATH_PRECNAME(function,suffix) args __THROW
                              ^
/usr/include/math.h:68:42: note: expanded from macro '__MATH_PRECNAME'
#define __MATH_PRECNAME(name,r) __CONCAT(name,r)
                                         ^
/usr/include/sys/cdefs.h:80:23: note: expanded from macro '__CONCAT'
#define __CONCAT(x,y)   x ## y
                        ^
/h2/sreepai/ws/gpuverify/bugle/src/include-blang/cuda_math_functions.h:160:19: note: previous declaration is here
__device__ double sin(double x);
                  ^
In file included from bug.cu:2:
In file included from /usr/include/math.h:71:
/usr/include/bits/mathcalls.h:68:13: error: exception specification in declaration does not match previous declaration
__MATHCALL (tan,, (_Mdouble_ __x));
            ^
/usr/include/math.h:55:25: note: expanded from macro '__MATHCALL'
  __MATHDECL (_Mdouble_,function,suffix, args)
                        ^
/usr/include/math.h:57:22: note: expanded from macro '__MATHDECL'
  __MATHDECL_1(type, function,suffix, args); \
                     ^
/usr/include/math.h:65:31: note: expanded from macro '__MATHDECL_1'
  extern type __MATH_PRECNAME(function,suffix) args __THROW
                              ^
/usr/include/math.h:68:42: note: expanded from macro '__MATH_PRECNAME'
#define __MATH_PRECNAME(name,r) __CONCAT(name,r)
                                         ^
/usr/include/sys/cdefs.h:80:23: note: expanded from macro '__CONCAT'
#define __CONCAT(x,y)   x ## y
                        ^
/h2/sreepai/ws/gpuverify/bugle/src/include-blang/cuda_math_functions.h:162:19: note: previous declaration is here
__device__ double tan(double x);
                  ^
In file included from bug.cu:2:
In file included from /usr/include/math.h:71:
/usr/include/bits/mathcalls.h:73:13: error: exception specification in declaration does not match previous declaration
__MATHCALL (cosh,, (_Mdouble_ __x));
            ^
/usr/include/math.h:55:25: note: expanded from macro '__MATHCALL'
  __MATHDECL (_Mdouble_,function,suffix, args)
                        ^
/usr/include/math.h:57:22: note: expanded from macro '__MATHDECL'
  __MATHDECL_1(type, function,suffix, args); \
                     ^
/usr/include/math.h:65:31: note: expanded from macro '__MATHDECL_1'
  extern type __MATH_PRECNAME(function,suffix) args __THROW
                              ^
/usr/include/math.h:68:42: note: expanded from macro '__MATH_PRECNAME'
#define __MATH_PRECNAME(name,r) __CONCAT(name,r)
                                         ^
/usr/include/sys/cdefs.h:80:23: note: expanded from macro '__CONCAT'
#define __CONCAT(x,y)   x ## y
                        ^
/h2/sreepai/ws/gpuverify/bugle/src/include-blang/cuda_math_functions.h:176:19: note: previous declaration is here
__device__ double cosh(double x);
...
fatal error: too many errors emitted, stopping now [-ferror-limit=]
20 errors generated.

file attachments

Closed Mar 3, 2015 at 4:51 PM by jketema
I fixed this particular issue for Linux systems. This should not cause any regressions on other systems, but expect those to be still broken. Other systems will be hard to fix without starting to override the system supplied headers in some way and distributing those with GPUVerify.

Mixed host/device code might still give errors on Linux due to other issues. Please file separate reports for those issues.

comments

allydonaldson wrote Sep 3, 2014 at 10:17 PM

Thanks for the report.

Is math.h a CUDA-specific, or the standard math.h file?

sreepathi wrote Sep 3, 2014 at 11:28 PM

This is the standard math.h file in /usr/include.

I think nvcc somehow ignores host-side code when processing device code.