     CudaSafeCall(cudaMalloc(&_R, s_V));
     CudaSafeCall(cudaMalloc(&_dsum, s_V));
     CudaSafeCall(cudaMalloc(&_rsum, s_V));
+
+    CudaSafeCall(cudaMemcpy(_elements, elements, s_Elements, cudaMemcpyHostToDevice));
+    CudaSafeCall(cudaMemcpy(_V, V, s_V, cudaMemcpyHostToDevice));
+    CudaSafeCall(cudaMemcpy(_nodes, nodes, s_Nodes, cudaMemcpyHostToDevice));
+
+    kernel_integration<<<elementblocks, threads>>>(ne, _elements, _nodes);
+    kernel_util_zero<<<nodeblocks, threads>>>(nn, _dsum);
+    kernel_util_zero<<<nodeblocks, threads>>>(nn, _rsum);
+    cudaDeviceSynchronize();
+    kernel_preprocess<<<elementblocks, threads>>>(ne, _elements, _V, _dsum, _rsum);
+    cudaDeviceSynchronize();
+    kernel_precond<<<nodeblocks, threads>>>(nn, _nodes, _dsum, _nsum, _R, _P, _V);
 
     return k;
 }

    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;
