Оптимизация выполнения времени ядра Cuda

Question

Я учащийся Cuda, и я хотел бы оптимизировать время выполнения моей функции ядра. В результате я понял короткую программу, вычисляющую разницу между двумя картинками. Поэтому я сравнил время выполнения между классическим исполнением процессора в C, и выполнение GPU в Cuda C.

Здесь вы можете найти код я говорю:

int *imgresult_data = (int *) malloc(width*height*sizeof(int)); 
int size = width*height; 

switch(computing_type) 
{ 

    case GPU: 

    HANDLE_ERROR(cudaMalloc((void**)&dev_data1, size*sizeof(unsigned char))); 
    HANDLE_ERROR(cudaMalloc((void**)&dev_data2, size*sizeof(unsigned char))); 
    HANDLE_ERROR(cudaMalloc((void**)&dev_data_res, size*sizeof(int))); 

    HANDLE_ERROR(cudaMemcpy(dev_data1, img1_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice)); 
    HANDLE_ERROR(cudaMemcpy(dev_data2, img2_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice)); 
    HANDLE_ERROR(cudaMemcpy(dev_data_res, imgresult_data, size*sizeof(int), cudaMemcpyHostToDevice)); 

    float time; 
    cudaEvent_t start, stop; 

    HANDLE_ERROR(cudaEventCreate(&start)); 
    HANDLE_ERROR(cudaEventCreate(&stop)); 
    HANDLE_ERROR(cudaEventRecord(start, 0)); 

    for(int m = 0; m < nb_loops ; m++) 
    { 
     diff<<<height, width>>>(dev_data1, dev_data2, dev_data_res); 
    } 

    HANDLE_ERROR(cudaEventRecord(stop, 0)); 
    HANDLE_ERROR(cudaEventSynchronize(stop)); 
    HANDLE_ERROR(cudaEventElapsedTime(&time, start, stop)); 

    HANDLE_ERROR(cudaMemcpy(imgresult_data, dev_data_res, size*sizeof(int), cudaMemcpyDeviceToHost)); 

    printf("Time to generate: %4.4f ms \n", time/nb_loops); 

    break; 

    case CPU: 

    clock_t begin = clock(), diff; 

    for (int z=0; z<nb_loops; z++) 
    { 
     // Apply the difference between 2 images 
     for (int i = 0; i < height; i++) 
     { 
      tmp = i*imgresult_pitch; 
      for (int j = 0; j < width; j++) 
      { 
       imgresult_data[j + tmp] = (int) img2_data[j + tmp] - (int) img1_data[j + tmp]; 
      } 
     } 
    } 
    diff = clock() - begin; 

    float msec = diff*1000/CLOCKS_PER_SEC; 
    msec = msec/nb_loops; 
    printf("Time taken %4.4f milliseconds", msec); 

    break; 
} 

И вот мое ядро функция:

__global__ void diff(unsigned char *data1 ,unsigned char *data2, int *data_res) 
{ 
    int row = blockIdx.x; 
    int col = threadIdx.x; 
    int v = col + row*blockDim.x; 

    if (row < MAX_H && col < MAX_W) 
    { 
     data_res[v] = (int) data2[v] - (int) data1[v]; 
    } 
} 

Я получил это время выполнения для каждого один

• CPU: 1,3210ms
• GPU: 0,3229ms

Интересно, почему результат GPU не так ниже, как это должно быть. Я новичок в Cuda, поэтому, пожалуйста, будьте всеобъемлющими, если есть некоторые классические ошибки.

EDIT1: Благодарим вас за отзыв. Я попытался удалить условие «if» из ядра, но это не сильно изменило время выполнения моей программы.

Однако, после установки профилировщика Cuda, он сказал мне, что мои потоки не запускались одновременно. Я не понимаю, почему у меня такое сообщение, но это кажется правдой, потому что у меня есть только 5 или 6 раз быстрее приложение с графическим процессором, чем с процессором. Это отношение должно быть больше, поскольку каждый поток должен обрабатывать один пиксель одновременно со всеми другими. Если у вас есть представление о том, что я делаю неправильно, это было бы полезно ...

Поток.

Answer 1

Возможно, есть другие проблемы с кодом, но вот что я вижу. Следующие строки __global__ void diff считаются не оптимален:

if (row < MAX_H && col < MAX_W) 
{ 
    data_res[v] = (int) data2[v] - (int) data1[v]; 
} 

Условные операторы в результате ядра в искривления дивергенции. Это означает, что if и else части внутри деформации выполняются последовательно, а не параллельно. Кроме того, как вы, возможно, поняли, if оценивает false только на границах. Для того, чтобы избежать расхождения и ненужных вычислений, разделить изображение на две части:

1. Центральная часть, где row < MAX_H && col < MAX_W всегда true. Создайте дополнительное ядро ​​для этой области. if здесь не нужно.

2. Области границы, которые будут использовать ваше ядро ​​diff.

Очевидно, что вы измените свой код, который вызывает ядра.

---

И на отдельной ноте:

1. GPU имеет пропускную-ориентированную архитектуру, но не латентности-ориентированный как центральный процессор. Это означает, что процессор может быть быстрее, чем CUDA, когда дело доходит до обработки небольших объемов данных. Вы пробовали использовать большие наборы данных?

2. CUDA Profiler - очень удобный инструмент, который скажет, что вы не оптимальны в коде.

Answer 2

Я не думаю, что вы правильно измеряете время, копия памяти занимает много времени в графическом процессоре, которую вы должны учитывать при измерении вашего времени.

Я вижу некоторые детали, которые вы можете проверить:

1. Я полагаю, вы используете MAX_H и MAX_H как константы, вы можете рассмотреть возможность сделать это с помощью cudaMemcpyToSymbol().

2. Не забудьте синхронизировать свои потоки, используя __syncthreads(), чтобы не возникало проблем между каждой итерацией цикла.

3. CUDA работает с перекосами, поэтому блок и количество потоков на блок работают лучше, чем кратные 8, но не более 512 потоков на блок, если ваше оборудование не поддерживает его. Ниже приведен пример использования 128 потоков на блок: < < < (cols * rows + 127)/128,128 >>>.

4. Помните также, чтобы освободить выделенную память в графическом процессоре и уничтожить созданные события.

5. В вашей функции ядра вы можете иметь одну переменную int v = threadIdx.x + blockIdx.x * blockDim.x.

6. Испытывали ли вы, помимо времени выполнения, что ваш результат верен? Я думаю, вы должны использовать cudaMallocPitch() и cudaMemcpy2D() во время работы с массивами из-за заполнения.