C vs OpenCL, как сравнить результаты измерения времени?

Итак, в другом посте я спросил об измерении времени C. Теперь я хочу знать, как сравнить результат «функции» C и «функции» OpenCL.

Это код хоста OpenCL и C

#define PROGRAM_FILE "sum.cl"
#define KERNEL_FUNC "float_sum"
#define ARRAY_SIZE 1000000


#include <stdio.h>
#include <stdlib.h>
#include <time.h>

#include <CL/cl.h>

int main()
{
    /* OpenCL Data structures */

    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_program program;
    cl_kernel kernel;    
    cl_command_queue queue;
    cl_mem vec_buffer, result_buffer;

    cl_event prof_event;;

    /* ********************* */

    /* C Data Structures / Data types */
    FILE *program_handle; //Kernel file handle
    char *program_buffer; //Kernel buffer

    float *vec, *non_parallel;
    float result[ARRAY_SIZE];

    size_t program_size; //Kernel file size

    cl_ulong time_start, time_end, total_time;

    int i;
    /* ****************************** */

    /* Errors */
    cl_int err;
    /* ****** */

    non_parallel = (float*)malloc(ARRAY_SIZE * sizeof(float));
    vec          = (float*)malloc(ARRAY_SIZE * sizeof(float));

    //Initialize the vector of floats
    for(i = 0; i < ARRAY_SIZE; i++)
    vec[i] = i + 1;

    /************************* C Function **************************************/
    clock_t start, end;

    start = clock();

    for( i = 0; i < ARRAY_SIZE; i++) 
    {
    non_parallel[i] = vec[i] * vec[i];
    }
    end = clock();
    printf( "Number of seconds: %f\n", (clock()-start)/(double)CLOCKS_PER_SEC );

    free(non_parallel);
    /***************************************************************************/




    clGetPlatformIDs(1, &platform, NULL);//Just want NVIDIA platform
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);

    // Context error?
    if(err)
    {
    perror("Cannot create context");
    return 1;
    }

    //Read the kernel file
    program_handle = fopen(PROGRAM_FILE,"r");
    fseek(program_handle, 0, SEEK_END);
    program_size = ftell(program_handle);
    rewind(program_handle);

    program_buffer = (char*)malloc(program_size + 1);
    program_buffer[program_size] = '\0';
    fread(program_buffer, sizeof(char), program_size, program_handle);
    fclose(program_handle);

    //Create the program
    program = clCreateProgramWithSource(context, 1, (const char**)&program_buffer, 
                    &program_size, &err);

    if(err)
    {
    perror("Cannot create program");
    return 1;
    }

    free(program_buffer);

    clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

    kernel = clCreateKernel(program, KERNEL_FUNC, &err);

    if(err)
    {
    perror("Cannot create kernel");
    return 1;
    }

    queue = clCreateCommandQueue(context, device, CL_QUEU_PROFILING_ENABLE, &err);

    if(err)
    {
    perror("Cannot create command queue");
    return 1;
    }

    vec_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                sizeof(float) * ARRAY_SIZE, vec, &err);
    result_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float)*ARRAY_SIZE, NULL, &err);

    if(err)
    {
    perror("Cannot create the vector buffer");
    return 1;
    }

    clSetKernelArg(kernel, 0, sizeof(cl_mem), &vec_buffer);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &result_buffer);

    size_t global_size = ARRAY_SIZE;
    size_t local_size = 0;

    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, &prof_event);

    clEnqueueReadBuffer(queue, result_buffer, CL_TRUE, 0, sizeof(float)*ARRAY_SIZE, &result, 0, NULL, NULL);
    clFinish(queue);



     clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_START,
           sizeof(time_start), &time_start, NULL);
     clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_END,
           sizeof(time_end), &time_end, NULL);
     total_time += time_end - time_start;

    printf("\nAverage time in nanoseconds = %lu\n", total_time/ARRAY_SIZE);



    clReleaseMemObject(vec_buffer);
    clReleaseMemObject(result_buffer);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(queue);
    clReleaseProgram(program);
    clReleaseContext(context);

    free(vec);

    return 0;
}

А ядро ​​такое:

__kernel void float_sum(__global float* vec,__global float* result){
    int gid = get_global_id(0);
    result[gid] = vec[gid] * vec[gid];
}

Теперь результаты таковы:

Количество секунд: 0,010000 ‹- Это для кода C

Среднее время в наносекундах = 140737284 ‹- Функция OpenCL

0,1407 секунд - это время выполнения ядра времени OpenCL, и это больше, чем функция C, это правильно? Потому что я думаю, что OpenCL должен быть быстрее, чем непараллельный алгоритм C...


person Community    schedule 14.04.2012    source источник
comment
Я очень удивлен этими результатами, тем более, что вы скорость OpenCL делите на размер массива. Вы уверены, что правильно синхронизируете коды? Вы используете Windows или Linux? Какой графический процессор вы используете?   -  person KLee1    schedule 15.04.2012
comment
Вы экспериментируете с этим конкретным примером? Если вы используете тип float4, вы можете произвести точечный продукт и суммировать 4 значения за одну операцию. Я ответил ниже, если вы ищете не такую ​​​​оптимизацию, а указатели opencl. Существует также формула суммы квадратов, которую вы можете использовать.   -  person mfa    schedule 26.10.2012


