如何克服堆栈大小警告？

Question

我想知道关于以下类型警告的最佳实践：

ptxas warning : Stack size for entry function '_Z11cuda_kernelv' cannot be statically determined

似乎将 virtual 关键字添加到 Internal 的析构函数中，即在以下程序中从 __device__ ~Internal(); 移动到 __device__ virtual ~Internal();：

template<typename T>
class Internal {
  T val;
public:
  __device__ Internal();
  __device__ virtual ~Internal();
  __device__ const T& get() const;
};

template<typename T>
__device__ Internal<T>::Internal(): val() {}
template<typename T>
__device__ Internal<T>::~Internal() {}
template<typename T>
__device__ const T& Internal<T>::get() const { return val; }


template<typename T>
class Wrapper {
  Internal<T> *arr;
public:
  __device__ Wrapper(size_t);
  __device__ virtual ~Wrapper();
};

template<typename T>
__device__ Wrapper<T>::Wrapper(size_t len): arr(nullptr) {
  printf("%s\n", __PRETTY_FUNCTION__);
  arr = new Internal<T>[len];
}

template<typename T>
__device__ Wrapper<T>::~Wrapper() {
  delete[] arr;
}

__global__ void cuda_kernel() {
  Wrapper<double> *wp = new Wrapper<double>(10);
  delete wp; 
}

int main() {
  cuda_kernel<<<1,1>>>();
  cudaDeviceSynchronize();
  return 0;
}

遇到上面的警告，我想知道在这种情况下我应该怎么做？

Answer 1

非常简短的回答是，对于这个特定的警告，您无能为力。

更详细：

此警告是汇编程序警告，而不是编译器警告
NVIDIA 工具链依赖大量汇编程序级优化来生成 SASS 硅片上运行的高性能机器代码。 NVIDIA 编译器发出 PTX 虚拟机语言，在组装时可以进行重大转换。这包括将单个静态赋值形式的编译器输出解析为静态寄存器赋值（并将寄存器溢出到本地内存）、函数的内联扩展以及静态编译堆栈预留的发射。所有这些都是潜在的性能优化操作。
这是来自汇编程序的信息性警告，告诉您在静态代码分析期间，汇编程序无法确定堆栈大小。
汇编程序发出此警告的最正常情况是在内核代码中检测到递归。您的用例显然是另一个。
警告来自汇编程序优化阶段。汇编程序让您知道 table 上还有潜在的性能改进机会，因为代码的编译器输出结构不允许它静态确定堆栈大小
汇编程序将使用的回退将更多样板 SASS 到 set-up 和 tear-down 内核将需要的每个线程堆栈运行。警告是让您知道发生了什么。

如何克服堆栈大小警告？

How to overcome Stack size warning?

c++

cuda

stack

ptxas