c - OpenCL 内核执行速度比单线程慢

Question

总而言之，我编写了一个非常简单的 OpenCL 内核，它使用简单的平均将 RGB 图像转换为灰度图像。

一些背景:

1. 图像以 24 位非填充内存块的形式存储在映射内存中
2. 输出数组存储在固定内存中(使用 clEnqueueMapBuffer 映射)并且为 8 bpp
3. 设备上分配了两个缓冲区(clCreateBuffer)，一个专门读(我们在内核启动前clWriteBuffer写入)，一个专门写(我们在内核完成后 clReadBuffer)

我在 1280x960 图像上运行它。该算法的串行版本平均为 60 毫秒，OpenCL 内核平均为 200 毫秒!!!我做错了什么，但我不知道如何进行，要优化什么。 (在没有内核调用的情况下对我的读/写进行计时，该算法在 15 毫秒内运行)

我正在附加内核设置(大小和参数)以及内核

---

编辑:所以我写了一个更笨的内核，里面没有全局内存访问，而且只有 150 毫秒……这仍然慢得离谱。我想也许我搞砸了全局内存读取，它们必须是 4 字节对齐的还是什么？不……

编辑 2: 从我的内核中删除所有参数使我的速度显着提高......我很困惑我认为因为我是 clEnqueueWriteBuffer 内核应该不从主机->设备和设备->主机进行内存传输....

编辑 3: 弄清楚了，但我仍然不明白为什么。如果有人可以解释，我很乐意将正确答案授予他们。问题是按值传递自定义结构。看起来我需要为它们分配一个全局内存位置并传递它们的 cl_mem

---

内核调用:

//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;
}

Answer 1

将值复制到将要创建的所有线程的开销可能是时间的可能原因；至于全局内存，在另一种情况下引用就足够了。只有 SDK 实现者才能准确回答.. :)