# Calling \_\_host\_\_ function from \_\_global\_\_ function **URL:** \ **Category:** CUDA Programming and Performance\ **Created:** [August 3, 2020, 7:51pm UTC](https://forums.developer.nvidia.com/t/calling-host-function-from-global-function/145231 "2020-08-03T19:51:45Z")\ **Posts on this page:** 6\ **Page:** 1
**Author:** ![VictoryCoorg](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@VictoryCoorg](https://forums.developer.nvidia.com/u/VictoryCoorg)\ **Post date:** [August 3, 2020, 7:51pm UTC](https://forums.developer.nvidia.com/t/calling-host-function-from-global-function/145231/1 "2020-08-03T19:51:45Z")
While building someone else’s PyTorch cpp/CUDA extension (SparseConvNet) on Windows, I am getting errors of the following type. My guess is that quite a number of people have built this library on Linux and probably aren’t getting these errors. Any idea why these errors may be getting reported only on Windows? Thanks. ```auto folderPath\SparseConvNet\sparseconvnet\SCN\CUDA/BatchNormalization.cu(49): error: calling a __host__ function("pow ") from a __global__ function("BatchNormalization_f_train ") is not allowed folderPath\SparseConvNet\sparseconvnet\SCN\CUDA/BatchNormalization.cu(49): error: identifier "pow " is undefined in device code folderPath\SparseConvNet\sparseconvnet\SCN\CUDA/BatchNormalization.cu(83): error: calling a __host__ function("pow ") from a __global__ function("BatchNormalization_f_test ") is not allowed folderPath\SparseConvNet\sparseconvnet\SCN\CUDA/BatchNormalization.cu(83): error: identifier "pow " is undefined in device code folderPath\SparseConvNet\sparseconvnet\SCN\CUDA/BatchNormalization.cu(49): error: calling a __host__ function("pow ") from a __global__ function("BatchNormalization_f_train ") is not allowed ``` ---
**Author:** ![VictoryCoorg](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@VictoryCoorg](https://forums.developer.nvidia.com/u/VictoryCoorg)\ **Post date:** [August 3, 2020, 10:14pm UTC](https://forums.developer.nvidia.com/t/calling-host-function-from-global-function/145231/2 "2020-08-03T22:14:04Z")
- Interestingly, “pow” happens to be listed among the available “Standard Functions” in the documentation at [Programming Guide :: CUDA Toolkit Documentation](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#standard-functions). The comment about the “Standard Functions” there says “The functions from this section can be used in both host and device code”. So, why is nvcc taking “pow” to be host only function? - Just as an experiment, I replaced “pow” by “powf” and the compiler did not generate errors. Any insights why this could be? Thanks. ---
**Author:** ![njuffa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@njuffa](https://forums.developer.nvidia.com/u/njuffa)\ **Post date:** [August 3, 2020, 10:45pm UTC](https://forums.developer.nvidia.com/t/calling-host-function-from-global-function/145231/3 "2020-08-03T22:45:24Z")
> [@VictoryCoorg](#): > > `pow` If I interpret the error message correctly, this is not a function signature supported by CUDA (or standard C++, best I know). Available choices are: pow (double, int) pow (double, double) pow (float, int) pow (float, float) If a function call in device code doesn’t match any of the available signatures and cannot be resolved to any particular existing signature (presumably because of ambiguity), the compiler may assume it refers to a host function. I am not entirely sure under which circumstances that happens. The first line in question seems to be: ` _saveInvStd = pow(_saveInvStd / nActive + eps, -0.5);` where `_saveInvStd` is of type `T`, `nActive` is of type `int` and `0.5` is of type `double`. So I am assuming you get this error message when you instantiate this with `T`=`float`, as that results in `pow (float, double)`. Here is what happens when I extract this into a simple standalone demo kernel: ``` #include #include __global__ void kernel (float f, double d, int i) { double r1, r2, r3, r4, r5, r6, r7, r8, r9; r1 = pow (f, i); r2 = pow (f, f); r3 = pow (d, i); r4 = pow (d, d); r5 = powf (f, i); r6 = powf (f, f); r7 = powf (d, i); r8 = powf (d, d); r9 = pow (f, d); printf ("%23.16e %23.16e %23.16e %23.16e %23.16e %23.16e %23.16e %23.16e\n", r1, r2, r3, r4, r5, r6, r7, r8); } int main (void) { kernel<<<1,1>>>(2,3,4); cudaDeviceSynchronize(); return EXIT_SUCCESS; } ``` The CUDA 9.2 toolchain on Windows complains: ``` pow.cu(15): error: more than one instance of overloaded function "pow" matches the argument list: function "pow(double, double)" function "pow(float, float)" function "pow(float, int)" argument types are: (float, double) ``` ---
**Author:** ![VictoryCoorg](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@VictoryCoorg](https://forums.developer.nvidia.com/u/VictoryCoorg)\ **Post date:** [August 3, 2020, 11:02pm UTC](https://forums.developer.nvidia.com/t/calling-host-function-from-global-function/145231/4 "2020-08-03T23:02:36Z")
Interesting… The actual call is pow(\_saveInvStd / nActive + eps, -0.5), which takes just two arguments. Perhaps, the compiler is seeing some other templated declaration of “pow” somewhere else that takes three parameters and is confusing it to be that instead of as calling the standard function. Wondering whether there is someway to indicate that what we want is the standard function “pow”… A namespace for anything like that… ---
**Author:** ![njuffa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@njuffa](https://forums.developer.nvidia.com/u/njuffa)\ **Post date:** [August 3, 2020, 11:10pm UTC](https://forums.developer.nvidia.com/t/calling-host-function-from-global-function/145231/5 "2020-08-03T23:10:01Z")
When instantiated with `T` = `float`, computing `pow(_saveInvStd / nActive + eps, -0.5)` turns into `pow (float/int + float, double)` ==\> `pow (float, double)`. This is not supported by the CUDA standard math library. I haven’t checked whether some host compilers have a prototype for and therefore accept `pow (float, double)`. If so, it would _seem to be_ an extension over the standard; at least I haven’t found anything in the C++ standard requiring such an overload. I have only very rudimentary knowledge of Python, but as far as I know it does not support a single-precision floating-point type, only double precision. Which probably means that in this context the code should never be instantiated with `T` = `float`. ---
**Author:** ![TomNVIDIA](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/tomnvidia/32/14181_2.png) [@TomNVIDIA](https://forums.developer.nvidia.com/u/TomNVIDIA)\ **Post date:** [October 12, 2021, 4:04pm UTC](https://forums.developer.nvidia.com/t/calling-host-function-from-global-function/145231/6 "2021-10-12T16:04:26Z")