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
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