4
\$\begingroup\$

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));
}
\$\endgroup\$
4
\$\begingroup\$

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;
}
\$\endgroup\$
3
\$\begingroup\$

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
   }
\$\endgroup\$

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Not the answer you're looking for? Browse other questions tagged or ask your own question.