Ответы (3)


Выполнение параллельного кода на GPU не обязательно быстрее, чем выполнение на CPU. Учтите, что вам также необходимо передавать данные в память графического процессора и из нее в дополнение к вычислениям.

В вашем примере вы передаете 2 * N элементов и выполняете операцию O (N) параллельно, что является очень неэффективным использованием графического процессора. Поэтому вполне вероятно, что процессор действительно быстрее для этого конкретного вычисления.

person Tudor    schedule 14.04.2012
comment
Существует ли способ сделать ядро ​​более эффективным? Я также думаю распараллелить данные... - person ; 14.04.2012
comment
@факундо: Не совсем. Боюсь, ваш единственный вариант — запустить более значительные вычисления на графическом процессоре. Например, вы можете попробовать реализовать умножение матриц. - person Tudor; 14.04.2012
comment
Это действительно не имеет никакого смысла. Для этого измерялось только время выполнения ядра. Не учитывает передачу данных. - person KLee1; 15.04.2012
comment
Скромный графический процессор часто побеждает высокопроизводительный процессор в этой задаче. время передачи не измерялось, но было всего 4М вверх/вниз. Пример немного надуман для начала. сумма первых N квадратов = (N * (N + 1) * (2N + 1)) / 6 - person mfa; 26.10.2012

Просто для других, обращающихся к ней за помощью: краткое введение в профилирование среды выполнения ядра с помощью OpenCL

Включить режим профилирования:

cmdQueue = clCreateCommandQueue(context, *devices, CL_QUEUE_PROFILING_ENABLE, &err);

Ядро профилирования:

cl_event prof_event; 
clEnqueueNDRangeKernel(cmdQueue, kernel, 1 , 0, globalWorkSize, NULL, 0, NULL, &prof_event);

Чтение данных профилирования в:

cl_ulong ev_start_time=(cl_ulong)0;     
cl_ulong ev_end_time=(cl_ulong)0;   

clFinish(cmdQueue);
err = clWaitForEvents(1, &prof_event);
err |= clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &ev_start_time, NULL);
err |= clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &ev_end_time, NULL);

Рассчитать время выполнения ядра:

float run_time_gpu = (float)(ev_end_time - ev_start_time)/1000; // in usec

Ваш подход с

total_time/ARRAY_SIZE

это не то, что вы хотите. Это даст вам время выполнения для каждого рабочего элемента.

Операции/время в наносекундах дадут вам GFLOPS (гига операций с плавающей запятой в секунду).

person jaba    schedule 25.10.2012

Это одна большая проблема с вашим приложением:

size_t global_size = ARRAY_SIZE;
size_t local_size = 0;

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

Вместо этого попробуйте следующее:

size_t global_size = ARRAY_SIZE / 250; //important: local_size*global_size must equal ARRAY_SIZE
size_t local_size = 250; //a power of 2 works well. 250 is close enough, yes divisible by your 1M array size

Теперь вы создаете большие группы, которые лучше насыщают АЛУ вашего графического оборудования. Ядро будет работать нормально в том виде, в котором оно есть сейчас, но есть и способы получить больше от части ядра.

Оптимизация ядра: передать ARRAY_SIZE в ядро ​​в качестве дополнительного параметра и использовать меньше групп более оптимального размера. Вы также избавитесь от необходимости, чтобы local_size*global_size был точно равен ARRAY_SIZE. Глобальный идентификатор рабочих элементов никогда не используется в этом ядре, и он не нужен, поскольку был передан общий размер.

__kernel void float_sum(__global float* vec,__global float* result,__global int count){
  int lId = get_local_id(0);
  int lSize = get_local_size(0);
  int grId = get_group_id(0);
  int totalOps = count/get_num_groups(0);
  int startIndex = grId * totalOps;
  int maxIndex = startIndex+totalOps;
  if(grId == get_num_groups(0)-1){
    endIndex = count;
  }
  for(int i=startIndex+lId;i<endIndex;i+=lSize){
    result[i] = vec[i] * vec[i];
  }
}

Теперь вы можете подумать, что для такого простого ядра очень много переменных. Помните, что каждое выполнение ядра будет выполнять несколько операций с данными, а не только одну. Используя приведенные ниже значения, на моем Radeon 5870 (20 вычислительных блоков) каждый рабочий элемент в конечном итоге вычисляет 781 или 782 значения в своем цикле for. Каждая группа будет вычислять 50000 единиц данных. Накладные расходы на переменные, которые я использую, намного меньше, чем накладные расходы на создание 4000 рабочих групп — или 1 миллиона.

size_t global_size = ARRAY_SIZE / numComputeUnits;
size_t local_size = 64; //also try other multiples of 16 or 64 for gpu; or multiples of your core-count for a cpu kernel

См. здесь о том, как получить значение для numComputeUnits

person mfa    schedule 26.10.2012