如何仅针对一项功能禁用 Cuda 主机设备警告？

Question

我的 Cuda 代码中有一个 C++14 模板，它是在 lambda 闭包上模板化的 __host__ 和 __device__，我收到警告：

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

但这是一个误报，因为它只是调用 __host__ 函数的模板的 __host__ 实例化，所以我希望仅针对这个模板定义抑制此警告。

我可以在模板前加上这个：

#pragma hd_warning_disable

并且警告消失了，但是，我担心我只想为这个模板函数抑制它，而不是编译单元的其余部分。我无法轻易将模板函数移动到编译单元的末尾。

我想要某种推送和弹出功能，但我在任何地方都找不到。

有没有办法在定义模板函数后使用 pragma 重新启用高清警告？

我试过：

#pragma hd_warning_enable

但这不起作用：

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

这里是演示问题的简单测试用例：

//#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 );
}

给出了这些警告：

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

Answer 1

不需要像 #pragma hd_warning_enable 这样的东西，因为 #pragma hd_warning_disable 只影响放在它前面的函数。似乎这在任何文档中都找不到，但下面的示例表明了这种行为。

旁注：还有 #pragma nv_exec_check_disable 并且流行的库已经迁移到该 pragma。参见例如this conversation 关于它。

#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

Answer 2

在某些情况下避免警告的另一种方法是创建较低级别的函数 constexpr 并使用 --expt-relaxed-constexpr 标志。

例如：

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 );
}

与未记录的 pragma 相比，其优点是 constexpr 的效果已被准确记录，尽管 nvcc 的确切行为仍未记录并且其符合标准不是 100 %。例如，上面的示例在 C++14 模式下适用于 nvcc 10.1.243，但不适用于 C++11。如果将函数的签名更改为 constexpr int constexprFunction(const Lambda lambda)，警告仍然会出现，因此 --expt-relaxed-constexpr 似乎并非在所有情况下都起作用。如果您添加 __device__ 函数，无论哪种方式都会导致错误：

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

如何仅针对一项功能禁用 Cuda 主机设备警告？

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

c++

cuda

closures

c++14

pragma