Trigonometric Functions and Integer Parameters

I’m stuck with the issue of integer and trigonometric functions.

__global__ void trigo(float f, int i) {

        sin(f);

        sin(i); //< comment this to have a compilable kernel

}

It seems to me that in C++ when you have “sin(a);” there is an implicit conversion of a as a float, but in a CUDA kernel you end up with a compilation error:

trigo.cu(3): error: calling a host function("std::sin ") from a device/global function(“trigo”) is not allowed

I wonder if this behavior is intended or is it a “bug”? Google didn’t help me a lot with this issue :-(

I am looking at the C++ standard [specifically, ISO/IEC 14882 Programming Languages – C++ (1998)] and do not see a function signature sin(int). Is this a recent addition to the C++ standard, or a vendor-specific extension that CUDA is exposed to because it incorporates the host system’s math.h ?

For now, simply cast to the desired type (either float or double). Note that sin(float) returns a result accurate to single precision, while sin(double) returns a result accurate to double precision, so you would want to cast based on the desired result precision.

I know that a cast would solve the issue, but this is an automatic generated code and I’d like to avoid hack the generator to make a special rules for Cuda :-(

I have established that math function signatures with integral arguments were added to C++ with the latest standard in September of 2011. I have raised the issue internally, but for now I would suggest changing the code generator so it accomodates compilers that do not yet support the 2011 C++ standard.

After looking at the preprocessed file I have this:

# 77 "/usr/include/c++/4.4/cmath" 3

namespace std __attribute__ ((__visibility__ ("default"))) {

...

using ::sin;

inline float

  sin(float __x)

  { return __builtin_sinf(__x); }

inline long double

  sin(long double __x)

  { return __builtin_sinl(__x); }

template<typename _Tp>

    inline typename __gnu_cxx::__enable_if<__is_integer<_Tp>::__value,

        double>::__type

    sin(_Tp __x)

    { return __builtin_sin(__x); }

...

I wonder if g++ doesn’t generate a sin(int) signature here which can cause the issue…

Our answers have crossed. Ok your point makes sense :)

Anyway I use gcc-4.4.6 which was released in April 2011, so before September 2011!

Prior to the new standard being issued, gcc may have supported this as a vendor-specific extension. To my knowledge, a majority of the material added to new language standards consists of codification of existing practices (such as vendor-specific extensions).

I would suggest filing a “request for enhancement” (RFE) through the bug reporting system to get these additional math function signatures added to CUDA. There is a link to the bug reporting form on the registered developer website. Thanks.