OpenCL 2.0 异构计算 [第三版] (中文)
  • Introduction
  • 序言
  • 第1章 简介异构计算
    • 1.1 关于异构计算
    • 1.2 本书目的
    • 1.3 并行思想
    • 1.4 并发和并行编程模型
    • 1.5 线程和共享内存
    • 1.6 消息通讯机制
    • 1.7 并行性的粒度
    • 1.8 使用OpenCL进行异构计算
    • 1.9 本书结构
  • 第2章 设备架构
    • 2.1 介绍
    • 2.2 硬件的权衡
    • 2.3 架构设计空间
    • 2.4 本章总结
  • 第3章 介绍OpenCL
    • 3.1 简介OpenCL
    • 3.2 OpenCL平台模型
    • 3.3 OpenCL执行模型
    • 3.4 内核和OpenCL编程模型
    • 3.5 OpenCL内存模型
    • 3.6 OpenCL运行时(例子)
    • 3.7 OpenCL C++ Wapper向量加法
    • 3.8 CUDA编程者使用OpenCL的注意事项
  • 第4章 OpenCL案例
    • 4.1 OpenCL实例
    • 4.2 直方图
    • 4.3 图像旋转
    • 4.4 图像卷积
    • 4.5 生产者-消费者
    • 4.6 基本功能函数
    • 4.7 本章总结
  • 第5章 OpenCL运行时和并发模型
    • 5.1 命令和排队模型
    • 5.2 多命令队列
    • 5.3 内核执行域:工作项、工作组和NDRange
    • 5.4 原生和内置内核
    • 5.5 设备端排队
    • 5.6 本章总结
  • 第6章 OpenCL主机端内存模型
    • 6.1 内存对象
    • 6.2 内存管理
    • 6.3 共享虚拟内存
    • 6.4 本章总结
  • 第7章 OpenCL设备端内存模型
    • 7.1 同步和交互
    • 7.2 全局内存
    • 7.3 常量内存
    • 7.4 局部内存
    • 7.5 私有内存
    • 7.6 统一地址空间
    • 7.7 内存序
    • 7.8 本章总结
  • 第8章 异构系统下解析OpenCL
    • 8.1 AMD FX-8350 CPU
    • 8.2 AMD RADEON R9 290X CPU
    • 8.3 OpenCL内存性能的考量
    • 8.4 本章总结
  • 第9章 案例分析:图像聚类
    • 9.1 图像聚类简介
    • 9.2 直方图的特性——CPU实现
    • 9.3 OpenCL实现
    • 9.4 性能分析
    • 9.5 本章总结
  • 第10章 OpenCL的分析和调试
    • 10.1 设置本章的原因
    • 10.2 使用事件分析OpenCL代码
    • 10.3 AMD CodeXL
    • 10.4 如何使用AMD CodeXL
    • 10.5 使用CodeXL分析内核
    • 10.6 使用CodeXL调试OpenCL内核
    • 10.7 使用printf调试
    • 10.8 本章总结
  • 第11章 高级语言映射到OpenCL2.0 —— 从编译器作者的角度
    • 11.1 简要介绍现状
    • 11.2 简单介绍C++ AMP
    • 11.3 编译器的目标 —— OpenCL 2.0
    • 11.4 C++ AMP与OpenCL对比
    • 11.5 C++ AMP的编译流
    • 11.6 编译之后的C++ AMP代码
    • 11.7 OpenCL 2.0提出共享虚拟内存的原因
    • 11.8 编译器怎样支持C++ AMP的线程块划分
    • 11.9 地址空间的推断
    • 11.10 优化数据搬运
    • 11.11 完整例子:二项式
    • 11.12 初步结果
    • 11.13 本章总结
  • 第12章 WebCL:使用OpenCL加速Web应用
    • 12.1 介绍WebCL
    • 12.2 如何使用WebCL编程
    • 12.3 同步机制
    • 12.4 WebCL的交互性
    • 12.5 应用实例
    • 12.6 增强安全性
    • 12.7 服务器端使用WebCL
    • 12.8 WebCL的状态和特性
  • 第13章 其他高级语言中OpenCL的使用
    • 13.1 本章简介
    • 13.2 越过C和C++
    • 13.3 Haskell中使用OpenCL
    • 13.4 本章总结
Powered by GitBook
On this page

Was this helpful?

  1. 第4章 OpenCL案例

4.4 图像卷积

Previous4.3 图像旋转Next4.5 生产者-消费者

Last updated 6 years ago

Was this helpful?

卷积在图像处理中经常用到,其会根据每个像素周围的像素点修改当前像素点的值。卷积核就是用来描述每个像素点如何被附近的像素点所影响。例如,模糊滤波中,使用平均的权重方式来进行计算,这样差异比较大的像素点会减少差异。对于相同的图像,如果想要做不同的操作,我们只需要变化滤波器即可,这样就能做锐化、模糊、边缘增强和图像压花。

卷积算法会遍历原始图像中的每一个像素点。对于每个原始像素点,滤波器中心点会为于像素点的上方,然后将中心点以及周围点的像素点与滤波器中对应的权重值相乘。然后将乘积的结果值相加后,产生出新的值作为输出。图4.3中就展示了该算法的具体过程。

图4.3 卷积滤波如何对原图像进行处理。

图4.4a是原始图,图4.4b中展示了原图经过模糊滤波后的结果,图4.4c展示了经过一个压花滤波器处理后的结果。

图4.4 不同卷积核对同一张图像进行处理:a)为原图;b)为模糊处理;c)为压花处理。

