[OpenCL] 读取一张raw图像(NV12格式)，转rgb格式并resize（双线性插值）操作，输出jpg图像，并统计GPU-kernel耗时。

已于 2024-07-08 16:07:44 修改
阅读量400 收藏
点赞数 7
文章标签： opencv c++ gpu算力 图像处理 计算机视觉
于 2024-07-08 16:06:08 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/AndreLe/article/details/140271235
版权

对于raw文件而言，读取其中的NV12格式的数据，需要利用二进制进行读取，再经过GPU处理后，保存输出的数据为jpg文件。

其中resize使用双线性插值法。

首先读取和保存图像接口：

1.data_io.h

 
#ifdef APPLE        //平台相关代码
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#endif
 
 
#include <opencv2/opencv.hpp>
#include <opencv2/highgui.hpp>
 
cl_int Load_image_from_NV12(const std::string &filename,cl_mem *imageObjects,cl_context context,int &img_h,int &img_w);
 
void Save_pic_by_opencv(const std::string &filename,unsigned char  *rd,int &img_h,int &img_w);

2.data_io.cpp

#include "data_io.h"
 
#include <iostream>
#include <fstream>
#include <sstream>
 
u_char * get_NV12_buffer(const std::string &filePath, int width, int height) 
{
    std::ifstream file(filePath, std::ios::binary | std::ios::ate);
    
    if (!file.is_open()) {
        std::cerr << "Error: Unable to open NV12 file " << filePath << std::endl;
        return NULL;
    }
 
    std::streamsize size = file.tellg();
    file.seekg(0, std::ios::beg);
 
    size=width*height*3/2;
    char *buffer=new char [size];
    
    if (!file.read(buffer, size)) {
        std::cerr << "Error: Unable to read NV12 data from file " << filePath << std::endl;
        return NULL;
    }
     
    return (u_char *)buffer;
}
 
cl_int Load_image_from_NV12(const std::string &filename,cl_mem *imageObjects,cl_context context,int &img_h,int &img_w)
{
 
    cl_int errNum;

    u_char * nv12data= get_NV12_buffer(filename, img_w, img_h);
     
    if(NULL==nv12data)
        return 1;
 
    int nv12_size=sizeof(u_char)*img_h*img_w*3/2;
    
    imageObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
                                    nv12_size, nv12data, &errNum);    
    return errNum;
}
 
void Save_pic_by_opencv(const std::string &filename,unsigned char  *rd,int &img_h,int &img_w)
{
   cv::Mat src_Image =cv::Mat(img_h, img_w, CV_8UC3, (void*)rd);
   cv::Mat dst_Image(img_h, img_w, CV_8UC3);
   cv::cvtColor(src_Image, dst_Image,cv::COLOR_RGB2BGR);
   cv::imwrite(filename, dst_Image);
}

3.main.cpp

#include <iostream>
#include <fstream>
#include <sstream>


#ifdef APPLE        //平台相关代码
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#endif

#include "data_io.h"



//在第一个平台中创建只包括GPU的上下文
cl_context CreateContext()
{
    cl_int errNum;
    cl_uint numPlatforms;
    cl_platform_id firstPlatformId;
    cl_context context = NULL;
    
    
    // 选择第一个平台
    errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
    if (errNum != CL_SUCCESS || numPlatforms <= 0)
    {
        std::cerr << "Failed to find any OpenCL platforms." << std::endl;
        return NULL;
    }
 
    //cout<<"numPlatforms="<<numPlatforms<<endl; 1

    // 接下来尝试通过GPU设备建立上下文
    cl_context_properties contextProperties[] =
    {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)firstPlatformId,
        0
    };
    context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
                                      NULL, NULL, &errNum);

    

    if (errNum != CL_SUCCESS)
    {
        std::cout << "Could not create GPU context, trying CPU..." << std::endl;
        context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU,
                                          NULL, NULL, &errNum);
        if (errNum != CL_SUCCESS)
        {
            std::cerr << "Failed to create an OpenCL GPU or CPU context." << std::endl;
            return NULL;
        }
    }
 
    return context;
}


