тяга динамического распределения памяти для массива

Question

У меня есть функтор, используемый тяги, когда мне нужно динамически указать его длину, как

struct func { 

     const int h; 

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

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

Я не знаю, как это сделать, потому что час должен быть известное число, но h неизвестно до времени выполнения.

Answer 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 или более поздней версии. Недостатком является то, что распределение памяти будет в куче рабочей среды в глобальной памяти, что ограничивает оптимизацию компилятора, а сами новые/бесплатные операторы очень медленны, поэтому, если это произойдет для каждого потока в запуске ядра, это будет стоить очень высокой производительности.

Альтернативы, если диапазон значений 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 фактически поддерживает динамическое распределение кучной памяти).