程序清单4.6中使用C/C++实现了一个串行的卷积操作。两层外部循环以遍历原始图像中所有的像素点。每一次滤波操作,每个原始点和其附近的点都要参与计算。需要注意的是,滤波器有可能访问到原始图像之外的区域。为了解决这个问题,我们在最内层循环中添加了四个显式的检查,当滤波器对应的坐标点位于原始图像之外,我们将使用与其最近的原始图像的边界值。

程序清单4.6 图像卷积的串行版本

OpenCL中,使用图像内存对象处理卷积操作要比使用数组内存对象更有优势。图像采样方式会自动对访问到图像之外的区域进行处理(类似于上节图像转换中所提到的),并且访问缓存中二维数据在硬件上也会有所优化(将会在第7章讨论)。

OpenCL上实现卷积操作几乎没有什么难度,并且写法类似于卷积操作的C版本。OpenCL版本中,我们为每一个输出的像素点创建了一个工作项,使用并行的方式将最外层两个循环去掉。那么每个工作项的任务就是完成最里面的两个循环,这两个循环完成的就是滤波操作。前面的例子中,读取源图像数据,需要配置一个OpenCL结构体,来指定数据的类型。本节的例子中,将继续使用read_imagef()。完整的内核代码将在代码清单4.7中展示。

kerenl void convolution( read_only image2d_t inputImage, write_only image2d_t outputImage, int rows, int cols, consetant float filter, int filterWidth, sampler_t sampler) { / Store each work-item's unique row and column */ int column = get_global_id(0); int row = get_global_id(1);

/* Half the width of the filter is needed for indexing

  • memory later */ int halfWidth = (int)(filterWidth / 2);

    /* All accesses to images return data as four-element vectors

  • (i.e., float4), although only the x component will contain

  • meaningful data in this code */ float4 sum = {0.0f,0.0f,0.0f,0.0f};

    / Iterator for the filter / int filterIdx = 0;

    /* Each work-item iterates around its local area on the basis of the

  • size of the filter*/ int2 coords; // Coordinates for accessing the image

    / Iterate the filter rows/ for (int i = -halfWidth; i <= halfWidth; i++) { coords.y = row + i; / Iterate over the filter columns / for (int j = -halfWidth; j <= halfWidth; j++) { coords.x - column + 1;

    /*Read a pixel from the image. A single-channel image

    • stores the pixel in the x coordinate of the reatured

    • vector. */

      pixel = read_imagef(inputImage, sampler, coords);

      sum.x += pixel.x * filter[filterIdx++];

    } }

    / Copy the data to the output image / coords.x = column; coords.y = row; write_imagef(outputImage, coords, sum); }

程序清单4.7 使用OpenCL C实现的图像卷积

访问图像总会返回具有四个值的向量(每个通道一个值)。前节例子中,我们加了.x来表示只需要返回值的第一个元素。本节例子中,我们会申明pixel(图像访问之后返回的值)和sum(存储结果数据,以拷贝到输出图像),其类型都为float4。当然,我们仅对x元素进行累加滤波像素值的计算(第45行)。

卷积核很适合放置到常量内存上,因为所有工作项所要用到的卷积核都是相同的。简单的添加关键字"__costant"在函数参数列表中(第7行),用于表示卷积核存放在常量内存中。

之前的例子中,我们是直接在内核内部创建了一个采样器。本节例子中,我们将使用主机端API创建一个采样器,并将其作为内核的一个参数传入。同样,本节例子中我们将使用C++ API(C++采样器构造函数需要相同的参数)。

主机端创建采样器的API如下所示:

cl_sampler clCreateSampler(
  cl_context context,
  cl_bool normalized_coords,
  cl_addressing_mode addressing_mode,
  cl_filter_mode filter_mode,
  cl_int *errcode_ret)

其C++ API如下所示:

cl::Sampler::Sampler(
  const Context &context,
  cl_bool normalized_coords,
  cl_addressing_mode addressing_mode,
  cl_filter_mode filter_mode,
  cl_int *err = NULL)

使用C++创建采样器的方式如下所示:

cl::Sampler sampler = new cl::Sampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST);

图像旋转例子中,采样器使用非标准化坐标。这里我们将展示不同于内部使用时的另外两个采样器参数:滤波模式将使用最近的像素点的值,而非进行差值后的值(CL_FILTER_NEAREST),并且寻址模式将在访问到图像区域之外时,将其值置为最接近的图像边界值(CL_ADDRESS_CLAMP_TO_EDGE)。(注:这里要注意一下CL和CLK的区别。CL开头的是使用在主机端API上,CLK则直接使用在OpenCL内核代码中。)

使用C++ API,创建二维图像使用image2D类进行创建,创建这个类需要一个ImageFormat对象作为参数。C API则不需要图像描述器的传入。

Image2D和ImageFormat的构造函数如下:

cl::Image2D::Image2D(
  Context& context,
  cl_mem_flags flags,
  ImageFormat format,
  ::size_t width,
  ::size_t height,
  ::size_t row_pitch = 0,
  void *host_ptr = NULL,
  cl_int *err = NULL)

cl::ImageFormat::ImageFormat(
  cl_channel_order order,
  cl_channel_type type)

这样我们就能创建卷积所需要的输入和输出图像,其调用方式如下所示:

cl::ImageFormat imageFormat = cl::ImageFormat(CL_R, CL_FLOAT);
cl::Image2D inputImage = cl::Image2D(context, CL_MEM_READ_ONLY, imageFormat, imageCols, imageRows);
cl::Image2D outputImage = cl::Image2D(context, CL_MEM_WRITE_ONLY, imageFormat, imageCols, imageRows);

使用C++ API实现图像卷积的主机端代码在代码清单4.8中展示。主程序中,使用了一个5x5的高斯模糊滤波核用于卷积处理。

程序清单4.8 图像卷积主机端完整代码