 }
 
 // Kernel de pré-condicionamento.
-__global__ void kernel_precond(int nn, node * nodes, float * rsum, float * dsum,
+__global__ void kernel_precond(int nn, node * nodes, float * dsum, float * rsum,
                                float * R, float * P, float * V) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i >= nn) return;

    P[i] = ri;
}

__global__ void kernel_iter_element(int ne) { }

// Função externa que processa o problema na GPU.
//    ne: número de elementos.

    float *p = (float*)malloc(nn*sizeof(float));
    float *u = (float*)malloc(nn*sizeof(float));

    // Pré-condicionamento.
    for (i = 0; i < nn; i++) {
        ri = nodes[i].calc ? rsum[i] - dsum[i]*V[i] : 0.0;
        p[i] = ri;
