Почему этот код в десять раз медленнее, чем у процессора?

Question

У меня проблема, которая сводится к выполнению некоторой арифметики для каждого элемента набора матриц. Я думал, что это похоже на то, какое вычисление может сильно повлиять на перенос на GPU. Тем не менее, мне удалось только замедлить вычисление в 10 раз!

Вот специфика моей тестовой системы:

• ОС: Windows 10
• CPU: Ядро i7-4700MQ @ 2,40 ГГц
• GPU: GeForce GT 750M (вычислительные возможности 3,0)
• CUDA SDK: v7.5

Нижеприведенный код выполняет эквивалентные вычисления в моем производственном коде, на процессоре и на графическом процессоре. Последний последовательно в десять раз медленнее на моей машине (процессор примерно 650 мс, GPU около 7 с).

Я попытался изменить размеры сетки и блока; Я увеличил и уменьшил размер массива, переданного на GPU; Я запустил его через визуальный профайлер; Я пробовал целочисленные данные, а не удваивал, но что бы я ни делал, версия графического процессора всегда значительно медленнее, чем эквивалент процессора.

Так почему же версия GPU настолько медленнее, и какие изменения, о которых я не упоминал выше, могу ли я попытаться улучшить ее производительность?

Вот моя командная строка: nvcc source.cu -o CPUSpeedTest.exe -arch=sm_30

А вот содержимое source.cu:

#include <iostream> 
#include <windows.h> 
#include <cuda_runtime_api.h> 

void AdjustArrayOnCPU(double factor1, double factor2, double factor3, double denominator, double* array, int arrayLength, double* curve, int curveLength) 
{ 
    for (size_t i = 0; i < arrayLength; i++) 
    { 
     double adjustmentFactor = factor1 * factor2 * factor3 * (curve[i]/denominator); 
     array[i] = array[i] * adjustmentFactor; 
    } 
} 

__global__ void CudaKernel(double factor1, double factor2, double factor3, double denominator, double* array, int arrayLength, double* curve, int curveLength) 
{ 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 

    if (idx < arrayLength) 
    { 
     double adjustmentFactor = factor1 * factor2 * factor3 * (curve[idx]/denominator); 
     array[idx] = array[idx] * adjustmentFactor; 
    } 
} 

void AdjustArrayOnGPU(double array[], int arrayLength, double factor1, double factor2, double factor3, double denominator, double curve[], int curveLength) 
{ 
    double *dev_row, *dev_curve; 

    cudaMalloc((void**)&dev_row, sizeof(double) * arrayLength); 
    cudaMalloc((void**)&dev_curve, sizeof(double) * curveLength); 

    cudaMemcpy(dev_row, array, sizeof(double) * arrayLength, cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_curve, curve, sizeof(double) * curveLength, cudaMemcpyHostToDevice); 

    CudaKernel<<<100, 1000>>>(factor1, factor2, factor3, denominator, dev_row, arrayLength, dev_curve, curveLength); 

    cudaMemcpy(array, dev_row, sizeof(double) * arrayLength, cudaMemcpyDeviceToHost); 

    cudaFree(dev_curve); 
    cudaFree(dev_row); 
} 

void FillArray(int length, double row[]) 
{ 
    for (size_t i = 0; i < length; i++) row[i] = 0.1 + i; 
} 

int main(void) 
{ 
    const int arrayLength = 10000; 

    double arrayForCPU[arrayLength], curve1[arrayLength], arrayForGPU[arrayLength], curve2[arrayLength];; 

    FillArray(arrayLength, curve1); 
    FillArray(arrayLength, curve2); 

    ///////////////////////////////////// CPU Version //////////////////////////////////////// 

    LARGE_INTEGER StartingTime, EndingTime, ElapsedMilliseconds, Frequency; 

    QueryPerformanceFrequency(&Frequency); 
    QueryPerformanceCounter(&StartingTime); 

    for (size_t iterations = 0; iterations < 10000; iterations++) 
    { 
     FillArray(arrayLength, arrayForCPU); 
     AdjustArrayOnCPU(1.0, 1.0, 1.0, 1.0, arrayForCPU, 10000, curve1, 10000); 
    } 

    QueryPerformanceCounter(&EndingTime); 

    ElapsedMilliseconds.QuadPart = EndingTime.QuadPart - StartingTime.QuadPart; 
    ElapsedMilliseconds.QuadPart *= 1000; 
    ElapsedMilliseconds.QuadPart /= Frequency.QuadPart; 
    std::cout << "Elapsed Milliseconds: " << ElapsedMilliseconds.QuadPart << std::endl; 

    ///////////////////////////////////// GPU Version //////////////////////////////////////// 

    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 

    cudaEventRecord(start); 

    for (size_t iterations = 0; iterations < 10000; iterations++) 
    { 
     FillArray(arrayLength, arrayForGPU); 
     AdjustArrayOnGPU(arrayForGPU, 10000, 1.0, 1.0, 1.0, 1.0, curve2, 10000); 
    } 

    cudaEventRecord(stop); 
    cudaEventSynchronize(stop); 

    float elapsedTime; 
    cudaEventElapsedTime(&elapsedTime, start, stop); 

    std::cout << "CUDA Elapsed Milliseconds: " << elapsedTime << std::endl; 

    cudaEventDestroy(start); 
    cudaEventDestroy(stop); 

    return 0; 
} 

А вот пример вывода CUDASpeedTest.exe

Elapsed Milliseconds: 565 
CUDA Elapsed Milliseconds: 7156.76 

Answer 1

Что следует, вероятно, смущает очевидность для большинства разработчиков, работающих с CUDA, но может быть полезной для других, таких как я, - которые являются новыми для этой технологии.

Код графического процессора в десять раз медленнее, чем эквивалент ЦП, поскольку код графического процессора демонстрирует идеальный шторм характеристик разрушения производительности.

Код GPU большую часть времени выделяет память на графическом процессоре, копируя данные на устройство, выполняя очень простой расчет (что в высшей степени быстро, независимо от типа процессора, на котором он работает), а затем копирование данных от устройства до хоста.

Как отмечено в комментариях, если существует верхняя граница размера обрабатываемых структур данных, тогда буфер на графическом процессоре может быть выделен ровно один раз и повторно использован. В приведенном выше коде это приводит к тому, что графический процессор переходит в режим работы процессора с 10: 1 до 4: 1.

Оставшееся несоответствие производительности связано с тем, что ЦП способен выполнять требуемые вычисления в последовательном порядке, миллионы раз за очень короткий промежуток времени из-за его простоты. В приведенном выше коде вычисление включает в себя считывание значения из массива, некоторого умножения и, наконец, назначение элементу массива. Что-то такое простое должно выполняться миллионы раз , прежде чем выгоды от этого параллельны, перевешивают необходимый штраф времени для передачи данных на GPU и обратно.В моей тестовой системе миллион элементов массива - это точка безубыточности, где GPU и CPU выполняют (приблизительно) столько же времени.