Производительность добавления 2D-массивов на GPU и CPU

Question

В настоящее время я экспериментирую с производительностью кода OpenCL с использованием графического процессора и с использованием C++ на процессоре. Я написал программы, которые вычисляют сумму z = x + y, где z, x и y - двумерные массивы (матрицы) для GPU и CPU. После тестирования этих программ я обнаружил, что процессор намного эффективнее вычисляет эту сумму, чем графический процессор из-за медленной передачи данных в шине PCI между GPU и CPU. Теперь я хочу определить, сколько еще потребуется больших сумм, чтобы сделать использование GPU более эффективным, чем процессор. Я планирую сделать это, увеличив сумму z = x + y до z = x + y + y + y + y + ... и так далее.

Можно ли сделать использование GPU более эффективным, чем процессор, просто увеличив количество сумм для этой конкретной проблемы?

Как и FYI: Я использую графическую карту nVIDIA GeForce GT 640 и i5 процессор Intel Core.

Любая помощь будет принята с благодарностью.

EDIT:

Ниже я прилагаю мой код на CPU:

int main(int argc, const char * argv[]) 
{ 

    //This value determines the size of the nxn (square array)    
    int n = 1000; 

    //Allocating the memory for the nxn arrays of floats. 
    float **x = (float**)malloc(sizeof(float*)*n); 
    float **y = (float**)malloc(sizeof(float*)*n); 
    float **z = (float**)malloc(sizeof(float*)*n); 


    //Initializing the arrays. 
    for(int i = 0; i<n; i++){ 
     x[i] = (float*)malloc(sizeof(float)*n); 
     y[i] = (float*)malloc(sizeof(float)*n); 
     z[i] = (float*)malloc(sizeof(float)*n); 

     for(int j = 0; j<n; j++){ 
      x[i][j] = i+j; 
      y[i][j] = i+j; 

     } 
    } 

    for(int i = 0; i<n; i++){ 
     for(int j = 0; j<n; j++){ 

      z[i][j] = x[i][j] + y[i][j]; 
      for(int k = 0; k < 100; k++){ 
       z[i][j] += y[i][j]; 
      } 
     } 
    } 

    return 0; 

} 

А вот C++ с использованием OpenCL: (используется для копирования данных и выполнить ядро ​​на GPU)

int n = 1000; 

for(int i = 0; i<n; i++) 
    { 
     //Writing the data from the host to the device 
     err = clEnqueueWriteBuffer(queue, d_xx, CL_TRUE, 0, sizeof(float)*n, h_xx[i], 0, NULL, NULL); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not write to buffer d_xx" << std::endl; 
      exit(1); 
     } 

     err = clEnqueueWriteBuffer(queue, d_yy, CL_TRUE, 0, sizeof(float)*n, h_yy[i], 0, NULL, NULL); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not write to buffer d_yy" << std::endl; 
      exit(1); 
     } 

     //Setting the Kernel Arguments 
     err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_xx); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not set kernel argument h_xx." << std::endl; 
      exit(1); 
     } 

     err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_yy); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not set kernel argument h_yy." << std::endl; 
      exit(1); 
     } 

     err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_zz); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not set kernel argument h_zz." << std::endl; 
     } 

     work_units_per_kernel = n; 

     //Executing the Kernel 
     err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &work_units_per_kernel, NULL, 0, NULL, NULL); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not execute kernel." << std::endl; 
      exit(1); 
     } 

     //Reading the Data from the Kernel 
     err = clEnqueueReadBuffer(queue, d_zz, CL_TRUE, 0, n*(sizeof(float)), h_zz[i], 0, NULL, NULL); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not read data from kernel." << std::endl; 
      exit(1); 
     } 

    } 

И, наконец, код ядра выполняется на GPU:

__kernel void arraysum(__global const float *d_aa, __global const float *d_bb, __global float *d_cc) 
{ 

    int i = get_global_id(0); 

    d_cc[i] = d_aa[i] + d_bb[i]; 


    for(int j = 0; j < 100; j++){ 
     d_cc[i] += d_bb[i]; 
    } 


} 

Answer 1

