OpenCL灰度Kernel仅读取首个像素问题求助

核心错误：图像尺寸参数传递错误 #

你在创建OpenCL图像对象时，错误地将int*类型的width和height指针直接传给了clCreateImage2D，而该函数要求的是size_t类型的图像宽高数值。这会导致设备上创建的图像尺寸被设置为指针的内存地址（一个超大的无符号整数），内核中get_image_width(inputImage)返回的是这个错误的超大值，所有线程的坐标(x,y)都远小于该值，读取时会因为内存布局错位，所有线程都读到同一块内存区域的数据，最终出现所有像素值相同的现象。

修复步骤 #

1. 修正图像对象创建代码 #

将width和height的取值转换为size_t类型后传入clCreateImage2D：

// 修正：取出width和height的实际值，转换为size_t类型
size_t img_width = static_cast<size_t>(*width);
size_t img_height = static_cast<size_t>(*height);

// 创建输入图像对象
cl_mem input_image = clCreateImage2D(*context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
                                    &input_format, img_width, img_height, 
                                    0, _input_data.data(), &err_num);
if(err_num != CL_SUCCESS){
    logger.log("Failed to create cl_image input_image mem object", Logger::LogLevel::ERROR);
}

// 创建输出图像对象
cl_mem output_image = clCreateImage2D(*context, CL_MEM_WRITE_ONLY, 
                                     &output_format, img_width, img_height, 
                                     0, nullptr, &err_num);
if(err_num != CL_SUCCESS){
    logger.log("Failed to create cl_image output_image mem object", Logger::LogLevel::ERROR);
}

2. 移除冗余的图像写入操作 #

你在创建输入图像时使用了CL_MEM_COPY_HOST_PTR标志，该标志已经会在创建时将主机端数据拷贝到设备内存，因此后续的clEnqueueWriteImage调用是冗余的，需要删除这部分代码。

3. 优化内核代码（可选但更规范） #

• 直接使用read_imagef返回的归一化像素值（0.0~1.0）计算灰度，无需乘以255；
• 输出图像为单通道CL_R，使用float1类型写入更匹配格式；
• 用全局工作尺寸判断线程范围，避免频繁调用get_image_width/get_image_height：

__kernel void grayscale(
    __read_only image2d_t inputImage,
    __write_only image2d_t outputImage)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if (x < get_global_size(0) && y < get_global_size(1)) {
        int2 coord = (int2)(x, y);
        float4 pixel = read_imagef(inputImage, 
                                   CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST, 
                                   coord);
        
        // 直接用归一化值计算灰度
        float gray = 0.299f * pixel.x + 0.587f * pixel.y + 0.114f * pixel.z;

        // 写入单通道灰度值
        write_imagef(outputImage, coord, (float1)(gray));
    }
}

4. 确保全局工作尺寸正确 #

执行内核时，global_work_size需要设置为图像的实际宽高：

size_t global_work_size[2] = {img_width, img_height};
err_num = clEnqueueNDRangeKernel(*command_queue, *kernel, 2, nullptr, 
                                global_work_size, nullptr, 0, nullptr, &kernel_event);

额外说明 #

• 输出图像格式使用CL_R + CL_FLOAT是可行的，但如果后续要转回8位灰度图，需要注意将浮点值（0.0_{1.0）转换为0}255的整数；
• 内核中的printf调试可以保留，但大量输出会影响性能，调试完成后建议移除。

内容的提问来源于stack exchange，提问作者Arief Kurniawan