Calculating neuron outputs and derivatives This function runs very often.  cudaMemcpy is at the start and works very slowly. How can I change this function to avoid this? I already have inputs in device memory.

void OpenNNL::calculateNeuronsOutputsAndDerivatives(double * inputs, double * deviceOutputs, double * deviceDerivatives)
{
    int inputsCount = _inputsCount;

    double * deviceTemp;
    double * deviceInputs;

    cudaCall(cudaMalloc ( (void**)&amp;deviceInputs, inputsCount*sizeof(double) ));

    cudaCall(cudaMemcpy ( deviceInputs, inputs, inputsCount*sizeof(double), cudaMemcpyDeviceToDevice ));

    for(int i=0;i&lt;_layersCount;i++)
    {
        cudaCall(cudaMalloc((void**)&amp;deviceTemp, _neuronsPerLayerCount[i]*inputsCount*sizeof(double)));

        dim3 threadsMul = dim3(BLOCK_SIZE, 1);
        int blocksCount = floor((double) _neuronsPerLayerCount[i]*inputsCount / threadsMul.x) + 1;
        dim3 blocksMul = dim3(blocksCount, 1);

        weighting&lt;&lt;&lt;blocksMul, threadsMul&gt;&gt;&gt;(deviceTemp, deviceInputs, _neuronsInputsWeights, _inputsInPreviousLayers[i], inputsCount, _neuronsPerLayerCount[i]);

        cudaCall(cudaFree(deviceInputs));

        cudaCall(cudaMalloc((void**)&amp;deviceInputs, _neuronsPerLayerCount[i]*sizeof(double)));

        dim3 threadsSum = dim3(BLOCK_SIZE, 1);
        blocksCount = floor((double) _neuronsPerLayerCount[i] / threadsSum.x) + 1;
        dim3 blocksSum = dim3(blocksCount, 1);

        calculateOutputsAndDerivatives&lt;&lt;&lt;blocksSum, threadsSum&gt;&gt;&gt;(deviceOutputs, deviceDerivatives, deviceInputs, deviceTemp, _neuronsBiases, inputsCount, _neuronsPerLayerCount[i], _neuronsInPreviousLayers[i]);

        inputsCount = _neuronsPerLayerCount[i];

        cudaCall(cudaFree(deviceTemp));
    }

    cudaCall(cudaFree(deviceInputs));
}
Try to minimaze memory allocations.

Allocate memory for deviceTemp and deviceInputs only once (in the constructor, for example):

cudaCall(cudaMalloc ( (void**)&amp;deviceInputs, some_big_value * sizeof(double) ));
cudaCall(cudaMalloc((void**)&amp;deviceTemp, some_big_value * sizeof(double)));


And in calculateNeuronsOutputsAndDerivatives, reallocate memory only if needed:

if (cur_deviceInputs_size &lt; inputsCount)
{
    cudaCall(cudaFree(deviceInputs));
    cudaCall(cudaMalloc ( (void**)&amp;deviceInputs, inputsCount*sizeof(double) ));
    cur_deviceInputs_size = inputsCount;
}