Для n = 1000 * 1000 вы дошли до того, что копирование, работа и копирование стоит того. Как отметил DarkZero, глобальная память НЕ оптимальна, поэтому, если вы можете кэшировать свою глобальную память в локальную память или память потоков и использовать локальные рабочие группы, это будет чрезвычайно полезно для CPU и графического процессора.

Начнем с ядра. d_cc ссылается 100 раз от Global Memory. Простым изменением в этом случае является кэширование глобальной памяти в поточную память, а затем в конце скопируйте локальную обратно в глобальную.

__kernel void arraysum(__global const float *d_aa, __global const float *d_bb, __global float *d_cc) 
{ 

    int i = get_global_id(0); 

    float t_d_cc = d_aa[i] + d_bb[i]; //make a thread only version of d_cc 

    for(int j = 0; j < 100; j++){ 
     t_d_cc += d_bb[i]; 
    } 

    d_cc[i] = t_d_cc; //copy the thread only back to global 
} 

Другим изменением, в зависимости от оборудования, является кеширование d_aa и d_bb в локальную память.Это позволяет OpenCL использовать пакетные копии из глобальной памяти. Это может быть немного сложнее, потому что каждое устройство OpenCL имеет разные размеры и кратность размеров локальной рабочей группы, которые можно использовать.

Например, у моего i5 максимальный размер рабочей группы 1024 и рабочая группа, кратная 1, поэтому для моих локальных рабочих групп может быть что угодно от 1 до 1024. Мой ATI-7970 имеет значения 256 и 64 соответственно, поэтому мой Местным рабочим группам должно быть 64, 128 и т. Д. Это гораздо более ограничительно.

__kernel void arraysum(__global const float *d_aa, 
         __local float *l_d_aa, 
         __global const float *d_bb, 
         __local float *l_d_bb, 
         __global float *d_cc, 
         __local float *l_d_cc) 
{ 

//In this example, the global_id(1) is the number of rows and global_id(0) is the columns 
//So when the kernel is called, the local work group size needs to be the size of the 
//number of columns 

int i = get_global_id(1)*get_global_size(0) + get_global_id(0); //Index of the row 
int j = get_local_id(0); 

l_d_aa[get_local_id(0)] = d_aa[i]; 
l_d_bb[get_local_id(0)] = d_bb[i]; 

read_mem_fence(CLK_LOCAL_MEM_FENCE); 

float l_d_cc[get_local_id(0)] = l_d_aa[get_local_id(0)] + l_d_bb[get_local_id(0)]; 

for(int j = 0; j < get_global_size(0); j++){ 
    l_d_cc[get_local_id(0)] += l_d_bb[j]; 
} 

d_cc[i] = l_d_cc[get_local_id(0)]; //copy the thread only back to global 

}

Я извиняюсь, если я получил алгоритм неправильно, но, надеюсь, он передает, как кэш глобальной памяти в локальную память. Опять же, на i5 размер локальной рабочей группы может быть от 1 до 1024, но ATI7970 ограничен размерами столбцов 64, 128 и т. Д.

Это концептуально сложнее, но производительность для OpenCL - это много, гораздо лучше при использовании этого подхода.

Сообщество, пожалуйста, не стесняйтесь очищать ядро.

Answer 2

Многие вещи замедляют вас:

1- Злоупотребление использованием глобальной памяти. Каждый доступ к глобальной памяти примерно в 400 раз медленнее, и вы используете только глобальную память (например, 200 операций чтения/записи). Глобальная память должна использоваться только для чтения в начале и писать в конце, а не как промежуточное значение.

2- Ваша длина N очень короткая. CPU завершится всего за 1000 инструкций, а все задержки в графическом процессоре намного медленнее. Поскольку копия размером 100 МБ намного эффективнее, чем 1-байтная копия, в операциях копирования есть накладные расходы.

3- Возможно, код ЦП оптимизируется компилятором на умножения, в то время как код GPU не может, поскольку он обращается к изменчивым переменным, таким как глобальные.

4 Память для чтения/записи на устройство очень дорого, если вы включите это в calc, CPU легко победит. Кроме того, создание буфера OpenCL и ядер очень дорого. Обратите внимание, что вы также используете блокирующие вызовы записи, это намного медленнее, чем неблокирующие вызовы.