Преобразовать из рекурсивной в итеративную функцию CUDA C++

Question

Я работаю над генетической программой, в которой перенесу часть тяжелой работы в CUDA. (Ранее просто OpenMP).

Он работает не очень быстро, и я получаю ошибку, связанную с рекурсией:

Размер стека для функции ввода '_Z9KScoreOnePdPiS_S_P9CPPGPNode' не может быть определен статически

Я добавил кусок логики, которая работает на CUDA. Я считаю, что этого достаточно, чтобы показать, как это работает. Я был бы рад услышать о других оптимизациях, которые я мог бы добавить, но я действительно хотел бы взять рекурсию, если она ускорит процесс.

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

__device__ double Fadd(double a, double b)   {
    return a + b;
};


__device__ double Fsubtract(double a, double b)   {
        return a - b;
};

__device__ double action (int fNo, double aa , double bb, double cc, double dd) {
    switch (fNo) {
    case 0 :
        return Fadd(aa,bb);
    case 1 :
        return Fsubtract(aa,bb);
    case 2 :
        return Fmultiply(aa,bb);
    case 3 :
        return Fdivide(aa,bb);
    default:
        return 0.0;
    }

}

__device__ double solve(int node,CPPGPNode * dev_m_Items,double * var_set) {
    if (dev_m_Items[node].is_terminal) {
        return var_set[dev_m_Items[node].tNo];
    } else {
        double values[4];
        for (unsigned int x = 0; x < 4; x++ ) {
            if (x < dev_m_Items[node].fInputs) {
                values[x] = solve(dev_m_Items[node].children[x],dev_m_Items,var_set);
            } else {
                values[x] = 0.0;
            }
        }
        return action(dev_m_Items[node].fNo,values[0],values[1],values[2],values[3]);
    }
}

__global__ void KScoreOne(double *scores,int * root_nodes,double * targets,double * cases,CPPGPNode * dev_m_Items) {
    int pid = blockIdx.x;

    // We only work if this node needs to be calculated
    if (root_nodes[pid] != -1) {
        for (unsigned int case_no = 0; case_no < FITNESS_CASES; case_no ++) {
            double result = solve(root_nodes[pid],dev_m_Items,&cases[case_no]);
            double target = targets[case_no];
            scores[pid] += abs(result - target);
        }
    }
}

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

Answer 1

Я решил эту проблему сейчас. Это был не совсем случай помещения рекурсивных аргументов в стек, но это была очень похожая система.

Как часть создания дерева узлов, я добавляю каждый узел в вектор. Теперь я решаю проблему в обратном порядке, используя http://en.wikipedia.org/wiki/Reverse_Polish_notation, что очень хорошо подходит, поскольку каждый узел содержит либо значение, либо функцию для выполнения.

Это также на ~20% быстрее, чем рекурсивная версия, так что я рад!