Trigonometric Functions and Integer Parameters
I'm stuck with the issue of integer and trigonometric functions.

[code]
__global__ void trigo(float f, int i) {
sin(f);
sin(i); //< comment this to have a compilable kernel
}
[/code]

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<int> ") 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'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<int> ") 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 :-(

#1
Posted 02/28/2012 04:48 PM   
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 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.

#2
Posted 02/28/2012 06:23 PM   
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 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 :-(

#3
Posted 02/28/2012 10:14 PM   
I have established that math function signatures with integral arguments were added to C++ with the latest standard in [b]September of 2011[/b]. 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.
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.

#4
Posted 02/28/2012 10:23 PM   
After looking at the preprocessed file I have this:

[code]
# 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); }

...
[/code]


I wonder if g++ doesn't generate a sin(int) signature here which can cause the issue...
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...

#5
Posted 02/28/2012 10:31 PM   
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!
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!

#6
Posted 02/28/2012 10:35 PM   
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.
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.

#7
Posted 02/28/2012 10:48 PM   
Scroll To Top