• OpenCL:图像处理基础note


    使用图像对象的理由

    虽然对于图像也可以把它的像素数据当做一般的缓存数据来处理,但是如果把它当做图像来处理有如下好处:

    1. 在GPU中,图像数据是保存在特殊的全局内存中,即纹理内存,它和一般的全局内存不相同,它是被缓存的,用于高速访问处理。GPU中有专门支持图像读写的硬件,使用内置读写函数可以充分发挥这个优势
    2. 只要OpenCL支持该图像格式,那么就可以不用考虑图像格式的前提下使用读写图像数据的函数
    3. 可以使用采样器来配置读取图像中数据的的方式
    4. OpenCL提供函数来获取图像相关信息,比如宽度等

    主机与内核的命名

    图像相关应用主要涉及两个数据类型,即图像对象采样器。图像对线用来保存主机和设备上的图像对象,而采样器则在设备接收数据时,说明如何读取这些颜色值

    图像对象 采样器
    主机 cl_mem (和一般缓存对象一样) cl_sampler
    设备 image2d_t或者image3d_t sampler_t

    主机编程的主要接口

    创建图像对象

    cl_mem clCreateImage2D()
    cl_mem clCreateImage3D()

    图像格式

    cl_image_format

    创建采样器

    cl_sampler clCreateSampler()

    设备上的接口

    图像数据

    由于很多设备都是将图像对象保存在特定的内存中,因此image2d_t以及image3d_t数据的前面通常会有read_only或者write_only等修饰符。图像不能即可写又可读
    图像数据作为内核的参数时,和普通的缓存对象不同,它不是指针参数,因为图像数据并不希望直接通过内存操作来访问数据。

    采样器

    根据khronos官网的说明,采样器其实就是一个uint类型(可以直接给内置cl读写函数传递标识位的组合来替代sampler_t,例如:float4 clr = read_imagef(img, CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST, (float2)(x, y));),通过bit-mode来说明坐标模式,超出边界的处理模式以及插值模式。在官方的sample中采用的是constant申明的模式,虽然也说了可以采用global声明,但这也许说明constant更好。
    可以通过设置参数clSetKernelArg的方式来给内核传递采样器。但是更为方便的方式是在cl文件中,在内核函数之前创建一个sampler_t对象,如下case:

    __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | 
          CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; 
    
    __kernel void simple_image(read_only image2d_t src_image,
                            write_only image2d_t dst_image) {
    
       /* Compute value to be subtracted from each pixel */
       uint offset = get_global_id(1) * 0x4000 + get_global_id(0) * 0x1000;
    
       /* Read pixel value */
       int2 coord = (int2)(get_global_id(0), get_global_id(1));
       uint4 pixel = read_imageui(src_image, sampler, coord);
    
       /* Subtract offset from pixel */
       pixel.x -= offset;
    
       /* Write new pixel value to output */
       write_imageui(dst_image, coord, pixel);
    }
    

    内置图像函数

    OpenCL的内置图像相关的函数主要包括三类:read_imageT读取类,write_imageT写入类以及get_image_X读取图像信息类。其中T表示数据类型,比如f,i,ui等,而X表示width,dim等图像信息

    read_image
    • read_image函数返回的都是四分量的数据,如果实际图片没有四个通道的数据,那么那些不含有的的通道对应的分量会被置为0,具体请看官网
    • 很多时候都会看到代码中用float4 read_imagef()去读图像数据为int或者uint类型的image,可能你会困惑为什么整形要用float去读?这是因为,这些图片在创建的时候就是以正则NORM化的数据类型,这个类型在cl_image_formatimage_channel_data_type中指定。以int8为例,image_channel_data_type可取的值如下表:参考
    image_channel_data_type 参数 含义 read_image
    CL_SNORM_INT8 每个通道都是正则化符号8位整数 使用read_imagef读取,返回的float4的每个分量的取值范围为[0.0, 1.0]
    CL_UNORM_INT8 每个通道都是正则化符号8位整数 使用read_imagef读取,返回的float4的每个分量的取值范围为[-1.0, 1.0]
    CL_SIGNED_INT8 每个通道都是8位有符号的整数 使用read_imagei读取,返回int4
    CL_USIGNED_INT8 每个通道都是8位无符号的整数 使用read_imageui读取,返回unsigned int4
    • 使用read_imageiread_imageui时,采样器的filter_mode必须设置为CLK_FILTER_NEAREST
    • 使用read_imagef时,如果坐标是整数,那么filte_mode必须设置为CLK_FILTER_NEAREST
    • 使用read_imageX时,X表示**,如果坐标是整数,那么采样器的坐标正则化必须设置为CLK_NORMALIZED_COORDS_FALSE,且address mode必须设置为CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP 或者 CLK_ADDRESS_NONE
  • 相关阅读:
    linux软件名规则
    给php开启mysql扩展
    centos6可用的Apache管理脚本
    Centos下设置redis开机自启动
    拆卸mysql
    如何判断是否在一个网路中
    linux 如何清理僵尸进程
    如何查找僵尸进程并Kill之,杀不掉的要查看父进程并杀之
    一张图告诉你php的命名空间和自动加载
    PHP的文件加载机制到底是什么目录
  • 原文地址:https://www.cnblogs.com/willhua/p/9672135.html
Copyright © 2020-2023  润新知