//在第一个设备上创建命令队列
cl_command_queue CreateCommandQueue(cl_context context, cl_device_id *device)
{
    cl_int errNum;
    cl_device_id *devices;
    cl_command_queue commandQueue = NULL;
    size_t deviceBufferSize = -1;
 
    // 首先获得设备的信息
    errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Failed call to clGetContextInfo(...,GL_CONTEXT_DEVICES,...)";
        return NULL;
    }
 
    if (deviceBufferSize <= 0)
    {
        std::cerr << "No devices available.";
        return NULL;
    }
 
    //为设备分配内存
    devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
    errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Failed to get device IDs";
        return NULL;
    }
 
    // 选择第一个设备并为其创建命令队列
    commandQueue = clCreateCommandQueue(context, devices[0], CL_QUEUE_PROFILING_ENABLE, NULL);
    //commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
    if (commandQueue == NULL)
    {
        std::cerr << "Failed to create commandQueue for device 0";
        return NULL;
    }
    
    //释放信息
    *device = devices[0];
    delete [] devices;
    return commandQueue;
}
 

 //  创建OpenCL程序对象
cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
    cl_int errNum;
    cl_program program;
 
    std::ifstream kernelFile(fileName, std::ios::in);
    if (!kernelFile.is_open())
    {
        std::cerr << "Failed to open file for reading: " << fileName << std::endl;
        return NULL;
    }
 
    std::ostringstream oss;
    oss << kernelFile.rdbuf();
 
    std::string srcStdStr = oss.str();
    const char *srcStr = srcStdStr.c_str();
    program = clCreateProgramWithSource(context, 1,
                                        (const char**)&srcStr,
                                        NULL, NULL);
    if (program == NULL)
    {
        std::cerr << "Failed to create CL program from source." << std::endl;
        return NULL;
    }
 
    errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    if (errNum != CL_SUCCESS)
    {
        // 输出错误信息
        char buildLog[16384];
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
                              sizeof(buildLog), buildLog, NULL);
 
        std::cerr << "Error in kernel: " << std::endl;
        std::cerr << buildLog;
        clReleaseProgram(program);
        return NULL;
    }
 
    return program;
}


//获取最接近的倍数
size_t RoundUp(int groupSize, int globalSize)
{
    int r = globalSize % groupSize;
    if(r == 0)
    {
         return globalSize;
    }
    else
    {
         return globalSize + groupSize - r;
    }
}


//清除资源
void Cleanup(cl_context context, cl_command_queue commandQueue,
             cl_program program, cl_kernel kernel, cl_mem imageObjects[2])
{
    for (int i = 0; i < 2; i++)
    {
        if (imageObjects[i] != 0)
            clReleaseMemObject(imageObjects[i]);
    }
    if (commandQueue != 0)
        clReleaseCommandQueue(commandQueue);
 
    if (kernel != 0)
        clReleaseKernel(kernel);
 
    if (program != 0)
        clReleaseProgram(program);
    

    if (context != 0)
        clReleaseContext(context);
 
}

