How to disable Cuda host device warning for just one function?

Question

I've got a C++14 template in my Cuda code that is templated over a lambda closure and is __host__ and __device__ and I get the warning:

warning: calling a __host__ function("Iter<(bool)1> ::Iter [subobject]")
         from a __host__ __device__ function("Horizontal::Horizontal")
         is not allowed

But this is a false positive, because it is only the __host__ instantiation of the template that calls the __host__ function, so I wish to suppress this warning for just this one template definition.

I can add this before the template:

#pragma hd_warning_disable

And the warning goes away, however, I'm concerned that I only want it suppressed for this one template function, not the remaining of the compilation unit. I cannot easily move the template function to the end of the compilation unit.

I'd like some kind of push and pop but I'm not finding this anywhere.

Is there a way to reenable hd warnings with a pragma after the template function is defined?

I tried:

#pragma hd_warning_enable

But that doesn't work:

test.cu:44:0: warning: ignoring #pragma 
hd_warning_enable  [-Wunknown-pragmas]
 #pragma hd_warning_enable

Here is simple test case to demonstrate the issue:

//#pragma hd_warning_disable

template<typename Lambda>
__host__ __device__
int hostDeviceFunction(const Lambda lambda)
{
    return lambda();
}


__device__
int deviceFunction()
{
    auto lambda = []() { return 0.0; };

    return hostDeviceFunction( lambda );
}

__host__
int hostFunction()
{
    auto lambda = []() { return 1.0; };

    return hostDeviceFunction( lambda );
}

Which gives these warnings:

test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed
test.cu(7): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction(Lambda) [with Lambda=lambda []()->double]" 
(24): here
test.cu(7): warning: calling a __host__ function(" const") from a __host__ __device__ function("hostDeviceFunction< ::> ") is not allowed

You could also consider restructuring your code to avoid the warning instead of suppressing it. *I* would, personally, try that option *first*. — Jesper Juhl, Apr 02 '19 at 18:21
The pragma only suppresses the warnings for the one function it's in front of. At least that's the case for me... — BlameTheBits, Apr 02 '19 at 20:05

BlameTheBits · Accepted Answer · 2019-11-07T16:32:24.310

There is no need for something like #pragma hd_warning_enable as #pragma hd_warning_disable only affects the function it is placed in front of. As it seems this is nothing one can find in any documentation but the example below suggests that behavior.

Sidenote: There is also #pragma nv_exec_check_disable and popular libraries have migrated to that pragma. See e.g. this conversation about it.

#include <iostream>
#include <cassert>

#pragma hd_warning_disable
//#pragma nv_exec_check_disable
template<typename Lambda>
__host__ __device__
int hostDeviceFunction1(const Lambda lambda)
{
    return lambda()*1.0;
}

__host__            
int hostFunction1()
{
    auto lambda = []() { return 1.0; };  
    return hostDeviceFunction1( lambda );
}

template<typename Lambda>
__host__ __device__
int hostDeviceFunction2(const Lambda lambda)
{                                       
    return lambda()*2.0;
}

__host__
int hostFunction2()
{
    auto lambda = []() { return 2.0; };  
    return hostDeviceFunction2( lambda );
}

int main()           
{ 
  std::cout << "hostFunction1: " << hostFunction1() << std::endl;
  assert(hostFunction1() == 1.0);

  std::cout << "hostFunction2: " << hostFunction2() << std::endl;
  assert(hostFunction2() == 4.0);

  return 0;
}

$ nvcc pragma_test.cu 
pragma_test.cu(24): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during instantiation of "int hostDeviceFunction2(Lambda) [with Lambda=lambda []()->double]" 
(31): here

If anyone finds the documentation about (any of) theses pragmas please tell us. There is also an [open question](https://stackoverflow.com/q/49328130/7968757) about it. — BlameTheBits, Apr 02 '19 at 20:48
I've confirmed that `#pragma nv_exec_check_disable` only affects the following function. — WilliamKF, Apr 02 '19 at 21:03
Be careful, both `nv_exec_check_disable` and `hd_warning_disable` pragmas makes my nvcc 11.2 generate *wrong* code. This is also reported here https://forums.developer.nvidia.com/t/pragma-hd-warning-disable-causes-nvcc-to-generate-incorrect-code-cuda-9-1/57755 . Given that, they are not useful for my case, having the warning displayed and working code seems preferable. Also, ` --diag_suppress=20014` does nothing in my case (warning still shown). — alfC, Oct 11 '21 at 10:39

score 2 · Answer 2 · answered Aug 24 '19 at 20:25

Another way to avoid the warning in some cases is to make the lower-level function constexpr and use the --expt-relaxed-constexpr flag.

For example:

template<typename Lambda>
constexpr int constexprFunction(Lambda&& lambda)
{
    return lambda();
}


__host__
int hostFunction()
{
    auto lambda = []() { return 1.0; };
    return constexprFunction( lambda );
}

__host__ __device__
int hostDeviceFunction()
{
    auto lambda = []() { return 0.0; };
    return constexprFunction( lambda );
}

Compared to the undocumented pragmas, this has the advantage that the effects of constexpr are exactly documented, although the exact behaviour of nvcc is still undocumented and its conformance to the standards is not 100%. For instance, the example above works with nvcc 10.1.243 in C++14 mode, but not in C++11. And if you change the signature of the function to constexpr int constexprFunction(const Lambda lambda), the warning still appears, so --expt-relaxed-constexpr does not seem to do its thing in all cases. And if you add a __device__ function, it results in an error either way:

__device__
int deviceFunction()
{
    auto lambda = []() { return 0.0; };
    return constexprFunction( lambda );
}

These seems to have changed with cuda 12/C++20. I can compile the device function without problems. — tommsch, Jul 26 '23 at 21:16

How to disable Cuda host device warning for just one function?

2 Answers2