CUDA __device__ 函数作为类成员：内联和性能？答案

【问题标题】：CUDA __device__ function as class member: Inlining and performance?CUDA __device__ 函数作为类成员：内联和性能？
【发布时间】：2014-12-30 01:56:03
【问题描述】：

我计划将我的计算划分为一个细粒度的函数/类框架，这些框架封装了某个部分。

类似这样，但有更多的类和通常更长的参数列表：

class Point{

  Coordinates thisPoint;
  Value getPointValue();
  Point getPoint(Offset offset); 
  Point getNumNeighbors();
  Point getNeighbor(int i);
  // many more

}

class Operator{

  void doOperation(Point p){
    // calls some of the functions in Point
  }

}

显然，这在任何面向对象的语言中都是一个很好的做法。但它旨在在 CUDA GPU 上运行。我不知道的是：当我将所有这些细粒度的函数限定为__device__ 并在内核中调用它们时——它们将如何实现？成员函数的调用会产生很大的开销，还是会被内联或以其他方式有效优化？通常，这些函数非常短，但会被调用很多次。

【问题讨论】：

我在所有 CUDA 应用程序中都遵循面向对象的模式，从未经历过显着的性能开销，但没有图表来支持我的陈述。我想开发您的应用程序然后对其进行分析，与 __device__ 函数不是成员函数的版本进行比较...实际上，我确实遇到了一个稍微相关的性能问题，请参阅更多信息 here 但这是由优化器引起的在非常特殊的情况下（包括调用__device__ 成员函数）无法正确解开循环。

标签： c++ cuda

【解决方案1】：

出于性能原因，GPU 编译器将积极地inline 函数。在这种情况下，对性能应该没有特别的影响。

如果无法内联函数，则会发生通常的性能开销，包括创建堆栈帧和调用函数 - 就像您在 CPU 调用非内联函数时观察到的那样。

如果您对特定示例有疑虑，可以创建一个简短的测试代码并使用cuobjdump -sass myexe 查看生成的汇编语言 (SASS)，并确定该函数是否已内联。

在作为类成员/方法的 __device__ 函数的内联中没有 general restrictions。

【讨论】：

所以在内联函数中使用this 指针时不会遇到问题？乍一看，这应该不是问题，但 C++ 的某些东西表现出没有经验的人不会预料到的行为。
我不知道 C++ 中关于 this 使用的任何 CUDA 限制。如果您对此有疑问，为什么不问一个新问题呢？更好的是，尝试一下。如果您遇到困难，请询问您遇到的困难。