int main()
{
    
   

    //initial varibles
    cl_int errNum=0;
    cl_device_id device = 0;
    cl_context context = 0;
    cl_program program = 0;
    cl_kernel kernel = 0;
    cl_command_queue commandQueue = 0;
    cl_sampler sampler = 0;
    cl_mem imageObjects[2] = { 0, 0 };
    
    int img_w = 3840;
    int img_h = 2176; 

    size_t read_size=img_w*img_h*3;

    size_t g_w=0;
    size_t g_h=0;
    
    std::string input_file,output_jpgfile;
    char *kernel_file=new char [100];
    char *kernel_name=new char [100];
   

    cl_float scale_w=0.5;
    cl_float scale_h=0.5;

    int new_w=0;
    int new_h=0;
        
    input_file="../data/test.raw";
    output_jpgfile="../data/05yuv2rgb_resize.jpg";
    strcpy(kernel_file,"yuv2rgb_resize.cl");
    strcpy(kernel_name,"kernel_yuv2rgb_resize");
         
     // 1.选择platform,创建contex上下文
    context = CreateContext();
    if (context == NULL)
    {
        std::cerr << "Failed to create OpenCL context." << std::endl;
        return 1;
    }
    
    
    // 2.创建命令队列
    commandQueue = CreateCommandQueue(context, &device);
    if (commandQueue == NULL)
    {
        std::cerr <<"CreateCommandQueue failed"<<std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        return 1;
    }
    
    
    // 3. 确保设备支持这种图像格式
    cl_bool imageSupport = CL_FALSE;
    clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool),
                    &imageSupport, NULL);
    if (imageSupport != CL_TRUE)
    {
        std::cerr << "OpenCL device does not support images." << std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        
        return 1;
    }
    
     //4.准备输入数据 
     
    errNum=Load_image_from_NV12(input_file,imageObjects,context,img_h,img_w);
    
    if (errNum != CL_SUCCESS)
    { 
        std::cerr << "Load_image failed." << std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        return 1;
    }
    
    // 5.创建输出的图像对象
   
    new_w=scale_w*img_w;
    new_h=scale_h*img_h;
    read_size=sizeof(u_char)*3*new_w*new_h;
    
    imageObjects[1] = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 
                                    read_size, NULL, &errNum);
    
     if (errNum != CL_SUCCESS)
    { 
        std::cerr << "create output imageobject failed." << std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        return 1;

    }

 
    //6.创建OpenCL-program对象
    
    program = CreateProgram(context, device, kernel_file);
    if (program == NULL)
    {
        std::cerr <<"CreateProgram failed"<<std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        return 1;
    }
    
    // 7.创建OpenCL核

    kernel = clCreateKernel(program, kernel_name, NULL);
    
    if (kernel == NULL)
    {
        std::cerr << "Failed to create kernel" << std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        return 1;
    }
 
  
    //8. 设定参数
 
   errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &imageObjects[0]);
   errNum |= clSetKernelArg(kernel, 1, sizeof(cl_int), &img_w);
   errNum |= clSetKernelArg(kernel, 2, sizeof(cl_int), &img_h);
   errNum |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &imageObjects[1]);
   errNum |= clSetKernelArg(kernel, 4, sizeof(cl_float), &scale_w);
   errNum |= clSetKernelArg(kernel, 5, sizeof(cl_float), &scale_h);

    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Error setting kernel arguments." << std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        return 1;
    }
    
    g_w=new_w;
    g_h=new_h;
 
    size_t localWorkSize[2] = { 16, 16 };
    size_t globalWorkSize[2] =  { RoundUp(localWorkSize[0], g_w),
                                  RoundUp(localWorkSize[1], g_h) };
    
    //9.启动内核，内核执行完成后，会将evt置为CL_SUCCESS/CL_COMPLETE
    cl_event evt;
    
    errNum = clEnqueueNDRangeKernel(commandQueue, kernel,  
                               2, 0, globalWorkSize, localWorkSize, 
                               0, NULL, &evt);  
     if (errNum != CL_SUCCESS)
    {
        std::cerr << "Error clEnqueueNDRangeKernel errNum=." <<errNum<< std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        return 1;
    }

    clFinish(commandQueue);
    clWaitForEvents(1, &evt);   //等待命令事件发生
    
    cl_ulong time_start= (cl_ulong)0;
    cl_ulong time_end= (cl_ulong)0;
    double total_time;
    
    clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
    clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);
    clReleaseEvent(evt);
    

    total_time =(double)( time_end - time_start);
    printf("\nExecution time in milliseconds = %f ms\n", (total_time / 1000000.0) );

    //10.读回数据
      
    u_char *read_data = new u_char [read_size];
    errNum =clEnqueueReadBuffer(commandQueue, imageObjects[1], 
                            CL_TRUE, 
                            0, 
                            read_size, 
                            read_data, 0, NULL, NULL);
    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Error reading result buffer enume." <<errNum<< std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        return 1;
    }
     
    if (read_data == NULL)
    {
        std::cerr << "Error reading result buffer null." <<errNum<< std::endl;
        Cleanup(context, commandQueue, program, kernel, imageObjects);
        return 1;
    }
    
    //显示图像

    Save_pic_by_opencv(output_jpgfile,read_data,new_h,new_w);

    Cleanup(context, commandQueue, program, kernel, imageObjects);
   
    return 0;
}

最后，GPU处理的kernel函数

4.yuv2rgb_resize.cl

