OpenCL ядрото се изпълнява по-бавно от единична нишка

Всичко, написах много просто OpenCL ядро, което трансформира RGB изображение в сива скала, използвайки просто осредняване.

Малко предистория:

1. Изображението се съхранява в картографирана памет като 24-битов, неподплатен блок памет
2. Изходният масив се съхранява в фиксирана памет (картирана с clEnqueueMapBuffer) и е 8 bpp
3. Има два буфера, разпределени на устройството (clCreateBuffer), единият е специално прочетен (в който ние clWriteBuffer влизаме преди стартирането на ядрото), а другият е специално записван (който ние clReadBuffer след като ядрото приключи)

Пускам това на изображение 1280x960. Серийната версия на алгоритъма е средно 60ms, OpenCL ядрото е средно 200ms!!! Правя нещо нередно, но нямам идея как да продължа, какво да оптимизирам. (Времето ми чете/записва без извикване на ядрото, алгоритъмът работи за 15 ms)

Прилагам настройката на ядрото (размери и аргументи), както и ядрото

РЕДАКТИРАНЕ: Затова написах още по-глупаво ядро, което няма глобален достъп до паметта вътре в него и беше само 150 ms... Това все още е абсурдно бавно. Помислих си, че може би бъркам с четенията на глобалната памет, те трябва да са подравнени в 4 байта или нещо подобно? Не...

Редактиране 2: Премахването на всички параметри от моето ядро ​​ми даде значително ускорение... Объркан съм, мислех си, че тъй като съм clEnqueueWriteBuffer, ядрото не трябва да извършва прехвърляне на памет от хост->устройство и устройство->хост....

Редактиране 3: Разбрах го, но все още не разбирам защо. Ако някой може да го обясни, ще се радвам да му дам правилния отговор. Проблемът беше предаването на персонализираните структури по стойност. Изглежда, че ще трябва да заделя глобална памет за тях и да предам техните cl_mems

Извикване на ядрото:

//Copy input to device
result = clEnqueueWriteBuffer(handles->queue, d_input_data, CL_TRUE, 0, h_input.widthStep*h_input.height, (void *)input->imageData, 0, 0, 0);
if(check_result(result, "opencl_rgb_to_gray", "Failed to write to input buffer on device!")) return 0;

//Set kernel arguments
result = clSetKernelArg(handles->current_kernel, 0, sizeof(OpenCLImage), (void *)&h_input);
if(check_result(result, "opencl_rgb_to_gray", "Failed to set input struct.")) return 0;
result = clSetKernelArg(handles->current_kernel, 1, sizeof(cl_mem), (void *)&d_input_data);
if(check_result(result, "opencl_rgb_to_gray", "Failed to set input data.")) return 0;
result = clSetKernelArg(handles->current_kernel, 2, sizeof(OpenCLImage), (void *)&h_output);
if(check_result(result, "opencl_rgb_to_gray", "Failed to set output struct.")) return 0;
result = clSetKernelArg(handles->current_kernel, 3, sizeof(cl_mem), (void *)&d_output_data);
if(check_result(result, "opencl_rgb_to_gray", "Failed to set output data.")) return 0;

//Determine run parameters
global_work_size[0] = input->width;//(unsigned int)((input->width / (float)local_work_size[0]) + 0.5);
global_work_size[1] = input->height;//(unsigned int)((input->height/ (float)local_work_size[1]) + 0.5);

printf("Global Work Group Size: %d %d\n", global_work_size[0], global_work_size[1]);

//Call kernel
result = clEnqueueNDRangeKernel(handles->queue, handles->current_kernel, 2, 0, global_work_size, local_work_size, 0, 0, 0);
if(check_result(result, "opencl_rgb_to_gray", "Failed to run kernel!")) return 0;

result = clFinish(handles->queue);
if(check_result(result, "opencl_rgb_to_gray", "Failed to finish!")) return 0;

//Copy output
result = clEnqueueReadBuffer(handles->queue, d_output_data, CL_TRUE, 0, h_output.widthStep*h_output.height, (void *)output->imageData, 0, 0, 0);
if(check_result(result, "opencl_rgb_to_gray", "Failed to write to output buffer on device!")) return 0;

Ядро:

typedef struct OpenCLImage_t
{
    int width;
    int widthStep;
    int height;
    int channels;
} OpenCLImage;

__kernel void opencl_rgb_kernel(OpenCLImage input, __global unsigned char*  input_data, OpenCLImage output, __global unsigned char * output_data)
{
    int pixel_x = get_global_id(0);
    int pixel_y = get_global_id(1);
    unsigned char * cur_in_pixel, *cur_out_pixel;
    float avg = 0;

    cur_in_pixel = (unsigned char *)(input_data + pixel_y*input.widthStep + pixel_x * input.channels);
    cur_out_pixel = (unsigned char *)(output_data + pixel_y*output.widthStep + pixel_x * output.channels);

    avg += cur_in_pixel[0];
    avg += cur_in_pixel[1];
    avg+= cur_in_pixel[2];
    avg /=3.0f;

    if(avg > 255.0)
        avg = 255.0;
    else if(avg < 0)
        avg = 0;

    *cur_out_pixel = avg;
}