Как можно оптимизировать этот кусочек CUDA-кода?

Question

[Robotex]

Взгляните на этот кусочек кода:

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));
    }

Эта функция запускается очень часто. И работает она очень медленно. Вы можете увидеть cudaMemcpy в начале функции. Как можно было бы переписать ее так, чтобы избежать копирования? Массив **inputs** уже находится в глобальной памяти.

Answer 1

[Robotex]

Вот здесь результаты работы профайлера: ubuntuone.com/41OVwiE7NEd3fI9jALsFi6
Но, увы, я не понимаю, что это значит (я в CUDA новичек). Я был бы рад, если бы кто-то смог мне пояснить это.

Answer 2

[Disasm]

Так не копируйте данные, передавайте их сразу на вход ядру weighting (если они не изменяются ядром, конечно).
Ещё было бы неплохо дожидаться окончания работы ядер.

Answer 3

[oleksandr_veles]

1.Я бы ещё избавился от cudaFree cudaMalloc во внутреннем цикле, хотя не знаю точно насколько они медленные.
2. Использовал бы одинарную точность, двойная на nvidia на порядок медленнее, тем более на мобильной карте.
3. Присоединяюсь к совету дожидаться окончания работы ядер.

Comments

[Robotex]

Что значит дожидаться окончания работы ядер?

[Robotex]

Избавился от cudaFree cudaMalloc — скорость сравнялась с процессорной.
Замена на одинарную точность дала уменьшение времени выполнения всего лишь на 12 секунд (с двойной — 1m28.486s, с одинарной — 1m16.102s, CPU — 1m38.294s)
Как заставить ее дожидаться окончания работы ядер?

[Disasm]

cudaThreadSynchronize();

[oleksandr_veles]

наверное можно cudaEvent* использовать, не помню точно, но синхронизация ядер повлияет только на правильность вычислений, а скорость скорее даже уменьшиться.
Трудно советовать не зная параметров задачи, но обычно нужно оптимизировать в сторону больше операций в потоке, меньше пересылок данных. Например сгруппировать, чтобы в потоке(в ядре) обрабатывать сразу несколько нейронов (4,16,32 и более, зависит от параметров обращения в память) и\или возиться с кэшированием данных.

[Robotex]

А разве ядро не ждет пока завершится предыдущее?

[oleksandr_veles]

Если не выставлена спец. переменная окружения и архитектура(железо+драйвер) позволяет, то не ждёт.

Answer 4

[Максим Милаков]

Используйте CUBLAS