#define BRD_CHECKS(x, y, X, Y, ybased, ubased,src_rgb)            \
    if ((x) >= 0 && (x) < (X) && (y) >= 0 && (y) < (Y)) {        \
        uchar yy = ybased[x + y * X];                             \
        uchar uu = ubased[(y/2) * X + (x / 2) * 2];                \
        uchar vv = ubased[(y/2) * X + (x / 2) * 2 + 1];            \
        uchar r = yy + 1.402 * (vv - 128);                       \
        uchar g = yy - 0.34413 * (uu - 128) - 0.71414 * (vv - 128);\
        uchar b = yy + 1.772 * (uu - 128);                        \
        src_rgb =(uchar3)(r,g,b);                              \
    }


__kernel void kernel_yuv2rgb_resize(const __global uchar *vinput,
        int input_w, int input_h,
        __global uchar *vout,
        float scale_x, float scale_y){
      
    const int w = get_global_id(0);
    const int h = get_global_id(1);
    
     int WIDTH = input_w;
     int HEIGHT = input_h;
    
     int out_w=input_w*scale_x;

    uchar *ybase = vinput;
    uchar *ubase = &vinput[WIDTH * HEIGHT];

    
    const float src_x = (float)(w)/scale_x;
    const float src_y = (float)(h)/scale_y ;
    const int x1 = convert_int_rtn(src_x);
    const int y1 = convert_int_rtn(src_y);
    const int x2 = x1 + 1;
    const int y2 = y1 + 1;
    
    int bh = h;
    do {
        
        float3 out = 0;
        //双线性插值计算
        {
            uchar3 src_reg = 0;
            BRD_CHECKS(x1, y1, WIDTH, HEIGHT,vinput,ubase, src_reg);
            //convert_float3函数用于类型转换
            out = out + convert_float3(src_reg) * ((x2 - src_x) * (y2 - src_y));
        }
        {
            uchar3 src_reg = 0;
            BRD_CHECKS(x2, y1, WIDTH, HEIGHT,vinput,ubase,src_reg);
            out = out + convert_float3(src_reg) * ((src_x - x1) * (y2 - src_y));
        }
        {
            uchar3 src_reg = 0;
            BRD_CHECKS(x1, y2, WIDTH, HEIGHT,vinput,ubase,src_reg);
            out = out + convert_float3(src_reg) * ((x2 - src_x) * (src_y - y1));
        }
        {
            uchar3 src_reg = 0;
            BRD_CHECKS(x2, y2, WIDTH, HEIGHT,vinput,ubase,src_reg);
            out = out + convert_float3(src_reg) * ((src_x - x1) * (src_y - y1));
        }

        float3 rgb = out;

        vstore3(convert_uchar3_sat(rgb), mad24(bh, out_w, w), vout);

    } while (0);
}

