Как повторно использовать функторы с данными-членами во многих выполнениях ядра в CUDA, чтобы улучшить использование памяти и уменьшить время копирования?

Question

Я перевожу программу c ++ 11, которая вычисляет силы контакта между парами частиц, в программу cuda. Все пары частиц не зависят друг от друга. Я использую функтор для расчета силы контакта. Этот функтор выполняет много вычислений и содержит много переменных-членов. Поэтому я пытаюсь повторно использовать функторы вместо создания одного нового функтора на пару частиц.

Поскольку функтор содержит виртуальные функции, клонирование функтора выполняется на устройстве, а не на хосте.

Я думаю о схеме, которая выглядит следующим образом:

1) Функторы клона М

2) Начать вычисление M пар частиц

3) Пара частиц M + 1 ожидает завершения одной пары частиц и затем повторно использует ее функтор

Однако, другие идеи также очень приветствуются.

Я сделал очень упрощенную версию программы. В этой программе воспроизведения переменная F не обязательно должна быть переменной-членом, но в реальной программе она должна быть. В реальной программе намного больше данных и пар частиц (N). N часто составляет несколько миллионов.

#include <stdio.h>

#define TPB 4 // realistic value = 128
#define N 10  // realistic value = 5000000
#define M 5   // trade of between copy time and parallel gain.
              // Realistic value somewhere around 1000 maybe

#define OPTION 1
// option 1: Make one functor per particle pair => works, but creates too many functor clones
// option 2: Only make one functor clone => no more thread independent member variables
// option 3: Make M clones which get reused => my suggestion, but I don't know how to program it

struct FtorBase
{
  __device__ virtual void execute(long i) = 0;

  __device__ virtual void show() = 0;
};

struct FtorA : public FtorBase
{

  __device__ void execute(long i) final
  {
    F = a*i;
  }

  __device__ void show() final
  {
    printf("F = %f\n", F);
  }

  double a;
  double F;
};

template <class T>
__global__ void cloneFtor(FtorBase** d_ftorBase, T ftor, long n_ftorClones)
{
  const long i = threadIdx.x + blockIdx.x * blockDim.x;

  if (i >= n_ftorClones) {
    return;
  }

  d_ftorBase[i] = new T(ftor);
}

struct ClassA
{
  typedef FtorA ftor_t;

  FtorBase** getFtor()
  {
    FtorBase** d_cmFtorBase;
    cudaMalloc(&d_cmFtorBase, N * sizeof(FtorBase*));

#if OPTION == 1 
    // option 1: Create one copy of the functor per particle pair
    printf("using option 1\n");
    cloneFtor<<<(N + TPB - 1) / TPB, TPB>>>(d_cmFtorBase, ftor_, N);
#elif OPTION == 2
    // option 2: Create just one copy of the functor
    printf("using option 2\n");
    cloneFtor<<<1, 1>>>(d_cmFtorBase, ftor_, 1);
#elif OPTION == 3
    // option 3: Create M functor clones
    printf("using option 3\n");
    printf("This option is not implemented. I don't know how to do this.\n");
    cloneFtor<<<(M + TPB - 1) / TPB, TPB>>>(d_cmFtorBase, ftor_, M);
#endif
    cudaDeviceSynchronize();

    return d_cmFtorBase;
  }

  ftor_t ftor_;
};


__global__ void cudaExecuteFtor(FtorBase** ftorBase)
{
  const long i = threadIdx.x + blockIdx.x * blockDim.x;

  if (i >= N) {
    return;
  }

#if OPTION == 1
  // option 1: One functor per particle was created
  ftorBase[i]->execute(i);
  ftorBase[i]->show();
#elif OPTION == 2
  // option 2: Only one single functor was created
  ftorBase[0]->execute(i);
  ftorBase[0]->show();
#elif OPTION == 3
  // option 3: Reuse the fuctors
  // I don't know how to do this
#endif
}

int main()
{
  ClassA* classA = new ClassA();
  classA->ftor_.a = .1;

  FtorBase** ftorBase = classA->getFtor();

  cudaExecuteFtor<<<(N + TPB - 1) / TPB, TPB>>>(ftorBase);
  cudaDeviceSynchronize();

  return 0;
}

Я проверяю вывод F, чтобы увидеть, является ли переменная-член независимой в каждом вызове. Как и ожидалось, при использовании разных функторов для каждой пары частиц (вариант 1) все значения F различны, а при использовании только одного функтора для всей программы (вариант 2) все значения F одинаковы.

using option 1
F = 0.800000
F = 0.900000
F = 0.000000
F = 0.100000
F = 0.200000
F = 0.300000
F = 0.400000
F = 0.500000
F = 0.600000
F = 0.700000

using option 2
F = 0.700000
F = 0.700000
F = 0.700000
F = 0.700000
F = 0.700000
F = 0.700000
F = 0.700000
F = 0.700000
F = 0.700000
F = 0.700000

Интересно, есть ли способ получить все различные значения F в этом примере воспроизведения, не беря N копий (вариант 3).

PS: я использую Ubuntu 18.04, nvcc 9.1 и графическую карту NVIDIA GeForce GTX 1060 Mobile (совместимость cuda 6.1).

UPDATE:

В предыдущем коде, который я представил, была проблема только в режиме отладки (компиляция с флагом -G), но не в версии выпуска. Я предполагаю, что компилятор оптимизировал printf("F = %f\n", F); до printf("F = %f\n", a*i);, так что проблема переменных, зависящих от потока, о чем этот вопрос, исчезла.

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