为数组推力动态内存分配答案

【问题标题】：thrust dynamic memory allocation for array为数组推力动态内存分配
【发布时间】：2013-12-27 13:47:05
【问题描述】：

我有一个推力使用的函子，我需要动态指定它的长度，比如

struct func { 

       const int h;

       func(const int _h): h(_h) {}

       __device__ __host__
       void operator()(int id) {
              double data[h];
      }
};

我不知道该怎么做，因为 h 必须是一个已知数字，但 h 直到运行时才知道。

【问题讨论】：

是否存在预期的、相对较小的 h 值集？
@JoeZ 这部分代码在设备上，不能使用 std::vector。只是指向设备内存的原始指针
@talonmies h 变化很大，在我的例子中它从 20 到 2000 不等。它是神经网络中隐藏神经元的数量。
啊，好的。我会删除我的评论。
@user2684645：乍一看，这似乎是一个非常糟糕的设计选择。每个线程都必须静态或动态分配data 的副本。每个线程 16kb 的本地或堆内存确实会影响性能。这真的是你想做的吗？

标签： c++ cuda thrust

【解决方案1】：

解决这个问题的明显方法是使用动态内存分配，所以函子变成了

   __device__ __host__
   void operator()(int id) {
        double *data  = new double[h];

        // functor code goes here

        // Heap memory has context scope, so delete is necessary to stop leaks
        delete[] data; 
   };

这适用于计算能力为 2.0 或更高版本的 GPU。缺点是内存分配将在全局 memoey 的运行时堆上，这限制了编译器的优化，并且 new/free 运算符本身非常慢，因此在内核启动中的每个线程都发生这种情况将消耗大量性能。

另一种选择，如果h 的值范围有限，请考虑将运算符代码中的 h 替换为模板参数，然后在已知情况下仅使用选择器，例如

   template<int j>
   __device__ __host__
   void guts(int id) {
       double data[j];
       // code here
   };

   __device__ __host__
   void guts_rt(int id) {
       double *data = new double[h];
       // code here
       delete[] data;
   };

   __device__ __host__
   void operator()(int id) {
       switch (h) {
           case 2:
           guts<2>(id);
           break;

           case 4:
           guts<4>(id);
           break;

           // As many as needed here

           default:
           guts_rt(id);
           break;
      }
  }

即。尽可能尝试使用硬编码数组（编译器可以优化），否则回退到动态解决方案（如果你的 GPU 确实支持堆内存的动态分配）。

【讨论】：

谢谢，很好的回答！