AndreLe
关注
OpenCL应用与双线性插值
SE_and_CS
03-28 3220
一．  CPU实现和GPU实现 本实验基于的测试图像是8000*1000大小的8位灰度钢板图像，缩小倍数为0.5。 CPU上算法实现的核心代码 void func(BYTE*src, int width, int height, int pitchSrc, BYTE *dst, int t_width, int t_height, int pitchDes) { float f
[OpenCL] 读取一张jpg图像，将其resize（双线性插值）操作，并做耗时统计。
AndreLe的博客
07-08 444
读取一张jpg图像，并且利用gpu-opencl对其进行resize操作，并将结果保存到jpg。其中resize使用双线性插值方法。
opencl:C++实现双线性插值图像缩放
10km的专栏
02-27 8512
用OpenCL实现图像缩放代码是比较简单的，因为OpenCL本身就支持双线性插值 下面是kernel代码(从Mali OpenCL SDK 抄来的：/samples/image_scaling/assets/image_scaling.cl) 非常简单只有4行// 定义采样器 // CLK_NORMALIZED_COORDS_TRUE指定使用归一化坐标 // CLK_ADDRESS_CLAMP
C++实现NV12转RGBA(不借助opencv)
最新发布
qq_44908396的博客
08-08 404
使用C++实现NV12转RGBA
C++——NV12(YUV)格式转换为RGB/RGBA格式
qq_45848817的博客
07-09 598
YUV图片的NV12格式转换为RGB/RGBA图片格式。
NV12_clip_resize.c
06-05
项目中需要对NV12灰度图像进行缩放和裁剪处理，调试通过，可以运行
nv12转rgb
dss875914213的博客
07-19 3421
nv12 格式 nv12 转 rgb 两种格式 代码 void NV12_T_RGB(unsigned int width, unsigned int height, unsigned char* Y, unsigned char* UV, unsigned char* rgb) { int r, g, b; int y, u, v; for (int i = 0; i < height; i++) { for (int j = 0; j < width; j++) { y =
nv12——resize
songyimin1208的博客
09-08 2435
#include #include #include #include #include typedef unsigned char uint8_t;   /**   * @param src input nv12 raw data array    * @param dst output nv12 raw data result,   * the memor
利用opencl，读取并resize一张raw图(NV12格式)，指定quality的JPEG压缩，输出jpg图像,并统计耗时
07-09
本资源采用opencl的C++绑定。 首先，读取一张raw图（NV12格式）， 其次，利用gpiu对其进行resize并以指定quality对其进行JPEG压缩， 最后，输出一张jpg图像到本地。
利用opencl，读取一张raw图（NV12格式），并对其进行指定quality的JPEG压缩，输出jpg图像到本地,并统计耗时
07-09
本资源采用opencl的C++绑定，实现了对nv12数据的JPEG压缩，可自定义quality。
YUV和NV12T转RGB
09-22
YUV和NV12T转RGB的一个小工具，可以在linux和android上边使用。
RAW格式原图例如bayer转RGB，那么输入就是raw格式的。
11-22
raw格式图片是相机直接拍摄出的一种格式，一般看到的图像是经过转换格式的， 在开发图像处理相关的产品时，需要测试图像转码，例如bayer转RGB，那么输入就是raw格式的。
YUV（NV12 NV21等） RGB（16，24，32位） 原始数据（无压缩）的图片和视频查看工具
10-11
可以直接打开 YUV NV21 NV12 YUV420 YUV444 YUV422 RGB16 RGB24 RGB32 RGBA等格式的从内存中直接保存的原始无压缩图片
NV12代码及缩放说明文档
07-15
libyuv 库基于Linux平台的NV12图片缩放demos
通过封装的opencv函数进行yuv.nv12到jpg的转换并画框显示
12-17
1、读取yuv.nv12文件转换成jpg格式文件输出； 2、通过封装opencv函数进行画框输出；
C语言利用opencl实现图像处理 包括yuv(NV12)转rgb、remap、resize、crop等，并做了耗时统计
07-08
读取raw文件（NV12格式），转rgb+resize操作，输出jpg图像。 token=6；读取raw文件（NV12格式），转rgb+remap+resize操作，输出jpg图像。 token=7；读取raw文件（NV12格式），转rgb+remap+crop操作，输出jpg图像。 ...
opencl GPU加速rgb24与yuv444p格式相互转换核函数参考
05-23
在图像和视频处理领域，颜色空间转换是一种常见的操作，用于适应不同的显示设备和处理需求。OpenCL（Open Computing Language）是一种跨平台的并行编程框架，特别适合利用GPU（Graphics Processing Unit）的强大计算...
opencl GPU加速rgb图像缩放 核函数参考
05-23
3. **图像缩放算法**：OpenCL中实现图像缩放可以采用多种算法，例如最近邻插值、双线性插值或更高级的三次卷积插值。这些算法决定了新像素的颜色是如何根据原图像像素来计算的。例如，最近邻插值简单快速，但可能会...
openc中的图像的扩展缩放resize、图像的平移及旋转warpAffine、getRotationMatrix2D
weixin_45670380的博客
10-28 451
图像的扩展缩放 dst=cv2.resize(src，dsize[,dst[,fx[,fy[,interpation]]]]) dsize 目标图像的大小，如果后面有了比例因子fx,fy，就可以为None dst 一般放到前面去 interpation插值方法可以有在缩放时我们推荐使用 cv2.INTER_AREA,在扩展时我们推荐使用 v2.INTER_CUBIC(慢) 和 v2.INTER_L...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值