OpenCL study - code03 rgb2gray

记录一下学习OpenCL的过程，下面更新关于rgb转换为gray的操作，需要OpenCL以及OpenCV

// grayscale.cl
__kernel void rgb_to_gray(__global const uchar* rgb_image,
                          __global uchar* gray_image,
                          const int width,
                          const int height)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int idx = y * width + x;

    if (x < width && y < height) {
        int rgb_index = idx * 3;
        uchar r = rgb_image[rgb_index + 0];
        uchar g = rgb_image[rgb_index + 1];
        uchar b = rgb_image[rgb_index + 2];

        uchar gray = (uchar)(0.299f * r + 0.587f * g + 0.114f * b);
        gray_image[idx] = gray;
    }
}

//
// Created by zixhu on 2025/7/26.
//

#ifndef COMPUTERVISION_RGBMAIN_H
#define COMPUTERVISION_RGBMAIN_H


#include <opencv2/opencv.hpp>
#include "../helper/opencl_helper.h"

int runRgb2gray() {
    // 1. 读取图像（BGR）
    cv::Mat img = cv::imread("../src/opencl/sources/img.png");
    if (img.empty()) {
        printf("Failed to load image.\n");
        return -1;
    }

    // 2. 转成 RGB（因为 OpenCL kernel 假设的是 RGB）
    cv::cvtColor(img, img, cv::COLOR_BGR2RGB);

    int width = img.cols;
    int height = img.rows;
    int size_rgb = width * height * 3;
    int size_gray = width * height;

    // 3. 初始化输出灰度图内存
    std::vector<uchar> gray_data(size_gray);

    // 4. 初始化 OpenCL
    OpenCLObjects ocl = init_opencl("grayscale.cl", "rgb_to_gray");
    cl_int err;

    // 5. 创建缓冲区
    cl_mem input_buf = clCreateBuffer(ocl.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                      size_rgb * sizeof(uchar), img.data, &err);
    CHECK_ERROR(err, "clCreateBuffer input");

    cl_mem output_buf = clCreateBuffer(ocl.context, CL_MEM_WRITE_ONLY,
                                       size_gray * sizeof(uchar), NULL, &err);
    CHECK_ERROR(err, "clCreateBuffer output");

    // 6. 设置 kernel 参数
    clSetKernelArg(ocl.kernel, 0, sizeof(cl_mem), &input_buf);
    clSetKernelArg(ocl.kernel, 1, sizeof(cl_mem), &output_buf);
    clSetKernelArg(ocl.kernel, 2, sizeof(int), &width);
    clSetKernelArg(ocl.kernel, 3, sizeof(int), &height);

    // 7. 启动 kernel
    size_t global_size[2] = { (size_t)width, (size_t)height };
    err = clEnqueueNDRangeKernel(ocl.queue, ocl.kernel, 2, NULL, global_size, NULL, 0, NULL, NULL);
    CHECK_ERROR(err, "clEnqueueNDRangeKernel");

    clFinish(ocl.queue);

    // 8. 拷贝结果回主机
    err = clEnqueueReadBuffer(ocl.queue, output_buf, CL_TRUE, 0,
                              size_gray * sizeof(uchar), gray_data.data(), 0, NULL, NULL);
    CHECK_ERROR(err, "clEnqueueReadBuffer");

    // 9. 构造灰度图并显示
    cv::Mat gray_img(height, width, CV_8UC1, gray_data.data());
    cv::imshow("Gray Image", gray_img);
    cv::waitKey(0);

    // 10. 清理
    clReleaseMemObject(input_buf);
    clReleaseMemObject(output_buf);
    release_opencl(&ocl);

    return 0;
}


#endif //COMPUTERVISION_RGBMAIN_H

//
// Created by zixhu on 2025/7/26.
//

#ifndef COMPUTERVISION_OPENCL_HELPER_H
#define COMPUTERVISION_OPENCL_HELPER_H


#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <OpenCL/opencl.h>

#define CHECK_ERROR(err, msg) \
    if (err != CL_SUCCESS) { \
        fprintf(stderr, "%s failed with error %d\n", msg, err); \
        exit(1); \
    }

typedef struct {
    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_command_queue queue;
    cl_program program;
    cl_kernel kernel;
} OpenCLObjects;

// 加载内核源码
char *read_source(const char *filename) {
    FILE *fp = fopen(filename, "r");
    if (!fp) {
        perror("Failed to open kernel file");
        exit(1);
    }
    fseek(fp, 0, SEEK_END);
    size_t size = ftell(fp);
    rewind(fp);
    char *source = (char *)malloc(size + 1);
    fread(source, 1, size, fp);
    source[size] = '\0';
    fclose(fp);
    return source;
}

// 初始化 OpenCL，并构建 kernel
OpenCLObjects init_opencl(const char *source_file, const char *kernel_name) {
    OpenCLObjects ocl;
    cl_int err;

    err = clGetPlatformIDs(1, &ocl.platform, NULL);
    CHECK_ERROR(err, "clGetPlatformIDs");

    err = clGetDeviceIDs(ocl.platform, CL_DEVICE_TYPE_DEFAULT, 1, &ocl.device, NULL);
    CHECK_ERROR(err, "clGetDeviceIDs");

    ocl.context = clCreateContext(NULL, 1, &ocl.device, NULL, NULL, &err);
    CHECK_ERROR(err, "clCreateContext");

    ocl.queue = clCreateCommandQueue(ocl.context, ocl.device, 0, &err);
    CHECK_ERROR(err, "clCreateCommandQueue");

    char *source = read_source(source_file);
    ocl.program = clCreateProgramWithSource(ocl.context, 1, (const char **)&source, NULL, &err);
    CHECK_ERROR(err, "clCreateProgramWithSource");

    err = clBuildProgram(ocl.program, 1, &ocl.device, NULL, NULL, NULL);
    if (err != CL_SUCCESS) {
        char log[4096];
        clGetProgramBuildInfo(ocl.program, ocl.device, CL_PROGRAM_BUILD_LOG, sizeof(log), log, NULL);
        fprintf(stderr, "Build Error:\n%s\n", log);
        exit(1);
    }

    ocl.kernel = clCreateKernel(ocl.program, kernel_name, &err);
    CHECK_ERROR(err, "clCreateKernel");

    free(source);
    return ocl;
}

void release_opencl(OpenCLObjects *ocl) {
    clReleaseKernel(ocl->kernel);
    clReleaseProgram(ocl->program);
    clReleaseCommandQueue(ocl->queue);
    clReleaseContext(ocl->context);
}

#endif //COMPUTERVISION_OPENCL_HELPER_H