如何禁用一个功能的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

警告消失了，然而，我担心我只想在这个模板函数中取消它，而不是编译单元的其余部分。我不能轻易地将模板函数移到编译单元的末尾。

我想要一些推送和弹出，但我找不到这个。

在定义模板函数后，有没有办法用杂注重新启用hd警告？

我试过了：

#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_disable只影响它前面的函数，所以不需要像#pragma hd_warning_enable这样的东西。这似乎在任何文档中都找不到，但下面的示例表明了这一行为。

旁注:还有#pragma nv_exec_check_disable，流行的库已经迁移到了那个编译指示。例如，请参阅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 );
}

与未记录的编译指示相比，这具有一个优势，即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 );
}