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**)&deviceInputs, inputsCount*sizeof(double) ));

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

for(int i=0;i<_layersCount;i++)
{
cudaCall(cudaMalloc((void**)&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<<<blocksMul, threadsMul>>>(deviceTemp, deviceInputs, _neuronsInputsWeights, _inputsInPreviousLayers[i], inputsCount, _neuronsPerLayerCount[i]);

cudaCall(cudaFree(deviceInputs));

cudaCall(cudaMalloc((void**)&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<<<blocksSum, threadsSum>>>(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**)&deviceInputs, some_big_value * sizeof(double) ));
cudaCall(cudaMalloc((void**)&deviceTemp, some_big_value * sizeof(double)));

And in calculateNeuronsOutputsAndDerivatives, reallocate memory only if needed:

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

There is some redundancy within the for loop which could be removed:

dim3 dimGrid = dim3(blocksCount, 1, 1);
dim3 dimBlock = dim3(BLOCK_SIZE, 1, 1);

for (int i = 0; i < _layersCount; i++)
{
cudaCall(cudaMalloc((void**)&deviceTemp, ... ));

size_t blocksCount = ...

weighting<<<dimGrid, dimBlock>>>(...);

// no need to call cudaFree on deviceInputs every iteration
// b/c cudaMalloc will be writing to the same location in device memory
cudaCall(cudaMalloc((void**)&deviceInputs, ... ));

blocksCount = ...

calcOutputsAndDer<<<dimGrid, dimBlock>>>(...);

// rest same as before
}