Ядро OpenCL выполняется медленнее, чем один поток

Все, я написал очень простое ядро ​​OpenCL, которое преобразует RGB-изображение в оттенки серого с помощью простого усреднения.

Немного фона:

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

Я запускаю это на изображении 1280x960. Серийная версия алгоритма в среднем 60 мс, ядро ​​OpenCL в среднем 200 мс!!! Я делаю что-то не так, но я понятия не имею, как действовать, что оптимизировать. (Рассчитывая время чтения/записи без вызова ядра, алгоритм работает за 15 мс)

Я прилагаю настройку ядра (размеры и аргументы), а также ядро

РЕДАКТИРОВАТЬ: Итак, я написал еще более тупое ядро, в котором нет доступа к глобальной памяти внутри него, и это было всего 150 мс... Это все еще смехотворно медленно. Я подумал, может быть, я ошибаюсь с чтением глобальной памяти, они должны быть выровнены по 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;
}