如何仅针对一个函数禁用Cuda主机设备警告？

Question

如何仅针对一个函数禁用Cuda主机设备警告？

3

我在我的CUDA代码中有一个针对lambda闭包的C++14模板，它是__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重新启用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

- WilliamKF

您也可以考虑重构代码以避免警告，而不是将其抑制。个人而言，我会首先尝试这个选项。 - Jesper Juhl

@JesperJuhl 是的，我尝试过了，但没有成功。 - WilliamKF

#pragma 只会抑制它前面的一个函数的警告。至少对我来说是这样... - BlameTheBits

2个回答

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

- Jakub Klinkovský

这些似乎在cuda 12/C++20中有所改变。我可以无问题地编译设备函数。 - tommsch

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- BlameTheBits · Accepted Answer

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

另外：还有#pragma nv_exec_check_disable，流行的库已经迁移到该指令。请参阅例如此对话。

#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